RDNA3.5 (gfx11 / gfx1151) MMQ prefill optimizations#32
Open
liangliangchang wants to merge 29 commits into
Open
RDNA3.5 (gfx11 / gfx1151) MMQ prefill optimizations#32liangliangchang wants to merge 29 commits into
liangliangchang wants to merge 29 commits into
Conversation
Instrumented CUDA kernels to log actual M×N×K dimensions to stderr for performance analysis and TFLOPS/bandwidth calculations. Changes: - quantize.cu: Log M, K, grid for quantize_mmq_q8_1 kernel - mmq.cu: Log M, N, K, ne dimensions for mul_mat_q (GEMM) kernel - mmvq.cu: Log M, K dimensions for mul_mat_vec_q (GEMV) kernel Usage: llama-bench ... 2>&1 | tee profile.log The logged dimensions are matched with rocprofv3 timing data to calculate: - TFLOPS: From actual M×N×K dimensions (2×M×N×K FLOPs) - Bandwidth: From actual data transfer patterns This instrumentation enabled accurate performance analysis showing: - GEMM kernels: 8.83-43.58 TFLOPS (varies by quantization type) - Quantize kernels: 284 GB/s bandwidth (data in L2 cache) - GEMV kernels: 0.05-5.52 TFLOPS (memory-bound decode) Co-Authored-By: Claude Sonnet 4 <noreply@anthropic.com>
Log GEMM as M=ne11, N=ne01, K=ne00 and GEMV as M=ne01, K=ne00 so instrumented shapes match actual matmul dimensions. Co-authored-by: Cursor <cursoragent@cursor.com>
On gfx115x, mmq_x=128 uses more than half of per-CU LDS and caps occupancy at one workgroup per CU. Select a smaller mmq_x that fits two workgroups without doubling M-tiles, improving P0 prefill TFLOPS ~25% on gfx1151. Co-authored-by: Cursor <cursoragent@cursor.com>
…ning. Restore WMMA granularity at mmq_x>=64, widen tile_y padding to cut LDS bank conflicts, and software-pipeline Q8_1 activation loads into registers during vec_dot to hide global memory latency on gfx115x. Co-authored-by: Cursor <cursoragent@cursor.com>
Enables MMQ_PROFILE_PHASES=1 to report load_tiles, y_act, and vec_dot cycle shares per launch, guiding RDNA3.5 prefill optimization work. Co-authored-by: Cursor <cursoragent@cursor.com>
…head. Add gfx115x load_tiles_q4_K_dm_rdna35 and vec_dot_q4_K_q8_1_mma with per-k dmA hoisting; restore pre-profiling tile loop when MMQ_PROFILE_PHASES is off so extra __syncthreads do not regress P0 TFLOPS. Co-authored-by: Cursor <cursoragent@cursor.com>
Split load_tiles_q4_K qs staging into load_tiles_q4_K_qs_wmma_rdna35 (one warp/row, coalesced qs global load) and mmq_q4_K_qs_lds_k helper that documents the pseudo-q8_1 LDS offsets used by load_ldmatrix. Perf (gfx1151, P0 128×12288×4096): neutral vs prior commit — 26.56 vs 26.51 TFLOPS, 485 vs 486 µs, pp128 ~1135 t/s (within run-to-run noise). Layout math is unchanged; this is groundwork for a Q4_K-native MMA tile. Co-authored-by: Cursor <cursoragent@cursor.com>
fprintf for MUL_MAT_Q, MUL_MAT_VEC_Q, QUANTIZE_MMQ_Q8_1, and MUL_MAT_Q_LAUNCH are off by default to avoid CPU overhead in llama-bench TTFT runs; set GGML_CUDA_MM_LOG=1 for profiling scripts. Co-authored-by: Cursor <cursoragent@cursor.com>
Hoist dmA*dsB into per-lane scl[] before mma on RDNA3.5 vec_dot_q4_K_q8_1_mma to separate scale multiply from C accumulation (ATT showed v_fma_mix hotspots). Prefill P0 Q4_K_M 128x12288x4096 (3x profile_mmq_prefill): median 26.59 TFLOPS / 484.7 us vs 26.56 / 485 us baseline; VGPR 224 (+8), LDS conflict ratio 0.348. Assisted-by: Auto Co-authored-by: Cursor <cursoragent@cursor.com>
clock64() can move backward across __syncthreads, which inflated phase cycle counts and produced garbage load_tiles percentages; discard small backward glitches and skip corrupt host-side totals. Co-authored-by: Cursor <cursoragent@cursor.com>
Use mmq_x=128 for small nrows_x grids instead of dual-WG LDS downsizing, and hoist dA in Q8_0/Q4_0 vec_dot on RDNA3.5 WMMA. Gate 128×32×4096 drops ~209µs to ~105µs with no P0 Q4_K regression. Co-authored-by: Cursor <cursoragent@cursor.com>
Precompute sclA = sc[k01/4]*d per k-slice in vec_dot_q6_K_q8_1_mma instead of reloading scales inside the j-loop. Q6_K FFN down 128×4096×12288 improves ~13.4 to ~15.2 TFLOPS with no P0 regression. Co-authored-by: Cursor <cursoragent@cursor.com>
Prefer ntx=1 tile width when batch ≤128 and LDS allows, improving Q6_K FFN down (~40% faster vs mmq_x=64 on gfx1151). Co-authored-by: Cursor <cursoragent@cursor.com>
…RDNA3.5 Prefetch B across j0 with batched WMMA/epilogue, pad tile_y loads for mmq_x<=64 software pipelining, add MMQ_PROFILE_VEC_SUB (off by default), and pick Q6_K mmq_x=128 only when LDS fits dual-WG occupancy (smpbo/2). Co-authored-by: Cursor <cursoragent@cursor.com>
Software register pipelining of tile_y loads showed no measurable gain (~4% y_act slice) and slightly regressed pp128 on gfx1151; keep padded tile_y LDS layout and the simple sequential load path. Co-authored-by: Cursor <cursoragent@cursor.com>
Load all B tiles per k-slice via load_generic (faster than load_ldmatrix from tile_y), interleave with A on the hot path, and merge the duplicate vec_dot profile loop into one path. Use fast process_tile when only MMQ_PROFILE_VEC_SUB is set. Co-authored-by: Cursor <cursoragent@cursor.com>
Revert the smpbo/2 gate added in 03927f1 for the Q6_K narrow-N launch heuristic. Q6 at mmq_x=128 (38 KiB, ntx=1) is ~40% faster than mmq_x=64 with ntx=2 even at 1 WG/CU; the dual-WG policy belongs on Q4_K only. Co-authored-by: Cursor <cursoragent@cursor.com>
Revert the d4511bd B-batch/load_generic vec_dot rewrite that regressed Q6 FFN down from ~13 to ~8 TFLOPS under rocprof trace (~1.5ms to ~1.1ms per launch). Keep the 4d9b940 j0 loop with load_ldmatrix for B; profile instrumentation remains on a separate g_mmq_profile early-return path only. Co-authored-by: Cursor <cursoragent@cursor.com>
Revert g_mmq_profile_outer/cudaMemcpyToSymbol and Q6 vec_dot global profile branches that regressed Q6_K FFN down ~15% under rocprof trace. Restore RDNA3.5 y-tile software pipeline and MMQ_PROFILE_PHASES-only instrumentation to match 4d9b940 prefill performance. Co-authored-by: Cursor <cursoragent@cursor.com>
Prefetch the next activation B tile while MMA runs on the current tile in the Q6_K narrow-N path, improving FFN-down kernel throughput without changing launch geometry. Co-authored-by: Cursor <cursoragent@cursor.com>
Drop the Q6-only mmq_x=128 upgrade so Q6_K FFN uses the same smpbo/2 mmq_x=64 ntx=2 launch as Q4_K, which measures faster on gfx115x prefill. Co-authored-by: Cursor <cursoragent@cursor.com>
Block quants (Q5_0, Q8_0) regressed when forced to mmq_x=64 for LDS dual-WG occupancy; keep mmq_x=128 for them while retaining mmq_x=64 for K-quants tuned on gfx115x (Q6_K FFN down). Co-authored-by: Cursor <cursoragent@cursor.com>
Keep gfx11 kernel tuning (dual-WG mmq_x, WMMA vec_dot, y-tile pipelining) while removing opt-in logging, phase profiling, and MMQ_FORCE_MMQ_X overrides. Co-authored-by: Cursor <cursoragent@cursor.com>
Stripping the profile branch caused HIP/clang to emit worse WMMA kernels (~12% pp regression). Keep a never-taken cold path and nullptr kernarg so the hot loop matches gfx11 codegen without runtime profiling hooks. Co-authored-by: Cursor <cursoragent@cursor.com>
Replace the dummy mmq_profile kernarg with MMQ_CODEGEN_SPLIT_COLD (threadIdx.z > 0) so HIP/clang keeps full WMMA unrolling on the hot path without an extra launch parameter. Add docs/backend/HIP-MMQ-compiler-codegen.md documenting the compiler issue, ASM evidence, and alternatives tested. Co-authored-by: Cursor <cursoragent@cursor.com>
Profiling removal left no callers; revert the define added for MMQ phase profiling. Co-authored-by: Cursor <cursoragent@cursor.com>
liangliangchang
changed the title
Re-add fprintf in the default switch case removed with profiling cleanup. Co-authored-by: Cursor <cursoragent@cursor.com>
mgehre-amd
reviewed
Jul 1, 2026
| #if defined(GGML_USE_HIP) | ||
| // ROCm/clang WMMA codegen workaround: an unprovable-false branch keeps full unroll/inline | ||
| // on the hot K-loop. blockDim.z is always 1 for mul_mat_q. See docs/backend/HIP-MMQ-compiler-codegen.md. | ||
| #define MMQ_CODEGEN_SPLIT_COLD (threadIdx.z > 0) |
There was a problem hiding this comment.
This approach looks brittle. Let's please investigate further what affects the unroll decision. Can you try with an explicit unroll pragma?
…ranch Replace MMQ_CODEGEN_SPLIT_COLD (threadIdx.z) with MMQ_HIP_TILE_BARRIER on HIP+RDNA3.5 to restore DS/WMMA/FMA_MIX scheduling without unreachable code. Co-authored-by: Cursor <cursoragent@cursor.com>
The tile-barrier fix is self-contained in mmq.cuh; the internal codegen bisection write-up is no longer needed in-tree. Co-authored-by: Cursor <cursoragent@cursor.com>
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
2 participants
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.
Summary
RDNA3.5 (
gfx11/gfx1151) MMQ prefill optimizations and a documented HIP/clang codegen workaround.MMQ optimizations (
mmq.cuh)mmq_xwhen LDS limits workgroups/CU; dual-WGmmq_x=64for narrow-N K-quant prefill (Q6_K FFN-down), gated to K-quants only.vec_dotfor WMMA;ldmatrix-friendly qs layout; precompute scale factors before WMMA accumulate.mmq_x=128for narrow-N prefill; batched B-tile loads (load_generic); pipelined B loads invec_dot; restoreload_ldmatrixpath.Q8_0on RDNA3.5.HIP compiler codegen workaround (required for performance)
Removing the cold/hot
if / elseloop structure caused HIP/clang ongfx1151to emit ~half the WMMA/unrolled vec_dot body (~12%pp128regression on Qwen2.5-0.5B Q4_K_M), even though only one path runs at runtime.Fix: keep a full duplicate cold loop behind
MMQ_CODEGEN_SPLIT_COLD (threadIdx.z > 0)on HIP. Launches useblockDim.z == 1, so the cold path never runs; the compiler cannot prove the branch dead and emits good hot-path ISA (same quality as the old profiling split, without a dummy kernarg).Full analysis, ASM metrics, and repro steps:
docs/backend/HIP-MMQ-compiler-codegen.md.Benchmarks (gfx1151, ROCm)
Run with
llama-bench(-p 128 -n 0 -ngl 999, flash-attn on). Rebuild MMQ aftermmq.cuhchanges:touch ggml/src/ggml-cuda/mmq.cuh && cmake --build build -j$(nproc).Files changed
ggml/src/ggml-cuda/mmq.cuh— RDNA3.5 MMQ optimizations +MMQ_CODEGEN_SPLIT_COLDdocs/backend/HIP-MMQ-compiler-codegen.md— compiler issue write-up (new)Test plan
GGML_HIP=ON, targetgfx11/gfx1151llama-benchon Q4_K_M and Q6_K models:pp128, narrow-N FFN-down shapesgfx11base; confirm no regression vs pre-strip branch with codegen splittemplate-instances/mmq-instance-q4_k.cuand verify hot-pathds_load_b128≈ 48 (not 24) permmq_x=128kernel body — see docAI usage disclosure
Used Cursor/AI for profiling analysis, ASM comparison, codegen-split investigation, and documentation.