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

fix ck tile build error

Browse Source
This commit is contained in:
Lin, Qun
2025-05-16 00:26:53 -05:00
committed by Qun Lin
parent 027d58110a
commit fb212bd921
3 changed files with 31 additions and 12 deletions
Download Patch File Download Diff File Expand all files Collapse all files

5
example/ck_tile/11_add_rmsnorm2d_rdquant/example_add_rmsnorm2d_rdquant_fwd.cpp
View File

@@ -186,7 +186,7 @@ bool run(const ck_tile::ArgParser& arg_parser)
// Rmsnorm2d
{
ck_tile::HostTensor<InvRmsDataType> invRms_host_ref({m});
ck_tile::HostTensor<ck_tile::null_type> unquant_y_host_ref({m, n}, {stride, 1});
// CAUSION: kernel use ComputeDataType version of x, but we use XDataType here for
// simplicity
ck_tile::reference_rmsnorm2d_fwd<XDataType,
@@ -194,7 +194,8 @@ bool run(const ck_tile::ArgParser& arg_parser)
ComputeDataType,
YDataType,
InvRmsDataType>(
x_host_ref, gamma_host, y_host, invRms_host_ref, epsilon);
x_host_ref, gamma_host, y_host, invRms_host_ref, unquant_y_host_ref, epsilon);
}
// yscale

21
include/ck_tile/ops/flatmm/pipeline/flatmm_pipeline_agmem_bgmem_creg_v1.hpp
View File

@@ -91,8 +91,14 @@ struct FlatmmPipelineAGmemBGmemCRegV1
constexpr index_t A_Buffer_Load_Inst_Num = kMPerBlock * kKPerBlock / BlockSize / KPerLoad;
constexpr index_t A_LDS_Read_Inst_Num = MIterPerWarp * KIterPerWarp;
constexpr index_t B_Buffer_Load_Inst_Num = NIterPerWarp * KIterPerWarp;
#endif
#if defined(USING_MFMA_16x16x32) && defined(ENABLE_FP8)
// constexpr index_t A_LDS_Read_Inst_Remain = A_LDS_Read_Inst_Num - A_Buffer_Load_Inst_Num;
ignore = A_Buffer_Load_Inst_Num;
ignore = A_LDS_Read_Inst_Num;
ignore = B_Buffer_Load_Inst_Num;
if constexpr (WG::kM == 16 && WG::kN == 16)
{
//#if defined(USING_MFMA_16x16x32) && defined(ENABLE_FP8)
static_for<0, A_Buffer_Load_Inst_Num, 1>{}([&](auto i) {
ignore = i;
__builtin_amdgcn_sched_group_barrier(0x100, 1, 0); // DS read
@@ -114,11 +120,11 @@ struct FlatmmPipelineAGmemBGmemCRegV1
__builtin_amdgcn_sched_group_barrier(0x200, 1, 0); // DS write
__builtin_amdgcn_sched_group_barrier(0x008, 4, 0); // MFMA
});
}
else if constexpr (WG::kM == 32 && WG::kN == 32)
{
#elif defined(USING_MFMA_32x32x16)
ignore = A_Buffer_Load_Inst_Num;
ignore = A_LDS_Read_Inst_Num;
ignore = B_Buffer_Load_Inst_Num;
//#elif defined(USING_MFMA_32x32x16)
#if 0
static_for<0,
A_LDS_Read_Inst_Num / 2 - A_Buffer_Load_Inst_Num - B_Buffer_Load_Inst_Num,
@@ -153,7 +159,8 @@ struct FlatmmPipelineAGmemBGmemCRegV1
});
__builtin_amdgcn_sched_group_barrier(0x008, 4, 0); // MFMA
#endif
#endif
}
//#endif
}
template <typename ADramBlockWindowTmp, typename BFlatBlockWindowTmp, typename AElementFunction>

17
include/ck_tile/ops/flatmm/pipeline/flatmm_pipeline_agmem_bgmem_creg_v1_policy.hpp
View File

@@ -19,7 +19,14 @@ struct UniversalFlatmmPipelineAgBgCrPolicy
CK_TILE_HOST_DEVICE static constexpr auto MakeALdsBlockDescriptor()
{
using namespace ck_tile;
#if defined(USING_MFMA_16x16x32) && defined(ENABLE_FP8)
// constexpr auto config = BlockFlatmm::BlockPolicy::template GetWarpGemmMWarpNWarp<Problem>();
//using WG = remove_cvref_t<decltype(config.template at<0>())>;
constexpr index_t MPerXdl = Problem::BlockGemmShape::WarpTile::at(I0);
constexpr index_t NPerXdl = Problem::BlockGemmShape::WarpTile::at(I1);
if constexpr (MPerXdl == 16 && NPerXdl == 16)
{
//#if defined(USING_MFMA_16x16x32) && defined(ENABLE_FP8)
/*reduce transform layers,compare with old ck*/
constexpr index_t MPerBlock = Problem::BlockGemmShape::kM;
constexpr index_t KPerBlock = Problem::BlockGemmShape::kK;
@@ -48,7 +55,10 @@ struct UniversalFlatmmPipelineAgBgCrPolicy
make_tuple(sequence<0>{}, sequence<1>{}));
return a_lds_block_desc;
#elif defined(USING_MFMA_32x32x16)
}
else
{
//#elif defined(USING_MFMA_32x32x16)
constexpr index_t kMPerBlock = Problem::BlockGemmShape::kM;
constexpr index_t kKPerBlock = Problem::BlockGemmShape::kK;
constexpr index_t kKPack = GetSmemPackA<Problem>();
@@ -67,7 +77,8 @@ struct UniversalFlatmmPipelineAgBgCrPolicy
make_tuple(sequence<0>{}, sequence<1>{}));
return a_lds_block_desc;
#endif
//#endif
}
/*xor*/
#if 0
constexpr index_t kMPerBlock = Problem::BlockGemmShape::kM;