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

Disable some tests

Browse Source
This commit is contained in:
Rostyslav Geyyer
2025-04-03 20:41:53 +00:00
parent e3386e4f66
commit 5d22aa296e
2 changed files with 37 additions and 34 deletions
Download Patch File Download Diff File Expand all files Collapse all files

66
test/mx_mfma_op/mx_mfma_op.cpp
View File

@@ -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,32 +125,34 @@ 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);
// }

5
test/mx_mfma_op/mx_mfma_op.hpp
View File

@@ -131,7 +131,7 @@ __device__ AFragT load_A_col_major(AType const* input_ptr)
// Register Mapping for 128x16 for FP4: || Register Mapping for 64x32 for FP4:
// Size | BLOCK_M | BLOCK_M | BLOCK_M | BLOCK_M | || Size | BLOCK_M | BLOCK_M | |
// M | 0 ... 15 | 0 ... 15 | 0 ... 15 | 0 ... 15 | Vector || N | 0 ... 31 | 0 ... 31 | Vector |
// M | 0 ... 15 | 0 ... 15 | 0 ... 15 | 0 ... 15 | Vector || M | 0 ... 31 | 0 ... 31 | Vector |
// Thread Id | 0 ... 15 | 16 ... 31 | 32 ... 47 | 48 ... 63 | Element || Thread Id | 0 ... 31 | 32 ... 63 | Element|
// Register Element |------------|-------------|------------|-------------|-----------|| Register Element |------------|-------------|--------|
// Reg 0 [0:7] | K0K1 | K32K33 | K64K65 | K96K97 | v[0] || Reg 0 [0:7] | K0K1 | K32K33 | v[0] |
@@ -1279,7 +1279,7 @@ struct TestMFMA
switch(init)
{
case 0:
a_m_k.GenerateTensorValue(GeneratorTensor_1<ADataType>{0.015625f});
a_m_k.GenerateTensorValue(GeneratorTensor_1<ADataType>{1.0f});
// NOTE: not all numbers are representable in FP8, BF8, etc.
b_n_k.GenerateTensorValue(GeneratorTensor_Sequential<BDataType, 1>{});
break;
@@ -1299,6 +1299,7 @@ struct TestMFMA
b_n_k.GenerateTensorValue(GeneratorTensor_4<BDataType>(1, 3));
break;
case 4:
// FP4 values case
a_m_k.GenerateTensorValue(GeneratorTensor_2<ADataType>{-4, 5});
b_n_k.GenerateTensorValue(GeneratorTensor_2<BDataType>{-4, 5});
break;