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

remove unnecessary hacky

Browse Source
This commit is contained in:
aska-0096
2025-05-09 16:07:22 +00:00
parent 86bff9c46e
commit 0987b0af44
4 changed files with 16 additions and 86 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
example/67_gemm_microscaling/gemm_mx_fp4.cpp
View File

@@ -27,7 +27,7 @@ constexpr ck::index_t KPerBlock = 256;
constexpr auto GemmSpec = ck::tensor_operation::device::GemmSpecialization::Default;
constexpr auto BlkGemmPSched = ck::BlockGemmPipelineScheduler::Intrawave;
constexpr auto BlkGemmPVer = ck::BlockGemmPipelineVersion::v1;
constexpr auto BlkGemmPVer = ck::BlockGemmPipelineVersion::v3;
// v3 should be performant one, However
// 1. some bug existed cause memory access fault in some cases, MNK=2k2k2k
// 2. Register spill observed, most likely unpack the e8m0 from single register then feed to

14
include/ck/tensor_operation/gpu/block/blockwise_gemm_mx_pipeline_xdlops_base.hpp
View File

@@ -58,7 +58,10 @@ struct BlockwiseGemmXdlops_mx_pipeline_base
//> store rows/cols into thread registers in chunks of 16
//> e.g. [k0,...,k15,k64,...,k79] or [k0,...,k15,k32,...,k47]
static constexpr index_t KThreadChunk = 16;
static constexpr index_t APackedSize = is_same_v<remove_cvref_t<ComputeTypeA>, f4x2_pk_t>
? 2
: 1;
static constexpr index_t KThreadChunk = 16 * APackedSize/ sizeof(ComputeTypeA);
static constexpr index_t KPerThread = KPerBlock / xdlops_gemm.K0PerXdlops;
static constexpr index_t KRepeat = KPerThread / KPack;
@@ -118,7 +121,7 @@ struct BlockwiseGemmXdlops_mx_pipeline_base
const auto xdlops_a_idx = xdlops_gemm.CalculateAThreadOriginDataIndex();
return make_tuple(0, waveId_m, xdlops_a_idx[I1], KThreadChunk * 2 * xdlops_a_idx[I0]);
return make_tuple(0, waveId_m, xdlops_a_idx[I1], KThreadChunk * xdlops_a_idx[I0]);
}
__device__ static auto CalculateBThreadOriginDataIndex()
@@ -129,8 +132,7 @@ struct BlockwiseGemmXdlops_mx_pipeline_base
const auto xdlops_b_idx = xdlops_gemm.CalculateBThreadOriginDataIndex();
return make_tuple(0, waveId_n, xdlops_b_idx[I1], KThreadChunk * 2 * xdlops_b_idx[I0]);
// [0-15] [0-3]
return make_tuple(0, waveId_n, xdlops_b_idx[I1], KThreadChunk * xdlops_b_idx[I0]);
}
template <index_t m0, index_t n0, index_t xdlops_i, index_t blk_i>
@@ -347,7 +349,7 @@ struct BlockwiseGemmXdlops_mx_pipeline_base
ComputeTypeA,
decltype(a_block_desc_m0_m1_m2_k),
decltype(a_thread_desc_),
Sequence<1, 1, 1, KThreadChunk * 2>,
Sequence<1, 1, 1, KThreadChunk>,
Sequence<0, 1, 2, 3>,
3,
A_K1,
@@ -357,7 +359,7 @@ struct BlockwiseGemmXdlops_mx_pipeline_base
ComputeTypeB,
decltype(b_block_desc_n0_n1_n2_k),
decltype(b_thread_desc_),
Sequence<1, 1, 1, KThreadChunk * 2>,
Sequence<1, 1, 1, KThreadChunk>,
Sequence<0, 1, 2, 3>,
3,
B_K1,

14
include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v3_mx.hpp
View File

@@ -423,7 +423,7 @@ struct BlockwiseGemmXdlops_pipeline_v3_mx<BlockGemmPipelineScheduler::Intrawave,
static_for<0, KRepeat, 1>{}([&](auto k) {
constexpr auto k_step = k * xdlops_gemm.KPerXdlops * (KPack / xdlops_gemm.K1PerXdlops);
static_for<0, MRepeat, 1>{}([&](auto m0) {
static_for<0, xdlops_gemm.K1PerXdlops / 2 / KThreadChunk, 1>{}([&](auto chunk) {
static_for<0, xdlops_gemm.K1PerXdlops / KThreadChunk, 1>{}([&](auto chunk) {
constexpr auto a_k_step_chunk =
k_step + chunk * KThreadChunk * xdlops_gemm.mfma_instr.num_input_blks;
a_thread_copy_.Run(a_block_desc_m0_m1_m2_k,
@@ -436,7 +436,7 @@ struct BlockwiseGemmXdlops_pipeline_v3_mx<BlockGemmPipelineScheduler::Intrawave,
});
static_for<0, NRepeat, 1>{}([&](auto n0) {
// read block data in chunks to assemble correct thread vectors
static_for<0, xdlops_gemm.K1PerXdlops / 2 / KThreadChunk, 1>{}([&](auto chunk) {
static_for<0, xdlops_gemm.K1PerXdlops / KThreadChunk, 1>{}([&](auto chunk) {
constexpr auto b_k_step_chunk =
k_step + chunk * KThreadChunk * xdlops_gemm.mfma_instr.num_input_blks;
b_thread_copy_.Run(b_block_desc_n0_n1_n2_k,
@@ -616,7 +616,7 @@ struct BlockwiseGemmXdlops_pipeline_v3_mx<BlockGemmPipelineScheduler::Intrawave,
k * xdlops_gemm.KPerXdlops * (KPack / xdlops_gemm.K1PerXdlops);
static_for<0, MRepeat, 1>{}([&](auto m0) {
static_for<0, xdlops_gemm.K1PerXdlops / 2 / KThreadChunk, 1>{}(
static_for<0, xdlops_gemm.K1PerXdlops / KThreadChunk, 1>{}(
[&](auto chunk) {
constexpr auto a_k_step_chunk =
k_step + chunk * KThreadChunk *
@@ -632,7 +632,7 @@ struct BlockwiseGemmXdlops_pipeline_v3_mx<BlockGemmPipelineScheduler::Intrawave,
});
static_for<0, NRepeat, 1>{}([&](auto n0) {
// read block data in chunks to assemble correct thread vectors
static_for<0, xdlops_gemm.K1PerXdlops / 2 / KThreadChunk, 1>{}(
static_for<0, xdlops_gemm.K1PerXdlops / KThreadChunk, 1>{}(
[&](auto chunk) {
constexpr auto b_k_step_chunk =
k_step + chunk * KThreadChunk *
@@ -656,7 +656,7 @@ struct BlockwiseGemmXdlops_pipeline_v3_mx<BlockGemmPipelineScheduler::Intrawave,
LoopFunc(I1, I0);
i += 2;
} while(i < (num_loop - 1));
} while(i < (num_loop - 2));
}
// tail
@@ -776,7 +776,7 @@ struct BlockwiseGemmXdlops_pipeline_v3_mx<BlockGemmPipelineScheduler::Intrawave,
k * xdlops_gemm.KPerXdlops * (KPack / xdlops_gemm.K1PerXdlops);
static_for<0, MRepeat, 1>{}([&](auto m0) {
static_for<0, xdlops_gemm.K1PerXdlops / 2 / KThreadChunk, 1>{}([&](auto chunk) {
static_for<0, xdlops_gemm.K1PerXdlops / KThreadChunk, 1>{}([&](auto chunk) {
constexpr auto a_k_step_chunk =
k_step + chunk * KThreadChunk * xdlops_gemm.mfma_instr.num_input_blks;
a_thread_copy_.Run(a_block_desc_m0_m1_m2_k,
@@ -789,7 +789,7 @@ struct BlockwiseGemmXdlops_pipeline_v3_mx<BlockGemmPipelineScheduler::Intrawave,
});
static_for<0, NRepeat, 1>{}([&](auto n0) {
// read block data in chunks to assemble correct thread vectors
static_for<0, xdlops_gemm.K1PerXdlops / 2 / KThreadChunk, 1>{}([&](auto chunk) {
static_for<0, xdlops_gemm.K1PerXdlops / KThreadChunk, 1>{}([&](auto chunk) {
constexpr auto b_k_step_chunk =
k_step + chunk * KThreadChunk * xdlops_gemm.mfma_instr.num_input_blks;
b_thread_copy_.Run(b_block_desc_n0_n1_n2_k,

72
include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v3r1.hpp
View File

@@ -314,78 +314,6 @@ struct ThreadwiseTensorSliceTransfer_v3r1
src_thread_scratch_tuple_(thread_scratch_id)
.template SetAsType<dst_vector_t>(src_data_idx_seq,
op_r_v.template AsType<dst_vector_t>()[I0]);
CK_PRINT<decltype(src_thread_scratch_tuple_(thread_scratch_id))>();
// using a = ck::StaticTensorTupleOfVectorBuffer<
// ck::AddressSpaceEnum::Vgpr,
// ck::f4x2_pk_t,
// 8,
// const ck::TensorDescriptor<
// ck::Tuple<ck::UnMerge<ck::Tuple<ck::integral_constant<int, 2>,
// ck::integral_constant<int, 1>,
// ck::integral_constant<int, 1>,
// ck::integral_constant<int, 8>>,
// false>,
// ck::PassThrough<ck::integral_constant<int, 2>>,
// ck::PassThrough<ck::integral_constant<int, 1>>,
// ck::Merge_v3_division_mod<ck::Tuple<ck::integral_constant<int, 1>,
// ck::integral_constant<int,
// 8>>>>,
// ck::Tuple<ck::Sequence<0>,
// ck::Sequence<1>,
// ck::Sequence<2>,
// ck::Sequence<3, 4>>,
// ck::Tuple<ck::Sequence<1, 2, 3, 4>,
// ck::Sequence<5>,
// ck::Sequence<6>,
// ck::Sequence<7>>,
// ck::Sequence<5, 6, 7>,
// ck::integral_constant<long, 16>>,
// true>;
#if 0
auto data_print = src_thread_scratch_tuple_(thread_scratch_id).data_;
printf("TID%03d src_thread_scratch_tuple_(thread_scratch_id).data_ (%dx%d): "
"0x%02x 0x%02x 0x%02x 0x%02x 0x%02x 0x%02x 0x%02x 0x%02x "
"0x%02x 0x%02x 0x%02x 0x%02x 0x%02x 0x%02x 0x%02x 0x%02x "
// "0x%02x 0x%02x 0x%02x 0x%02x 0x%02x 0x%02x 0x%02x 0x%02x "
// "0x%02x 0x%02x 0x%02x 0x%02x 0x%02x 0x%02x 0x%02x 0x%02x "
"\n",
get_thread_local_1d_id(),
static_cast<int>(data_print.s_per_v),
static_cast<int>(data_print.num_of_v_),
*reinterpret_cast<const uint8_t*>(&data_print[Number<0>{}]),
*reinterpret_cast<const uint8_t*>(&data_print[Number<1>{}]),
*reinterpret_cast<const uint8_t*>(&data_print[Number<2>{}]),
*reinterpret_cast<const uint8_t*>(&data_print[Number<3>{}]),
*reinterpret_cast<const uint8_t*>(&data_print[Number<4>{}]),
*reinterpret_cast<const uint8_t*>(&data_print[Number<5>{}]),
*reinterpret_cast<const uint8_t*>(&data_print[Number<6>{}]),
*reinterpret_cast<const uint8_t*>(&data_print[Number<7>{}]),
*reinterpret_cast<const uint8_t*>(&data_print[Number<8>{}]),
*reinterpret_cast<const uint8_t*>(&data_print[Number<9>{}]),
*reinterpret_cast<const uint8_t*>(&data_print[Number<10>{}]),
*reinterpret_cast<const uint8_t*>(&data_print[Number<11>{}]),
*reinterpret_cast<const uint8_t*>(&data_print[Number<12>{}]),
*reinterpret_cast<const uint8_t*>(&data_print[Number<13>{}]),
*reinterpret_cast<const uint8_t*>(&data_print[Number<14>{}]),
*reinterpret_cast<const uint8_t*>(&data_print[Number<15>{}])
// *reinterpret_cast<const uint8_t*>(&data_print[Number<16>{}]),
// *reinterpret_cast<const uint8_t*>(&data_print[Number<17>{}]),
// *reinterpret_cast<const uint8_t*>(&data_print[Number<18>{}]),
// *reinterpret_cast<const uint8_t*>(&data_print[Number<19>{}]),
// *reinterpret_cast<const uint8_t*>(&data_print[Number<20>{}]),
// *reinterpret_cast<const uint8_t*>(&data_print[Number<21>{}]),
// *reinterpret_cast<const uint8_t*>(&data_print[Number<22>{}]),
// *reinterpret_cast<const uint8_t*>(&data_print[Number<23>{}]),
// *reinterpret_cast<const uint8_t*>(&data_print[Number<24>{}]),
// *reinterpret_cast<const uint8_t*>(&data_print[Number<25>{}]),
// *reinterpret_cast<const uint8_t*>(&data_print[Number<26>{}]),
// *reinterpret_cast<const uint8_t*>(&data_print[Number<27>{}]),
// *reinterpret_cast<const uint8_t*>(&data_print[Number<28>{}]),
// *reinterpret_cast<const uint8_t*>(&data_print[Number<29>{}]),
// *reinterpret_cast<const uint8_t*>(&data_print[Number<30>{}]),
// *reinterpret_cast<const uint8_t*>(&data_print[Number<31>{}])
);
#endif
constexpr auto move_on_dim = [&]() constexpr
{