From 0987b0af441311af201ec3530ab21bbb6c2b7a00 Mon Sep 17 00:00:00 2001 From: aska-0096 Date: Fri, 9 May 2025 16:07:22 +0000 Subject: [PATCH] remove unnecessary hacky --- example/67_gemm_microscaling/gemm_mx_fp4.cpp | 2 +- ...blockwise_gemm_mx_pipeline_xdlops_base.hpp | 14 ++-- .../blockwise_gemm_pipeline_xdlops_v3_mx.hpp | 14 ++-- .../threadwise_tensor_slice_transfer_v3r1.hpp | 72 ------------------- 4 files changed, 16 insertions(+), 86 deletions(-) diff --git a/example/67_gemm_microscaling/gemm_mx_fp4.cpp b/example/67_gemm_microscaling/gemm_mx_fp4.cpp index 3b84d12c3b..01d57d0471 100644 --- a/example/67_gemm_microscaling/gemm_mx_fp4.cpp +++ b/example/67_gemm_microscaling/gemm_mx_fp4.cpp @@ -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 diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_mx_pipeline_xdlops_base.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_mx_pipeline_xdlops_base.hpp index 52b66703dd..aa5a899779 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_mx_pipeline_xdlops_base.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_gemm_mx_pipeline_xdlops_base.hpp @@ -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, 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 @@ -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, diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v3_mx.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v3_mx.hpp index ffb97c2297..e9ddf8346a 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v3_mx.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v3_mx.hpp @@ -423,7 +423,7 @@ struct BlockwiseGemmXdlops_pipeline_v3_mx{}([&](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{}([&](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{}([&](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{}([&](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{}([&](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{}([&](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, diff --git a/include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v3r1.hpp b/include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v3r1.hpp index e6c3b9a76c..1592d77286 100644 --- a/include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v3r1.hpp +++ b/include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v3r1.hpp @@ -314,78 +314,6 @@ struct ThreadwiseTensorSliceTransfer_v3r1 src_thread_scratch_tuple_(thread_scratch_id) .template SetAsType(src_data_idx_seq, op_r_v.template AsType()[I0]); - CK_PRINT(); - // using a = ck::StaticTensorTupleOfVectorBuffer< - // ck::AddressSpaceEnum::Vgpr, - // ck::f4x2_pk_t, - // 8, - // const ck::TensorDescriptor< - // ck::Tuple, - // ck::integral_constant, - // ck::integral_constant, - // ck::integral_constant>, - // false>, - // ck::PassThrough>, - // ck::PassThrough>, - // ck::Merge_v3_division_mod, - // ck::integral_constant>>>, - // ck::Tuple, - // ck::Sequence<1>, - // ck::Sequence<2>, - // ck::Sequence<3, 4>>, - // ck::Tuple, - // ck::Sequence<5>, - // ck::Sequence<6>, - // ck::Sequence<7>>, - // ck::Sequence<5, 6, 7>, - // ck::integral_constant>, - // 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(data_print.s_per_v), - static_cast(data_print.num_of_v_), - *reinterpret_cast(&data_print[Number<0>{}]), - *reinterpret_cast(&data_print[Number<1>{}]), - *reinterpret_cast(&data_print[Number<2>{}]), - *reinterpret_cast(&data_print[Number<3>{}]), - *reinterpret_cast(&data_print[Number<4>{}]), - *reinterpret_cast(&data_print[Number<5>{}]), - *reinterpret_cast(&data_print[Number<6>{}]), - *reinterpret_cast(&data_print[Number<7>{}]), - *reinterpret_cast(&data_print[Number<8>{}]), - *reinterpret_cast(&data_print[Number<9>{}]), - *reinterpret_cast(&data_print[Number<10>{}]), - *reinterpret_cast(&data_print[Number<11>{}]), - *reinterpret_cast(&data_print[Number<12>{}]), - *reinterpret_cast(&data_print[Number<13>{}]), - *reinterpret_cast(&data_print[Number<14>{}]), - *reinterpret_cast(&data_print[Number<15>{}]) - // *reinterpret_cast(&data_print[Number<16>{}]), - // *reinterpret_cast(&data_print[Number<17>{}]), - // *reinterpret_cast(&data_print[Number<18>{}]), - // *reinterpret_cast(&data_print[Number<19>{}]), - // *reinterpret_cast(&data_print[Number<20>{}]), - // *reinterpret_cast(&data_print[Number<21>{}]), - // *reinterpret_cast(&data_print[Number<22>{}]), - // *reinterpret_cast(&data_print[Number<23>{}]), - // *reinterpret_cast(&data_print[Number<24>{}]), - // *reinterpret_cast(&data_print[Number<25>{}]), - // *reinterpret_cast(&data_print[Number<26>{}]), - // *reinterpret_cast(&data_print[Number<27>{}]), - // *reinterpret_cast(&data_print[Number<28>{}]), - // *reinterpret_cast(&data_print[Number<29>{}]), - // *reinterpret_cast(&data_print[Number<30>{}]), - // *reinterpret_cast(&data_print[Number<31>{}]) - ); -#endif constexpr auto move_on_dim = [&]() constexpr {