From 89a1103ce9a29236d7c3d6ddfb1d9b5becb46fcb Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Sun, 2 Aug 2026 12:57:27 -0700 Subject: [PATCH] cuda: enable the Hopper wgmma prefill path by default Inverts the runtime gate from opt-in (GGML_HOPPER_Q1) to opt-out (GGML_HOPPER_Q1_DISABLE) for builds configured with GGML_CUDA_HOPPER_Q1. The build-time flag is unchanged and still off by default, since it requires a CUTLASS checkout. Validated internally: prefill improves for the shapes this path serves and decode is unaffected. A control arm with the path disabled matches baseline, so the sm_90a plus CUTLASS build costs nothing by itself. The per-128-K activation scale is coarser than q8_1's per-32, so results are not bit-identical. The deviation was measured against a same-path control; it is small and the tail is bounded. --- ggml/src/ggml-cuda/mmq-hopper-q1.cu | 19 +++++++++++++------ 1 file changed, 13 insertions(+), 6 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq-hopper-q1.cu b/ggml/src/ggml-cuda/mmq-hopper-q1.cu index 9aa63ec56271..b9e5d78c36eb 100644 --- a/ggml/src/ggml-cuda/mmq-hopper-q1.cu +++ b/ggml/src/ggml-cuda/mmq-hopper-q1.cu @@ -1,8 +1,15 @@ // Hopper (sm_90a) wgmma MMQ path for Q1_0: dequant-in-SMEM + int8 wgmma with exact per-block scaling. -// Experimental opt-in path (env GGML_HOPPER_Q1) targeting large-batch prefill on sm_90a. -// Activations are quantized fp32 -> int8 with a per-128-K absmax scale (coarser than q8_1's per-32; -// flagged for KLD validation). Dispatched only when M,N,K % 128 == 0 and cc >= 900; otherwise the -// caller falls through to the standard MMQ path. +// Active by default in builds configured with GGML_CUDA_HOPPER_Q1; set GGML_HOPPER_Q1_DISABLE to +// fall back to the standard MMQ path. Dispatched only when M,N,K % 128 == 0 and cc >= 900 (Hopper +// only: not Ada, not Blackwell); every other shape falls through to standard MMQ. +// +// Activations are quantized fp32 -> int8 with a per-128-K absmax scale, coarser than q8_1's +// per-32, so this path is NOT bit-identical to standard MMQ. The deviation has been measured +// against the same build with the path disabled, using a same-path control to establish the +// floor: it is small, the tail is bounded, and only a small fraction of tokens change their +// argmax. Figures are recorded in the internal notes. If you need bit-exact output, set +// GGML_HOPPER_Q1_DISABLE. + #include "common.cuh" #include @@ -442,8 +449,8 @@ bool ggml_cuda_mul_mat_q1_hopper(ggml_backend_cuda_context & ctx, const ggml_tensor * src1, ggml_tensor * dst) { #if defined(GGML_USE_HOPPER_Q1) - static const bool enabled = getenv("GGML_HOPPER_Q1") != nullptr; - if (!enabled) { + static const bool disabled = getenv("GGML_HOPPER_Q1_DISABLE") != nullptr; + if (disabled) { return false; } const int cc = ggml_cuda_info().devices[ctx.device].cc;