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

test generalization to handle FP16 shuffle better

Browse Source
This commit is contained in:
Zoltan Lakatos
2025-06-30 15:49:00 +00:00
parent 6ba1dc66ac
commit 07948332e9
6 changed files with 51 additions and 56 deletions
Download Patch File Download Diff File Expand all files Collapse all files

72
test/gemm_add/CMakeLists.txt
View File

@@ -1,51 +1,51 @@
# Implements test instances for MultipleD with xdl and wmma support.
# add_gtest_executable(test_gemm_add_xdl test_gemm_add_xdl.cpp)
# if(result EQUAL 0)
# target_link_libraries(test_gemm_add_xdl PRIVATE utility device_gemm_add_instance)
# endif()
add_gtest_executable(test_gemm_add_xdl test_gemm_add_xdl.cpp)
if(result EQUAL 0)
target_link_libraries(test_gemm_add_xdl PRIVATE utility device_gemm_add_instance)
endif()
# add_gtest_executable(test_gemm_add_relu_xdl test_gemm_add_relu_xdl.cpp)
# if(result EQUAL 0)
# target_link_libraries(test_gemm_add_relu_xdl PRIVATE utility device_gemm_add_instance device_gemm_add_relu_instance)
# endif()
add_gtest_executable(test_gemm_add_relu_xdl test_gemm_add_relu_xdl.cpp)
if(result EQUAL 0)
target_link_libraries(test_gemm_add_relu_xdl PRIVATE utility device_gemm_add_instance device_gemm_add_relu_instance)
endif()
# add_gtest_executable(test_gemm_add_silu_xdl test_gemm_add_silu_xdl.cpp)
# if(result EQUAL 0)
# target_link_libraries(test_gemm_add_silu_xdl PRIVATE utility device_gemm_add_instance device_gemm_add_silu_instance)
# endif()
add_gtest_executable(test_gemm_add_silu_xdl test_gemm_add_silu_xdl.cpp)
if(result EQUAL 0)
target_link_libraries(test_gemm_add_silu_xdl PRIVATE utility device_gemm_add_instance device_gemm_add_silu_instance)
endif()
# add_gtest_executable(test_gemm_add_fastgelu_xdl test_gemm_add_fastgelu_xdl.cpp)
# if(result EQUAL 0)
# target_link_libraries(test_gemm_add_fastgelu_xdl PRIVATE utility device_gemm_add_instance device_gemm_add_fastgelu_instance)
# endif()
add_gtest_executable(test_gemm_add_fastgelu_xdl test_gemm_add_fastgelu_xdl.cpp)
if(result EQUAL 0)
target_link_libraries(test_gemm_add_fastgelu_xdl PRIVATE utility device_gemm_add_instance device_gemm_add_fastgelu_instance)
endif()
# add_gtest_executable(test_gemm_fastgelu_wmma test_gemm_fastgelu_wmma.cpp)
# if(result EQUAL 0)
# target_link_libraries(test_gemm_fastgelu_wmma PRIVATE utility device_gemm_fastgelu_instance)
# endif()
add_gtest_executable(test_gemm_fastgelu_wmma test_gemm_fastgelu_wmma.cpp)
if(result EQUAL 0)
target_link_libraries(test_gemm_fastgelu_wmma PRIVATE utility device_gemm_fastgelu_instance)
endif()
# add_gtest_executable(test_gemm_add_fastgelu_wmma test_gemm_add_fastgelu_wmma.cpp)
# if(result EQUAL 0)
# target_link_libraries(test_gemm_add_fastgelu_wmma PRIVATE utility device_gemm_add_fastgelu_instance)
# endif()
add_gtest_executable(test_gemm_add_fastgelu_wmma test_gemm_add_fastgelu_wmma.cpp)
if(result EQUAL 0)
target_link_libraries(test_gemm_add_fastgelu_wmma PRIVATE utility device_gemm_add_fastgelu_instance)
endif()
# add_gtest_executable(test_gemm_add_add_fastgelu_wmma test_gemm_add_add_fastgelu_wmma.cpp)
# if(result EQUAL 0)
# target_link_libraries(test_gemm_add_add_fastgelu_wmma PRIVATE utility device_gemm_add_add_fastgelu_instance)
# endif()
add_gtest_executable(test_gemm_add_add_fastgelu_wmma test_gemm_add_add_fastgelu_wmma.cpp)
if(result EQUAL 0)
target_link_libraries(test_gemm_add_add_fastgelu_wmma PRIVATE utility device_gemm_add_add_fastgelu_instance)
endif()
# add_gtest_executable(test_gemm_add_multiply_wmma test_gemm_add_multiply_wmma.cpp)
# if(result EQUAL 0)
# target_link_libraries(test_gemm_add_multiply_wmma PRIVATE utility device_gemm_add_multiply_instance)
# endif()
add_gtest_executable(test_gemm_add_multiply_wmma test_gemm_add_multiply_wmma.cpp)
if(result EQUAL 0)
target_link_libraries(test_gemm_add_multiply_wmma PRIVATE utility device_gemm_add_multiply_instance)
endif()
add_gtest_executable(test_gemm_multiply_add_wmma test_gemm_multiply_add_wmma.cpp)
if(result EQUAL 0)
target_link_libraries(test_gemm_multiply_add_wmma PRIVATE utility device_gemm_multiply_add_instance)
endif()
# add_gtest_executable(test_gemm_bilinear_wmma test_gemm_bilinear_wmma.cpp)
# if(result EQUAL 0)
# target_link_libraries(test_gemm_bilinear_wmma PRIVATE utility device_gemm_bilinear_instance)
# endif()
add_gtest_executable(test_gemm_bilinear_wmma test_gemm_bilinear_wmma.cpp)
if(result EQUAL 0)
target_link_libraries(test_gemm_bilinear_wmma PRIVATE utility device_gemm_bilinear_instance)
endif()

2
test/gemm_add/test_gemm_add_fastgelu_wmma.cpp
View File

@@ -32,4 +32,4 @@ using KernelTypes = ::testing::Types<std::tuple<F16, F16, F32, F16, F16, Row, Ro
std::tuple<F16, F16, F32, F16, F16, Col, Col, Row, Row>>;
TYPED_TEST_SUITE(TestGemmAddFastgelu, KernelTypes);
TYPED_TEST(TestGemmAddFastgelu, Test_BF16FP16) { this->Run(); }
TYPED_TEST(TestGemmAddFastgelu, Test_FP16FP16) { this->Run(); }

3
test/gemm_add/test_gemm_add_multiply_wmma.cpp
View File

@@ -35,4 +35,5 @@ using KernelTypes =
std::tuple<F16, F16, F32, F16, F16, F16, Col, Row, Row, Row, Row>>;
TYPED_TEST_SUITE(TestGemmAddMultiply, KernelTypes);
TYPED_TEST(TestGemmAddMultiply, Test_BF16FP16) { this->Run(); }
// Due to F16 shuffle data type tests has to run with limited K size. Change instances to FP32?
TYPED_TEST(TestGemmAddMultiply, Test) { this->Run({{16, 32, 64}, {2048, 1024, 256}}); }

7
test/gemm_add/test_gemm_bilinear_wmma.cpp
View File

@@ -32,11 +32,8 @@ class TestGemmBilinear : public ::testing::Test
ELayout>;
public:
void Run()
void Run(TestMatrixSizes const& lengths)
{
std::vector<std::vector<ck::index_t>> lengths = {
{16, 32, 64}, {512, 2048, 4096}, {2048, 1024, 16}};
bool all_success = true;
for(auto length : lengths)
@@ -69,4 +66,4 @@ using KernelTypes = ::testing::Types<std::tuple<F16, F16, F32, F16, F16, Row, Ro
std::tuple<I8, I8, I32, I8, I8, Col, Col, Row, Row>>;
TYPED_TEST_SUITE(TestGemmBilinear, KernelTypes);
TYPED_TEST(TestGemmBilinear, Test) { this->Run(); }
TYPED_TEST(TestGemmBilinear, Test) { this->Run({{16, 32, 64}, {512, 2048, 4096}, {2048, 1024, 16}}); }

20
test/gemm_add/test_gemm_common.hpp
View File

@@ -16,6 +16,11 @@ using F16 = ck::half_t;
using F32 = float;
using F8 = ck::f8_t;
// M, N, K
using TestMatrixSizes = std::vector<std::vector<ck::index_t>>;
static const TestMatrixSizes DefaultTestMatrixSizes = {{16, 32, 64}, {512, 2048, 4096}, {2048, 1024, 16}};
template <typename Tuple>
class TestGemmCommon : public ::testing::Test
{
@@ -32,11 +37,8 @@ class TestGemmCommon : public ::testing::Test
virtual ProfileCall GetImpl() = 0;
void Run()
void Run(const TestMatrixSizes& lengths = DefaultTestMatrixSizes)
{
std::vector<std::vector<ck::index_t>> lengths = {
{16, 32, 64}, {2048, 4096, 8192}, {2048, 1024, 16}};
bool all_success = true;
for(auto length : lengths)
@@ -75,11 +77,8 @@ class TestGemmD0Common : public ::testing::Test
virtual ProfileCall GetImpl() = 0;
void Run()
void Run(const TestMatrixSizes& lengths = DefaultTestMatrixSizes)
{
std::vector<std::vector<ck::index_t>> lengths = {
{16, 32, 64}, {512, 2048, 4096}, {2048, 1024, 16}};
bool all_success = true;
for(auto length : lengths)
@@ -121,11 +120,8 @@ class TestGemmD0D1Common : public ::testing::Test
virtual ProfileCall GetImpl() = 0;
void Run()
void Run(const TestMatrixSizes& lengths = DefaultTestMatrixSizes)
{
std::vector<std::vector<ck::index_t>> lengths = {
{16, 32, 64}, {512, 2048, 4096}, {2048, 1024, 16}};
bool all_success = true;
for(auto length : lengths)

3
test/gemm_add/test_gemm_multiply_add_wmma.cpp
View File

@@ -36,4 +36,5 @@ using KernelTypes = ::testing::Types<
std::tuple<F16, F16, F32, F16, F16, F16, Row, Row, Row, Row, Row>>;
TYPED_TEST_SUITE(TestGemmMultiplyAdd, KernelTypes);
TYPED_TEST(TestGemmMultiplyAdd, Test) { this->Run(); }
// Due to F16 shuffle data type tests has to run with limited K size. Change instances to FP32?
TYPED_TEST(TestGemmMultiplyAdd, Test) { this->Run({{16, 32, 64}, {2048, 1024, 256}}); }