diff --git a/example/01_gemm/run_gemm_example_streamk_v2.inc b/example/01_gemm/run_gemm_example_streamk_v2.inc index 6527472113..b4c37bb4e3 100644 --- a/example/01_gemm/run_gemm_example_streamk_v2.inc +++ b/example/01_gemm/run_gemm_example_streamk_v2.inc @@ -12,14 +12,14 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config) using namespace ck::literals; - auto M = problem_size.M; - auto N = problem_size.N; - auto K = problem_size.K; - auto StrideA = problem_size.StrideA; - auto StrideB = problem_size.StrideB; - auto StrideC = problem_size.StrideC; - // auto Grid_size = problem_size.Grid_size; - // auto Streamk_sel = problem_size.Streamk_sel; + auto M = problem_size.M; + auto N = problem_size.N; + auto K = problem_size.K; + auto StrideA = problem_size.StrideA; + auto StrideB = problem_size.StrideB; + auto StrideC = problem_size.StrideC; + auto Grid_size = problem_size.Grid_size; + auto Streamk_sel = problem_size.Streamk_sel; auto f_host_tensor_descriptor = [](std::size_t row, std::size_t col, std::size_t stride, auto layout) { @@ -51,20 +51,20 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config) return static_cast(stride); }; - /**auto f_get_default_streamk_policy = [](ck::index_t streamk_sel) { + auto f_get_default_streamk_policy = [](ck::index_t streamk_sel) { if(streamk_sel == -1) { return static_cast(4); } else return static_cast(streamk_sel); - };**/ + }; StrideA = f_get_default_stride(M, K, StrideA, ALayout{}); StrideB = f_get_default_stride(K, N, StrideB, BLayout{}); StrideC = f_get_default_stride(M, N, StrideC, CLayout{}); - // Streamk_sel = f_get_default_streamk_policy(Streamk_sel); + Streamk_sel = f_get_default_streamk_policy(Streamk_sel); Tensor a_m_k(f_host_tensor_descriptor(M, K, StrideA, ALayout{})); Tensor b_k_n(f_host_tensor_descriptor(K, N, StrideB, BLayout{})); @@ -148,10 +148,11 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config) StrideA, StrideB, StrideC, + Streamk_sel, + Grid_size, a_element_op, b_element_op, - c_element_op, - problem_size.NumSKBlocks); // TODO: add NumSKBlocks arg into Problem? + c_element_op); // TODO: add NumSKBlocks arg into Problem? if(!gemm.IsSupportedArgument(argument)) { diff --git a/include/ck/tensor_operation/gpu/device/device_gemm_streamk_v2.hpp b/include/ck/tensor_operation/gpu/device/device_gemm_streamk_v2.hpp index 8ceb770fc7..1a4d684f14 100644 --- a/include/ck/tensor_operation/gpu/device/device_gemm_streamk_v2.hpp +++ b/include/ck/tensor_operation/gpu/device/device_gemm_streamk_v2.hpp @@ -20,19 +20,21 @@ template struct DeviceGemm_Streamk_V2 : public BaseOperator { - virtual std::unique_ptr MakeArgumentPointer(const void* p_a, - const void* p_b, - void* p_c, - ck::index_t M, - ck::index_t N, - ck::index_t K, - ck::index_t StrideA, - ck::index_t StrideB, - ck::index_t StrideC, - AElementwiseOperation a_element_op, - BElementwiseOperation b_element_op, - CElementwiseOperation c_element_op, - index_t NumSKBlocks = 0) = 0; + virtual std::unique_ptr + MakeArgumentPointer(const void* p_a, + const void* p_b, + void* p_c, + ck::index_t M, + ck::index_t N, + ck::index_t K, + ck::index_t StrideA, + ck::index_t StrideB, + ck::index_t StrideC, + ck::index_t Streamk_sel, + ck::index_t Grid_size, + AElementwiseOperation a_element_op, + BElementwiseOperation b_element_op, + CElementwiseOperation c_element_op) = 0; virtual std::unique_ptr MakeInvokerPointer() = 0; }; diff --git a/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_streamk_v3.hpp b/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_streamk_v3.hpp index ccffbb1b44..b3cc5df8b7 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_streamk_v3.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_streamk_v3.hpp @@ -159,16 +159,21 @@ struct DeviceGemm_Xdl_CShuffle_Streamk_V3 : public DeviceGemm_Streamk_V2(num_cu), - static_cast(occupancy), - NumSKBlocks}; + return Argument{p_a, p_b, p_c, M, N, K, StrideA, StrideB, StrideC, streamk_sel, Grid_size}; } static auto MakeInvoker() { return Invoker{}; } @@ -726,239 +721,12 @@ struct DeviceGemm_Xdl_CShuffle_Streamk_V3 : public DeviceGemm_Streamk_V2; - calculate_grid_size(kernel); - } - // Tail number could be One to Seven - else if constexpr(BlkGemmPipelineVer == BlockGemmPipelineVersion::v2) - { - - if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::One) - { - // printf("Case 2\n"); - const auto kernel = kernel_gemm_xdl_cshuffle_v3; - calculate_grid_size(kernel); - } - else if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::Full) - { - // printf("Case 3\n"); - const auto kernel = kernel_gemm_xdl_cshuffle_v3; - calculate_grid_size(kernel); - } - - if constexpr(GridwiseGemm::BlockwiseGemmPipe::PrefetchStages > 2) - { - if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::Two) - { - // printf("Case 4\n"); - const auto kernel = - kernel_gemm_xdl_cshuffle_v3; - calculate_grid_size(kernel); - } - } - - if constexpr(GridwiseGemm::BlockwiseGemmPipe::PrefetchStages > 3) - { - if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::Three) - { - // printf("Case 5\n"); - const auto kernel = - kernel_gemm_xdl_cshuffle_v3; - calculate_grid_size(kernel); - } - } - - if constexpr(GridwiseGemm::BlockwiseGemmPipe::PrefetchStages > 4) - { - if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::Four) - { - // printf("Case 6\n"); - const auto kernel = - kernel_gemm_xdl_cshuffle_v3; - calculate_grid_size(kernel); - } - } - - if constexpr(GridwiseGemm::BlockwiseGemmPipe::PrefetchStages > 5) - { - if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::Five) - { - // printf("Case 7\n"); - const auto kernel = - kernel_gemm_xdl_cshuffle_v3; - calculate_grid_size(kernel); - } - } - - if constexpr(GridwiseGemm::BlockwiseGemmPipe::PrefetchStages > 6) - { - if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::Six) - { - // printf("Case 8\n"); - const auto kernel = - kernel_gemm_xdl_cshuffle_v3; - calculate_grid_size(kernel); - } - } - - if constexpr(GridwiseGemm::BlockwiseGemmPipe::PrefetchStages > 7) - { - if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::Seven) - { - // printf("Case 9\n"); - const auto kernel = - kernel_gemm_xdl_cshuffle_v3; - calculate_grid_size(kernel); - } - } - } - // Tail number could be Odd or Even - else if constexpr(BlkGemmPipelineVer == BlockGemmPipelineVersion::v4) - { - - if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::Odd) - { - // printf("Case 10\n"); - const auto kernel = - kernel_gemm_xdl_cshuffle_v3_2lds; - calculate_grid_size(kernel); - } - else - { - // printf("Case 11\n"); - const auto kernel = - kernel_gemm_xdl_cshuffle_v3_2lds; - calculate_grid_size(kernel); - } - } - else - { - - if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::Odd) - { - // printf("Case 12\n"); - const auto kernel = kernel_gemm_xdl_cshuffle_v3; - calculate_grid_size(kernel); - } - else - { - // printf("Case 13\n"); - const auto kernel = kernel_gemm_xdl_cshuffle_v3; - calculate_grid_size(kernel); - } - } - } - else - { - // printf("not main k block loop\n"); - // Tail number always 1 - if constexpr(BlkGemmPipelineVer == BlockGemmPipelineVersion::v1 || - BlkGemmPipelineVer == BlockGemmPipelineVersion::v2 || - BlkGemmPipelineVer == BlockGemmPipelineVersion::v3 || - BlkGemmPipelineVer == BlockGemmPipelineVersion::v5) - { - - // printf("Case 14\n"); - const auto kernel = kernel_gemm_xdl_cshuffle_v3; - calculate_grid_size(kernel); - } - else if constexpr(BlkGemmPipelineVer == BlockGemmPipelineVersion::v4) - { - - // printf("Case 15\n"); - const auto kernel = kernel_gemm_xdl_cshuffle_v3_2lds; - calculate_grid_size(kernel); - } - } - // printf("num_cu: %u\n", static_cast(num_cu)); - // printf("occupancy: %u\n", static_cast(occupancy)); return std::make_unique(static_cast(p_a), static_cast(p_b), static_cast(p_c), @@ -968,9 +736,8 @@ struct DeviceGemm_Xdl_CShuffle_Streamk_V3 : public DeviceGemm_Streamk_V2(num_cu), - static_cast(occupancy), - static_cast(NumSKBlocks)); + streamk_sel, + Grid_size); } // polymorphic diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_streamk_v3.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_streamk_v3.hpp index c8f5be068a..9a944e11a6 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_streamk_v3.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_streamk_v3.hpp @@ -513,18 +513,16 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3 index_t StrideA_, index_t StrideB_, index_t StrideC_, - uint32_t num_cu_, - uint32_t occupancy_, - uint32_t num_sk_blocks_) + index_t Streamk_sel_, + index_t Grid_size_) : M{M_}, N{N_}, K{K_}, StrideA{StrideA_}, StrideB{StrideB_}, StrideC{StrideC_}, - num_cu{num_cu_}, - occupancy{occupancy_}, - num_sk_blocks{num_sk_blocks_}, + Streamk_sel{Streamk_sel_}, + Grid_size{Grid_size_}, MPadded{CalculateMPadded(M_)}, NPadded{CalculateNPadded(N_)}, KRead{CalculateKRead(K_, 1)}, @@ -590,35 +588,26 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3 index_t StrideA_, index_t StrideB_, index_t StrideC_, - uint32_t num_cu_, - uint32_t occupancy_, - uint32_t num_sk_blocks_) - : Problem{M_, - N_, - K_, - StrideA_, - StrideB_, - StrideC_, - num_cu_, - occupancy_, - num_sk_blocks_}, + index_t Streamk_sel_, + index_t Grid_size_) + : Problem{M_, N_, K_, StrideA_, StrideB_, StrideC_, Streamk_sel_, Grid_size_}, p_a_grid{p_a_grid_}, p_b_grid{p_b_grid_}, p_c_grid{p_c_grid_}, block_2_ctile_map_streamk( - M_, N_, AK0Number * CalculateKPadded(K_, 1), num_cu_, occupancy_, num_sk_blocks_) - + M_, N_, AK0Number * CalculateKPadded(K_, 1), Grid_size_, Streamk_sel_) { } const ADataType* p_a_grid; const BDataType* p_b_grid; CDataType* p_c_grid; - BlockToCTileMap_GemmStreamK + BlockToCTileMap_GemmStreamK_v2 block_2_ctile_map_streamk; }; @@ -1250,11 +1239,12 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3 }(); return c_partial_acc_block_m_n; } - using Block2CTileMap_streamk = BlockToCTileMap_GemmStreamK; + using Block2CTileMap_streamk = BlockToCTileMap_GemmStreamK_v2; template ( p_c_grid, c_grid_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize()); - // lds max alignment - constexpr auto max_lds_align = math::lcm(AK1Number, BK1Number); - - // A matrix in LDS memory, dst of blockwise copy - constexpr auto a_block_desc_ak0_m_ak1 = GetABlockDescriptor_AK0PerBlock_MPerBlock_AK1(); - - // B matrix in LDS memory, dst of blockwise copy - constexpr auto b_block_desc_bk0_n_bk1 = GetBBlockDescriptor_BK0PerBlock_NPerBlock_BK1(); - - // LDS allocation for A and B: be careful of alignment - constexpr auto a_block_space_size_aligned = math::integer_least_multiple( - a_block_desc_ak0_m_ak1.GetElementSpaceSize(), max_lds_align); - - // Cast after lds - auto a_block_buf = make_dynamic_buffer( - static_cast(p_shared), a_block_desc_ak0_m_ak1.GetElementSpaceSize()); - - auto b_block_buf = make_dynamic_buffer( - static_cast(p_shared) + - a_block_space_size_aligned * sizeof(ADataType) / sizeof(BDataType), - b_block_desc_bk0_n_bk1.GetElementSpaceSize()); - - constexpr auto a_block_slice_copy_step = make_multi_index(KPerBlock / AK1Number, 0, 0); - constexpr auto b_block_slice_copy_step = make_multi_index(KPerBlock / BK1Number, 0, 0); - - // Blockwise GEMM pipeline - static_assert(std::is_default_constructible_v); - auto blockwise_gemm_pipeline = BlockwiseGemmPipe{}; - auto c_thread_buf = blockwise_gemm_pipeline.GetCThreadBuffer(); - uint32_t* p_semaphore = reinterpret_cast( reinterpret_cast(p_workspace) + block_2_ctile_map_streamk.get_workspace_size_for_acc(sizeof(AccDataType))); @@ -1336,21 +1295,14 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3 static_cast(block_idx) >= block_2_ctile_map_streamk.dp_start_block_idx && static_cast(block_idx) < block_2_ctile_map_streamk.reduction_start_block_idx; - is_reduction_block = - static_cast(block_idx) >= block_2_ctile_map_streamk.reduction_start_block_idx; - is_padding_block = - static_cast(block_idx) >= block_2_ctile_map_streamk.sk_num_blocks && - static_cast(block_idx) < block_2_ctile_map_streamk.dp_start_block_idx; - if(is_padding_block) - { - return; - } block_2_ctile_map_streamk.get_block_itr(block_idx, iter_start, iter_end); - uint32_t total_iter_length = iter_end - iter_start; + num_k_block_main_loop = iter_end - iter_start; if constexpr(Block2CTileMap_streamk::ReductionStrategy == StreamKReductionStrategy::Reduction) { + is_reduction_block = static_cast(block_idx) >= + block_2_ctile_map_streamk.reduction_start_block_idx; if(is_reduction_block) { // descriptors @@ -1557,7 +1509,7 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3 { uint32_t current_iter_length = __builtin_amdgcn_readfirstlane(block_2_ctile_map_streamk.get_current_iter_length( - iter_start, iter_end, total_iter_length)); + iter_start, iter_end, num_k_block_main_loop)); uint32_t tile_idx, iter_offset; block_2_ctile_map_streamk.get_tile_idx_with_offset(iter_end - 1, tile_idx, iter_offset); iter_offset = __builtin_amdgcn_readfirstlane(iter_offset - current_iter_length + 1); @@ -1573,6 +1525,15 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3 const index_t k0_block_data_idx_on_grid = __builtin_amdgcn_readfirstlane(iter_offset * AK0Number); + // lds max alignment + constexpr auto max_lds_align = math::lcm(AK1Number, BK1Number); + + // A matrix in LDS memory, dst of blockwise copy + constexpr auto a_block_desc_ak0_m_ak1 = GetABlockDescriptor_AK0PerBlock_MPerBlock_AK1(); + + // B matrix in LDS memory, dst of blockwise copy + constexpr auto b_block_desc_bk0_n_bk1 = GetBBlockDescriptor_BK0PerBlock_NPerBlock_BK1(); + // A matrix blockwise copy auto a_blockwise_copy = ThreadGroupTensorSliceTransfer_v4r1, + Sequence<0, 1, 2>, ABlockTransferSrcVectorDim, 2, ABlockTransferSrcScalarPerVector, @@ -1618,7 +1579,7 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3 decltype(b_grid_desc_bk0_n_bk1), decltype(b_block_desc_bk0_n_bk1), BBlockTransferSrcAccessOrder, - Sequence<1, 0, 2>, + Sequence<0, 1, 2>, BBlockTransferSrcVectorDim, 2, BBlockTransferSrcScalarPerVector, @@ -1635,7 +1596,30 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3 make_multi_index(0, 0, 0), ck::tensor_operation::element_wise::PassThrough{}); - const index_t num_k_block_main_loop = current_iter_length; + // LDS allocation for A and B: be careful of alignment + constexpr auto a_block_space_size_aligned = math::integer_least_multiple( + a_block_desc_ak0_m_ak1.GetElementSpaceSize(), max_lds_align); + + // Cast after lds + auto a_block_buf = make_dynamic_buffer( + static_cast(p_shared), a_block_desc_ak0_m_ak1.GetElementSpaceSize()); + + auto b_block_buf = make_dynamic_buffer( + static_cast(p_shared) + + a_block_space_size_aligned * sizeof(ADataType) / sizeof(BDataType), + b_block_desc_bk0_n_bk1.GetElementSpaceSize()); + + constexpr auto a_block_slice_copy_step = make_multi_index(KPerBlock / AK1Number, 0, 0); + constexpr auto b_block_slice_copy_step = make_multi_index(KPerBlock / BK1Number, 0, 0); + + // Blockwise GEMM pipeline + static_assert(std::is_default_constructible_v); + auto blockwise_gemm_pipeline = BlockwiseGemmPipe{}; + auto c_thread_buf = blockwise_gemm_pipeline.GetCThreadBuffer(); + + num_k_block_main_loop = __builtin_amdgcn_readfirstlane( + (a_grid_desc_ak0_m_ak1.GetLength(I0) * a_grid_desc_ak0_m_ak1.GetLength(I2)) / + KPerBlock); /**if (threadIdx.x == 0 && threadIdx.y == 0 ){ printf("num_k_block_main_loop: %d\n", num_k_block_main_loop); }**/ @@ -1828,7 +1812,107 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3 make_multi_index(0, 0, 0, 0), c_element_op}; - constexpr auto mxdlperwave_forward_step = + // space filling curve for threadwise C in VGPR + constexpr auto sfc_c_vgpr = + SpaceFillingCurve, + Sequence<0, 1, 2, 3, 4, 5, 6, 7>, + Sequence>{}; + + // space filling curve for shuffled blockwise C in global mem + constexpr auto sfc_c_global = + SpaceFillingCurve, + Sequence<0, 2, 1, 3>, + Sequence<1, + CShuffleMXdlPerWavePerShuffle * MWave * MPerXdl, + 1, + CShuffleNXdlPerWavePerShuffle * NWave * NPerXdl>>{}; + + constexpr index_t num_access = sfc_c_vgpr.GetNumOfAccess(); + + static_assert(num_access == sfc_c_global.GetNumOfAccess(), "wrong!"); + + static_for<0, num_access, 1>{}([&](auto access_id) { + // make sure it's safe to write to LDS + block_sync_lds(); + + // each thread write its data from VGPR to LDS + c_thread_copy_vgpr_to_lds.Run(c_m0_n0_m1_n1_m2_m3_m4_n2_thread_desc, + sfc_c_vgpr.GetIndexTupleOfNumber(access_id), + c_thread_buf, + c_block_desc_m0_n0_m1_n1_m2_m3_m4_n2, + c_block_buf); + + // make sure it's safe to read from LDS + block_sync_lds(); + c_block_copy_lds_to_global.SetSrcSliceOrigin( + c_block_desc_mblock_mpershuffle_nblock_npershuffle, make_tuple(0, 0, 0, 0)); + + if(is_dp_block) + { + // each block copy its data from LDS to global + c_block_copy_lds_to_global.template Run( + c_block_desc_mblock_mpershuffle_nblock_npershuffle, + c_block_buf, + c_grid_desc_mblock_mperblock_nblock_nperblock, + c_grid_buf); + } + else if(is_sk_block) + { + if constexpr(Block2CTileMap_streamk::ReductionStrategy == + StreamKReductionStrategy::Atomic) + { + // each block copy its data from LDS to global + c_block_copy_lds_to_global + .template Run( + c_block_desc_mblock_mpershuffle_nblock_npershuffle, + c_block_buf, + c_grid_desc_mblock_mperblock_nblock_nperblock, + c_grid_buf); + } + else if constexpr(Block2CTileMap_streamk::ReductionStrategy == + StreamKReductionStrategy::Reduction) + { + // constexpr offset + c_block_copy_lds_to_partial_acc.SetSrcSliceOrigin( + c_block_desc_mblock_mpershuffle_nblock_npershuffle, + make_tuple(0, 0, 0, 0)); + + c_block_copy_lds_to_partial_acc.SetDstSliceOrigin( + c_block_desc_mshuffle_mpershuffle_nshuffle_npershuffle, + make_tuple(MXdlPerWave, 0, NXdlPerWave, 0)); + + c_block_copy_lds_to_partial_acc + .template Run( + c_block_desc_mblock_mpershuffle_nblock_npershuffle, + c_block_buf, + c_block_desc_mshuffle_mpershuffle_nshuffle_npershuffle, + c_partial_acc_buf); + } + } + + if constexpr(access_id < num_access - 1) + { + constexpr auto c_global_step = sfc_c_global.GetForwardStep(access_id); + + // move on C + c_block_copy_lds_to_global.MoveDstSliceWindow( + c_grid_desc_mblock_mperblock_nblock_nperblock, c_global_step); + } + }); + /**constexpr auto mxdlperwave_forward_step = make_multi_index(0, CShuffleMXdlPerWavePerShuffle * MWave * MPerXdl, 0, 0); constexpr auto nxdlperwave_forward_step = make_multi_index(0, 0, 0, CShuffleNXdlPerWavePerShuffle * NWave * NPerXdl); @@ -1940,9 +2024,9 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3 c_grid_desc_mblock_mperblock_nblock_nperblock, mxdlperwave_forward_step); } - }); + });**/ - /**if constexpr(Block2CTileMap_streamk::ReductionStrategy == + if constexpr(Block2CTileMap_streamk::ReductionStrategy == StreamKReductionStrategy::Reduction) { if(is_sk_block) @@ -1951,7 +2035,7 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3 workgroup_barrier wg_barrier(p_semaphore); wg_barrier.inc(tile_idx); } - }**/ + } } // exit condition @@ -1966,7 +2050,7 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3 } // make sure next loop LDS is ready for use block_sync_lds(); - if constexpr(Block2CTileMap_streamk::ReductionStrategy == + /**if constexpr(Block2CTileMap_streamk::ReductionStrategy == StreamKReductionStrategy::Reduction) { if(is_sk_block) @@ -1975,7 +2059,7 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3 workgroup_barrier wg_barrier(p_semaphore); wg_barrier.inc(tile_idx); } - } + }**/ } } @@ -2021,9 +2105,8 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3 Block2CTileMap_streamk block_2_ctile_map_streamk(problem.M, problem.N, AK0Number * problem.KPadded, - problem.num_cu, - problem.occupancy, - problem.num_sk_blocks); + problem.Grid_size, + problem.Streamk_sel); auto block_idx = get_block_1d_id(); is_sk_block = static_cast(block_idx) < block_2_ctile_map_streamk.sk_num_blocks;