CUDA: refactor MMQ kernel configuration - #24127
Conversation
f28de39 to
3c84aa4
Compare
Performance check NVIDIA
Performance check AMD
On my NVIDIA hardware I am seeing no changes to performance beyond statistical fluctuations. On my AMD hardware however the performance is changing and quite frankly I don't understand why, the only thing that should have really changed is that there are now some |
|
@JohannesGaessler There are some failing tests on DGX Spark with MXFP4 and NVFP: |
|
Should be fixed now, the Blackwell config was wrong but in such a way that did not consistently result in incorrect outputs. |
|
Im still way behind on things since moving, but hope to be able to take a real look at this this coming weekend. |
IMbackK
left a comment
There was a problem hiding this comment.
Overall this is a huge improvement and i am exited at the possibility this opens up wrt per arch tuneing.
I ran test-backend-ops perf and a couple of models by this and dident find anything unusal performance wise on cdna1 and rdna3 but do think its pretty easy for this to have some regressions on some arch by accident.
I would be good if someone with a rdna1 device could check if the change to the compiler register allocation has a detrimental effect for some reason.
| } | ||
| } | ||
| return ret; |
There was a problem hiding this comment.
Is there a valid usecase for this to not assert at this point?
There was a problem hiding this comment.
No, I think raising an error here would make sense.
| 1, 1, 0, 0, 0, | ||
| 1, 1, 0, 0, 0, | ||
| use_stream_k, src1_ncols}; | ||
| false, src1_ncols}; |
There was a problem hiding this comment.
When stream_k is used changed, on pupose?
There was a problem hiding this comment.
The semantic meaning of the boolean is different. Previously it was whether or not to use stream-k, now it is whether or not it is safe to use stream-k. However, after #24216 it should now always be safe to use stream-k and the boolean can be removed.
59a1e08 to
af72f0d
Compare
|
@ggml-cuda I've rebased on top of the latest master commit and removed legacy code related to stream-k configuration. Can I please get a second review? |
|
Meant to ping @ggml-org/ggml-cuda . |
…#24127) Upstream PR ggml-org#24127 replaced the per-arch MMQ macro/ternary helpers with per-arch config tables keyed by (type, J, fallback), renamed mmq_x/mmq_y to J/I, and made __launch_bounds__ mandatory. It deleted all six functions the old gfx1151 patch edited, so that patch was dropped on this rebase rather than replayed; this re-authors it against the new architecture. amd_wmma_available() still covers all of RDNA3, so without a dispatch branch gfx1151 silently inherits the rdna4 table (nthreads=256, occupancy=2, I=128 for every type) - the same tile shape this fork has been beating since April. - mmq-config-rdna3_5.cuh: generated from mmq-config-rdna4.cuh with nthreads 256 -> 128 and I 128 -> 64. sram_layout, K_vram and stream_k are carried over unchanged for all 21 types; only tile shape is retuned. Adds a J=48 fallback=true row per type: rdna4 thins fallback specializations to powers of 2, so without it the MoE cap below would silently settle on 32 whenever ne01 % 128 != 0. - mmq.cuh: host and device dispatch, both ahead of the AMD_WMMA branch. - mmq.cuh: MoE J cap in mul_mat_q_switch_J. args.expert_bounds still reaches the J-selection loop, so the old dense/MoE split survives as a bound on the search rather than an mmq_x_max override. I and nthreads are not independent: the MMA write-back needs I == nwarps*16 on wave32 (rdna4 is 128 == 8*16, this table is 64 == 4*16). The old patch's mmq_y=64 / nwarps=4 pair was obeying the same constraint. occupancy=2 is unvalidated - the knob did not exist when this tuning was last measured. It gives 2x rdna4's per-thread register budget but only 256 threads resident per CU. occupancy 2 vs 4 is the first sweep if the bench comes in flat. mmq-table-check.cpp verifies the table is well-formed and dispatches as intended without ROCm: it reuses the real CASE macro so the upstream static_asserts fire, checks every rdna4-reachable entry has an rdna3_5 counterpart, checks I == nwarps*16, and checks no input drives J_best to 0. That proves the table is correct, not that it is fast. NOT BENCHED. Never run on gfx1151. Gated on the two-run A/B (port off, then port on) in strix-halo/mmq-rdna3_5-config-table.md - this rebase also carries 234 upstream commits and the ROCm 7.14.0 switch, so the re-port is not attributable without it. Assisted-by: Claude Fable 5
…24546) Upstream PR ggml-org#24546 (ravel7524, open) sizes routed-MoE MMQ N-tiles from the typical expert width instead of the worst case. That is the same problem the static J_max=48 cap from Findings ggml-org#8/ggml-org#9 solves, approached from the other end: the cap clamps the search ceiling, the picker fixes the search objective. This commit swaps the cap for the picker so the two can be A/B'd on gfx1151. The PR does not apply as written - it patches mmq_x/mmq_x_max/ mmq_get_granularity_host in mul_mat_q_case, all deleted by PR ggml-org#24127, the same refactor that forced the Finding ggml-org#5 -> ggml-org#9 re-port. Re-authored onto mul_mat_q_switch_J's config-table search and scoped to RDNA3.5: gfx1151 is the only arch we can measure, and narrowing the predicate keeps the A/B single-variable. Upstream's predicate already covers RDNA3.5 via GGML_CUDA_CC_IS_RDNA3, and the author confirmed on 2026-07-04 that Strix Halo should benefit, but nobody has ever tested it there. On the routed-MoE path ncols_max is ne12 - the worst case of one expert receiving every token - so the search never reaches ntiles==1 and simply runs to whatever ceiling it is given. The picker instead tiles from ncols_dst/nchannels_x, the width a typical expert covers. launch_mul_mat_q still builds the grid from ncols_max, so worst-case coverage is unchanged. - mmq.cuh: replace the RDNA3.5 MoE J_max=48 clamp with the typical-width picker; J_max is now a uniform 128 and the search terminates on its own. - mmq-table-check.cpp: pick_J mirrors the selection loop, so update it in lockstep or the check silently becomes a rubber stamp. Adds the routed-MoE arg shape from mmq.cu and a crossover case. For Qwen 3.6 35B-A3B (256 experts, 8 active) at the production ub=2048 the typical width is 2048*8/256 = 64, so this is a 48 -> 64 change at our operating point - narrow enough that it may land in noise. Verified host-side without ROCm: MoE J=64, dense J=128, no aborts across all 21 types. Above ub=4096 the typical width reaches 128, the picker disengages and the search runs to J=128 where the cap held 48. Findings ggml-org#5/ggml-org#8 only ever measured the typical=64 regime, so that is untested on this chip; production ub=2048 cannot reach it. See strix-halo/mmq-moe-ncols-picker.md for the bench plan and the keep/revert criteria. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Upstream landed native Q2_0 (type 42) with full Vulkan support (ggml-org#25430) and refactored the CUDA MMQ kernel configuration (ggml-org#24127). Reconciled feat/q2_0-gpu against both: - Vulkan Q2_0: took upstream's implementation (identical QK2_0=64 block format; our branch had cherry-picked the same PR). Dropped our redundant shader/pipeline copies. - CUDA Q2_0: dropped entirely (MMQ port collided with the ggml-org#24127 refactor and the half-kept dequant path would have routed small-batch matmuls to a mmvq kernel that no longer exists). Q2_0 now runs on Vulkan (BC-250) only, matching upstream; the 3090 has no Q2_0 path, same as mainline. Re-porting the CUDA kernel is a deliberate follow-up if a House config ever needs Q2_0 on CUDA. Preserved fork features through the merge: mtp_dev pin, hy_v3 (.bias suffix + per-layer n_ff_exp_impl), nemotron-h MTP filter, ssm d_state 96 (CUDA+Vulkan), vk-uma-mem clamp, disk-cache, rpc-cache(-preflight), decode-timing instrumentation. Assisted-by: Claude Fable 5 <noreply@anthropic.com>
…, OpenCL Q6_K/Adreno, CORS, checkpoint min-step, prompt cache refactor, MoE expert API stays) Upstream highlights since 6be7459: - model: DFlash speculative with KV rotation (ggml-org#25823) - model: Hy3 (hy_v3) with MTP speculative decoding (ggml-org#25395) - model: DeepseekV4 with fused hyper-connection ops (ggml-org#25585) - ggml: 0.17.0, LIGHTNING_INDEXER, out_prod, f16 set_rows - vulkan: Q2_0 support, native e2m1/e4m3 conversions, transfer-queue race fix - CUDA: MMQ kernel config refactor (ggml-org#24127), tighter MMQ src1 buffer for fp4 (ggml-org#25613), CUDA graphs on Volta/Turing, MoE gate/up dedup, CUDA Virtual Devices - ROCm: hexagon L2 cache rework, native fp4, FP16/INT8 coopmat on AMD - SYCL: Battlemage flash attention via oneDNN XMX, XIELU op, fp16 conv2d_dw - OpenCL: Q6_K GEMM/GEMV fix, ragged-tile MoE prefill FP16, Adreno vectorized LD/ST, A7x optimizations, ABS op - kleidiai: SME2 f32 kernel, SME vs SME2 dispatch - server: refactor prompt cache state ownership (ggml-org#25649) - new server_prompt_cache_state separates prompt metadata from KV data - server: evict checkpoints within min-step (ggml-org#25472) - server: text-only slot save/restore with mtmd (ggml-org#25076) - server: --cors-* options (ggml-org#25655) - server: refactored server_stream (ggml-org#25541) - server: respect min-step when splitting prompt batches (ggml-org#25420) - server: move chat-template thinking probe inside init try/catch (ggml-org#24093) - common: auto-download dflash/eagle3 HF sidecars (ggml-org#25811), drop --stdin mutual-exclusion, align tokenize usage - conversion: BitNetForCausalLM, dflash tokenizer fix, split MTP export for HY V3 - llama-quant: exclude i32 ffn_gate_tid2eid routing table, allow manual tensor types with --pure - llama-batch: fix allowed decreasing pos in a seq (ggml-org#25449), n_keep_tail in split_equal for recurrent - llama: refactor fused ops (ggml-org#24646), TP fix for Phi3/Bert/Plamo2/3/ChatGLM - ui: agentic content UX, reasoning effort on mobile add sheet, MCP panel fixes, thinking menu fix - vendor: BoringSSL 0.20250713.0 - tests: actually exercise test-recurrent-state-rollback, ds_v4_hc sentinel init, export-graph-ops graceful exit CachyLLama preservation work (conflict resolution): 1. tools/server/server-task.h: Accept upstream's server_prompt refactor (no data member, clear() method). Move our t_last_used field from server_prompt to server_prompt_cache_state (where it now lives after the refactor). server_prompt_cache_state already has the size() method, so our old size() on server_prompt is no longer needed. 2. tools/server/server-context.cpp (create_checkpoint): Take upstream's min-step eviction pre-filter as the FIRST pass, then keep our existing highest-pos_min eviction as the capacity overflow fallback. These are complementary: min-step removes redundant checkpoints from the same task; highest-pos_min keeps the rec-window-friendly checkpoints when at cap. 3. tools/server/server-context.cpp (handle_completions_impl): Keep our std::vector<server_task> tasks batching for multi-prompt requests and per-user concurrency check, AND take upstream's res->set_req(&req) for spipe ownership transfer. 4. tools/server/server-task.cpp: Fix references to entry.tokens -> entry.prompt.tokens, entry.checkpoints -> entry.prompt.checkpoints, entry.n_tokens() -> entry.prompt.n_tokens(). Update find_eviction_candidate return type from list<server_prompt>::iterator to list<server_prompt_cache_state>::iterator. 5. ggml/src/ggml-cuda/mmq.cuh + new mmq-config-rdna3_5.cuh: Upstream's massive MMQ refactor moved per-architecture config into separate files but did NOT add RDNA3.5 (gfx1150/1/2/3, Strix Halo). Create mmq-config-rdna3_5.cuh (231 CASE entries) derived from rdna2 with nthreads=128 (4 warps) and I=48 (smaller X tile) matching our original Strix Halo tuning. Wire into both host and device dispatch paths before the RDNA4 / RDNA2 fallback. 6. README.md and AGENTS.md: Keep CachyLLama-specific links and project context where upstream added parallel content. Verified: - cmake --build builds clean (Release, CPU-only) - llama-server starts, --help shows all CachyLLama flags preserved: --cache-ssd-hot-ram, --cache-ssd-warm-ram, --cache-ssd-system-prompts, --cache-ssd-system-max-days, --cache-ssd-no-fsync, --cache-ssd-max-conversations, --max-concurrent-per-user - /expert-stats and /expert-tracking endpoints preserved - 55/58 tests pass; 3 failures unrelated to merge: - test-tokenizers-ggml-vocabs: missing model downloads - test-jinja-py: missing jinja2 Python module - test-quant-type-selection: snapshot mismatch on upstream's new MXFP4_MOE heuristic Custom CachyLLama files untouched (no upstream conflicts): - common/kv-ssd-cache.{cpp,h}, common/kv-ssd-posix.h, common/kv-ssd-system-cache.{cpp,h} - common/kv_page_manager.{cpp,h} - tools/server/server-context-page-manager.{cpp,h} - tools/server/server-context-ssd-cache.{cpp,h} - test_kv_page_manager.cpp, tests/test-ssd-cache-caps.cpp - STRIX_HALO_NOTES.md, docs/development/user-isolation-design.md - .github/workflows/build-cpu.yml, build-cuda-windows.yml, build-vulkan.yml
Upstream's MMQ refactor (ggml-org#24127) added `static_assert((I_) % 32 == 0, "bad I")` to the CASE macro. The rdna3_5 config introduced in cf4f42e used I=48, which fails the assertion (48 % 32 == 16 != 0) and breaks HIP builds for gfx1150/1/2/3 (Strix Halo). 48 was never a valid choice even without the static_assert. With nwarps=4 and the AMD MMA write_back path, each warp covers rows_per_warp=16 (hardcoded for AMD_MFMA/AMD_WMMA) and tile_C::I=16, so total rows = nwarps*16 = 64. I=48 would have left warp 3 writing rows 48-63 out of bounds. 64 is the smallest valid I for nwarps=4, satisfies I % 32 == 0, keeps the "smaller X tile" intent vs rdna2/rdna4's I=128, and matches the project notes in STRIX_HALO_NOTES.md ("use 64 as the safe final value"). All 232 CASE entries + the default return at the bottom updated. Fixes fewtarius/llama-ai#7
* CUDA: refactor MMQ kernel configuration * fix Blackwell config * remove legacy code
The Strix Halo row claimed "Merged upstream" but `mmq-config-rdna3_5.cuh` does not exist on `upstream/master` - gaetan-puleo's tables never landed in ggml-org/llama.cpp. The follow-up `71d1e8f2f` (I=48 -> I=64) is now linked into that entry instead of being orphaned. Added the missing Vulkan APU `nodes_per_submit` auto-lower row (`1c19480da`) so the table reflects what CachyLLama actually carries. Reworded the "drop on merge" guidance: CachyLLama focuses downstream, so in-tree copies of unupstreamed third-party patches stay until they land upstream cleanly and we rebase the local additions on top. Flagged upstream ggml-org#24127 as a watch item for future Strix Halo bumps - its `static_assert((I_) % 32 == 0)` is why `71d1e8f2f` exists.
…#24127) Upstream PR ggml-org#24127 replaced the per-arch MMQ macro/ternary helpers with per-arch config tables keyed by (type, J, fallback), renamed mmq_x/mmq_y to J/I, and made __launch_bounds__ mandatory. It deleted all six functions the old gfx1151 patch edited, so that patch was dropped on this rebase rather than replayed; this re-authors it against the new architecture. amd_wmma_available() still covers all of RDNA3, so without a dispatch branch gfx1151 silently inherits the rdna4 table (nthreads=256, occupancy=2, I=128 for every type) - the same tile shape this fork has been beating since April. - mmq-config-rdna3_5.cuh: generated from mmq-config-rdna4.cuh with nthreads 256 -> 128 and I 128 -> 64. sram_layout, K_vram and stream_k are carried over unchanged for all 21 types; only tile shape is retuned. Adds a J=48 fallback=true row per type: rdna4 thins fallback specializations to powers of 2, so without it the MoE cap below would silently settle on 32 whenever ne01 % 128 != 0. - mmq.cuh: host and device dispatch, both ahead of the AMD_WMMA branch. - mmq.cuh: MoE J cap in mul_mat_q_switch_J. args.expert_bounds still reaches the J-selection loop, so the old dense/MoE split survives as a bound on the search rather than an mmq_x_max override. I and nthreads are not independent: the MMA write-back needs I == nwarps*16 on wave32 (rdna4 is 128 == 8*16, this table is 64 == 4*16). The old patch's mmq_y=64 / nwarps=4 pair was obeying the same constraint. occupancy=2 is unvalidated - the knob did not exist when this tuning was last measured. It gives 2x rdna4's per-thread register budget but only 256 threads resident per CU. occupancy 2 vs 4 is the first sweep if the bench comes in flat. mmq-table-check.cpp verifies the table is well-formed and dispatches as intended without ROCm: it reuses the real CASE macro so the upstream static_asserts fire, checks every rdna4-reachable entry has an rdna3_5 counterpart, checks I == nwarps*16, and checks no input drives J_best to 0. That proves the table is correct, not that it is fast. NOT BENCHED. Never run on gfx1151. Gated on the two-run A/B (port off, then port on) in strix-halo/mmq-rdna3_5-config-table.md - this rebase also carries 234 upstream commits and the ROCm 7.14.0 switch, so the re-port is not attributable without it. Assisted-by: Claude Fable 5
…24546) Upstream PR ggml-org#24546 (ravel7524, open) sizes routed-MoE MMQ N-tiles from the typical expert width instead of the worst case. That is the same problem the static J_max=48 cap from Findings ggml-org#8/ggml-org#9 solves, approached from the other end: the cap clamps the search ceiling, the picker fixes the search objective. This commit swaps the cap for the picker so the two can be A/B'd on gfx1151. The PR does not apply as written - it patches mmq_x/mmq_x_max/ mmq_get_granularity_host in mul_mat_q_case, all deleted by PR ggml-org#24127, the same refactor that forced the Finding ggml-org#5 -> ggml-org#9 re-port. Re-authored onto mul_mat_q_switch_J's config-table search and scoped to RDNA3.5: gfx1151 is the only arch we can measure, and narrowing the predicate keeps the A/B single-variable. Upstream's predicate already covers RDNA3.5 via GGML_CUDA_CC_IS_RDNA3, and the author confirmed on 2026-07-04 that Strix Halo should benefit, but nobody has ever tested it there. On the routed-MoE path ncols_max is ne12 - the worst case of one expert receiving every token - so the search never reaches ntiles==1 and simply runs to whatever ceiling it is given. The picker instead tiles from ncols_dst/nchannels_x, the width a typical expert covers. launch_mul_mat_q still builds the grid from ncols_max, so worst-case coverage is unchanged. - mmq.cuh: replace the RDNA3.5 MoE J_max=48 clamp with the typical-width picker; J_max is now a uniform 128 and the search terminates on its own. - mmq-table-check.cpp: pick_J mirrors the selection loop, so update it in lockstep or the check silently becomes a rubber stamp. Adds the routed-MoE arg shape from mmq.cu and a crossover case. For Qwen 3.6 35B-A3B (256 experts, 8 active) at the production ub=2048 the typical width is 2048*8/256 = 64, so this is a 48 -> 64 change at our operating point - narrow enough that it may land in noise. Verified host-side without ROCm: MoE J=64, dense J=128, no aborts across all 21 types. Above ub=4096 the typical width reaches 128, the picker disengages and the search runs to J=128 where the cap held 48. Findings ggml-org#5/ggml-org#8 only ever measured the typical=64 regime, so that is untested on this chip; production ub=2048 cannot reach it. See strix-halo/mmq-moe-ncols-picker.md for the bench plan and the keep/revert criteria. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* CUDA: refactor MMQ kernel configuration * fix Blackwell config * remove legacy code
* CUDA: refactor MMQ kernel configuration * fix Blackwell config * remove legacy code
Upstream's MMQ refactor (ggml-org#24127) added `static_assert((I_) % 32 == 0, "bad I")` to the CASE macro. The rdna3_5 config introduced in cf4f42e used I=48, which fails the assertion (48 % 32 == 16 != 0) and breaks HIP builds for gfx1150/1/2/3 (Strix Halo). 48 was never a valid choice even without the static_assert. With nwarps=4 and the AMD MMA write_back path, each warp covers rows_per_warp=16 (hardcoded for AMD_MFMA/AMD_WMMA) and tile_C::I=16, so total rows = nwarps*16 = 64. I=48 would have left warp 3 writing rows 48-63 out of bounds. 64 is the smallest valid I for nwarps=4, satisfies I % 32 == 0, keeps the "smaller X tile" intent vs rdna2/rdna4's I=128, and matches the project notes in STRIX_HALO_NOTES.md ("use 64 as the safe final value"). All 232 CASE entries + the default return at the bottom updated. Fixes fewtarius/llama-ai#7
Upstream's MMQ refactor (ggml-org#24127) added `static_assert((I_) % 32 == 0, "bad I")` to the CASE macro. The rdna3_5 config introduced in cf4f42e used I=48, which fails the assertion (48 % 32 == 16 != 0) and breaks HIP builds for gfx1150/1/2/3 (Strix Halo). 48 was never a valid choice even without the static_assert. With nwarps=4 and the AMD MMA write_back path, each warp covers rows_per_warp=16 (hardcoded for AMD_MFMA/AMD_WMMA) and tile_C::I=16, so total rows = nwarps*16 = 64. I=48 would have left warp 3 writing rows 48-63 out of bounds. 64 is the smallest valid I for nwarps=4, satisfies I % 32 == 0, keeps the "smaller X tile" intent vs rdna2/rdna4's I=128, and matches the project notes in STRIX_HALO_NOTES.md ("use 64 as the safe final value"). All 232 CASE entries + the default return at the bottom updated. Fixes fewtarius/llama-ai#7
The Strix Halo row claimed "Merged upstream" but `mmq-config-rdna3_5.cuh` does not exist on `upstream/master` - gaetan-puleo's tables never landed in ggml-org/llama.cpp. The follow-up `71d1e8f2f` (I=48 -> I=64) is now linked into that entry instead of being orphaned. Added the missing Vulkan APU `nodes_per_submit` auto-lower row (`1c19480da`) so the table reflects what CachyLLama actually carries. Reworded the "drop on merge" guidance: CachyLLama focuses downstream, so in-tree copies of unupstreamed third-party patches stay until they land upstream cleanly and we rebase the local additions on top. Flagged upstream ggml-org#24127 as a watch item for future Strix Halo bumps - its `static_assert((I_) % 32 == 0)` is why `71d1e8f2f` exists.
…#24127) Upstream PR ggml-org#24127 replaced the per-arch MMQ macro/ternary helpers with per-arch config tables keyed by (type, J, fallback), renamed mmq_x/mmq_y to J/I, and made __launch_bounds__ mandatory. It deleted all six functions the old gfx1151 patch edited, so that patch was dropped on this rebase rather than replayed; this re-authors it against the new architecture. amd_wmma_available() still covers all of RDNA3, so without a dispatch branch gfx1151 silently inherits the rdna4 table (nthreads=256, occupancy=2, I=128 for every type) - the same tile shape this fork has been beating since April. - mmq-config-rdna3_5.cuh: generated from mmq-config-rdna4.cuh with nthreads 256 -> 128 and I 128 -> 64. sram_layout, K_vram and stream_k are carried over unchanged for all 21 types; only tile shape is retuned. Adds a J=48 fallback=true row per type: rdna4 thins fallback specializations to powers of 2, so without it the MoE cap below would silently settle on 32 whenever ne01 % 128 != 0. - mmq.cuh: host and device dispatch, both ahead of the AMD_WMMA branch. - mmq.cuh: MoE J cap in mul_mat_q_switch_J. args.expert_bounds still reaches the J-selection loop, so the old dense/MoE split survives as a bound on the search rather than an mmq_x_max override. I and nthreads are not independent: the MMA write-back needs I == nwarps*16 on wave32 (rdna4 is 128 == 8*16, this table is 64 == 4*16). The old patch's mmq_y=64 / nwarps=4 pair was obeying the same constraint. occupancy=2 is unvalidated - the knob did not exist when this tuning was last measured. It gives 2x rdna4's per-thread register budget but only 256 threads resident per CU. occupancy 2 vs 4 is the first sweep if the bench comes in flat. mmq-table-check.cpp verifies the table is well-formed and dispatches as intended without ROCm: it reuses the real CASE macro so the upstream static_asserts fire, checks every rdna4-reachable entry has an rdna3_5 counterpart, checks I == nwarps*16, and checks no input drives J_best to 0. That proves the table is correct, not that it is fast. NOT BENCHED. Never run on gfx1151. Gated on the two-run A/B (port off, then port on) in strix-halo/mmq-rdna3_5-config-table.md - this rebase also carries 234 upstream commits and the ROCm 7.14.0 switch, so the re-port is not attributable without it. Assisted-by: Claude Fable 5
…24546) Upstream PR ggml-org#24546 (ravel7524, open) sizes routed-MoE MMQ N-tiles from the typical expert width instead of the worst case. That is the same problem the static J_max=48 cap from Findings ggml-org#8/ggml-org#9 solves, approached from the other end: the cap clamps the search ceiling, the picker fixes the search objective. This commit swaps the cap for the picker so the two can be A/B'd on gfx1151. The PR does not apply as written - it patches mmq_x/mmq_x_max/ mmq_get_granularity_host in mul_mat_q_case, all deleted by PR ggml-org#24127, the same refactor that forced the Finding ggml-org#5 -> ggml-org#9 re-port. Re-authored onto mul_mat_q_switch_J's config-table search and scoped to RDNA3.5: gfx1151 is the only arch we can measure, and narrowing the predicate keeps the A/B single-variable. Upstream's predicate already covers RDNA3.5 via GGML_CUDA_CC_IS_RDNA3, and the author confirmed on 2026-07-04 that Strix Halo should benefit, but nobody has ever tested it there. On the routed-MoE path ncols_max is ne12 - the worst case of one expert receiving every token - so the search never reaches ntiles==1 and simply runs to whatever ceiling it is given. The picker instead tiles from ncols_dst/nchannels_x, the width a typical expert covers. launch_mul_mat_q still builds the grid from ncols_max, so worst-case coverage is unchanged. - mmq.cuh: replace the RDNA3.5 MoE J_max=48 clamp with the typical-width picker; J_max is now a uniform 128 and the search terminates on its own. - mmq-table-check.cpp: pick_J mirrors the selection loop, so update it in lockstep or the check silently becomes a rubber stamp. Adds the routed-MoE arg shape from mmq.cu and a crossover case. For Qwen 3.6 35B-A3B (256 experts, 8 active) at the production ub=2048 the typical width is 2048*8/256 = 64, so this is a 48 -> 64 change at our operating point - narrow enough that it may land in noise. Verified host-side without ROCm: MoE J=64, dense J=128, no aborts across all 21 types. Above ub=4096 the typical width reaches 128, the picker disengages and the search runs to J=128 where the cap held 48. Findings ggml-org#5/ggml-org#8 only ever measured the typical=64 regime, so that is untested on this chip; production ub=2048 cannot reach it. See strix-halo/mmq-moe-ncols-picker.md for the bench plan and the keep/revert criteria. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* CUDA: refactor MMQ kernel configuration * fix Blackwell config * remove legacy code
* CUDA: refactor MMQ kernel configuration * fix Blackwell config * remove legacy code
On master it is not possible to configure the CUDA MMQ kernel as a function of batch size and data type. This PR fixes that with a general refactor of the MMQ kernel that resembles more the mma FA kernel with a table of parameters rather than a bunch of macros and functions that return hard-coded values per architecture. Also I moved a lot of the code out of the main
mmq.cuhfile since it is pretty overloaded on master with 4k LoC. I removed the variable namesmmq_xandmmq_yand replaced them withJandI(same as the FA kernels) to avoid confusion with thexandydata pointers. There are no (intentional) functional changes from this PR other than:src0->ne[1]direction as well as the fallback version with those checks (something like ~5% end-to-end performance difference). However, for the fallback case it should be fine to compile fewer template specializations; a lot of them are just there to make pp snappier for short prompts where the number of tokens is not necessarily cleanly divided by e.g. 64 or 128. So for the fallback case I reduced the template specializations to only powers of 2. Longer-term we can consider adding a compilation option likeGGML_CUDA_FULLas an opt-in for template specializations that are rarely useful but blow up the compilation time.__launch_bounds__is optional, with this PR it becomes mandatory in the configuration. This should only affect RDNA1 where a targeted occupancy of 2 is now given.Going forward this PR will enable:
__dp4awith 4 byte loads, to be replaced with 16 byte loads (~10% end-to-end speedup for e.g. P40) that can to a large degree re-use the SRAM layout for tensor cores. Originally I was going to do this transition first so code is being removed before the refactor but this triggered performance regressions for some combinations of GPUs and data types. So I'm taking a more granular approach where I will do the transition piece-by-piece; the refactor in this PR still has some WIPs and inconsistencies that I will gradually phase out.Requirements