Skip to content

RDNA3.5 (gfx11 / gfx1151) MMQ prefill optimizations #32

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
liangliangchang wants to merge 29 commits into
base: gfx11
Choose a base branch
from
+618 −35
Open

RDNA3.5 (gfx11 / gfx1151) MMQ prefill optimizations #32

Changes from all commits
Commits
Show all changes
29 commits
Select commit Hold shift + click to select a range
a70d5dd
Add kernel dimension logging for profiling analysis
liangliangchang Jun 23, 2026
3c63ba2
Fix M/N/K dimension logging for MMQ and MMVQ profiling.
liangliangchang Jun 23, 2026
6051c1f
Improve RDNA3.5 MMQ occupancy by reducing mmq_x when LDS limits WGs/CU.
liangliangchang Jun 23, 2026
36249b0
Improve RDNA3.5 MMQ prefill via LDS layout and activation load pipeli…
liangliangchang Jun 23, 2026
7248ec9
Add opt-in HIP MMQ phase profiling to split dequant vs MMA time.
liangliangchang Jun 23, 2026
7ed8d1e
Improve RDNA3.5 Q4_K MMQ dequant/vec_dot and fix phase profiling over…
liangliangchang Jun 23, 2026
fb63447
Refactor RDNA3.5 Q4_K qs load for WMMA ldmatrix layout.
liangliangchang Jun 24, 2026
5924b0a
Gate MMQ kernel dimension logs behind GGML_CUDA_MM_LOG.
liangliangchang Jun 24, 2026
afa1455
cuda: precompute Q4_K vec_dot scale factors before WMMA accumulate
liangliangchang Jun 24, 2026
46b38f2
cuda: fix MMQ phase profiler clock64 deltas on HIP
liangliangchang Jun 24, 2026
bee2b6d
cuda: optimize RDNA3.5 tiny-M MMQ for gate Q8_0 prefill
liangliangchang Jun 25, 2026
6e777ce
cuda: hoist Q6_K vec_dot scales on RDNA3.5 WMMA
liangliangchang Jun 25, 2026
4d9b940
cuda: use mmq_x=128 for Q6_K narrow-N prefill on RDNA3.5
liangliangchang Jun 25, 2026
03927f1
cuda: Q6_K MMQ B-prefetch, y-tile pipeline, and vec_dot profiling on …
liangliangchang Jun 30, 2026
cf83d46
cuda: drop RDNA3.5 activation y-pipeline from MMQ
liangliangchang Jun 30, 2026
d4511bd
cuda: batch Q6_K B loads with load_generic on RDNA3.5 WMMA
liangliangchang Jun 30, 2026
eb2a7de
cuda: restore Q6_K mmq_x=128 for narrow-N FFN down on RDNA3.5
liangliangchang Jun 30, 2026
49b036d
cuda: restore Q6_K vec_dot load_ldmatrix path on RDNA3.5
liangliangchang Jun 30, 2026
94e9d1d
cuda: remove always-on MMQ profiling overhead on RDNA3.5 hot path
liangliangchang Jun 30, 2026
272c12e
cuda: pipeline Q6_K B-tile loads in RDNA3.5 WMMA vec_dot
liangliangchang Jun 30, 2026
4648486
cuda: use dual-WG mmq_x=64 for Q6_K narrow-N on RDNA3.5
liangliangchang Jun 30, 2026
35bd178
cuda: limit RDNA3.5 dual-WG mmq_x downgrade to K-quants
liangliangchang Jun 30, 2026
55945ef
cuda: strip profiling hooks from RDNA3.5 MMQ optimizations
liangliangchang Jun 30, 2026
df6dae2
cuda: restore HIP codegen split for RDNA3.5 MMQ without profiling
liangliangchang Jun 30, 2026
1c862d5
cuda: use threadIdx.z codegen split for RDNA3.5 MMQ
liangliangchang Jun 30, 2026
6044323
cuda: drop unused cudaMemcpyToSymbol HIP shim
liangliangchang Jun 30, 2026
3579242
cuda: restore mmq_x_best diagnostic on MMQ launch abort
liangliangchang Jun 30, 2026
b33c2de
cuda: fix RDNA3.5 MMQ regression with tile barriers instead of dead b…
liangliangchang Jul 7, 2026
d51d5b3
docs: remove HIP-MMQ-compiler-codegen investigation notes
liangliangchang Jul 7, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Loading
Loading