[CK] enable new gfx1250-strict target. - #12203
Merged
Merged
Conversation
github-actions
Bot
added
project: composablekernel
ck: convolution
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Contributor
There was a problem hiding this comment.
🟡 Changes recommended
Several strict-target WMMA paths remain selectable but silently produce unchanged or zero accumulators.
Get a fresh assessment by requesting another Copilot review.
Pull request overview
Enables CK compilation for gfx1250-strict while avoiding unsupported conversion and WMMA builtins.
Changes:
- Adds strict-target architecture detection and CMake handling.
- Introduces scalar fallbacks for packed MX conversions.
- Adjusts kernels, instances, and tests for strict-target capabilities.
File summaries
| File | Description |
|---|---|
projects/composablekernel/CMakeLists.txt |
Maps the strict target to gfx1250. |
projects/composablekernel/include/ck/ck.hpp |
Adds strict architecture aliases. |
projects/composablekernel/include/ck/utility/amd_wmma.hpp |
Restricts WMMA builtin use. |
projects/composablekernel/include/ck/utility/e4m3.hpp |
Extends gfx125 handling. |
projects/composablekernel/include/ck/utility/e5m3.hpp |
Extends gfx125 handling. |
projects/composablekernel/include/ck/utility/mxf4_utils.hpp |
Restricts FP4 fast paths. |
projects/composablekernel/include/ck/utility/mxf6_utils.hpp |
Restricts FP6 fast paths. |
projects/composablekernel/include/ck/utility/mxf8_utils.hpp |
Restricts FP8 fast paths. |
projects/composablekernel/include/ck/utility/mxfp_utils.hpp |
Adds non-strict architecture detection. |
projects/composablekernel/include/ck/utility/scaled_type_convert.hpp |
Limits native scaled conversions. |
projects/composablekernel/include/ck/utility/type_convert.hpp |
Selects strict-compatible fallbacks. |
projects/composablekernel/include/ck_tile/core/arch/amd_cluster_load.hpp |
Enables cluster loads for gfx125 variants. |
projects/composablekernel/include/ck_tile/core/config.hpp |
Registers strict as gfx125. |
projects/composablekernel/include/ck_tile/core/numeric/e4m3.hpp |
Extends gfx125 conversion handling. |
projects/composablekernel/include/ck_tile/core/numeric/e5m3.hpp |
Extends gfx125 conversion handling. |
projects/composablekernel/include/ck_tile/core/numeric/float8_ext.hpp |
Adds strict FP8 conversion fallbacks. |
projects/composablekernel/include/ck_tile/core/numeric/mxfp_scale.hpp |
Adds strict packed-scale selection. |
projects/composablekernel/include/ck_tile/core/numeric/pk_f6.hpp |
Adds strict FP6 conversion fallback. |
projects/composablekernel/include/ck_tile/core/numeric/pk_fp4.hpp |
Adds strict FP4 conversion fallback. |
projects/composablekernel/include/ck_tile/core/tensor/buffer_view.hpp |
Extends gfx125 buffer handling. |
projects/composablekernel/include/ck_tile/ops/flatmm/pipeline/mx_flatmm_pipeline_agmem_bgmem_creg_v1_policy.hpp |
Extends gfx125 tile distributions. |
projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v3.hpp |
Disables the 8-warp schedule for strict. |
projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl_8bit_traits.hpp |
Adds decomposed strict WMMA handling. |
projects/composablekernel/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_weight_kernel.hpp |
Applies gfx125 stream-K restrictions. |
projects/composablekernel/library/src/tensor_operation_instance/gpu/CMakeLists.txt |
Excludes strict from MHA instances. |
projects/composablekernel/library/src/tensor_operation_instance/gpu/gemm_multiply_multiply/device_gemm_multiply_multiply_wmma_c_shuffle_f8_f8_bf16_mk_nk_mn.cpp |
Restricts unsupported WMMA tiles. |
projects/composablekernel/library/src/tensor_operation_instance/gpu/gemm_multiply_multiply/device_gemm_multiply_multiply_wmma_c_shuffle_f8_f8_f16_mk_nk_mn.cpp |
Restricts unsupported WMMA tiles. |
projects/composablekernel/library/src/tensor_operation_instance/gpu/gemm_multiply_multiply/device_gemm_multiply_multiply_wmma_c_shuffle_i8_i8_bf16_mk_nk_mn.cpp |
Restricts unsupported WMMA tiles. |
projects/composablekernel/library/src/tensor_operation_instance/gpu/gemm_multiply_multiply/device_gemm_multiply_multiply_wmma_c_shuffle_i8_i8_f16_mk_nk_mn.cpp |
Restricts unsupported WMMA tiles. |
projects/composablekernel/library/src/tensor_operation_instance/gpu/gemm_universal/device_gemm_wmma_universal_bf16_i4_bf16/device_gemm_wmma_universal_bf16_i4_bf16_km_nk_mn.hpp |
Restricts problematic gfx125 instances. |
projects/composablekernel/library/src/tensor_operation_instance/gpu/gemm_universal/device_gemm_wmma_universal_bf16_i4_bf16/device_gemm_wmma_universal_bf16_i4_bf16_mk_nk_mn.hpp |
Restricts problematic gfx125 instances. |
projects/composablekernel/library/src/tensor_operation_instance/gpu/gemm_universal/device_gemm_wmma_universal_f16_f8_f16/device_gemm_wmma_universal_f16_f8_f16_km_kn_mn.hpp |
Restricts problematic gfx125 instances. |
projects/composablekernel/library/src/tensor_operation_instance/gpu/gemm_universal/device_gemm_wmma_universal_f16_f8_f16/device_gemm_wmma_universal_f16_f8_f16_km_nk_mn.hpp |
Restricts problematic gfx125 instances. |
projects/composablekernel/library/src/tensor_operation_instance/gpu/gemm_universal/device_gemm_wmma_universal_f8_f8_bf16/device_gemm_wmma_universal_f8_f8_bf16_mk_kn_mn.hpp |
Restricts unsupported gfx125 tiles. |
projects/composablekernel/library/src/tensor_operation_instance/gpu/gemm_universal/device_gemm_wmma_universal_f8_f8_bf16/device_gemm_wmma_universal_f8_f8_bf16_mk_nk_mn.hpp |
Restricts unsupported gfx125 tiles. |
projects/composablekernel/library/src/tensor_operation_instance/gpu/quantization/gemm/device_gemm_quantization_wmma_c_shuffle_i8_i8_i8_instance.hpp |
Restricts unsupported gfx125 tiles. |
projects/composablekernel/test/CMakeLists.txt |
Excludes strict from SMFMAC tests. |
projects/composablekernel/test/ck_tile/multicast_load/test_cluster_load_async_to_lds.cpp |
Enables cluster tests for gfx125 variants. |
projects/composablekernel/test/data_type/CMakeLists.txt |
Excludes unsupported strict data-type tests. |
projects/composablekernel/example/ck_tile/42_mx_gemm/run_mx_gemm.inc |
Reformats a scale-buffer call. |
Review details
Suppressed comments (3)
projects/composablekernel/include/ck/utility/amd_wmma.hpp:1374
- This makes the scale16 operation a silent no-op on
gfx1250-strict, althoughGetMfmastill returnswmma_scale16_f32_16x16x128_f8f6f4_gfx125unconditionally atxdlops_gemm.hpp:2093-2105and the strict build includes its correctness tests. Route strict builds to a supported implementation or mark/filter this operation as unavailable instead of leaving the accumulator unchanged.
#if defined(__gfx1250__)
projects/composablekernel/include/ck/utility/amd_wmma.hpp:1429
- For
gfx1250-strict, the 32x16 FP4 operation now compiles as a no-op that leavesreg_cunchanged. It remains publicly callable throughmfma_type(xdlops_gemm.hpp:1509-1543), andtest_mx_coexecutionis still built for everygfx125*target and calls this exact operation. Please remove/filter this operation for strict or implement it using supported instructions.
#if defined(__gfx1250__)
projects/composablekernel/include/ck/utility/amd_wmma.hpp:1483
- The scale16 32x16 FP4 path is also still exposed and exercised by
test_mx_coexecutionongfx1250-strict, but this guard makesRunleave the accumulator untouched. Please provide a supported decomposition/fallback or make the operation unavailable for strict builds so callers cannot receive silently incorrect results.
#if defined(__gfx1250__)
- Files reviewed: 40/40 changed files
- Comments generated: 2
- Review effort level: Balanced
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
andriy-ca
requested changes
Sep 17, 2026
andriy-ca
left a comment
andriy-ca
left a comment
Contributor
There was a problem hiding this comment.
Could you please document which defines or their combinations are expected to be defined under gfx1250-strict? What do we expect to be defined for A0? For B0?
Contributor
Author
Yes, so the gfx1250 = B0, gfx1250_strict = A0, gfx125 = A0 + B0. |
…a passing test on gfx1250 emulator
assistant-librarian Bot
pushed a commit
to ROCm/composable_kernel
that referenced
this pull request
Sep 18, 2026
[CK] enable new gfx1250-strict target. ## Motivation These changes will enable building CK for the new gfx1250-strict target by limiting access to some of the unsupported builtins. I have verified that the code builds successfully for the gfx1250-strict with a rocm/compiler installation that supports it. JIRA ID : AICK-2239 ## Technical Details <!-- Explain the changes along with any relevant GitHub links. --> ## Test Plan <!-- Explain any relevant testing done to verify this PR. --> ## Test Result <!-- Briefly summarize test outcomes. --> ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
These changes will enable building CK for the new gfx1250-strict target by limiting access to some of the unsupported builtins. I have verified that the code builds successfully for the gfx1250-strict with a rocm/compiler installation that supports it.
JIRA ID : AICK-2239
Technical Details
Test Plan
Test Result
Submission Checklist