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

Enable tests

Browse Source
This commit is contained in:
Rostyslav Geyyer
2025-04-30 16:33:36 +00:00
parent 5d22aa296e
commit 840527ec41
1 changed files with 46 additions and 48 deletions
Download Patch File Download Diff File Expand all files Collapse all files

94
test/mx_mfma_op/mx_mfma_op.cpp
View File

@@ -20,9 +20,9 @@ using ck::type_convert;
template <typename AType, typename BType, typename CType, ck::MFMA_F8F6F4 mfma>
bool run_mfma_test(ck::index_t init)
{
using ALayout = ck::tensor_layout::gemm::ColumnMajor;
using ALayout = ck::tensor_layout::gemm::RowMajor;
using BLayout = ck::tensor_layout::gemm::ColumnMajor;
using CLayout = ck::tensor_layout::gemm::ColumnMajor;
using CLayout = ck::tensor_layout::gemm::RowMajor;
using AccType = float; // only MFMA_F32 instructions supported
using CPUAccType = AccType;
@@ -52,19 +52,19 @@ bool run_mfma_test(ck::index_t init)
return pass;
}
// TEST(MFMA, FP8MFMA16x16x128)
// {
// auto AB_init = 7;
// auto pass = run_mfma_test<f8_t, f8_t, half_t, ck::MFMA_F8F6F4::F32_16x16x128>(AB_init);
// EXPECT_TRUE(pass);
// }
TEST(MFMA, FP8MFMA16x16x128)
{
auto AB_init = 7;
auto pass = run_mfma_test<f8_t, f8_t, half_t, ck::MFMA_F8F6F4::F32_16x16x128>(AB_init);
EXPECT_TRUE(pass);
}
// TEST(MFMA, FP8MFMA32x32x64)
// {
// auto AB_init = 7;
// auto pass = run_mfma_test<f8_t, f8_t, float, ck::MFMA_F8F6F4::F32_32x32x64>(AB_init);
// EXPECT_TRUE(pass);
// }
TEST(MFMA, FP8MFMA32x32x64)
{
auto AB_init = 7;
auto pass = run_mfma_test<f8_t, f8_t, float, ck::MFMA_F8F6F4::F32_32x32x64>(AB_init);
EXPECT_TRUE(pass);
}
TEST(MFMA, FP4MFMA16x16x128)
{
@@ -73,12 +73,12 @@ TEST(MFMA, FP4MFMA16x16x128)
EXPECT_TRUE(pass);
}
// TEST(MFMA, FP4MFMA32x32x64)
// {
// auto AB_init = 4;
// auto pass = run_mfma_test<f4x2_pk_t, f4x2_pk_t, half_t,
// ck::MFMA_F8F6F4::F32_32x32x64>(AB_init); EXPECT_TRUE(pass);
// }
TEST(MFMA, FP4MFMA32x32x64)
{
auto AB_init = 4;
auto pass = run_mfma_test<f4x2_pk_t, f4x2_pk_t, half_t, ck::MFMA_F8F6F4::F32_32x32x64>(AB_init);
EXPECT_TRUE(pass);
}
/**
* @brief Run the test for the given MX MFMA instruction
@@ -125,34 +125,32 @@ bool run_mxmfma_test(ck::index_t init)
return pass;
}
// TEST(MXMFMA, MXFP8MFMA16x16x128)
// {
// auto AB_init = 7;
// auto pass = run_mxmfma_test<f8_t, f8_t, float,
// ck::MFMA_F8F6F4::SCALE_F32_16x16x128>(AB_init); EXPECT_TRUE(pass);
// }
TEST(MXMFMA, MXFP8MFMA16x16x128)
{
auto AB_init = 7;
auto pass = run_mxmfma_test<f8_t, f8_t, float, ck::MFMA_F8F6F4::SCALE_F32_16x16x128>(AB_init);
EXPECT_TRUE(pass);
}
// TEST(MXMFMA, MXFP8MFMA32x32x64)
// {
// auto AB_init = 7;
// auto pass = run_mxmfma_test<f8_t, f8_t, half_t,
// ck::MFMA_F8F6F4::SCALE_F32_32x32x64>(AB_init); EXPECT_TRUE(pass);
// }
TEST(MXMFMA, MXFP8MFMA32x32x64)
{
auto AB_init = 7;
auto pass = run_mxmfma_test<f8_t, f8_t, half_t, ck::MFMA_F8F6F4::SCALE_F32_32x32x64>(AB_init);
EXPECT_TRUE(pass);
}
// TEST(MXMFMA, MXFP4MFMA16x16x128)
// {
// auto AB_init = 4;
// auto pass =
// run_mxmfma_test<f4x2_pk_t, f4x2_pk_t, float,
// ck::MFMA_F8F6F4::SCALE_F32_16x16x128>(AB_init);
// EXPECT_TRUE(pass);
// }
TEST(MXMFMA, MXFP4MFMA16x16x128)
{
auto AB_init = 4;
auto pass =
run_mxmfma_test<f4x2_pk_t, f4x2_pk_t, float, ck::MFMA_F8F6F4::SCALE_F32_16x16x128>(AB_init);
EXPECT_TRUE(pass);
}
// TEST(MXMFMA, MXFP4MFMA32x32x64)
// {
// auto AB_init = 4;
// auto pass =
// run_mxmfma_test<f4x2_pk_t, f4x2_pk_t, half_t,
// ck::MFMA_F8F6F4::SCALE_F32_32x32x64>(AB_init);
// EXPECT_TRUE(pass);
// }
TEST(MXMFMA, MXFP4MFMA32x32x64)
{
auto AB_init = 4;
auto pass =
run_mxmfma_test<f4x2_pk_t, f4x2_pk_t, half_t, ck::MFMA_F8F6F4::SCALE_F32_32x32x64>(AB_init);
EXPECT_TRUE(pass);
}