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

Disable all tests and instances with f8 on gfx11

Browse Source 
Even though f8_f16 and f16_f8 don't require f8 WMMA instructions,
gfx11 still lacks hardware instructions for fast f8->f32 conversion.
This commit is contained in:
Anton Gorenko
2025-05-23 15:17:36 +05:00
parent 668914cd58
commit 2679c0aa2b
5 changed files with 12 additions and 13 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
library/include/ck/library/tensor_operation_instance/gpu/gemm_universal.hpp
View File

@@ -160,7 +160,7 @@ struct DeviceOperationInstanceFactory<
}
}
#endif
#if(defined(CK_ENABLE_BF16) && defined(CK_ENABLE_FP8))
#if(defined(CK_ENABLE_BF16) && defined(CK_ENABLE_FP8) && defined(CK_USE_WMMA_FP8))
if constexpr(is_same_v<ADataType, f8_t> && is_same_v<BDataType, f8_t> &&
is_same_v<CDataType, bhalf_t>)
{
@@ -185,8 +185,7 @@ struct DeviceOperationInstanceFactory<
op_ptrs);
}
}
#endif
#if(defined(CK_ENABLE_BF16) && defined(CK_ENABLE_FP8))
if constexpr(is_same_v<ADataType, bhalf_t> && is_same_v<BDataType, pk_i4_t> &&
is_same_v<CDataType, bhalf_t>)
{
@@ -204,7 +203,7 @@ struct DeviceOperationInstanceFactory<
}
}
#endif
#if(defined(CK_ENABLE_FP16) && defined(CK_ENABLE_FP8))
#if(defined(CK_ENABLE_FP16) && defined(CK_ENABLE_FP8) && defined(CK_USE_WMMA_FP8))
if constexpr(is_same_v<ADataType, f8_t> && is_same_v<BDataType, half_t> &&
is_same_v<CDataType, half_t>)
{

4
library/include/ck/library/tensor_operation_instance/gpu/gemm_universal_wmma.inc
View File

@@ -155,7 +155,7 @@ void add_device_gemm_wmma_universal_bf16_bf16_bf16_km_nk_mn_comp_mnkpadding_inst
DeviceGemmV2<Col, Col, Row, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough>>>&
instances);
#endif
#if(defined(CK_ENABLE_BF16) && defined(CK_ENABLE_FP8))
#if(defined(CK_ENABLE_BF16) && defined(CK_ENABLE_FP8) && defined(CK_USE_WMMA_FP8))
void add_device_gemm_wmma_universal_f8_f8_bf16_mk_kn_mn_comp_default_instances(
std::vector<std::unique_ptr<
DeviceGemmV2<Row, Row, Row, F8, F8, BF16, PassThrough, PassThrough, PassThrough>>>&
@@ -215,7 +215,7 @@ void add_device_gemm_wmma_universal_bf16_i4_bf16_km_nk_mn_comp_default_instances
DeviceGemmV2<Col, Col, Row, BF16, I4, BF16, PassThrough, PassThrough, PassThrough>>>&
instances);
#endif
#if(defined(CK_ENABLE_FP16) && defined(CK_ENABLE_FP8))
#if(defined(CK_ENABLE_FP16) && defined(CK_ENABLE_FP8) && defined(CK_USE_WMMA_FP8))
void add_device_gemm_wmma_universal_f8_f16_f16_mk_kn_mn_comp_default_instances(
std::vector<std::unique_ptr<
DeviceGemmV2<Row, Row, Row, F8, F16, F16, PassThrough, PassThrough, PassThrough>>>&

4
test/gemm_universal/test_gemm_universal_wmma_bf16.cpp
View File

@@ -59,7 +59,7 @@ using KernelTypes_MK_KN = ::testing::Types<
using KernelTypes_MK_NK = ::testing::Types<
// ADataType, BDataType, ComputeDataType, CDataType
#if defined(CK_ENABLE_FP8)
#if defined(CK_ENABLE_FP8)
std::tuple< BF16, I4, BF16, BF16>,
#endif
std::tuple< BF16, BF16, BF16, BF16>
@@ -72,7 +72,7 @@ using KernelTypes_KM_KN = ::testing::Types<
using KernelTypes_KM_NK = ::testing::Types<
// ADataType, BDataType, ComputeDataType, CDataType
#if defined(CK_ENABLE_FP8)
#if defined(CK_ENABLE_FP8)
std::tuple< BF16, I4, BF16, BF16>,
#endif
std::tuple< BF16, BF16, BF16, BF16>

8
test/gemm_universal/test_gemm_universal_wmma_fp16.cpp
View File

@@ -56,7 +56,7 @@ class TestGemmUniversal_FP16_KM_NK
// clang-format off
using KernelTypes_MK_KN = ::testing::Types<
// ADataType, BDataType, ComputeDataType, CDataType
#if defined(CK_ENABLE_FP8)
#if defined(CK_ENABLE_FP8) && defined(CK_USE_WMMA_FP8)
std::tuple< F8, F16, F16, F16>,
std::tuple< F16, F8, F16, F16>,
#endif
@@ -65,7 +65,7 @@ using KernelTypes_MK_KN = ::testing::Types<
using KernelTypes_MK_NK = ::testing::Types<
// ADataType, BDataType, ComputeDataType, CDataType
#if defined(CK_ENABLE_FP8)
#if defined(CK_ENABLE_FP8) && defined(CK_USE_WMMA_FP8)
std::tuple< F8, F16, F16, F16>,
std::tuple< F16, F8, F16, F16>,
std::tuple< F16, I4, F16, F16>,
@@ -75,7 +75,7 @@ using KernelTypes_MK_NK = ::testing::Types<
using KernelTypes_KM_NK = ::testing::Types<
// ADataType, BDataType, ComputeDataType, CDataType
#if defined(CK_ENABLE_FP8)
#if defined(CK_ENABLE_FP8) && defined(CK_USE_WMMA_FP8)
std::tuple< F8, F16, F16, F16>,
std::tuple< F16, F8, F16, F16>,
std::tuple< F16, I4, F16, F16>,
@@ -85,7 +85,7 @@ using KernelTypes_KM_NK = ::testing::Types<
using KernelTypes_KM_KN = ::testing::Types<
// ADataType, BDataType, ComputeDataType, CDataType
#if defined(CK_ENABLE_FP8)
#if defined(CK_ENABLE_FP8) && defined(CK_USE_WMMA_FP8)
std::tuple< F8, F16, F16, F16>,
std::tuple< F16, F8, F16, F16>,
#endif

2
test/gemm_universal/test_gemm_universal_wmma_fp8.cpp
View File

@@ -7,7 +7,7 @@
#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
#include "test_gemm_universal_util.hpp"
#if CK_USE_WMMA_FP8
#if defined(CK_USE_WMMA_FP8)
using F8 = ck::f8_t;
using BF16 = ck::bhalf_t;