cuda: enable the Hopper wgmma Q1_0/Q2_0 prefill path by default

[bri-prism]

What

Flips the runtime gate on the Hopper (sm_90a) wgmma prefill path from opt-in to opt-out. In a build configured with GGML_CUDA_HOPPER_Q1, the path is now active unless GGML_HOPPER_Q1_DISABLE is set.

The build-time option is deliberately left off by default, because it requires a CUTLASS checkout via GGML_CUDA_CUTLASS_DIR and would otherwise break configuration for everyone else. See "Scope" below for what that means in practice.

Why

The path had never been built or exercised. It turns out to work, and to be worth having on for the shapes it serves.

Validated internally on Hopper with a three-arm comparison: baseline build, this build with the path off, and this build with the path on. The path-off arm matches baseline closely on every measurement, so the sm_90a plus CUTLASS build costs nothing by itself and the improvement is attributable to the path rather than to the build. Prefill improves; decode is unaffected, which is the expected behaviour for a prefill-only path and a useful check that nothing leaked into the wrong dispatch.

The gain is larger for the wider of the two low-bit types, consistent with it having more unpack work for int8 wgmma to absorb, and shrinks as the token count grows, so standard MMQ closes the gap on long prompts. Numbers are in the internal notes rather than here.

Numerics

This path is not bit-identical to standard MMQ. It quantizes activations to int8 with a coarser per-128-K absmax scale where q8_1 uses per-32, and the original source comment flagged it for validation. That validation has now been done, using the same binary with only the env flag differing, against a same-path control to establish the floor.

The deviation is small and the tail is bounded. Perplexity moves well inside its error bar, and a small fraction of tokens change their argmax. It is acceptable for the prefill gain, but it is a behaviour change rather than a pure optimization, which is why the opt-out exists and why the measured figures are recorded in the source comment for whoever reads this next.

Scope and what is not changed

• GGML_CUDA_HOPPER_Q1 still defaults off at configure time. Shipping this in release binaries is a separate CI question, since the build image would need CUTLASS. Happy to do that as a follow-up if we want it in prebuilt artifacts.
• Hardware gating is unchanged: Hopper only, explicitly excluding Ada and Blackwell, which lack wgmma and would need the tcgen05 path.
• Shape gating is unchanged: supported weight types, F32 activations, contiguous, and M, N and K all divisible by 128. Everything else falls through to standard MMQ as before.

Test plan

• Patched translation unit compiles clean at arch=compute_90a,code=[compute_90a,sm_90a] with CUTLASS includes.
• Path confirmed to fire by profiler kernel trace: the wgmma kernel, its activation quantizer and its repack all appear when enabled, and the standard MMQ variant disappears, so it displaces work rather than running alongside.
• test-backend-ops test -b CUDA0, zero failures in both arms.
• Perplexity and logit-KLD with a same-path control.
• Not yet re-run end to end after the gate inversion itself. The change is a two-line boolean inversion and the enabled path is byte-for-byte the code that was validated, but a confirming run on the built branch would be reasonable before merge.