diff --git a/test/ck_tile/batched_gemm/CMakeLists.txt b/test/ck_tile/batched_gemm/CMakeLists.txt index 6f29225291..6509113a09 100644 --- a/test/ck_tile/batched_gemm/CMakeLists.txt +++ b/test/ck_tile/batched_gemm/CMakeLists.txt @@ -2,5 +2,66 @@ # SPDX-License-Identifier: MIT if(GPU_TARGETS MATCHES "gfx9|gfx11|gfx12") - add_gtest_executable(test_ck_tile_batched_gemm test_batched_gemm.cpp) + + add_custom_target(test_ck_tile_batched_gemm) + + add_gtest_executable(test_ck_tile_batched_gemm_f16 test_batched_gemm_f16.cpp) + add_gtest_executable(test_ck_tile_batched_gemm_bf16 test_batched_gemm_bf16.cpp) + + add_dependencies(test_ck_tile_batched_gemm + test_ck_tile_batched_gemm_f16 + test_ck_tile_batched_gemm_bf16) + + # FP8 / BF8 batched GEMM relies on the fp8/bf8 MFMA or WMMA paths. + # - gfx908 / gfx90a use the scalar f32 fallback in + # include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma_impl.hpp. + # - gfx942 / gfx950 use the native v_mfma_f32_32x32x16_{fp8,bf8}_{fp8,bf8} instructions. + # - gfx12 uses the native v_wmma_f32_16x16x16_{fp8,bf8}_{fp8,bf8}_w32_gfx12 instructions + # (see include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl_8bit_traits.hpp). + # - gfx11 has no native fp8/bf8 WMMA instruction and no software fallback, so the + # warp-gemm silently returns CVecType{0} and the kernel produces all-zero output. + # Skip the test there until an fp8/bf8 path is implemented for gfx11. + # + # On gfx950 / gfx12 the FP8/BF8 hardware uses the OCP encoding (bias 15 for + # E5M2, bias 7 for E4M3). On gfx94x it uses FNUZ (bias 16 / 8). The device + # code in include/ck_tile/core/numeric/float8.hpp picks the encoding from + # CK_TILE_USE_OCP_FP8, which in config.hpp defaults to 1 only inside + # __HIP_DEVICE_COMPILE__. For the host-side reference_gemm and tensor fill + # to use the same encoding as the device kernel (otherwise A/B and the + # reference output are produced in FNUZ while the device interprets the + # same bytes as OCP, yielding a 4x error from the off-by-one bias on each + # operand), propagate -DCK_TILE_USE_OCP_FP8 to host code when the global + # CK_USE_OCP_FP8 toggle is on (set by the top-level CMakeLists.txt for + # gfx950 / gfx12). + if(GPU_TARGETS MATCHES "gfx9|gfx12") + set(CK_TILE_BATCHED_GEMM_F8_COMPILE_OPTIONS) + if(CK_USE_OCP_FP8) + list(APPEND CK_TILE_BATCHED_GEMM_F8_COMPILE_OPTIONS -DCK_TILE_USE_OCP_FP8) + endif() + add_gtest_executable(test_ck_tile_batched_gemm_f8 test_batched_gemm_f8.cpp) + add_gtest_executable(test_ck_tile_batched_gemm_bf8 test_batched_gemm_bf8.cpp) + target_compile_options(test_ck_tile_batched_gemm_f8 + PRIVATE ${CK_TILE_BATCHED_GEMM_F8_COMPILE_OPTIONS}) + target_compile_options(test_ck_tile_batched_gemm_bf8 + PRIVATE ${CK_TILE_BATCHED_GEMM_F8_COMPILE_OPTIONS}) + if(TARGET test_ck_tile_batched_gemm_f8) + add_dependencies(test_ck_tile_batched_gemm test_ck_tile_batched_gemm_f8) + endif() + if(TARGET test_ck_tile_batched_gemm_bf8) + add_dependencies(test_ck_tile_batched_gemm test_ck_tile_batched_gemm_bf8) + endif() + endif() + + # INT8 batched GEMM relies on the int32 MFMA / WMMA paths. + # On gfx908 / gfx90a there is no native v_mfma_i32_32x32x16_i8 instruction and + # the warp-gemm impl in include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma_impl.hpp + # falls back to v_mfma_f32_32x32x2f32, writing fp32 bit patterns into an int32_t + # accumulator and ultimately into the int32_t C buffer (see how the existing + # test_ck_tile_gemm_pipeline_compv3 INT8 cases are restricted to gfx94|gfx95). + if(GPU_TARGETS MATCHES "gfx94|gfx95|gfx11|gfx12") + add_gtest_executable(test_ck_tile_batched_gemm_int8 test_batched_gemm_int8.cpp) + if(TARGET test_ck_tile_batched_gemm_int8) + add_dependencies(test_ck_tile_batched_gemm test_ck_tile_batched_gemm_int8) + endif() + endif() endif() diff --git a/test/ck_tile/batched_gemm/test_batched_gemm_bf16.cpp b/test/ck_tile/batched_gemm/test_batched_gemm_bf16.cpp new file mode 100644 index 0000000000..0305fe779e --- /dev/null +++ b/test/ck_tile/batched_gemm/test_batched_gemm_bf16.cpp @@ -0,0 +1,37 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include + +#include "gtest/gtest.h" + +#include "ck_tile/host.hpp" +#include "test_batched_gemm_util.hpp" + +using F16 = ck_tile::half_t; +using F32 = float; +using BF16 = ck_tile::bf16_t; + +using Row = ck_tile::tensor_layout::gemm::RowMajor; +using Col = ck_tile::tensor_layout::gemm::ColumnMajor; + +// clang-format off +using KernelTypes = ::testing::Types< + // ALayout, BLayout, CLayout, ADataType, BDataType, AccDataType, CDataType + std::tuple< Row, Row, Row, BF16, BF16, F32, F16>, + std::tuple< Col, Row, Row, BF16, BF16, F32, F16>, + std::tuple< Row, Col, Row, BF16, BF16, F32, F16>, + std::tuple< Col, Col, Row, BF16, BF16, F32, F16> + >; +// clang-format on + +template +class TestCkTileBatchedGemmBF16 : public TestCkTileBatchedGemm +{ +}; + +TYPED_TEST_SUITE(TestCkTileBatchedGemmBF16, KernelTypes); + +#define TEST_CKTILE_BGEMM_SUITE_NAME TestCkTileBatchedGemmBF16 + +#include "test_batched_gemm_ut_cases.inc" diff --git a/test/ck_tile/batched_gemm/test_batched_gemm_bf8.cpp b/test/ck_tile/batched_gemm/test_batched_gemm_bf8.cpp new file mode 100644 index 0000000000..0b14c25a60 --- /dev/null +++ b/test/ck_tile/batched_gemm/test_batched_gemm_bf8.cpp @@ -0,0 +1,37 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include + +#include "gtest/gtest.h" + +#include "ck_tile/host.hpp" +#include "test_batched_gemm_util.hpp" + +using F16 = ck_tile::half_t; +using F32 = float; +using BF8 = ck_tile::bf8_t; + +using Row = ck_tile::tensor_layout::gemm::RowMajor; +using Col = ck_tile::tensor_layout::gemm::ColumnMajor; + +// clang-format off +using KernelTypes = ::testing::Types< + // ALayout, BLayout, CLayout, ADataType, BDataType, AccDataType, CDataType + std::tuple< Row, Row, Row, BF8, BF8, F32, F16>, + std::tuple< Col, Row, Row, BF8, BF8, F32, F16>, + std::tuple< Row, Col, Row, BF8, BF8, F32, F16>, + std::tuple< Col, Col, Row, BF8, BF8, F32, F16> + >; +// clang-format on + +template +class TestCkTileBatchedGemmBF8 : public TestCkTileBatchedGemm +{ +}; + +TYPED_TEST_SUITE(TestCkTileBatchedGemmBF8, KernelTypes); + +#define TEST_CKTILE_BGEMM_SUITE_NAME TestCkTileBatchedGemmBF8 + +#include "test_batched_gemm_ut_cases.inc" diff --git a/test/ck_tile/batched_gemm/test_batched_gemm.cpp b/test/ck_tile/batched_gemm/test_batched_gemm_f16.cpp similarity index 82% rename from test/ck_tile/batched_gemm/test_batched_gemm.cpp rename to test/ck_tile/batched_gemm/test_batched_gemm_f16.cpp index 9e28c96da8..825bb3d302 100644 --- a/test/ck_tile/batched_gemm/test_batched_gemm.cpp +++ b/test/ck_tile/batched_gemm/test_batched_gemm_f16.cpp @@ -24,6 +24,13 @@ using KernelTypes = ::testing::Types< >; // clang-format on -TYPED_TEST_SUITE(TestCkTileBatchedGemm, KernelTypes); +template +class TestCkTileBatchedGemmF16 : public TestCkTileBatchedGemm +{ +}; + +TYPED_TEST_SUITE(TestCkTileBatchedGemmF16, KernelTypes); + +#define TEST_CKTILE_BGEMM_SUITE_NAME TestCkTileBatchedGemmF16 #include "test_batched_gemm_ut_cases.inc" diff --git a/test/ck_tile/batched_gemm/test_batched_gemm_f8.cpp b/test/ck_tile/batched_gemm/test_batched_gemm_f8.cpp new file mode 100644 index 0000000000..cbaaa80dd5 --- /dev/null +++ b/test/ck_tile/batched_gemm/test_batched_gemm_f8.cpp @@ -0,0 +1,42 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include + +#include "gtest/gtest.h" + +#include "ck_tile/host.hpp" +#include "test_batched_gemm_util.hpp" + +using BF16 = ck_tile::bf16_t; +using F16 = ck_tile::half_t; +using F32 = float; +using F8 = ck_tile::fp8_t; + +using Row = ck_tile::tensor_layout::gemm::RowMajor; +using Col = ck_tile::tensor_layout::gemm::ColumnMajor; + +// clang-format off +using KernelTypes = ::testing::Types< + // ALayout, BLayout, CLayout, ADataType, BDataType, AccDataType, CDataType + std::tuple< Row, Row, Row, F8, F8, F32, F16>, + std::tuple< Col, Row, Row, F8, F8, F32, F16>, + std::tuple< Row, Col, Row, F8, F8, F32, F16>, + std::tuple< Col, Col, Row, F8, F8, F32, F16>, + std::tuple< Row, Row, Row, F8, F8, F32, BF16>, + std::tuple< Col, Row, Row, F8, F8, F32, BF16>, + std::tuple< Row, Col, Row, F8, F8, F32, BF16>, + std::tuple< Col, Col, Row, F8, F8, F32, BF16> + >; +// clang-format on + +template +class TestCkTileBatchedGemmF8 : public TestCkTileBatchedGemm +{ +}; + +TYPED_TEST_SUITE(TestCkTileBatchedGemmF8, KernelTypes); + +#define TEST_CKTILE_BGEMM_SUITE_NAME TestCkTileBatchedGemmF8 + +#include "test_batched_gemm_ut_cases.inc" diff --git a/test/ck_tile/batched_gemm/test_batched_gemm_int8.cpp b/test/ck_tile/batched_gemm/test_batched_gemm_int8.cpp new file mode 100644 index 0000000000..dde39b5992 --- /dev/null +++ b/test/ck_tile/batched_gemm/test_batched_gemm_int8.cpp @@ -0,0 +1,36 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include + +#include "gtest/gtest.h" + +#include "ck_tile/host.hpp" +#include "test_batched_gemm_util.hpp" + +using INT8 = ck_tile::int8_t; +using INT32 = ck_tile::int32_t; + +using Row = ck_tile::tensor_layout::gemm::RowMajor; +using Col = ck_tile::tensor_layout::gemm::ColumnMajor; + +// clang-format off +using KernelTypes = ::testing::Types< + // ALayout, BLayout, CLayout, ADataType, BDataType, AccDataType, CDataType + std::tuple< Row, Row, Row, INT8, INT8, INT32, INT32>, + std::tuple< Col, Row, Row, INT8, INT8, INT32, INT32>, + std::tuple< Row, Col, Row, INT8, INT8, INT32, INT32>, + std::tuple< Col, Col, Row, INT8, INT8, INT32, INT32> + >; +// clang-format on + +template +class TestCkTileBatchedGemmInt8 : public TestCkTileBatchedGemm +{ +}; + +TYPED_TEST_SUITE(TestCkTileBatchedGemmInt8, KernelTypes); + +#define TEST_CKTILE_BGEMM_SUITE_NAME TestCkTileBatchedGemmInt8 + +#include "test_batched_gemm_ut_cases.inc" diff --git a/test/ck_tile/batched_gemm/test_batched_gemm_ut_cases.inc b/test/ck_tile/batched_gemm/test_batched_gemm_ut_cases.inc index 8c07b78ce4..bcda4b766b 100644 --- a/test/ck_tile/batched_gemm/test_batched_gemm_ut_cases.inc +++ b/test/ck_tile/batched_gemm/test_batched_gemm_ut_cases.inc @@ -21,7 +21,7 @@ struct StrideConfig int batchStrideC; }; -TYPED_TEST(TestCkTileBatchedGemm, Basic) +TYPED_TEST(TEST_CKTILE_BGEMM_SUITE_NAME, Basic) { std::vector gemmParams{{256, 256, 256, 1}, {256, 256, 256, 2}, diff --git a/test/ck_tile/batched_gemm/test_batched_gemm_util.hpp b/test/ck_tile/batched_gemm/test_batched_gemm_util.hpp index 557dd295cf..ef28d8da5a 100644 --- a/test/ck_tile/batched_gemm/test_batched_gemm_util.hpp +++ b/test/ck_tile/batched_gemm/test_batched_gemm_util.hpp @@ -1,7 +1,11 @@ // Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT #pragma once +#include +#include +#include #include +#include #include #include "ck_tile/core.hpp" @@ -261,9 +265,23 @@ class TestCkTileBatchedGemm : public ::testing::Test ck_tile::reference_batched_gemm( a_m_k, b_n_k, c_m_n_host_ref); - constexpr double rtol = 2e-3; - constexpr double atol = 2e-3; - pass = ck_tile::check_err( + using ComputeType = + std::conditional_t; + + constexpr ck_tile::index_t kbatch = 1; + + const float max_accumulated_value = std::accumulate( + c_m_n_host_ref.mData.begin(), c_m_n_host_ref.mData.end(), 0.0f, [](float acc, auto v) { + return std::max(acc, std::abs(static_cast(v))); + }); + + const auto rtol = ck_tile::get_relative_threshold( + ck_tile::integer_divide_ceil(K, kbatch)); + const auto atol = ck_tile::get_absolute_threshold( + static_cast(max_accumulated_value > 0 ? max_accumulated_value : 1.0f) / kbatch, + ck_tile::integer_divide_ceil(K, kbatch)); + + pass = ck_tile::check_err( c_m_n_dev_result, c_m_n_host_ref, "Error: Incorrect results!", rtol, atol); EXPECT_TRUE(pass); } diff --git a/test/ck_tile/grouped_gemm/CMakeLists.txt b/test/ck_tile/grouped_gemm/CMakeLists.txt index 8a5c8d33ba..efe6cd7007 100644 --- a/test/ck_tile/grouped_gemm/CMakeLists.txt +++ b/test/ck_tile/grouped_gemm/CMakeLists.txt @@ -4,18 +4,64 @@ if(GPU_TARGETS MATCHES "gfx9|gfx11|gfx12") add_custom_target(test_ck_tile_grouped_gemm) - + add_gtest_executable(test_ck_tile_grouped_gemm_f16 test_grouped_gemm_f16.cpp) add_gtest_executable(test_ck_tile_grouped_gemm_bf16 test_grouped_gemm_bf16.cpp) - if(GPU_TARGETS MATCHES "gfx1250") - add_gtest_executable(test_ck_tile_grouped_gemm_f8 test_grouped_gemm_f8.cpp) - endif() - + add_dependencies(test_ck_tile_grouped_gemm test_ck_tile_grouped_gemm_f16 test_ck_tile_grouped_gemm_bf16) - if(GPU_TARGETS MATCHES "gfx1250") - add_dependencies(test_ck_tile_grouped_gemm - test_ck_tile_grouped_gemm_f8) + + # FP8 / BF8 grouped GEMM relies on the fp8/bf8 MFMA or WMMA paths. + # - gfx908 / gfx90a use the scalar f32 fallback in + # include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma_impl.hpp. + # - gfx942 / gfx950 use the native v_mfma_f32_32x32x16_{fp8,bf8}_{fp8,bf8} instructions. + # - gfx12 uses the native v_wmma_f32_16x16x16_{fp8,bf8}_{fp8,bf8}_w32_gfx12 instructions + # (see include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl_8bit_traits.hpp). + # - gfx11 has no native fp8/bf8 WMMA instruction and no software fallback, so the + # warp-gemm silently returns CVecType{0} and the kernel produces all-zero output. + # Skip the test there until an fp8/bf8 path is implemented for gfx11. + # + # On gfx950 / gfx12 the FP8/BF8 hardware uses the OCP encoding (bias 15 for + # E5M2, bias 7 for E4M3). On gfx94x it uses FNUZ (bias 16 / 8). The device + # code in include/ck_tile/core/numeric/float8.hpp picks the encoding from + # CK_TILE_USE_OCP_FP8, which in config.hpp defaults to 1 only inside + # __HIP_DEVICE_COMPILE__. For the host-side reference_gemm and tensor fill + # to use the same encoding as the device kernel (otherwise A/B and the + # reference output are produced in FNUZ while the device interprets the + # same bytes as OCP, yielding a 4x error from the off-by-one bias on each + # operand), propagate -DCK_TILE_USE_OCP_FP8 to host code when the global + # CK_USE_OCP_FP8 toggle is on (set by the top-level CMakeLists.txt for + # gfx950 / gfx12). + if(GPU_TARGETS MATCHES "gfx9|gfx12") + set(CK_TILE_GROUPED_GEMM_F8_COMPILE_OPTIONS) + if(CK_USE_OCP_FP8) + list(APPEND CK_TILE_GROUPED_GEMM_F8_COMPILE_OPTIONS -DCK_TILE_USE_OCP_FP8) + endif() + add_gtest_executable(test_ck_tile_grouped_gemm_f8 test_grouped_gemm_f8.cpp) + add_gtest_executable(test_ck_tile_grouped_gemm_bf8 test_grouped_gemm_bf8.cpp) + target_compile_options(test_ck_tile_grouped_gemm_f8 + PRIVATE ${CK_TILE_GROUPED_GEMM_F8_COMPILE_OPTIONS}) + target_compile_options(test_ck_tile_grouped_gemm_bf8 + PRIVATE ${CK_TILE_GROUPED_GEMM_F8_COMPILE_OPTIONS}) + if(TARGET test_ck_tile_grouped_gemm_f8) + add_dependencies(test_ck_tile_grouped_gemm test_ck_tile_grouped_gemm_f8) + endif() + if(TARGET test_ck_tile_grouped_gemm_bf8) + add_dependencies(test_ck_tile_grouped_gemm test_ck_tile_grouped_gemm_bf8) + endif() + endif() + + # INT8 grouped GEMM relies on the int32 MFMA / WMMA paths. + # On gfx908 / gfx90a there is no native v_mfma_i32_32x32x16_i8 instruction and + # the warp-gemm impl in include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma_impl.hpp + # falls back to v_mfma_f32_32x32x2f32, writing fp32 bit patterns into an int32_t + # accumulator and ultimately into the int32_t C buffer (this is why the + # test_ck_tile_batched_gemm_int8 cases are also restricted to gfx94|gfx95|gfx11|gfx12). + if(GPU_TARGETS MATCHES "gfx94|gfx95|gfx11|gfx12") + add_gtest_executable(test_ck_tile_grouped_gemm_int8 test_grouped_gemm_int8.cpp) + if(TARGET test_ck_tile_grouped_gemm_int8) + add_dependencies(test_ck_tile_grouped_gemm test_ck_tile_grouped_gemm_int8) + endif() endif() endif() diff --git a/test/ck_tile/grouped_gemm/test_grouped_gemm_bf8.cpp b/test/ck_tile/grouped_gemm/test_grouped_gemm_bf8.cpp new file mode 100644 index 0000000000..e86abeb02c --- /dev/null +++ b/test/ck_tile/grouped_gemm/test_grouped_gemm_bf8.cpp @@ -0,0 +1,42 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include + +#include "gtest/gtest.h" + +#include "ck_tile/host.hpp" +#include "test_grouped_gemm_util.hpp" + +using F16 = ck_tile::half_t; +using F32 = float; +using BF8 = ck_tile::bf8_t; +using Row = ck_tile::tensor_layout::gemm::RowMajor; +using Col = ck_tile::tensor_layout::gemm::ColumnMajor; +using True = ck_tile::bool_constant; +using False = ck_tile::bool_constant; + +// clang-format off +using KernelTypes = ::testing::Types< + // ALayout, BLayout, CLayout, ADataType, BDataType, AccDataType, CDataType, Persistent + std::tuple< Row, Col, Row, BF8, BF8, F32, F16, True>, + std::tuple< Row, Col, Row, BF8, BF8, F32, F16, False>, + std::tuple< Col, Col, Row, BF8, BF8, F32, F16, True>, + std::tuple< Col, Col, Row, BF8, BF8, F32, F16, False>, + std::tuple< Row, Row, Row, BF8, BF8, F32, F16, True>, + std::tuple< Row, Row, Row, BF8, BF8, F32, F16, False>, + std::tuple< Col, Row, Row, BF8, BF8, F32, F16, True>, + std::tuple< Col, Row, Row, BF8, BF8, F32, F16, False> + >; +// clang-format on + +template +class TestCkTileGroupedGemmBF8 : public TestCkTileGroupedGemm +{ +}; + +TYPED_TEST_SUITE(TestCkTileGroupedGemmBF8, KernelTypes); + +#define TEST_CKTILE_GGEMM_SUITE_NAME TestCkTileGroupedGemmBF8 + +#include "test_grouped_gemm_ut_cases.inc" diff --git a/test/ck_tile/grouped_gemm/test_grouped_gemm_f8.cpp b/test/ck_tile/grouped_gemm/test_grouped_gemm_f8.cpp index 554b9e063f..29aac818ec 100644 --- a/test/ck_tile/grouped_gemm/test_grouped_gemm_f8.cpp +++ b/test/ck_tile/grouped_gemm/test_grouped_gemm_f8.cpp @@ -8,9 +8,10 @@ #include "ck_tile/host.hpp" #include "test_grouped_gemm_util.hpp" -using F8 = ck_tile::fp8_t; using BF16 = ck_tile::bf16_t; +using F16 = ck_tile::half_t; using F32 = float; +using F8 = ck_tile::fp8_t; using Row = ck_tile::tensor_layout::gemm::RowMajor; using Col = ck_tile::tensor_layout::gemm::ColumnMajor; using True = ck_tile::bool_constant; @@ -19,22 +20,32 @@ using False = ck_tile::bool_constant; // clang-format off using KernelTypes = ::testing::Types< // ALayout, BLayout, CLayout, ADataType, BDataType, AccDataType, CDataType, Persistent - std::tuple< Col, Col, Row, F8, F8, F32, BF16, True>, - std::tuple< Col, Col, Row, F8, F8, F32, BF16, False>, - std::tuple< Row, Row, Row, F8, F8, F32, BF16, True>, - std::tuple< Row, Row, Row, F8, F8, F32, BF16, False>, - std::tuple< Col, Row, Row, F8, F8, F32, BF16, True>, - std::tuple< Col, Row, Row, F8, F8, F32, BF16, False> + std::tuple< Row, Col, Row, F8, F8, F32, F16, True>, + std::tuple< Row, Col, Row, F8, F8, F32, F16, False>, + std::tuple< Col, Col, Row, F8, F8, F32, F16, True>, + std::tuple< Col, Col, Row, F8, F8, F32, F16, False>, + std::tuple< Row, Row, Row, F8, F8, F32, F16, True>, + std::tuple< Row, Row, Row, F8, F8, F32, F16, False>, + std::tuple< Col, Row, Row, F8, F8, F32, F16, True>, + std::tuple< Col, Row, Row, F8, F8, F32, F16, False>, + std::tuple< Row, Col, Row, F8, F8, F32, BF16, True>, + std::tuple< Row, Col, Row, F8, F8, F32, BF16, False>, + std::tuple< Col, Col, Row, F8, F8, F32, BF16, True>, + std::tuple< Col, Col, Row, F8, F8, F32, BF16, False>, + std::tuple< Row, Row, Row, F8, F8, F32, BF16, True>, + std::tuple< Row, Row, Row, F8, F8, F32, BF16, False>, + std::tuple< Col, Row, Row, F8, F8, F32, BF16, True>, + std::tuple< Col, Row, Row, F8, F8, F32, BF16, False> >; // clang-format on template -class TestCkTileGroupedGemmF16 : public TestCkTileGroupedGemm +class TestCkTileGroupedGemmF8 : public TestCkTileGroupedGemm { }; -TYPED_TEST_SUITE(TestCkTileGroupedGemmF16, KernelTypes); +TYPED_TEST_SUITE(TestCkTileGroupedGemmF8, KernelTypes); -#define TEST_CKTILE_GGEMM_SUITE_NAME TestCkTileGroupedGemmF16 +#define TEST_CKTILE_GGEMM_SUITE_NAME TestCkTileGroupedGemmF8 #include "test_grouped_gemm_ut_cases.inc" diff --git a/test/ck_tile/grouped_gemm/test_grouped_gemm_int8.cpp b/test/ck_tile/grouped_gemm/test_grouped_gemm_int8.cpp new file mode 100644 index 0000000000..453a7c6b9e --- /dev/null +++ b/test/ck_tile/grouped_gemm/test_grouped_gemm_int8.cpp @@ -0,0 +1,41 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include + +#include "gtest/gtest.h" + +#include "ck_tile/host.hpp" +#include "test_grouped_gemm_util.hpp" + +using INT8 = ck_tile::int8_t; +using INT32 = ck_tile::int32_t; +using Row = ck_tile::tensor_layout::gemm::RowMajor; +using Col = ck_tile::tensor_layout::gemm::ColumnMajor; +using True = ck_tile::bool_constant; +using False = ck_tile::bool_constant; + +// clang-format off +using KernelTypes = ::testing::Types< + // ALayout, BLayout, CLayout, ADataType, BDataType, AccDataType, CDataType, Persistent + std::tuple< Row, Col, Row, INT8, INT8, INT32, INT32, True>, + std::tuple< Row, Col, Row, INT8, INT8, INT32, INT32, False>, + std::tuple< Col, Col, Row, INT8, INT8, INT32, INT32, True>, + std::tuple< Col, Col, Row, INT8, INT8, INT32, INT32, False>, + std::tuple< Row, Row, Row, INT8, INT8, INT32, INT32, True>, + std::tuple< Row, Row, Row, INT8, INT8, INT32, INT32, False>, + std::tuple< Col, Row, Row, INT8, INT8, INT32, INT32, True>, + std::tuple< Col, Row, Row, INT8, INT8, INT32, INT32, False> + >; +// clang-format on + +template +class TestCkTileGroupedGemmInt8 : public TestCkTileGroupedGemm +{ +}; + +TYPED_TEST_SUITE(TestCkTileGroupedGemmInt8, KernelTypes); + +#define TEST_CKTILE_GGEMM_SUITE_NAME TestCkTileGroupedGemmInt8 + +#include "test_grouped_gemm_ut_cases.inc" diff --git a/test/ck_tile/grouped_gemm/test_grouped_gemm_util.hpp b/test/ck_tile/grouped_gemm/test_grouped_gemm_util.hpp index 16c7dfcb29..54872af25d 100644 --- a/test/ck_tile/grouped_gemm/test_grouped_gemm_util.hpp +++ b/test/ck_tile/grouped_gemm/test_grouped_gemm_util.hpp @@ -267,13 +267,14 @@ class TestCkTileGroupedGemm : public ::testing::Test // Calculate thresholds const auto rtol = ck_tile::get_relative_threshold( ck_tile::integer_divide_ceil(K, kbatch)); - auto atol = ck_tile::get_absolute_threshold( - max_accumulated_value / kbatch, ck_tile::integer_divide_ceil(K, kbatch)); + const float safe_max = max_accumulated_value > 0 ? max_accumulated_value : 1.0f; + auto atol = ck_tile::get_absolute_threshold( + safe_max / kbatch, ck_tile::integer_divide_ceil(K, kbatch)); // Calculate error due to split_k accumulation const auto rtol_split_k = ck_tile::get_relative_threshold(kbatch); - auto atol_split_k = ck_tile::get_absolute_threshold( - max_accumulated_value, kbatch); + auto atol_split_k = + ck_tile::get_absolute_threshold(safe_max, kbatch); // Add extra tolerance for BF16 to account for hardware vs software conversion differences // Hardware __bf16 conversion and software float_to_bf16 can differ by up to 1 ULP