[rocm-libraries] ROCm/rocm-libraries#9359 (commit 0cb05ba)

fix(ck-tile): Fix compiler issue
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

## Motivation

Bug causes miopen_composable_kernel_ut to fail due to numerical
correctness mismatches in Composable Kernel
GEMM/MX/blockscale/contraction paths. The failures are not build
failures; the tests run, but device output does not match CPU/reference
output, including large mismatches and NaN outputs.

## Technical Details

__builtin_amdgcn_sched_barrier(0) with mask 0 = "no instruction of any
kind may be reordered across this point, in either direction". It lowers
to no hardware instruction — purely a compile-time scheduling boundary.

## Test Plan

Confirm with the Compiler team

## Test Result

TBD

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

JIRA ID : AICK-1556

include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_blockscale_b_preshuffle_v1.hpp

@@ -520,6 +520,7 @@ struct BlockwiseGemmXdlops_pipeline_blockscale_bpreshuffle_v1<BlockGemmPipelineS
                         });
                     });
 
+                    __builtin_amdgcn_sched_barrier(0);
                     block_sync_lds();
 
                     static_ford<Sequence<MRepeat, KRepeat, KGroup>>{}([&](auto mkk) {

include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_moe_blockscale_b_preshuffle_v3.hpp

@@ -679,6 +679,7 @@ struct BlockwiseGemmXdlops_pipeline_moe_blockscale_bpreshuffle_v3<
 
                         if constexpr(m0.value == (MRepeat - 2))
                         {
+                            __builtin_amdgcn_sched_barrier(0);
                             block_sync_lds();
 
                             static_ford<Sequence<KRepeat, KGroup>>{}([&](auto kk) {

include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v3_mx_bpreshuffle.hpp

@@ -478,6 +478,7 @@ struct BlockwiseGemmXdlops_pipeline_v3_mx_bprehuffle<BlockGemmPipelineScheduler:
 
         // Global prefetch 1
         a_blockwise_copy.Run(a_grid_desc, a_grid_buf, a_block_desc, a_block_bufs(I0));
+        __builtin_amdgcn_sched_barrier(0);
         b_blockwise_copy.Run(
             b_grid_desc, b_grid_buf, b_block_desc, b_block_origin_idx, b_thread_bufs(I0));
 
@@ -556,6 +557,7 @@ struct BlockwiseGemmXdlops_pipeline_v3_mx_bprehuffle<BlockGemmPipelineScheduler:
 
         // Global prefetch 2
         a_blockwise_copy.Run(a_grid_desc, a_grid_buf, a_block_desc, a_block_bufs(I1));
+        __builtin_amdgcn_sched_barrier(0);
         a_blockwise_copy.MoveSrcSliceWindow(a_grid_desc, a_block_copy_step);
 
         // Initialize C
@@ -733,6 +735,7 @@ struct BlockwiseGemmXdlops_pipeline_v3_mx_bprehuffle<BlockGemmPipelineScheduler:
                                                  a_grid_buf,
                                                  a_block_desc,
                                                  a_block_bufs(scale_comp_buf));
+                            __builtin_amdgcn_sched_barrier(0);
                             a_blockwise_copy.MoveSrcSliceWindow(a_grid_desc, a_block_copy_step);
                         }