From 33ae2b65f791294a98736734c041b545c8faa780 Mon Sep 17 00:00:00 2001 From: Aleksander Dudek Date: Tue, 16 Sep 2025 10:48:58 +0000 Subject: [PATCH] [CK_TILE] Add supported grouped and batched GEMM test cases --- .../batched_gemm/test_batched_gemm.cpp | 27 +++++++- .../grouped_gemm/test_grouped_gemm.cpp | 68 +++++++++++++++---- 2 files changed, 80 insertions(+), 15 deletions(-) diff --git a/test/ck_tile/batched_gemm/test_batched_gemm.cpp b/test/ck_tile/batched_gemm/test_batched_gemm.cpp index 9bcf082e18..54befea08f 100644 --- a/test/ck_tile/batched_gemm/test_batched_gemm.cpp +++ b/test/ck_tile/batched_gemm/test_batched_gemm.cpp @@ -8,8 +8,13 @@ #include "ck_tile/host.hpp" #include "test_batched_gemm_util.hpp" -using F16 = ck_tile::half_t; -using F32 = float; +using INT8 = ck_tile::int8_t; +using INT32 = ck_tile::int32_t; +using F16 = ck_tile::half_t; +using F32 = float; +using F8 = ck_tile::fp8_t; +using BF8 = ck_tile::bf8_t; +using BF16 = ck_tile::bf16_t; using Row = ck_tile::tensor_layout::gemm::RowMajor; using Col = ck_tile::tensor_layout::gemm::ColumnMajor; @@ -20,7 +25,23 @@ using KernelTypes = ::testing::Types< std::tuple< Row, Row, Row, F16, F16, F32, F16>, std::tuple< Col, Row, Row, F16, F16, F32, F16>, std::tuple< Row, Col, Row, F16, F16, F32, F16>, - std::tuple< Col, Col, Row, F16, F16, F32, F16> + std::tuple< Col, Col, Row, F16, F16, F32, F16>, + 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>, + 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>, + 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, 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 diff --git a/test/ck_tile/grouped_gemm/test_grouped_gemm.cpp b/test/ck_tile/grouped_gemm/test_grouped_gemm.cpp index 7d71f9f927..c458fee0a4 100644 --- a/test/ck_tile/grouped_gemm/test_grouped_gemm.cpp +++ b/test/ck_tile/grouped_gemm/test_grouped_gemm.cpp @@ -8,8 +8,13 @@ #include "ck_tile/host.hpp" #include "test_grouped_gemm_util.hpp" +using INT8 = ck_tile::int8_t; +using INT32 = ck_tile::int32_t; using F16 = ck_tile::half_t; using F32 = float; +using F8 = ck_tile::fp8_t; +using BF16 = ck_tile::bf16_t; +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; @@ -17,18 +22,57 @@ 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, F16, F16, F32, F16, True>, - std::tuple< Row, Col, Row, F16, F16, F32, F16, False>, - std::tuple< Row, Col, Row, F16, F16, F32, F16, True>, - std::tuple< Row, Col, Row, F16, F16, F32, F16, False>, - - std::tuple< Col, Col, Row, F16, F16, F32, F16, True>, - std::tuple< Col, Col, Row, F16, F16, F32, F16, False>, - std::tuple< Row, Row, Row, F16, F16, F32, F16, True>, - std::tuple< Row, Row, Row, F16, F16, F32, F16, False>, - std::tuple< Col, Row, Row, F16, F16, F32, F16, True>, - std::tuple< Col, Row, Row, F16, F16, F32, F16, False> + // ALayout, BLayout, CLayout, ADataType, BDataType, AccDataType, CDataType, Persistent + std::tuple< Row, Col, Row, F16, F16, F32, F16, True>, + std::tuple< Row, Col, Row, F16, F16, F32, F16, False>, + std::tuple< Row, Col, Row, F16, F16, F32, F16, True>, + std::tuple< Row, Col, Row, F16, F16, F32, F16, False>, + std::tuple< Col, Col, Row, F16, F16, F32, F16, True>, + std::tuple< Col, Col, Row, F16, F16, F32, F16, False>, + std::tuple< Row, Row, Row, F16, F16, F32, F16, True>, + std::tuple< Row, Row, Row, F16, F16, F32, F16, False>, + std::tuple< Col, Row, Row, F16, F16, F32, F16, True>, + std::tuple< Col, Row, Row, F16, F16, F32, F16, False>, + std::tuple< Row, Col, Row, BF16, BF16, F32, F16, True>, + std::tuple< Row, Col, Row, BF16, BF16, F32, F16, False>, + std::tuple< Row, Col, Row, BF16, BF16, F32, F16, True>, + std::tuple< Row, Col, Row, BF16, BF16, F32, F16, False>, + std::tuple< Col, Col, Row, BF16, BF16, F32, F16, True>, + std::tuple< Col, Col, Row, BF16, BF16, F32, F16, False>, + std::tuple< Row, Row, Row, BF16, BF16, F32, F16, True>, + std::tuple< Row, Row, Row, BF16, BF16, F32, F16, False>, + std::tuple< Col, Row, Row, BF16, BF16, F32, F16, True>, + std::tuple< Col, Row, Row, BF16, BF16, F32, F16, False>, + std::tuple< Row, Col, Row, BF8, BF8, F32, F16, True>, + std::tuple< Row, Col, Row, BF8, BF8, F32, F16, False>, + 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>, + std::tuple< Row, Col, Row, F8, F8, F32, F16, True>, + std::tuple< Row, Col, Row, F8, F8, F32, F16, 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, INT8, INT8, INT32, INT32, True>, + std::tuple< Row, Col, Row, INT8, INT8, INT32, INT32, False>, + 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