spec(KERNEL-QUANT-CIQ-GEMM-ROCM): scope a cooperative activation share for the WMMA kernels - #3035
Open
joral wants to merge 3 commits into
Open
spec(KERNEL-QUANT-CIQ-GEMM-ROCM): scope a cooperative activation share for the WMMA kernels#3035joral wants to merge 3 commits into
joral wants to merge 3 commits into
Conversation
…asure, reject (mudler#3032) Issue mudler#3032's same-tool rocprofv3 trace found KQuantGemmKWmmaQ4K/Q6K 4.9x-10.6x slower per-kernel than llama.cpp's mul_mat_q on the identical tensor-core mechanism, with llama.cpp launching double the warps per block and 16-48x fewer, bigger blocks. Each warp here already owns a fully independent output tile with its own shared-memory slice, so widening the block from 4 to 8 warps was the obvious first lever to test: does packing more of that already-independent work into fewer, bigger launches close any of the gap. Templated both kernels on WarpsPerBlock and added an 8-warp instantiation behind VT_ROCM_QUANT_WMMA_WIDE=1 (default off, same same-binary-A/B posture VT_ROCM_QUANT_WMMA and VT_ROCM_Q6K_SMALL_PRIVATE already ship with). Hardware-verified correct on both configs: ctest -R rocm|cross_device, 46/46 cases, 84066/84066 assertions, zero regression. The answer is no. Op-level A/B (quant-gemm-bench, RX 9060 XT, best-of-4, idle host) across the six Q4_K/Q6_K prefill shapes: geomean -4.3%, a net regression. Each warp shares no loaded or dequantized data with any other warp in its block, so widening only spreads the same LDS-capped occupancy budget over fewer, bigger blocks rather than doing more work per warp -- the axis this measurement rules out, not the one that explains llama.cpp's actual advantage. The sharper, still-open hypothesis this points at -- cross-warp data reuse in llama.cpp's own kernel -- is recorded in the spec's Owed section, unread and unconfirmed, as the next traceable step. The toggle stays in the tree, default off, as a ready-made A/B for re-checking this specific axis on different hardware or a future toolchain revision, rather than making the next person re-derive it from scratch. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-sonnet-5 [Claude Code]
…thesis by reading llama.cpp's source The wide-block experiment's spec entry left this as an unread inference. Reading ggml-cuda/mmq.cuh, mmq-config-rdna4.cuh, and mmq-load-tiles.cuh confirms it: llama.cpp's Q4_K block cooperatively loads and shares one 128x128 tile across all warps, while this kernel gives each warp its own independent 16x16 tile with no sharing at all, even between warps that share the same M-rows. That is the mechanism gap the rejected wider-block experiment could not reach. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-sonnet-5 [Claude Code]
…e for the WMMA kernels (mudler#3034) The wider-block A/B (mudler#3032/mudler#3033) measured a net regression, and reading llama.cpp's mul_mat_q source found the actual mechanism it uses: one cooperatively-loaded shared tile per block, not just more warps. Tracing this kernel's own tile/warp index math (tile = blockIdx.x * WarpsPerBlock + threadIdx.y; it = tile / n_tiles; jt = tile % n_tiles) shows every warp in one block already shares the same activation rows (it), since n_tiles is always far larger than WarpsPerBlock. Each warp reads those identical rows independently via its own load_matrix_sync call, every superblock -- the precise, confirmed redundancy, and exactly why widening the block alone (mudler#3033) did not help: it made that same redundant pattern wider without adding any sharing to fix it. This spec scopes a smaller, targeted fix instead of matching llama.cpp's full 128x128 cooperative tile: reuse the already-landed 8-warp block (kQuantWmmaWideWarpsPerBlock, mudler#3033, currently unused on its own), and add cooperative activation staging so the shared superblock's BlockQ8_K bytes are loaded into shared memory once per block instead of once per warp. LDS budget checked: ~55 KiB total against the 64 KiB limit. Pull request shape: separate spec and implementation pull requests, same recorded preference this row already uses. This pull request lands the spec only. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-sonnet-5 [Claude Code]
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
1 participant
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.
Issue #3034. Row:
KERNEL-QUANT-CIQ-GEMM-ROCM(alreadyACTIVE). One newspec section, no product code.
The wider-block A/B (#3032/#3033) measured a net regression (geomean
-4.3%), and reading llama.cpp's
mul_mat_qsource (confirmed, notinferred:
ggml-cuda/mmq.cuh,mmq-config-rdna4.cuh,mmq-load-tiles.cuh) found the actual mechanism: one cooperatively-loadedshared tile per block, not just more warps.
Tracing this kernel's own tile/warp index math (
tile = blockIdx.x * WarpsPerBlock + threadIdx.y;it = tile / n_tiles;jt = tile % n_tiles)shows every warp in one block already shares the same activation rows
(
it), sincen_tilesis always far larger thanWarpsPerBlock. Eachwarp reads those identical rows independently via its own
load_matrix_synccall, every superblock — the precise, confirmedredundancy, and exactly why widening the block alone (#3033) did not
help: it made that same redundant pattern wider without adding any
sharing to fix it.
This spec scopes a smaller, targeted fix rather than matching
llama.cpp's full 128x128 cooperative tile: reuse the already-landed
8-warp block (
kQuantWmmaWideWarpsPerBlock, #3033, currently unused onits own), and add cooperative activation staging so the shared
superblock's
BlockQ8_Kbytes load into shared memory once per blockinstead of once per warp. LDS budget checked: ~55 KiB total against the
64 KiB limit — no change to the existing weight-side staging needed to
fit it.
Deliberately scoped smaller than a full retile so the measurement
isolates one variable: if this wins, sharing was the lever the
wider-block-alone experiment could not reach; if it does not, the
redundant reads were likely already served by cache and the remaining
gap points elsewhere (raw WMMA instruction throughput, shared-memory
bank conflicts) — recorded either way, not assumed.
Pull request shape: separate spec and implementation, same recorded
preference this row already uses. This pull request lands the spec
only; the implementation follows as its own PR.
FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-sonnet-5 [Claude Code]