ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-17 17:19:12 +00:00
Code Issues Packages Projects Releases Wiki Activity

[rocm-libraries] ROCm/rocm-libraries#6521 (commit 222385c)

Browse Source 
feat(composablekernel): More data type tests for ck tile
 batched grouped gemm (#6521)

## Motivation

Add more data type tests for ck tile batched and grouped gemm.

ISSUE ID : AICK-228

## Technical Details

Making sure the following are covered: fp16, bf16, fp8, bf8, int8.

## Test Plan

Added tests must pass.

## Test Result

Tests passing locally on gfx942.

## Submission Checklist

- [ ] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
This commit is contained in:
aledudek
2026-07-16 22:21:19 +00:00
committed by assistant-librarian[bot]
parent 973a0bfd81
commit e44760f220
13 changed files with 407 additions and 28 deletions
Download Patch File Download Diff File Expand all files Collapse all files

63
test/ck_tile/batched_gemm/CMakeLists.txt
View File

@@ -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()

37
test/ck_tile/batched_gemm/test_batched_gemm_bf16.cpp Normal file
View File

@@ -0,0 +1,37 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
#include <tuple>
#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 <typename Tuple>
class TestCkTileBatchedGemmBF16 : public TestCkTileBatchedGemm<Tuple>
{
};
TYPED_TEST_SUITE(TestCkTileBatchedGemmBF16, KernelTypes);
#define TEST_CKTILE_BGEMM_SUITE_NAME TestCkTileBatchedGemmBF16
#include "test_batched_gemm_ut_cases.inc"

37
test/ck_tile/batched_gemm/test_batched_gemm_bf8.cpp Normal file
View File

@@ -0,0 +1,37 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
#include <tuple>
#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 <typename Tuple>
class TestCkTileBatchedGemmBF8 : public TestCkTileBatchedGemm<Tuple>
{
};
TYPED_TEST_SUITE(TestCkTileBatchedGemmBF8, KernelTypes);
#define TEST_CKTILE_BGEMM_SUITE_NAME TestCkTileBatchedGemmBF8
#include "test_batched_gemm_ut_cases.inc"

9
test/ck_tile/batched_gemm/test_batched_gemm.cpp → test/ck_tile/batched_gemm/test_batched_gemm_f16.cpp
View File

@@ -24,6 +24,13 @@ using KernelTypes = ::testing::Types<
>;
// clang-format on
TYPED_TEST_SUITE(TestCkTileBatchedGemm, KernelTypes);
template <typename Tuple>
class TestCkTileBatchedGemmF16 : public TestCkTileBatchedGemm<Tuple>
{
};
TYPED_TEST_SUITE(TestCkTileBatchedGemmF16, KernelTypes);
#define TEST_CKTILE_BGEMM_SUITE_NAME TestCkTileBatchedGemmF16
#include "test_batched_gemm_ut_cases.inc"

42
test/ck_tile/batched_gemm/test_batched_gemm_f8.cpp Normal file
View File

@@ -0,0 +1,42 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
#include <tuple>
#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 <typename Tuple>
class TestCkTileBatchedGemmF8 : public TestCkTileBatchedGemm<Tuple>
{
};
TYPED_TEST_SUITE(TestCkTileBatchedGemmF8, KernelTypes);
#define TEST_CKTILE_BGEMM_SUITE_NAME TestCkTileBatchedGemmF8
#include "test_batched_gemm_ut_cases.inc"

36
test/ck_tile/batched_gemm/test_batched_gemm_int8.cpp Normal file
View File

@@ -0,0 +1,36 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
#include <tuple>
#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 <typename Tuple>
class TestCkTileBatchedGemmInt8 : public TestCkTileBatchedGemm<Tuple>
{
};
TYPED_TEST_SUITE(TestCkTileBatchedGemmInt8, KernelTypes);
#define TEST_CKTILE_BGEMM_SUITE_NAME TestCkTileBatchedGemmInt8
#include "test_batched_gemm_ut_cases.inc"

2
test/ck_tile/batched_gemm/test_batched_gemm_ut_cases.inc
View File

@@ -21,7 +21,7 @@ struct StrideConfig
int batchStrideC;
};
TYPED_TEST(TestCkTileBatchedGemm, Basic)
TYPED_TEST(TEST_CKTILE_BGEMM_SUITE_NAME, Basic)
{
std::vector<GemmParams> gemmParams{{256, 256, 256, 1},
{256, 256, 256, 2},

24
test/ck_tile/batched_gemm/test_batched_gemm_util.hpp
View File

@@ -1,7 +1,11 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
#pragma once
#include <algorithm>
#include <cmath>
#include <numeric>
#include <sstream>
#include <type_traits>
#include <gtest/gtest.h>
#include "ck_tile/core.hpp"
@@ -261,9 +265,23 @@ class TestCkTileBatchedGemm : public ::testing::Test
ck_tile::reference_batched_gemm<ADataType, BDataType, AccDataType, CDataType>(
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<sizeof(ADataType) < sizeof(BDataType), ADataType, BDataType>;
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<float>(v)));
});
const auto rtol = ck_tile::get_relative_threshold<ComputeType, CDataType, AccDataType>(
ck_tile::integer_divide_ceil(K, kbatch));
const auto atol = ck_tile::get_absolute_threshold<ComputeType, CDataType, AccDataType>(
static_cast<double>(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);
}

62
test/ck_tile/grouped_gemm/CMakeLists.txt
View File

@@ -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()

42
test/ck_tile/grouped_gemm/test_grouped_gemm_bf8.cpp Normal file
View File

@@ -0,0 +1,42 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
#include <tuple>
#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<true>;
using False = ck_tile::bool_constant<false>;
// 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 <typename Tuple>
class TestCkTileGroupedGemmBF8 : public TestCkTileGroupedGemm<Tuple>
{
};
TYPED_TEST_SUITE(TestCkTileGroupedGemmBF8, KernelTypes);
#define TEST_CKTILE_GGEMM_SUITE_NAME TestCkTileGroupedGemmBF8
#include "test_grouped_gemm_ut_cases.inc"

31
test/ck_tile/grouped_gemm/test_grouped_gemm_f8.cpp
View File

@@ -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<true>;
@@ -19,22 +20,32 @@ using False = ck_tile::bool_constant<false>;
// 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 <typename Tuple>
class TestCkTileGroupedGemmF16 : public TestCkTileGroupedGemm<Tuple>
class TestCkTileGroupedGemmF8 : public TestCkTileGroupedGemm<Tuple>
{
};
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"

41
test/ck_tile/grouped_gemm/test_grouped_gemm_int8.cpp Normal file
View File

@@ -0,0 +1,41 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
#include <tuple>
#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<true>;
using False = ck_tile::bool_constant<false>;
// 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 <typename Tuple>
class TestCkTileGroupedGemmInt8 : public TestCkTileGroupedGemm<Tuple>
{
};
TYPED_TEST_SUITE(TestCkTileGroupedGemmInt8, KernelTypes);
#define TEST_CKTILE_GGEMM_SUITE_NAME TestCkTileGroupedGemmInt8
#include "test_grouped_gemm_ut_cases.inc"

9
test/ck_tile/grouped_gemm/test_grouped_gemm_util.hpp
View File

@@ -267,13 +267,14 @@ class TestCkTileGroupedGemm : public ::testing::Test
// Calculate thresholds
const auto rtol = ck_tile::get_relative_threshold<ComputeType, CDataType, AccDataType>(
ck_tile::integer_divide_ceil(K, kbatch));
auto atol = ck_tile::get_absolute_threshold<ComputeType, CDataType, AccDataType>(
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<ComputeType, CDataType, AccDataType>(
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<CDataType, CDataType, CDataType>(kbatch);
auto atol_split_k = ck_tile::get_absolute_threshold<CDataType, CDataType, CDataType>(
max_accumulated_value, kbatch);
auto atol_split_k =
ck_tile::get_absolute_threshold<CDataType, CDataType, CDataType>(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