From 485088bdd35a667f99d21c547f5481159c7e0777 Mon Sep 17 00:00:00 2001 From: danielhanchen Date: Sat, 5 Sep 2026 23:41:42 -0700 Subject: [PATCH 1/5] CUDA: MMQ tile configuration and FP4 accumulation for GB10 (sm_121) Dense prompt matmuls on GB10 run the MMQ mainloop at one block per SM: the 128-row tile needs 61.5 KB of shared memory at J=128 and the kernel asks for 256 threads, so nothing can overlap the register-staged global loads or the two barriers per K step. Halving the tile to 64 rows with 128 threads and occupancy 2 makes two blocks resident and is worth 9 to 13 percent on the K-quants at prompt-sized batches. mul_mat_id does not want this: each expert only holds ncols_max*n_expert_used/n_experts columns, so halving the tile height only doubles the number of row tiles that pay for a mostly empty J tile. The tile table therefore keeps the stock 128-row tile in its fallback == true entries and mul_mat_id asks for that config on Blackwell; requesting the bounds-checked variant cannot change results, it only enables the src0->ne[1] range check. The block-scaled FP4 MMA already takes C as an accumulator input, but the vec_dot zeroed a temporary tile per instruction and added it to the running sum afterwards. Accumulating into the sum directly removes 805 FADD and all of the register spilling from the NVFP4 J=128 kernel. Both changes are confined to the Blackwell configuration and to code under BLACKWELL_MMA_AVAILABLE, so no other GPU sees a different kernel. --- ggml/src/ggml-cuda/mmq-config-blackwell.cuh | 82 +++++++++++++++------ ggml/src/ggml-cuda/mmq-vec-dot.cuh | 8 +- ggml/src/ggml-cuda/mmq.cu | 11 ++- 3 files changed, 73 insertions(+), 28 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq-config-blackwell.cuh b/ggml/src/ggml-cuda/mmq-config-blackwell.cuh index 9fbe32b6972b..f9689c7d77dd 100644 --- a/ggml/src/ggml-cuda/mmq-config-blackwell.cuh +++ b/ggml/src/ggml-cuda/mmq-config-blackwell.cuh @@ -1,37 +1,75 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_blackwell(ggml_type type, int J, bool fallback) { + CASE(GGML_TYPE_Q4_K, 128, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_K, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_K, 128, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_K, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_K, 128, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_K, 128, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_K, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_K, 128, 2, 64, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_K, 128, 2, 64, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_K, 128, 2, 64, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_K, 128, 2, 64, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + + CASE(GGML_TYPE_Q5_K, 128, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q5_K, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q5_K, 128, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q5_K, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q5_K, 128, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q5_K, 128, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q5_K, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q5_K, 128, 2, 64, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q5_K, 128, 2, 64, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q5_K, 128, 2, 64, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q5_K, 128, 2, 64, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + + CASE(GGML_TYPE_Q6_K, 128, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q6_K, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q6_K, 128, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q6_K, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q6_K, 128, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q6_K, 128, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q6_K, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q6_K, 128, 2, 64, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q6_K, 128, 2, 64, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q6_K, 128, 2, 64, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q6_K, 128, 2, 64, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_MXFP4, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true); CASE(GGML_TYPE_MXFP4, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true); CASE(GGML_TYPE_MXFP4, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true); CASE(GGML_TYPE_MXFP4, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true); CASE(GGML_TYPE_MXFP4, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true); - CASE(GGML_TYPE_MXFP4, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_MXFP4, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_MXFP4, 256, 1, 128, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_MXFP4, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_MXFP4, 256, 1, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_MXFP4, 256, 1, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_MXFP4, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_MXFP4, 256, 1, 128, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_MXFP4, 256, 1, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_MXFP4, 256, 1, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_MXFP4, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + + CASE(GGML_TYPE_MXFP4, 128, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 128, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 128, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 128, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 128, 2, 64, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 128, 2, 64, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 128, 2, 64, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 128, 2, 64, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); CASE(GGML_TYPE_NVFP4, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true); CASE(GGML_TYPE_NVFP4, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true); CASE(GGML_TYPE_NVFP4, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true); CASE(GGML_TYPE_NVFP4, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true); CASE(GGML_TYPE_NVFP4, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true); - CASE(GGML_TYPE_NVFP4, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_NVFP4, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_NVFP4, 256, 1, 128, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_NVFP4, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_NVFP4, 256, 1, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_NVFP4, 256, 1, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_NVFP4, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_NVFP4, 256, 1, 128, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_NVFP4, 256, 1, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_NVFP4, 256, 1, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_NVFP4, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + + CASE(GGML_TYPE_NVFP4, 128, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 128, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 128, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 128, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 128, 2, 64, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 128, 2, 64, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 128, 2, 64, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 128, 2, 64, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); return ggml_cuda_mmq_get_config_ampere(type, J, fallback); } diff --git a/ggml/src/ggml-cuda/mmq-vec-dot.cuh b/ggml/src/ggml-cuda/mmq-vec-dot.cuh index d573433865f8..2a1cd12c09b4 100644 --- a/ggml/src/ggml-cuda/mmq-vec-dot.cuh +++ b/ggml/src/ggml-cuda/mmq-vec-dot.cuh @@ -1236,14 +1236,12 @@ template static __device__ __forceinline_ #pragma unroll for (int n = 0; n < ntx; ++n) { + // The block-scaled MMA takes C as an accumulator input, so accumulate into the running sum directly. + static_assert(sizeof(tile_C) == tile_C::ne * sizeof(float), "tile_C must be a plain float array"); + tile_C & C = *reinterpret_cast(sum + (j0 / tile_C::J + n) * tile_C::ne); #pragma unroll for (int frag = 0; frag < nfrags; ++frag) { - tile_C C = {}; mma_block_scaled_fp4(C, A[n][frag], B[frag], scaleA[n][frag], scaleB[frag]); -#pragma unroll - for (int l = 0; l < tile_C::ne; ++l) { - sum[(j0 / tile_C::J + n) * tile_C::ne + l] += C.x[l]; - } } } } diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 707437ea3e52..168d800ee311 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -126,7 +126,16 @@ void ggml_cuda_mul_mat_q( const int64_t s03 = src0->nb[3] / ts_src0; const int64_t s3 = dst->nb[3] / ts_dst; - const bool fallback = ne01 % 128 != 0; + // The Blackwell MMQ tile table uses a 64-row / occupancy-2 tile for the K-quants and the FP4 + // types, which is a win for dense mul_mat but a loss for mul_mat_id: the MoE launch picks its J + // from ncols_max = the total token count while each expert only holds + // ncols_max*n_expert_used/n_experts columns, so halving I only doubles the number of row tiles + // that pay for the mostly empty J tile. The table therefore keeps the stock 128-row tile in its + // fallback == true entries and mul_mat_id selects those. Requesting the bounds-checked variant + // is always safe: `fallback` only enables the src0->ne[1] range check in the tile loads and the + // write-back, so this cannot change results, and it is gated on Blackwell so no other GPU sees + // a different kernel. + const bool fallback = ne01 % 128 != 0 || (ids && blackwell_mma_available(cc)); const bool use_native_fp4 = blackwell_mma_available(cc) && (src0->type == GGML_TYPE_MXFP4 || src0->type == GGML_TYPE_NVFP4); const size_t y_block_size = use_native_fp4 ? sizeof(block_fp4_mmq) : sizeof(block_q8_1_mmq); From 92db8b7f7c49b1760b98c5a441e19402af8f356e Mon Sep 17 00:00:00 2001 From: danielhanchen Date: Sat, 5 Sep 2026 23:59:55 -0700 Subject: [PATCH 2/5] CUDA: MMQ, carry the fallback tile choice from the caller instead of re-deriving it mul_mat_q_case re-derived the fallback flag from args.nrows_x % 128, so the flag that ggml_cuda_mul_mat_q computes (and that already sizes the src1 padding through ggml_cuda_mmq_get_J_max) never reached the kernel selection. It matched by construction while the only input was ne01, but any other reason to want the bounds-checked configuration was silently dropped, including the mul_mat_id case added in the previous commit. Carry the decision in mmq_args so the configuration the padding was sized for is the configuration that runs. No behaviour change for the ne01-only rule. --- ggml/src/ggml-cuda/mmq.cu | 4 ++-- ggml/src/ggml-cuda/mmq.cuh | 6 +++++- 2 files changed, 7 insertions(+), 3 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 168d800ee311..127cff5d557d 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -180,7 +180,7 @@ void ggml_cuda_mul_mat_q( ne00, ne01, ne1, s01, ne11, s1, ne02, ne12, s02, s12, s2, ne03, ne13, s03, s13, s3, - ne1}; + ne1, fallback}; ggml_cuda_mul_mat_q_switch_type(ctx, args, stream); return; } @@ -260,7 +260,7 @@ void ggml_cuda_mul_mat_q( ne00, ne01, ne_get_rows, s01, ne_get_rows, s1, ne02, ne02, s02, s12, s2, ne03, ne13, s03, s13, s3, - ne12}; + ne12, fallback}; ggml_cuda_mul_mat_q_switch_type(ctx, args, stream); } diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 2eb15fdfad93..ee6f38c7c4dc 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -1375,6 +1375,10 @@ struct mmq_args { int64_t nchannels_x; int64_t nchannels_y; int64_t stride_channel_x; int64_t stride_channel_y; int64_t stride_channel_dst; int64_t nsamples_x; int64_t nsamples_y; int64_t stride_sample_x; int64_t stride_sample_y; int64_t stride_sample_dst; int64_t ncols_max; + // Whether to select the bounds-checked ("fallback") tile configuration. Decided by the caller in + // ggml_cuda_mul_mat_q so that the config chosen here matches the one the src1 padding was sized + // for; it is set whenever src0->ne[1] % 128 != 0, and additionally for mul_mat_id on Blackwell. + bool fallback; }; static size_t mmq_get_nbytes_shared(const ggml_cuda_mmq_config & config, const int cc) { @@ -1551,7 +1555,7 @@ void mul_mat_q_switch_J(ggml_backend_cuda_context & ctx, const mmq_args & args, template void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) { - if (args.nrows_x % 128 == 0) { + if (!args.fallback) { constexpr bool fallback = false; mul_mat_q_switch_J(ctx, args, stream); } else { From 585f9c9eb4747f8984902dacf9fd86bba473ffcd Mon Sep 17 00:00:00 2001 From: danielhanchen Date: Tue, 8 Sep 2026 18:41:33 -0700 Subject: [PATCH 3/5] cuda: trim comments in the MMQ tile configuration change --- ggml/src/ggml-cuda/mmq-vec-dot.cuh | 2 +- ggml/src/ggml-cuda/mmq.cu | 12 +++--------- ggml/src/ggml-cuda/mmq.cuh | 4 +--- 3 files changed, 5 insertions(+), 13 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq-vec-dot.cuh b/ggml/src/ggml-cuda/mmq-vec-dot.cuh index 2a1cd12c09b4..95dbd88b9fbf 100644 --- a/ggml/src/ggml-cuda/mmq-vec-dot.cuh +++ b/ggml/src/ggml-cuda/mmq-vec-dot.cuh @@ -1236,7 +1236,7 @@ template static __device__ __forceinline_ #pragma unroll for (int n = 0; n < ntx; ++n) { - // The block-scaled MMA takes C as an accumulator input, so accumulate into the running sum directly. + // block-scaled MMA takes C as an accumulator input: accumulate into the running sum. static_assert(sizeof(tile_C) == tile_C::ne * sizeof(float), "tile_C must be a plain float array"); tile_C & C = *reinterpret_cast(sum + (j0 / tile_C::J + n) * tile_C::ne); #pragma unroll diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 127cff5d557d..f092b7f172c2 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -126,15 +126,9 @@ void ggml_cuda_mul_mat_q( const int64_t s03 = src0->nb[3] / ts_src0; const int64_t s3 = dst->nb[3] / ts_dst; - // The Blackwell MMQ tile table uses a 64-row / occupancy-2 tile for the K-quants and the FP4 - // types, which is a win for dense mul_mat but a loss for mul_mat_id: the MoE launch picks its J - // from ncols_max = the total token count while each expert only holds - // ncols_max*n_expert_used/n_experts columns, so halving I only doubles the number of row tiles - // that pay for the mostly empty J tile. The table therefore keeps the stock 128-row tile in its - // fallback == true entries and mul_mat_id selects those. Requesting the bounds-checked variant - // is always safe: `fallback` only enables the src0->ne[1] range check in the tile loads and the - // write-back, so this cannot change results, and it is gated on Blackwell so no other GPU sees - // a different kernel. + // mul_mat_id takes the fallback entries to keep the stock 128-row tile: J comes from ncols_max + // but an expert holds only ncols_max*n_expert_used/n_experts of it, so a 64-row tile just doubles + // the row tiles paying for a mostly empty J. fallback only adds the src0->ne[1] range check. const bool fallback = ne01 % 128 != 0 || (ids && blackwell_mma_available(cc)); const bool use_native_fp4 = blackwell_mma_available(cc) && (src0->type == GGML_TYPE_MXFP4 || src0->type == GGML_TYPE_NVFP4); diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index ee6f38c7c4dc..d75a351f3cfa 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -1375,9 +1375,7 @@ struct mmq_args { int64_t nchannels_x; int64_t nchannels_y; int64_t stride_channel_x; int64_t stride_channel_y; int64_t stride_channel_dst; int64_t nsamples_x; int64_t nsamples_y; int64_t stride_sample_x; int64_t stride_sample_y; int64_t stride_sample_dst; int64_t ncols_max; - // Whether to select the bounds-checked ("fallback") tile configuration. Decided by the caller in - // ggml_cuda_mul_mat_q so that the config chosen here matches the one the src1 padding was sized - // for; it is set whenever src0->ne[1] % 128 != 0, and additionally for mul_mat_id on Blackwell. + // set by ggml_cuda_mul_mat_q, not re-derived, so it matches the config src1 padding was sized for bool fallback; }; From d0a262e1a60ff7ad4ef55177abd47e2177db7db6 Mon Sep 17 00:00:00 2001 From: danielhanchen Date: Tue, 8 Sep 2026 21:40:57 -0700 Subject: [PATCH 4/5] CUDA: keep the GB10 MMQ tiles off every other Blackwell part The swept tiles were written into mmq-config-blackwell.cuh, so every sm_12x device got a table measured on one DGX Spark. GB10 is one SM cluster against unified LPDDR5X; a discrete Blackwell part has neither, and nothing in the sweep was measured on one. The tiles move to mmq-config-dgx-spark.cuh and mmq-config-blackwell.cuh returns to master byte for byte. Types the sweep did not cover fall through to the Blackwell table, so this narrows the blast radius without narrowing coverage. Dispatch is an exact equality against GGML_CUDA_CC_DGX_SPARK, not a >= range, following mmvq.cu:111, which tests the same ggml_cuda_highest_compiled_arch the existing dispatch uses. fattn.cu:194 is a >= threshold and is deliberately not the pattern copied here: a range would catch every future sm_12x part with a table measured on none of them. The host and device selectors mirror each other, or the host would size shared memory for one config while the kernel compiled another. Verified by SASS rather than by reading. Compiling the q4_k, q5_k, q6_k and q8_0 MMQ instantiations for sm_75, sm_80, sm_90, sm_120 and sm_121a and disassembling each cubin, against UPSTREAM MASTER as the reference: sm_75, sm_80, sm_90, sm_120: match master for all four types sm_121a: differs for q4_k, q5_k, q6_k; matches master for q8_0 which is exactly the intended shape, since q8_0 is not in the Spark table. The residual on every matching cell is 32 lines carrying one constant, 0x3bf to 0x3ca, an assert's __LINE__ shifted by the 11 lines this adds above it, materialised as MOV on sm_75, IMAD.MOV.U32 on sm_90 and HFMA2 on sm_80 and sm_120. No instruction, register allocation or scheduling differs, so the claim is "unchanged apart from an assert's line number" rather than byte-identical. ggml-cuda builds green for sm_121a. --- ggml/src/ggml-cuda/mmq-config-blackwell.cuh | 82 ++++++--------------- ggml/src/ggml-cuda/mmq-config-dgx-spark.cuh | 70 ++++++++++++++++++ ggml/src/ggml-cuda/mmq.cu | 11 ++- ggml/src/ggml-cuda/mmq.cuh | 13 +++- 4 files changed, 111 insertions(+), 65 deletions(-) create mode 100644 ggml/src/ggml-cuda/mmq-config-dgx-spark.cuh diff --git a/ggml/src/ggml-cuda/mmq-config-blackwell.cuh b/ggml/src/ggml-cuda/mmq-config-blackwell.cuh index f9689c7d77dd..9fbe32b6972b 100644 --- a/ggml/src/ggml-cuda/mmq-config-blackwell.cuh +++ b/ggml/src/ggml-cuda/mmq-config-blackwell.cuh @@ -1,75 +1,37 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_blackwell(ggml_type type, int J, bool fallback) { - CASE(GGML_TYPE_Q4_K, 128, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q4_K, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q4_K, 128, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q4_K, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q4_K, 128, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q4_K, 128, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q4_K, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q4_K, 128, 2, 64, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q4_K, 128, 2, 64, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q4_K, 128, 2, 64, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q4_K, 128, 2, 64, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - - CASE(GGML_TYPE_Q5_K, 128, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q5_K, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q5_K, 128, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q5_K, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q5_K, 128, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q5_K, 128, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q5_K, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q5_K, 128, 2, 64, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q5_K, 128, 2, 64, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q5_K, 128, 2, 64, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q5_K, 128, 2, 64, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - - CASE(GGML_TYPE_Q6_K, 128, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q6_K, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q6_K, 128, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q6_K, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q6_K, 128, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q6_K, 128, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q6_K, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q6_K, 128, 2, 64, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q6_K, 128, 2, 64, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q6_K, 128, 2, 64, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q6_K, 128, 2, 64, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_MXFP4, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true); CASE(GGML_TYPE_MXFP4, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true); CASE(GGML_TYPE_MXFP4, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true); CASE(GGML_TYPE_MXFP4, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true); CASE(GGML_TYPE_MXFP4, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true); - - CASE(GGML_TYPE_MXFP4, 128, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_MXFP4, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_MXFP4, 128, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_MXFP4, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_MXFP4, 128, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_MXFP4, 128, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_MXFP4, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_MXFP4, 128, 2, 64, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_MXFP4, 128, 2, 64, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_MXFP4, 128, 2, 64, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_MXFP4, 128, 2, 64, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 256, 1, 128, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 256, 1, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 256, 1, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 256, 1, 128, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 256, 1, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 256, 1, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); CASE(GGML_TYPE_NVFP4, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true); CASE(GGML_TYPE_NVFP4, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true); CASE(GGML_TYPE_NVFP4, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true); CASE(GGML_TYPE_NVFP4, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true); CASE(GGML_TYPE_NVFP4, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true); - - CASE(GGML_TYPE_NVFP4, 128, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_NVFP4, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_NVFP4, 128, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_NVFP4, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_NVFP4, 128, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_NVFP4, 128, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_NVFP4, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_NVFP4, 128, 2, 64, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_NVFP4, 128, 2, 64, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_NVFP4, 128, 2, 64, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); - CASE(GGML_TYPE_NVFP4, 128, 2, 64, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 256, 1, 128, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 256, 1, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 256, 1, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 256, 1, 128, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 256, 1, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 256, 1, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); return ggml_cuda_mmq_get_config_ampere(type, J, fallback); } diff --git a/ggml/src/ggml-cuda/mmq-config-dgx-spark.cuh b/ggml/src/ggml-cuda/mmq-config-dgx-spark.cuh new file mode 100644 index 000000000000..2861f06d9143 --- /dev/null +++ b/ggml/src/ggml-cuda/mmq-config-dgx-spark.cuh @@ -0,0 +1,70 @@ +// GB10 (sm_121a) ONLY. These tiles were swept on a DGX Spark, which is one SM cluster against +// unified LPDDR5X, and nothing here was measured on a discrete Blackwell part: an RTX 50-series +// card or an RTX PRO 6000 has its own SM count, its own GDDR7 bandwidth and its own occupancy +// limits, so a config tuned for the Spark is a guess there. Anything this table does not name +// falls through to the stock Blackwell config, so the override is exactly the swept entries. +static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_dgx_spark(ggml_type type, int J, bool fallback) { + CASE(GGML_TYPE_Q4_K, 128, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_K, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_K, 128, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_K, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_K, 128, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_K, 128, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_K, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_K, 128, 2, 64, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_K, 128, 2, 64, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_K, 128, 2, 64, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_K, 128, 2, 64, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + + CASE(GGML_TYPE_Q5_K, 128, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q5_K, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q5_K, 128, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q5_K, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q5_K, 128, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q5_K, 128, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q5_K, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q5_K, 128, 2, 64, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q5_K, 128, 2, 64, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q5_K, 128, 2, 64, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q5_K, 128, 2, 64, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + + CASE(GGML_TYPE_Q6_K, 128, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q6_K, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q6_K, 128, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q6_K, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q6_K, 128, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q6_K, 128, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q6_K, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q6_K, 128, 2, 64, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q6_K, 128, 2, 64, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q6_K, 128, 2, 64, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q6_K, 128, 2, 64, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, true, false); + + + CASE(GGML_TYPE_MXFP4, 128, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 128, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 128, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 128, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 128, 2, 64, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 128, 2, 64, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 128, 2, 64, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_MXFP4, 128, 2, 64, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + + + CASE(GGML_TYPE_NVFP4, 128, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 128, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 128, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 128, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 128, 2, 64, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 128, 2, 64, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 128, 2, 64, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + CASE(GGML_TYPE_NVFP4, 128, 2, 64, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, false); + + return ggml_cuda_mmq_get_config_blackwell(type, J, fallback); +} diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index f092b7f172c2..829eab18a0f0 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -126,10 +126,13 @@ void ggml_cuda_mul_mat_q( const int64_t s03 = src0->nb[3] / ts_src0; const int64_t s3 = dst->nb[3] / ts_dst; - // mul_mat_id takes the fallback entries to keep the stock 128-row tile: J comes from ncols_max - // but an expert holds only ncols_max*n_expert_used/n_experts of it, so a 64-row tile just doubles - // the row tiles paying for a mostly empty J. fallback only adds the src0->ne[1] range check. - const bool fallback = ne01 % 128 != 0 || (ids && blackwell_mma_available(cc)); + // GB10 only, matching the tile table it belongs to: mul_mat_id takes the fallback entries to + // keep the stock 128-row tile, because J comes from ncols_max but an expert holds only + // ncols_max*n_expert_used/n_experts of it, so a 64-row tile just doubles the row tiles paying + // for a mostly empty J. fallback only adds the src0->ne[1] range check. + const bool spark = GGML_CUDA_CC_IS_NVIDIA(cc) + && ggml_cuda_highest_compiled_arch(cc) == GGML_CUDA_CC_DGX_SPARK; + const bool fallback = ne01 % 128 != 0 || (ids && spark); const bool use_native_fp4 = blackwell_mma_available(cc) && (src0->type == GGML_TYPE_MXFP4 || src0->type == GGML_TYPE_NVFP4); const size_t y_block_size = use_native_fp4 ? sizeof(block_fp4_mmq) : sizeof(block_q8_1_mmq); diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index d75a351f3cfa..9dbae780fc52 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -216,6 +216,7 @@ struct ggml_cuda_mmq_config { #include "mmq-config-pascal.cuh" #include "mmq-config-ampere.cuh" #include "mmq-config-blackwell.cuh" +#include "mmq-config-dgx-spark.cuh" #include "mmq-config-cdna.cuh" #include "mmq-config-rdna2.cuh" @@ -241,6 +242,12 @@ static __host__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config(const ggml_type ty } return ggml_cuda_mmq_get_config_rdna2(type, J, fallback); } + // Ahead of the Blackwell branch, and an exact match rather than a range: the swept tiles + // below are a GB10 measurement and every other sm_12x part has its own SM count and memory. + // highest_compiled_arch, like mmvq.cu, so the host picks the table the device was built with. + if (GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) == GGML_CUDA_CC_DGX_SPARK) { + return ggml_cuda_mmq_get_config_dgx_spark(type, J, fallback); + } if (blackwell_mma_available(cc)) { return ggml_cuda_mmq_get_config_blackwell(type, J, fallback); } @@ -264,7 +271,11 @@ static constexpr __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config(ggml_t return ggml_cuda_mmq_get_config_rdna2(type, J, fallback); #endif // CDNA #else -#ifdef BLACKWELL_MMA_AVAILABLE +#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ == GGML_CUDA_CC_DGX_SPARK + // Must mirror the host selector above, or the host would size shared memory for one config + // while the kernel compiled another. + return ggml_cuda_mmq_get_config_dgx_spark(type, J, fallback); +#elif defined(BLACKWELL_MMA_AVAILABLE) return ggml_cuda_mmq_get_config_blackwell(type, J, fallback); #elif __CUDA_ARCH__ >= GGML_CUDA_CC_VOLTA return ggml_cuda_mmq_get_config_ampere(type, J, fallback); From daccdb2732be9639e4213d4b9d1251b3778b0ab8 Mon Sep 17 00:00:00 2001 From: danielhanchen Date: Wed, 9 Sep 2026 06:34:56 -0700 Subject: [PATCH 5/5] CUDA: keep the GB10 FP4 accumulator off every other Blackwell part The gate narrowing in d0a262e1 was entirely in the tile configuration selector. The other half of this PR, the FP4 accumulator change, lives in ggml_cuda_mmq_vec_dot_fp4_fp4_mma, which is compiled for every architecture defining BLACKWELL_MMA_AVAILABLE, i.e. __CUDA_ARCH__ in [1200, 1300). That is all of sm_120 as well as sm_121, so an RTX 50-series card was still getting it. Measured with cuobjdump -sass over all 23 MMQ translation units at 75-real;80-real;90-real;120a-real;121a-real, against the merge base, with each differing line classified rather than filtered (the +11 __LINE__ shift this PR introduces in mmq.cuh is compared on the instruction encoding word, since ptxas renders the same constant as MOV on sm_75 and HFMA2.MMA on sm_80): arch before the narrowing at d0a262e1 with this commit sm_75 0 of 23 0 of 23 0 of 23 sm_80 0 of 23 0 of 23 0 of 23 sm_90 0 of 23 0 of 23 0 of 23 sm_120a 5 of 23 2 of 23 0 of 23 sm_121a 5 of 23 5 of 23 5 of 23 The two that remained on sm_120a were mmq-instance-mxfp4 (139680 SASS lines) and mmq-instance-nvfp4 (127968). GB10 loses nothing. A strict comparison of d0a262e1 against this commit is byte-identical on sm_121a for all 23 translation units, so the Spark runs the same machine code either way. The change being gated is bit-neutral, so this is about not shipping unmeasured code generation rather than about results: the merge base plus that hunk alone, against the merge base, is bit-identical on all 48 (type, shape) MUL_MAT cases including MXFP4 and NVFP4. It is gated anyway, because PTX leaves mma's accumulation order unspecified and the hunk was swept on a DGX Spark and nowhere else, which is the same reason the tile table is gated. --- ggml/src/ggml-cuda/mmq-vec-dot.cuh | 18 +++++++++++++++++- 1 file changed, 17 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/mmq-vec-dot.cuh b/ggml/src/ggml-cuda/mmq-vec-dot.cuh index 95dbd88b9fbf..f149f9131ba8 100644 --- a/ggml/src/ggml-cuda/mmq-vec-dot.cuh +++ b/ggml/src/ggml-cuda/mmq-vec-dot.cuh @@ -1236,13 +1236,29 @@ template static __device__ __forceinline_ #pragma unroll for (int n = 0; n < ntx; ++n) { - // block-scaled MMA takes C as an accumulator input: accumulate into the running sum. +#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ == GGML_CUDA_CC_DGX_SPARK + // GB10 only, like the tile table this PR adds. Block-scaled MMA takes C as an + // accumulator input, so accumulating into the running sum removes a zeroed temporary + // per instruction, 805 FADD and all of the register spilling. It was swept on a DGX + // Spark and nowhere else, and PTX leaves mma's accumulation order unspecified, so it + // stays off every other Blackwell part for the same reason the tiles do. static_assert(sizeof(tile_C) == tile_C::ne * sizeof(float), "tile_C must be a plain float array"); tile_C & C = *reinterpret_cast(sum + (j0 / tile_C::J + n) * tile_C::ne); #pragma unroll for (int frag = 0; frag < nfrags; ++frag) { mma_block_scaled_fp4(C, A[n][frag], B[frag], scaleA[n][frag], scaleB[frag]); } +#else +#pragma unroll + for (int frag = 0; frag < nfrags; ++frag) { + tile_C C = {}; + mma_block_scaled_fp4(C, A[n][frag], B[frag], scaleA[n][frag], scaleB[frag]); +#pragma unroll + for (int l = 0; l < tile_C::ne; ++l) { + sum[(j0 / tile_C::J + n) * tile_C::ne + l] += C.x[l]; + } + } +#endif // __CUDA_ARCH__ == GGML_CUDA_CC_DGX_SPARK } } }