ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-06-30 11:47:48 +00:00
Code Issues Packages Projects Releases Wiki Activity

fixed synchronization issue in block gemm pipeline v1 that caused b_scale to fail

Browse Source
This commit is contained in:
Kevin Abraham
2025-10-30 21:16:02 +00:00
committed by Ahamed
parent 833ae1d051
commit 4aeb73c616
1 changed files with 8 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_wmmaops_v1.hpp
View File

@@ -289,6 +289,10 @@ struct BlockwiseGemmWmmaops_pipeline_v1<BlockGemmPipelineScheduler::Intrawave,
block_sync_lds();
b_scale_struct.template GlobalLoad<0>((i + 2) % num_loop_per_scale == 0);
if constexpr(ck::is_same<BScaleStruct, Empty>::value == false)
{
block_sync_lds();
}
a_blockwise_copy.RunWrite(a_block_desc, a_block_buf);
b_blockwise_copy.RunWrite(b_block_desc, b_block_buf);
@@ -631,6 +635,10 @@ struct BlockwiseGemmWmmaops_pipeline_v1<BlockGemmPipelineScheduler::Interwave,
blockwise_gemm_func();
b_scale_struct.template GlobalLoad<0>((i + 2) % num_loop_per_scale == 0);
if constexpr(ck::is_same<BScaleStruct, Empty>::value == false)
{
block_sync_lds();
}
a_blockwise_copy.RunWrite(a_block_desc, a_block_buf);
b_blockwise_copy.RunWrite(b_block_desc, b_block_buf);