Repository navigation
fix(ck,ck_tile): microscale pipeline B and BQ layout mismatch & bug in mxfp conversion with float scale - #11560
Merged
Merged
Conversation
✅ 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. |
… FP8
fix clang format
fix clang format in ck examples and host utility headers
clang-format-18 requires a space before an empty function body following a
constructor initializer list. These nine files predate that and were failing
the repo-wide format CI stage.
Co-authored-by: Cursor <cursoragent@cursor.com>
revert clang-format spacing to match CI's clang-format-18
CI formats an empty function body following a constructor initializer list
as `){};` with no space, which is what develop uses. The locally installed
clang-format 18.1.8 produces `) {};` instead. Earlier commits followed the
local binary, so each one added more files to the CI format failure (1 file
at build #2, 2 at #3, 11 at #4). This restores develop's spacing on all 11
affected files; the ck_tile file keeps its gfx125 functional changes.
Co-authored-by: Cursor <cursoragent@cursor.com>
JiaLuo-CAN
force-pushed
the
users/jialuo/ck/aick_1884
branch
from
September 1, 2026 23:02
1491815 to
8791408
Compare
Contributor
There was a problem hiding this comment.
🔵 Needs a closer look
The hardware-specific lane mapping and tile distributions require gfx1250 validation, but no test results were provided.
Pull request overview
Fixes gfx1250 microscale GEMM scale-layout alignment and float-scale conversion correctness.
Changes:
- Aligns BQuant distribution with B’s warp/thread-raked layout.
- Corrects per-lane FP4 and FP8 scaling.
- Uses native packed FP4 vector conversions on gfx1250.
File summaries
| File | Description |
|---|---|
gemm_microscale_pipeline_ag_bg_cr_policy.hpp |
Mirrors B access patterns for BQuant. |
unary_element_wise_operation.hpp |
Uses gfx1250 FP4 vector conversions. |
pk_fp4.hpp |
Adds per-lane FP4 scale selection. |
float8_ext.hpp |
Corrects per-lane FP8 scale selection. |
Review details
- Files reviewed: 4/4 changed files
- Comments generated: 0
- 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
reviewed
Sep 2, 2026
Contributor
Author
|
Some changes according to the review comments:
|
assistant-librarian Bot
pushed a commit
to ROCm/composable_kernel
that referenced
this pull request
Sep 11, 2026
fix(ck,ck_tile): microscale pipeline B and BQ layout mismatch & bug in mxfp conversion with float scale (#11560) ## Motivation Jira ID: AICK-1884 two bugs found in tests: test_tile_gemm_quant_bquant_microscale_rcr_1d_64 test_tile_gemm_quant_bquant_microscale_crr_1d_64 test_tile_gemm_quant_bquant_microscale_rrr_1d_64 test_tile_gemm_quant_bquant_microscale_ccr_1d_64 test_tile_gemm_quant_bquant_microscale_rcr_1d_128 test_tile_gemm_quant_bquant_microscale_crr_1d_128 test_tile_gemm_quant_bquant_microscale_rrr_1d_128 test_tile_gemm_quant_bquant_microscale_ccr_1d_128 Two root causes: 1. B is 16-bit and B is ColumnMajor (8 failures) 2. B is PkFP4 (10 failures; every PkFP4 case fails) ## Technical Details fix for 1: BQuant tile distribution is hardcoded thread-raked while B on gfx1250 is warp-raked. Ask BQuant to check B's layout. fix for 2: this is more general, the gfx1250 mx fp4 conversion use packed 4 scale. The pipeline is using legacy way of passing float scale. Fix the fp4 conversion with float scale. (fp8 has the same issue and similar fix.) **Caveat for fix2:** The fix is for correctness. Using float scale for mx fp4 conversion is only for backward compatibility and is not an efficient approach. Fast approach/optimization is to use packed scale in the kernel/pipeline to better utilize the native conversion instruction.
Contributor
TheRock Submodule Bump ActivityNewest first
|
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
Jira ID: AICK-1884
two bugs found in tests:
test_tile_gemm_quant_bquant_microscale_rcr_1d_64
test_tile_gemm_quant_bquant_microscale_crr_1d_64
test_tile_gemm_quant_bquant_microscale_rrr_1d_64
test_tile_gemm_quant_bquant_microscale_ccr_1d_64
test_tile_gemm_quant_bquant_microscale_rcr_1d_128
test_tile_gemm_quant_bquant_microscale_crr_1d_128
test_tile_gemm_quant_bquant_microscale_rrr_1d_128
test_tile_gemm_quant_bquant_microscale_ccr_1d_128
Two root causes:
Technical Details
fix for 1: BQuant tile distribution is hardcoded thread-raked while B on gfx1250 is warp-raked. Ask BQuant to check B's layout.
fix for 2: this is more general, the gfx1250 mx fp4 conversion use packed 4 scale. The pipeline is using legacy way of passing float scale. Fix the fp4 conversion with float scale. (fp8 has the same issue and similar fix.)
Caveat for fix2: The fix is for correctness. Using float scale for mx fp4 conversion is only for backward compatibility and is not an efficient approach. Fast approach/optimization is to use packed scale in the kernel/pipeline to better utilize the native conversion instruction.
Test Plan
Test Result
Submission Checklist