fix(ck-tile): correct compv3 intrawave 8-warp schedule on gfx1250 (wave32) by ozturkosu · Pull Request #11175 · ROCm/rocm-libraries

ozturkosu · 2026-08-24T01:38:53Z

Summary

Fixes the compv3 intrawave GEMM pipeline producing numerically wrong results on gfx1250 (MI400, RDNA4 / WMMA / wave32) for 8-warp block arrangements (2x4x1 / 4x2x1). max relative error 0.14–0.87 vs an fp32 CPU reference (not tolerance). Fixes #11161.

Root cause

The NumWarps == 8 special-cased hot-loop/tail schedule in BaseGemmPipelineAgBgCrCompV3 (BlockHasHotloop / GetBlockLoopTailNum / TailHandler) was written for wave64 512-thread blocks (gfx9xx / MFMA).

On gfx1250 an 8-warp block is wave32 = 256 threads — the same thread count as a 4-warp wave64 block — so it must follow the standard (≤4-warp) schedule.

The wave64 8-warp path mis-selects the loop tail (TailNumber::Even for even loop counts), which makes the intrawave RUN path execute one extra block_gemm on a non-existent K-tile → garbage accumulation. (For num_loop == 3 the 8-warp BlockHasHotloop uses > 3 instead of > 2, skipping the hot loop entirely.)

Fix

Add an arch-guarded compile-time predicate and switch the three checks on it:

#if defined(__gfx1250__) || defined(__GFX12__)
    static constexpr bool Use8WarpSchedule = false;
#else
    static constexpr bool Use8WarpSchedule = (Problem::BlockGemmShape::NumWarps == 8);
#endif

On gfx942 / gfx950 (wave64): Use8WarpSchedule == (NumWarps == 8) — byte-for-byte identical, zero regression.
On gfx1250 / GFX12 (wave32): the 8-warp special case is disabled, so those blocks use the standard schedule.

All users of these functions are CK_TILE_DEVICE and TailHandler::scenarios[] compiles in the same device pass, so the __gfx1250__/__GFX12__ guard is host/device consistent (verified hipcc --offload-arch=gfx1250 predefines both __gfx1250__ and __GFX12__).

Validation (on-device, gfx1250 MI400, GPU output vs fp32 CPU reference)

config	before	after
compv3+intrawave 2x4x1 (8-warp)	MISMATCH 0.326	VERIFY rel=5.08e-04
compv3+intrawave 4x2x1 (8-warp)	wrong	VERIFY rel=5.08e-04
compv3+intrawave 1x8x1 / 8x1x1 (8-warp)	wrong	VERIFY rel=5.08e-04
compv3+intrawave 2x2x1 / 1x4x1 / 4x1x1 (4-warp control)	correct	VERIFY (no regression)

Validated across num_loop = 3, 4, 5, 10, 11, 12 (odd + even, hot-loop and no-hot-loop) and both fp16 (~5e-4) and bf16 (~3.8e-3, normal tolerance).

Notes

This is the real fix for #11161. An interim gate that simply drops these kernels from the gfx1250 sweep was applied to the gfx1250 enablement PRs (#10921 / #11048 / #10928); once this lands, that gate can be removed so the (now-correct) 8-warp compv3 + intrawave kernels re-enter the sweep.

…ve32) The compv3 intrawave GEMM pipeline produced wrong results on gfx1250 (MI400, WMMA/wave32) for 8-warp block arrangements (2x4x1 / 4x2x1) -- max_rel 0.14-0.87 vs an fp32 CPU reference. Root cause: the NumWarps==8 special-cased hot-loop/tail schedule in BaseGemmPipelineAgBgCrCompV3 (BlockHasHotloop / GetBlockLoopTailNum / TailHandler) was written for wave64 512-thread blocks (gfx9xx / MFMA). On gfx1250 an 8-warp block is wave32 = 256 threads -- the same thread count as a 4-warp wave64 block -- so it must follow the standard (<=4-warp) schedule. The wave64 8-warp path mis-selects the loop tail (TailNumber::Even for even loop counts), making the intrawave RUN path execute one extra block_gemm on a non-existent K-tile -> garbage accumulation. Fix: add an arch-guarded compile-time predicate Use8WarpSchedule (false on __gfx1250__/__GFX12__, unchanged NumWarps==8 on wave64) and switch the three checks on it. Byte-identical on gfx942/gfx950 (wave64) -- zero regression; on gfx1250/GFX12 the 8-warp blocks use the standard schedule. Fixes #11161.

therock-pr-bot · 2026-08-24T01:41:22Z

✅ All Policy Checks Passed

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	⚠️ Warning	Error: Source/code files changed without an accompanying unit test. Expected: add at least one test file named like `test_<name>.py` / `test_<name>.cpp` (or `<name>_test.`). Current:* code file(s) changed: `projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v3.hpp`; no test file found
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—

🎉 All policy checks passed!

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

therock-pr-bot · 2026-08-24T01:41:24Z

🎉 All checks passed! This PR is ready for review.

ozturkosu · 2026-08-24T05:31:39Z

Full-sweep on-device validation (gfx1250, MI400)

Re-ran the complete default_config sweep with this fix, interim gate off so the previously-failing compv3 + intrawave + 8-warp kernels are actually generated and tested. GPU output vs an fp32 CPU reference, 3 tile-divisible shapes (768³ / 1536³ / 3072³), (ROCm 10.1):

op	total (1,824 kernels × 3)	the `compv3+intrawave+8-warp` class (was 552 wrong)
`gemm_multi_abd`	5,472 / 5,472 pass, 0 fail	552 / 552 correct
`gemm_universal`	5,472 / 5,472 pass, 0 fail	552 / 552 correct

Before this fix those 552 kernels were wrong (max_rel 0.14–0.87); now they all verify (~5e-4), and the full sweep has zero failures. No regression on ≤4-warp compv3 or on compv4 / mem / interwave.

The gfx11 stage failed with 'no ROCm-capable device is detected' / 'invalid device ordinal' after the first streamk test -- the amdgpu driver dropped off the CI node partway through ctest, so every remaining GPU test aborted. This is infrastructure, not a code change: test_ck_tile_gemm_pipeline_compv3_wmma (the direct test of the touched file) passed, and on gfx11 the __gfx1250__/__GFX12__ guard added here never fires, so the emitted code is byte-identical to develop. Empty commit to re-run CI on a healthy node.

ozturkosu · 2026-08-24T08:31:25Z

I dug into the red CI here — the failure is infrastructure, not this change — and I re-validated the fix end-to-end on MI400.

CI: the gfx11 failure is a node GPU-drop, not this PR

The Build CK and run Tests on gfx11 stage died because the runner lost its GPU partway through ctest:

_amdgpu_device_initialize: amdgpu_query_gpu_info_init failed
HIP runtime error: no ROCm-capable device is detected   /   invalid device ordinal

The first streamk test ran ~10 s, hit invalid device ordinal, and then every remaining GPU test failed instantly (0–1 ms) with "no ROCm-capable device is detected" — the amdgpu driver fell off the runner mid-run. Three things confirm it isn't this change:

test_ck_tile_gemm_pipeline_compv3_wmma — the direct test of the file I touched — passed before the GPU vanished.
On gfx11 the __gfx1250__ / __GFX12__ guard I added never fires, so the emitted code is byte-identical to develop there.
The sibling gfx1250-enablement PRs (which don't touch this file) are green on gfx11.

I re-triggered CI (empty commit e573f68); the GitHub-Actions checks are re-running green and Math CI is re-queued.

On-device re-validation on MI400 (gfx1250, Helios, ROCm 10.1)

GPU output verified against an fp32 CPU reference. I ran with the interim gate off so the previously-broken compv3 + intrawave + 8-warp kernels are actually emitted and tested.

Before/after on the exact failing class (same kernel, same shapes):

kernel	with this fix	fix reverted (old wave64 8-warp path)
`compv3_intrawave_128x128x64_2x4x1_16x16x32`	VERIFY, rel ~5e-4	MISMATCH rel 0.33–0.86
`compv3_intrawave_128x128x64_4x2x1_16x16x32`	VERIFY, rel ~5e-4	MISMATCH rel 0.33–0.86

The mismatches land exactly on the even-tail loop counts — matching the root cause (the wave64 8-warp path mis-selects TailNumber and runs an extra block_gemm on a non-existent K-tile).

Full default_config sweep, 3 tile-divisible shapes (768³ / 1536³ / 3072³), gate off:

op	measurements	pass	fail	previously-wrong 8-warp class
`gemm_universal`	5,472	5,472	0	1,824 / 1,824 now correct
`gemm_multi_abd`	5,472	5,472	0	1,824 / 1,824 now correct

max_rel median ~4–5e-4 (bit-accurate), zero failures. The 8-warp kernels that returned garbage on develop (#11161) all verify now, and there's no regression on ≤4-warp compv3 or on compv4 / mem / interwave.

ozturkosu · 2026-08-24T08:32:31Z

Full `default_config` sweep on MI400 (gfx1250) — clean

To make sure this fix holds across the whole config space (not just the handful of shapes in the repro), I ran the complete default_config sweep on-device on Helios (gfx1250, wave32, ROCm 10.1). Each kernel is compiled as a runtime .so for gfx1250, launched on the GPU, and its output verified against an fp32 NumPy CPU reference. I ran with the interim gate off, so the compv3 + intrawave + 8-warp kernels that were wrong on develop (#11161) are actually generated and checked.

3 tile-divisible shapes per kernel: 768³ / 1536³ / 3072³ (divisible by every tile size, so pad=false kernels never hit the non-divisibility path).

op	measurements	pass	fail	previously-broken 8-warp class
`gemm_universal`	5,472	5,472	0	1,824 / 1,824 now correct
`gemm_multi_abd`	5,472	5,472	0	1,824 / 1,824 now correct

max_rel: universal median 5.08e-4 (max 5.11e-4), multi-ABD median 4.03e-4 (max 4.67e-4) — bit-accurate for fp16.

Every 2x4x1 / 4x2x1 (8-warp) kernel that returned garbage on develop now verifies, and there's no regression anywhere else (≤4-warp compv3, compv4, mem, interwave all stay clean). Zero failures across 10,944 measurements.

…gfx1250-fix

…n gfx1250 (#11161)" This reverts commit 837dca8. The compv3+intrawave 8-warp (2x4x1 / 4x2x1) kernels are re-enabled on gfx1250 so the #11161 wrong-results defect is reproducible in this PR. Real pipeline fix is tracked in #11175.

ozturkosu · 2026-08-25T15:51:01Z

Independent on-hardware verification (gfx1250 MI400, native example)

Verified this fix on a real gfx1250 box (heliosr-1b114-a07-1, ROCm 10.1 / clang23) via the plain ck_tile example tile_example_gemm_universal — no Dispatcher / Tile-Engine harness, just the base GEMM binary against an fp32 CPU reference. Fresh ROCm/composable_kernel clone; only include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v3.hpp patched with this PR's change; identical binary rebuilt and rerun.

The gfx1250 default config GemmConfigComputeV3_WMMA is compv3 + intrawave + 4x2x1 (8 warps, 256 threads), so it exercises exactly the fixed path.

Before the patch (768³, fp16, rcr):

shape: tile_gemm_shape_128x128x32x8_4x2x1_16x16x32   blocks: {256,1,1}
max err: 63.5, 89.11183% wrong values
The CPU verification result is: fail

After the patch (same kernel, same shape, only the header changed):

shape: tile_gemm_shape_128x128x32x8_4x2x1_16x16x32   blocks: {256,1,1}
The CPU verification result is: correct

The 4-warp 2x2x1 control passes both before and after (no regression on the standard path). The arch-guarded Use8WarpSchedule predicate does exactly what the summary claims: on gfx1250 the wave64 8-warp schedule is disabled and the 8-warp block correctly follows the standard ≤4-warp schedule.

LGTM — the fix resolves #11161 on-device.

…gfx1250-fix

jefyang1 · 2026-08-26T16:38:33Z

I reproduced the issue with tile_example_gemm_universal and this PR fixed it.

fix(ck-tile): correct compv3 intrawave 8-warp schedule on gfx1250 (wave32) (#11175) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ## Summary Fixes the `compv3` intrawave GEMM pipeline producing **numerically wrong results** on **gfx1250 (MI400, RDNA4 / WMMA / wave32)** for **8-warp block arrangements** (`2x4x1` / `4x2x1`). max relative error 0.14–0.87 vs an fp32 CPU reference (not tolerance). Fixes #11161. ## Root cause The `NumWarps == 8` special-cased hot-loop/tail schedule in `BaseGemmPipelineAgBgCrCompV3` (`BlockHasHotloop` / `GetBlockLoopTailNum` / `TailHandler`) was written for **wave64 512-thread** blocks (gfx9xx / MFMA). On gfx1250 an 8-warp block is **wave32 = 256 threads** — the same thread count as a 4-warp wave64 block — so it must follow the standard (≤4-warp) schedule. The wave64 8-warp path mis-selects the loop tail (`TailNumber::Even` for even loop counts), which makes the intrawave RUN path execute **one extra `block_gemm` on a non-existent K-tile** → garbage accumulation. (For `num_loop == 3` the 8-warp `BlockHasHotloop` uses `> 3` instead of `> 2`, skipping the hot loop entirely.) ## Fix Add an arch-guarded compile-time predicate and switch the three checks on it: ```cpp #if defined(__gfx1250__) || defined(__GFX12__) static constexpr bool Use8WarpSchedule = false; #else static constexpr bool Use8WarpSchedule = (Problem::BlockGemmShape::NumWarps == 8); #endif ``` - On **gfx942 / gfx950** (wave64): `Use8WarpSchedule == (NumWarps == 8)` — **byte-for-byte identical, zero regression**. - On **gfx1250 / GFX12** (wave32): the 8-warp special case is disabled, so those blocks use the standard schedule. All users of these functions are `CK_TILE_DEVICE` and `TailHandler::scenarios[]` compiles in the same device pass, so the `__gfx1250__`/`__GFX12__` guard is host/device consistent (verified `hipcc --offload-arch=gfx1250` predefines both `__gfx1250__` and `__GFX12__`). ## Validation (on-device, gfx1250 MI400, GPU output vs fp32 CPU reference) | config | before | after | |--------|--------|-------| | compv3+intrawave **2x4x1** (8-warp) | MISMATCH 0.326 | **VERIFY rel=5.08e-04** | | compv3+intrawave **4x2x1** (8-warp) | wrong | **VERIFY rel=5.08e-04** | | compv3+intrawave **1x8x1 / 8x1x1** (8-warp) | wrong | **VERIFY rel=5.08e-04** | | compv3+intrawave **2x2x1 / 1x4x1 / 4x1x1** (4-warp control) | correct | **VERIFY** (no regression) | Validated across `num_loop` = 3, 4, 5, 10, 11, 12 (odd + even, hot-loop and no-hot-loop) and both fp16 (~5e-4) and bf16 (~3.8e-3, normal tolerance). ## Notes This is the **real** fix for #11161. An interim gate that simply drops these kernels from the gfx1250 sweep was applied to the gfx1250 enablement PRs (#10921 / #11048 / #10928); once this lands, that gate can be removed so the (now-correct) 8-warp `compv3 + intrawave` kernels re-enter the sweep.

…n gfx1250 (#11161)" This reverts commit 837dca8. The compv3+intrawave 8-warp (2x4x1 / 4x2x1) kernels are re-enabled on gfx1250 so the #11161 wrong-results defect is reproducible in this PR. Real pipeline fix is tracked in #11175.

JIRA ID : AICK-2101 JIRA ID : AICK-2133 ## Summary This PR enables the **multi-ABD** (`a2_b2_d2`) GEMM TE→dispatcher bridge (#9305) on **gfx1250 / MI400**. It was split out of #10921 so the other three ops (`gemm_universal`, `grouped_gemm`, `gemm_multi_d`) weren't held back. Comparison of this PR's original file set vs `develop` (branch head `56e7bc3`): **So the substantive multi-ABD gfx1250 enablement is already on `develop`; the remaining unique delta in this PR is a single config file (`example_problems.json`).** ## Motivation (why the enablement was needed) gfx1250 runs **WMMA**, not MFMA: the valid fp16/bf16 warp tile is **16x16x32**, whereas the CDNA MFMA **32x32x16** tile in `default_ci_config.json` instantiates zero runnable kernels there. ## Supported surface (gfx1250) The `a2_b2_d2` variant is intentionally narrow — the codegen, ctypes lib, and runner are all **fp16-only** and **rcrr-only** (`expand_sweep` rejects other dtypes/layouts). So **fp16 × rcrr is the complete supported matrix**. ## Validation (on-device, MI400 / gfx1250) Old-TE has no gfx1250 back end, so this is **correctness-only** (GPU output vs an fp32 CPU reference), not a performance baseline. Full `default_config` sweep of the fp16/rcrr surface, 3 tile-divisible shapes (768³ / 1536³ / 3072³): | op | supported matrix | kernels | result rows | fail | |----|------------------|--------:|------------:|-----:| | `gemm_multi_abd` | fp16 × rcrr (only) | 1,640 | 4,920 | **0** | Every kernel in the supported surface is correct on gfx1250 (rel err ≈ 4e-4). ## Note on #11161 (compv3 + intrawave + 8-warp) The full sweep initially surfaced 552 failures, all in the `compv3 + intrawave + 8-warp (2x4x1 / 4x2x1)` class — the **#11161** defect (a wave64 assumption in the `comp_v3` pipeline, wrong on gfx1250 wave32). An interim gate excludes that class so the supported surface reports clean; the real fix, **#11175** (`Use8WarpSchedule`), is already on `develop`. ## Notes / dependencies - Built on #10921's shared gfx1250 foundation (now on `develop` via the sibling series). - Supersedes the earlier closed attempt **#10917**. - Sibling gfx1250 bridge PRs: #10921 (universal / grouped / multi-D), #10927 (grouped quant), #10928 (batched GEMM + contraction), #11043 (rowcol / tensor quant), #10934 (contraction multi-ABD). --------- Co-authored-by: Muhammed Emin Ozturk <3836908+ozturkosu@users.noreply.github.com>

Narrowing WARP_SUPPORTED_COMBINATIONS["gfx1250"] from nine warp maps to four was a regression. That table is shared by seven instance builders, and one of the maps removed, 4x2x1, is plain GEMM's default gfx1250 WMMA configuration -- issue #11161 reported it returning wrong results, PR #11175 fixed it, and the fix was verified correct on device. Measured on the predicate: validate_warp_configuration(4, 2, 1, "gfx1250") develop ACCEPT after my change REJECT <- deleted another operator's working default The table is now byte-identical to develop again. The restriction itself is still needed, but only here: these two grouped quant bridges do not emit a loadable kernel entry for >4-warp blocks on gfx1250, so those kernels abort at launch with "cannot find symbol" -- 3,220 rows in a GPU sweep. That is a property of this bridge's code generation, not of the hardware, which is precisely why it does not belong in a hardware capability table. It now lives in validate_gemm_rowcol_tensor_quant(), alongside the warp-tile rule that was already scoped to the four quant validators. warp_k > 1 keeps its own rule rather than riding on the warps-per-block cap: [1,2,2] has a product of 4 and would otherwise pass. Verified: plain GEMM's 4x2x1 default is accepted again; this operator still rejects 4x2x1, 2x4x1, 4x4x1 and 1x2x2; 32x32x32 is still rejected; the legal 16x16x128 tile is still accepted; and across gfx90a, gfx942 and gfx950 there are zero decision changes.

…and 8-wave compv3, and offered compv5 (#13609) JIRA ID : AICK-2558 ## Why On gfx1250 the GEMM sweep's gates disagreed with what the pipelines support, in both directions: - **TDM kernels were refused problems they run.** `comp_tdm` and `comp_tdm_v2` declare `skipCheckValidLaunchParams`, so `UniversalGemmKernel::IsSupportedArgument` accepts every problem for them: the Tensor Data Mover clips each box to the tensor's real extent and moves A, B and C at width 1. The dispatcher modelled their A/B global vector widths as per-thread load widths, and both the Python pairing and `GeneratedTileKernelInstance::supports()` applied the unpadded tile-divisibility rule. An unpadded TDM kernel, the only kind codegen emits, was therefore refused for every problem whose contiguous extents are not tile and vector multiples. - **The padded-TDM rejection reason was wrong.** It said kPad's right-pad transforms inflate the extents TDM clips against. TDM clips every box to the pre-pad extent (`tdm_real_lengths`), so kPad is a no-op for it. Padded TDM stays rejected, for the real reason: each padded kernel would duplicate the unpadded one, which already serves every shape. - **compv3 8-wave kernels were hidden.** `expand_sweep` skipped compv3 intrawave kernels with 8 waves (2x4x1, 4x2x1) on gfx1250, because the pipeline's 8-warp schedule, written for wave64, gave wrong results on wave32 (#11161). The schedule is now keyed on `CK_TILE_USE_WMMA`, so wave32 builds take the standard schedule (#11175), and a wave32 device build without `CK_TILE_USE_WMMA` fails a `static_assert` instead. The gate only hid correct kernels. - **compv5 kernels were offered although none compiles.** `GemmPipelineAgBgCrCompV5` sizes its C LDS view with `BlockGemm::Traits::KPack`, but its block GEMM, `BlockGemmARegBRegCRegV1`, defines only `KPackA` and `KPackB`. `gfx1250_pipeline_reject_reason` judged only the gfx1250-only pipelines and the tdm epilogue, and `gemm_utils._gfx1250_pipeline_supported` returned early for every other pipeline before consulting it. The block GEMM is the same on other targets, where compv5 does not compile either; this PR rejects it only on gfx1250 and leaves other targets unchanged. ## What changes - `dispatcher/codegen/codegen_common.py`: the tdm epilogue's A/B/C widths are 1; new `gemm_pipeline_skips_launch_checks(pipeline)`, the Python mirror of `skipCheckValidLaunchParams`; corrected `TDM_PAD_REJECT_REASON`; `gfx1250_pipeline_reject_reason` rejects compv5 for gfx1250 (`GFX1250_COMPV5_REJECT_REASON`). - `tile_engine/ops/gemm/gemm_vector_fallback.py`: `_tile_fits` skips the tile rule for those pipelines. - `dispatcher/include/ck_tile/dispatcher/backends/generated_tile_backend.hpp`: new trait `pipeline_skips_launch_checks`; `supports()` skips the tile check for pipelines that set `skipCheckValidLaunchParams`. - `dispatcher/python/gemm_utils.py`: the compv3 8-wave gate is removed, and `_gfx1250_pipeline_supported` passes every pipeline through `gfx1250_pipeline_reject_reason`, so codegen, `arch_filter` and the sweep agree. - `tile_engine/ops/gemm/gemm_validation_utils.py`: its copy of the padded-TDM reason. Every other pipeline keeps both checks, and targets other than gfx1250 are unchanged.

github-actions Bot added the project: composablekernel label Aug 24, 2026

ozturkosu self-assigned this Aug 24, 2026

ozturkosu requested a lite review from Copilot August 24, 2026 01:57

This comment was marked as off-topic.

Sign in to view

assistant-librarian Bot added the organization: ROCm label Aug 24, 2026

Merge branch 'develop' into users/muozturk/ck/compv3-intrawave-8warp-…

1455405

…gfx1250-fix

ozturkosu marked this pull request as ready for review August 24, 2026 08:33

ozturkosu requested a review from a team as a code owner August 24, 2026 08:33

ozturkosu added 2 commits August 24, 2026 01:44

Merge branch 'develop' into users/muozturk/ck/compv3-intrawave-8warp-…

156b899

…gfx1250-fix

Merge branch 'develop' into users/muozturk/ck/compv3-intrawave-8warp-…

648aa24

…gfx1250-fix

ozturkosu requested a review from jefyang1 August 24, 2026 16:25

Merge branch 'develop' into users/muozturk/ck/compv3-intrawave-8warp-…

ae1c413

…gfx1250-fix

ozturkosu mentioned this pull request Aug 25, 2026

compv3 intrawave GEMM pipeline returns wrong results on gfx1250 (MI400) for 8-warp blocks (2x4x1/4x2x1) #11161

Closed

ozturkosu requested review from DDEle and kensclin August 25, 2026 15:48

Merge branch 'develop' into users/muozturk/ck/compv3-intrawave-8warp-…

2f8c6e6

…gfx1250-fix

illsilin approved these changes Aug 26, 2026

View reviewed changes

illsilin enabled auto-merge (squash) August 26, 2026 18:59

kensclin reviewed Aug 26, 2026

View reviewed changes

Comment thread projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v3.hpp Resolved

illsilin merged commit c0cb6df into develop Aug 27, 2026
131 of 146 checks passed

illsilin deleted the users/muozturk/ck/compv3-intrawave-8warp-gfx1250-fix branch August 27, 2026 16:34

This was referenced Aug 31, 2026

feat(ck-tile): enable grouped rowcolquant/tensorquant GEMM bridges on gfx1250 #11043

Merged

feat(ck-tile): enable grouped AQuant/ABQuant/BQuant GEMM bridges on gfx1250 (MI400) #10927

Merged

ozturkosu mentioned this pull request Sep 10, 2026

feat(ck-tile): grouped rowcolquant/tensorquant GEMM bridges on gfx1250 [Experimental- Dont Review] #11904

Closed

9 tasks done

andriy-ca mentioned this pull request Oct 14, 2026

fix(dispatcher): gfx1250 GEMM sweep gates refused valid TDM problems and 8-wave compv3, and offered compv5 #13609

Merged

Navigation Menu

Repository navigation

fix(ck-tile): correct compv3 intrawave 8-warp schedule on gfx1250 (wave32) - #11175

fix(ck-tile): correct compv3 intrawave 8-warp schedule on gfx1250 (wave32)#11175
illsilin merged 7 commits into
developfrom
users/muozturk/ck/compv3-intrawave-8warp-gfx1250-fix

ozturkosu commented Aug 24, 2026 •

edited by afagaj

Loading

Uh oh!

therock-pr-bot Bot commented Aug 24, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Aug 24, 2026

Uh oh!

This comment was marked as off-topic.

Uh oh!

ozturkosu commented Aug 24, 2026 •

edited

Loading

Uh oh!

ozturkosu commented Aug 24, 2026

Uh oh!

ozturkosu commented Aug 24, 2026

Uh oh!

ozturkosu commented Aug 25, 2026

Uh oh!

jefyang1 commented Aug 26, 2026

Uh oh!

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

Conversation

ozturkosu commented Aug 24, 2026 • edited by afagaj Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Root cause

Fix

Validation (on-device, gfx1250 MI400, GPU output vs fp32 CPU reference)

Notes

Related

Uh oh!

therock-pr-bot Bot commented Aug 24, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

✅ All Policy Checks Passed

Uh oh!

therock-pr-bot Bot commented Aug 24, 2026

Uh oh!

This comment was marked as off-topic.

Uh oh!

ozturkosu commented Aug 24, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Full-sweep on-device validation (gfx1250, MI400)

Uh oh!

ozturkosu commented Aug 24, 2026

CI: the gfx11 failure is a node GPU-drop, not this PR

On-device re-validation on MI400 (gfx1250, Helios, ROCm 10.1)

Uh oh!

ozturkosu commented Aug 24, 2026

Full default_config sweep on MI400 (gfx1250) — clean

Uh oh!

ozturkosu commented Aug 25, 2026

Independent on-hardware verification (gfx1250 MI400, native example)

Uh oh!

jefyang1 commented Aug 26, 2026

Uh oh!

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

ozturkosu commented Aug 24, 2026 •

edited by afagaj

Loading

therock-pr-bot Bot commented Aug 24, 2026 •

edited

Loading

ozturkosu commented Aug 24, 2026 •

edited

Loading

Full `default_config` sweep on MI400 (gfx1250) — clean