diff --git a/test/ck_tile/gemm/test_gemm_pipeline_kernel_types.hpp b/test/ck_tile/gemm/test_gemm_pipeline_kernel_types.hpp index 7f5d41a99a..fe80b6f22f 100644 --- a/test/ck_tile/gemm/test_gemm_pipeline_kernel_types.hpp +++ b/test/ck_tile/gemm/test_gemm_pipeline_kernel_types.hpp @@ -79,23 +79,10 @@ using KernelTypesMemWmma = ::testing::Types< >; using KernelTypesCompV3 = ::testing::Types< -<<<<<<< HEAD -<<<<<<< HEAD -<<<<<<< HEAD std::tuple< Row, Row, Col, F16, F16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3> //std::tuple< Row, Row, Row, F16, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, //std::tuple< Row, Row, Row, BF16, BF16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, //std::tuple< Row, Row, Row, BF16, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, -======= - std::tuple< Row, Row, Col, F16, F16, F32, F16, I128, I128, I64, I32, I32, I16, Intrawave, CompV3>//, - //std::tuple< Row, Row, Row, F8, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, - //std::tuple< Row, Col, Row, F16, F16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, - //std::tuple< Row, Col, Row, F8, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, - //std::tuple< Col, Row, Row, F16, F16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, - //std::tuple< Col, Row, Row, F8, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, - //std::tuple< Col, Col, Row, F16, F16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, - //std::tuple< Col, Col, Row, F8, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, ->>>>>>> 0a9ceadca ([CK_TILE] working version) //std::tuple< Row, Row, Row, INT8, INT8, INT32, INT32, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, //std::tuple< Row, Row, Row, F8, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, //std::tuple< Row, Row, Row, F8, BF8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, @@ -132,36 +119,6 @@ using KernelTypesCompV3 = ::testing::Types< //std::tuple< Col, Col, Row, F8, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, //std::tuple< Col, Col, Row, BF8, BF8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, //std::tuple< Col, Col, Row, BF8, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3> -======= - std::tuple< Row, Row, Col, F16, F16, F32, F16, I128, I128, I64, I32, I32, I16, Intrawave, CompV3>, - std::tuple< Row, Row, Col, F8, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, -======= - std::tuple< Row, Row, Col, F16, F16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, ->>>>>>> 05ce4e524 ([CK_TILE] TEST vector stores c col layout part1) - std::tuple< Row, Col, Col, F16, F16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, - std::tuple< Col, Row, Col, F16, F16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, -<<<<<<< HEAD - std::tuple< Col, Row, Col, F8, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, - std::tuple< Col, Col, Col, F16, F16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, - std::tuple< Col, Col, Col, F8, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, - std::tuple< Row, Row, Col, INT8, INT8, INT32, INT32, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, - std::tuple< Row, Col, Col, INT8, INT8, INT32, INT32, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, - std::tuple< Col, Row, Col, INT8, INT8, INT32, INT32, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, - std::tuple< Col, Col, Col, INT8, INT8, INT32, INT32, I256, I256, I64, I32, I32, I16, Intrawave, CompV3> ->>>>>>> 12c48382b ([CK_TILE] working version and tests) -======= - std::tuple< Col, Col, Col, F16, F16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3> - // - //std::tuple< Row, Row, Col, F8, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, - //std::tuple< Row, Col, Col, F8, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, - //std::tuple< Col, Row, Col, F8, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, - //std::tuple< Col, Col, Col, F8, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, - // - //std::tuple< Row, Row, Col, INT8, INT8, INT32, INT32, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, - //std::tuple< Row, Col, Col, INT8, INT8, INT32, INT32, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, - //std::tuple< Col, Row, Col, INT8, INT8, INT32, INT32, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, - //std::tuple< Col, Col, Col, INT8, INT8, INT32, INT32, I256, I256, I64, I32, I32, I16, Intrawave, CompV3> ->>>>>>> 05ce4e524 ([CK_TILE] TEST vector stores c col layout part1) >; using KernelTypesCompV3Wmma = ::testing::Types< diff --git a/test/ck_tile/gemm/test_gemm_pipeline_util.hpp b/test/ck_tile/gemm/test_gemm_pipeline_util.hpp index db520bbcb6..108676023a 100644 --- a/test/ck_tile/gemm/test_gemm_pipeline_util.hpp +++ b/test/ck_tile/gemm/test_gemm_pipeline_util.hpp @@ -318,7 +318,7 @@ class TestCkTileGemmPipeline : public ::testing::Test } } - template + template void RunSingle(const int M, const int N, const int K, @@ -334,58 +334,12 @@ class TestCkTileGemmPipeline : public ::testing::Test ck_tile::index_t stride_C = ck_tile::get_default_stride(M, N, StrideC, is_row_major(CLayout{})); - auto f_host_tensor_descriptor = [](std::size_t row, - std::size_t col, - std::size_t stride, - auto layout) { - if constexpr(std::is_same_v) - { - return ck_tile::HostTensorDescriptor({row, col}, {stride, 1_uz}); - } - else - { - return ck_tile::HostTensorDescriptor({col, row}, {stride, 1_uz}); - } - }; - - auto f_get_default_stride = - [](std::size_t row, std::size_t col, std::size_t stride, auto layout) { - if(stride == 0) - { - // give a chance if stride is zero, return a default packed stride - if constexpr(std::is_same_v) - { - return col; - } - else - { - return row; - } - } - else - return stride; - }; - - ck_tile::index_t stride_A = f_get_default_stride(M, K, StrideA, ALayout{}); - ck_tile::index_t stride_B = f_get_default_stride(K, N, StrideB, BLayout{}); - ck_tile::index_t stride_C = f_get_default_stride(M, N, StrideC, CLayout{}); - - ck_tile::HostTensor a_m_k(f_host_tensor_descriptor(M, K, stride_A, ALayout{})); - ck_tile::HostTensor b_k_n(f_host_tensor_descriptor(K, N, stride_B, BLayout{})); + ck_tile::HostTensor a_m_k( + ck_tile::host_tensor_descriptor(M, K, stride_A, is_row_major(ALayout{}))); + ck_tile::HostTensor b_k_n( + ck_tile::host_tensor_descriptor(K, N, stride_B, is_row_major(BLayout{}))); ck_tile::HostTensor c_m_n_dev_result( -<<<<<<< HEAD ck_tile::host_tensor_descriptor(M, N, stride_C, is_row_major(CLayout{}))); -======= - f_host_tensor_descriptor(M, N, stride_C, CLayout{})); ->>>>>>> 05ce4e524 ([CK_TILE] TEST vector stores c col layout part1) - - //std::cout << "a_m_k: "; - //a_m_k.print_first_n(std::cout) << '\n'; - //std::cout << "b_k_n: "; - //b_k_n.print_first_n(std::cout) << '\n'; - //std::cout << "c_m_n_dev_result: "; - //c_m_n_dev_result.print_first_n(std::cout) << '\n'; ck_tile::FillUniformDistributionIntegerValue{-5, 5, 11939}(a_m_k); ck_tile::FillUniformDistributionIntegerValue{-5, 5, 11940}(b_k_n); @@ -421,7 +375,7 @@ class TestCkTileGemmPipeline : public ::testing::Test stride_B, stride_C}; - invoke_gemm(args, ck_tile::stream_config{nullptr, false, 2}); + invoke_gemm(args, ck_tile::stream_config{nullptr, false}); c_m_n_dev_buf.FromDevice(c_m_n_dev_result.data()); bool pass = true; @@ -433,13 +387,6 @@ class TestCkTileGemmPipeline : public ::testing::Test ck_tile::reference_gemm( a_m_k, b_k_n, c_m_n_host_ref); - //std::cout << "a_m_k: "; - //a_m_k.print_first_n(std::cout) << '\n'; - //std::cout << "b_k_n: "; - //b_k_n.print_first_n(std::cout) << '\n'; - //std::cout << "c_m_n_dev_result: "; - //c_m_n_dev_result.print_first_n(std::cout) << '\n'; - const float max_accumulated_value = *std::max_element(c_m_n_host_ref.mData.begin(), c_m_n_host_ref.mData.end()); const auto rtol_atol = calculate_rtol_atol(