From a6028c883b4b8adbf0a575005bce8a300a0b116a Mon Sep 17 00:00:00 2001
From: Rostyslav Geyyer <46627076+geyyer@users.noreply.github.com>
Date: Wed, 15 Jul 2026 17:34:24 +0000
Subject: [PATCH] [rocm-libraries] ROCm/rocm-libraries#9359 (commit 0cb05ba)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

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
---
 ...ockwise_gemm_pipeline_xdlops_blockscale_b_preshuffle_v1.hpp | 1 +
 ...ise_gemm_pipeline_xdlops_moe_blockscale_b_preshuffle_v3.hpp | 1 +
 .../block/blockwise_gemm_pipeline_xdlops_v3_mx_bpreshuffle.hpp | 3 +++
 3 files changed, 5 insertions(+)

diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_blockscale_b_preshuffle_v1.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_blockscale_b_preshuffle_v1.hpp
index 214e21385f..c0cc81c42a 100644
--- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_blockscale_b_preshuffle_v1.hpp
+++ b/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) {
diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_moe_blockscale_b_preshuffle_v3.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_moe_blockscale_b_preshuffle_v3.hpp
index aef501dc86..56785a46d7 100644
--- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_moe_blockscale_b_preshuffle_v3.hpp
+++ b/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) {
diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v3_mx_bpreshuffle.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v3_mx_bpreshuffle.hpp
index 90d7b831b4..a28a36dfd9 100644
--- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v3_mx_bpreshuffle.hpp
+++ b/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);
                         }