fix(ck-tile): Fix compiler issue - #9359
Conversation
❌ PR Check — Action Required
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
There was a problem hiding this comment.
Pull request overview
This PR introduces explicit AMDGPU compiler scheduling barriers in several Composable Kernel XDLops GEMM pipeline implementations to prevent problematic instruction reordering that can lead to numerical correctness mismatches (including large errors and NaNs) observed in miopen_composable_kernel_ut.
Changes:
- Insert
__builtin_amdgcn_sched_barrier(0)after specific prefetch/copy operations to enforce a compile-time scheduling boundary. - Add the same scheduling barrier before certain
block_sync_lds()synchronization points in blockscale-related pipelines.
Reviewed changes
Copilot reviewed 3 out of 3 changed files in this pull request and generated no comments.
| File | Description |
|---|---|
| projects/composablekernel/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v3_mx_bpreshuffle.hpp | Adds scheduling barriers around global prefetch/copy steps to prevent compiler reordering across critical boundaries. |
| projects/composablekernel/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_moe_blockscale_b_preshuffle_v3.hpp | Adds a scheduling barrier immediately before an LDS block sync in the hot loop path. |
| projects/composablekernel/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_blockscale_b_preshuffle_v1.hpp | Adds a scheduling barrier before block_sync_lds() in the blockscale preshuffle pipeline sequence. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
shumway
left a comment
shumway
left a comment
There was a problem hiding this comment.
Have you checked for performance regressions? Being correct is more important than performance, but if there's a significant regressions we should understand that and work on that, too.
Compiler team confirmed that "It lowers to no hardware instruction — purely a compile-time scheduling boundary.", so it has no effect on performance. |
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
## 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
fix(ck-tile): Fix compiler issue ## 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
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
JIRA ID : AICK-1556