Skip to content
Draft
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
22 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
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
9 changes: 9 additions & 0 deletions ggml/src/ggml-cuda/common.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -1639,3 +1639,12 @@ static __inline__ void ggml_cuda_kernel_launch(Kernel kernel, const ggml_cuda_ke
CUDA_CHECK(cudaGetLastError());
}

// Opt-in stderr logging for MMQ/MMVQ kernel dimensions (set GGML_CUDA_MM_LOG=1).
static inline bool ggml_cuda_mm_log_enabled() {
static int enabled = -1;
if (enabled < 0) {
enabled = getenv("GGML_CUDA_MM_LOG") ? 1 : 0;
}
return enabled;
}

15 changes: 15 additions & 0 deletions ggml/src/ggml-cuda/mmq.cu
Original file line number Diff line number Diff line change
Expand Up @@ -157,6 +157,15 @@ void ggml_cuda_mul_mat_q(
ne02, ne12, s02, s12, s2,
ne03, ne13, s03, s13, s3,
use_stream_k, ne1};

// Log kernel dimensions for profiling analysis (GGML_CUDA_MM_LOG=1).
// dst[ne01, ne11] = src0[ne00, ne01] @ src1[ne10, ne11]; ne00==ne10 (K)
// GEMM C[M,N] = W[N,K] @ X[K,M]: M=ne11, N=ne01, K=ne00
if (ggml_cuda_mm_log_enabled()) {
fprintf(stderr, "[MUL_MAT_Q] M=%ld N=%ld K=%ld ne00=%ld ne01=%ld ne1=%ld ne11=%ld ne02=%ld ne12=%ld ne03=%ld ne13=%ld type=%d\n",
ne11, ne01, ne00, ne00, ne01, ne1, ne11, ne02, ne12, ne03, ne13, src0->type);
}

ggml_cuda_mul_mat_q_switch_type(ctx, args, stream);
return;
}
Expand Down Expand Up @@ -219,6 +228,12 @@ void ggml_cuda_mul_mat_q(
ne03, ne13, s03, s13, s3,
use_stream_k, ne12};

// Log kernel dimensions for profiling analysis (batched/MoE path, GGML_CUDA_MM_LOG=1).
if (ggml_cuda_mm_log_enabled()) {
fprintf(stderr, "[MUL_MAT_Q_MoE] M=%ld N=%ld K=%ld ne_get_rows=%ld ne00=%ld ne01=%ld ne02=%ld ne12=%ld type=%d\n",
ne11_flat, ne01, ne00, ne_get_rows, ne00, ne01, ne02, ne12, src0->type);
}

ggml_cuda_mul_mat_q_switch_type(ctx, args, stream);
}

Expand Down
Loading
Loading