ggml : add SM75 prefill experiment controls

Assisted-by: Codex

docs/sm75-ninfer-prefill-experiments.md

@@ -0,0 +1,118 @@
+# SM75 prefill experiments
+
+## Base and scope
+
+Base: `xrip/llama.cpp-avx1-numa-sm75`, public `master`, commit `af084752abb72652d51785602ef9536f631fec9d`.
+
+Idea source: `SirDebugALot/ninfer-3090`, commit `0d0ad44d8dac9c08c2e28b7c8c86527e10ba30a4`.
+
+This is a small experiment patch, not a port of the SM86 kernels and not a measured speed improvement. It uses the fork's existing cuBLAS path and MMQ specializations. It adds no BF16 MMA, asynchronous copy, new weight format, or persistent unpacked model copy.
+
+Important finding: the base already has a cuBLAS override for four FFN shapes in `ggml-cuda.cu`, before `ggml_cuda_should_use_mmq()`. The default mode below preserves that override. Do not describe the base as always using MMQ on SM75.
+
+## Controls
+
+Environment variables are read once per process. Set them before starting the server or benchmark. Restart the process for each case. CUDA Graphs can retain the selected kernels.
+
+| Control | Default | Meaning |
+| --- | --- | --- |
+| `GGML_CUDA_SM75_PREFILL_MODE` | `auto` | `auto`, `mmq`, or `cublas`; only the dense SM75 prefill override |
+| `GGML_CUDA_SM75_PREFILL_MIN_BATCH` | `1024` | Minimum N in experimental `cublas` mode; integer from 512 to 1048576 |
+| `GGML_CUDA_SM75_PREFILL_MAX_SCRATCH_MIB` | `512` | Conversion-scratch estimate cap in experimental `cublas` mode; integer from 1 to 65536 |
+| `GGML_CUDA_SM75_MMQ_MAX_J` | `0` | Limit the search over existing MMQ token tiles: 0, 32, 64, 96, or 128; 0 means unchanged |
+| `GGML_CUDA_SM75_PREFILL_TRACE` | `0` | 1 logs dispatch decisions and selected MMQ tiles; turn it off for timing |
+| CMake `GGML_CUDA_SM75_Q8_0_FAST_HALF` | `OFF` | Compile the experimental Q8 V to FP16 conversion for the SM75 device target |
+
+Invalid values produce a warning and use that setting's default. An empty value is invalid; use `unset` to remove a variable.
+
+### Dense cuBLAS policy
+
+All modes retain the existing checks for SM75, F32 activations/output, contiguous inputs, device-resident weights, and single-channel/single-sample matrices. The earlier MMVF/MMF/MMVQ decisions are unchanged. `MUL_MAT_ID` is not modified.
+
+`auto` is the exact old override: Q4_K or IQ4_XS; K/M equal to 4096/12288, 12288/4096, 5120/17408, or 17408/5120; N between 1024 and 2048 inclusive. The new scratch cap and minimum-batch setting do not alter this baseline mode.
+
+`mmq` disables this one FFN cuBLAS override and falls through to the existing backend selector. It is not a global force flag and cannot override `GGML_CUDA_FORCE_CUBLAS` or an earlier fallback.
+
+`cublas` allows Q4_K, Q5_K, Q6_K, and IQ4_XS outside the old shape list. It requires K >= 256 and divisible by 256, M >= 128 and divisible by 128, N >= the selected minimum and divisible by 128, and a fitting scratch estimate. Other shapes fall through to the existing selector. Small batches below 512 do not enter this experiment.
+
+The estimate is `4 * (M*K + K*N + M*N)` bytes, using the dimensions passed to this CUDA operation. This conservatively includes FP32 conversion overrides. Products are checked before multiplication. For example, K=5120, M=17408, N=2048 gives an estimate of 516 MiB, so the default 512 MiB cap rejects it. Raising the cap to 768 MiB admits that shape, but requires separate VRAM validation.
+
+This is not a total-VRAM guarantee. It excludes cuBLAS internal workspace, allocator cache, model/KV storage, other live tensors, and concurrent operations. No memory availability probe or OOM recovery is added.
+
+Compile-time force flags keep their existing priority. Existing fusions which call MMQ directly are not rerouted by this policy. Use trace output or a profiler to confirm which operations actually changed. Do not interpret the absence of a speed change as a measurement of a path that did not run.
+
+### MMQ tile experiment
+
+The cap applies only to SM75 Q4_K/IQ4_XS dense operations with no expert IDs, one channel/sample, complete 128-row tiles, and N >= 512 divisible by 128. All other operations retain the normal search up to J=128.
+
+The patch changes only the host search range. It keeps the original configuration-validity and shared-memory checks. If no valid configuration fits the cap, it repeats the original search. It does not change I, thread counts, launch bounds, weight layout, Stream-K, fixup logic, or buffer padding. `ggml_cuda_mmq_get_J_max()` is deliberately unchanged, so a smaller launch does not shrink a buffer needed by another path.
+
+The cap is not a promise to use exactly that width. It selects the best candidate from the existing configurations up to that width. This is not a new SM75-tuned configuration table or an autotuner.
+
+### Q8 V conversion experiment
+
+The helper constructs two FP16 bit patterns for `1152 + q`, then subtracts 1152 with `__hsub2_rn`. It covers every signed INT8 code, including -128. The existing half2 scale multiplication stays separate, with the same operands and rounding stage as before.
+
+The existing Q8 loads, block offsets, lengths, and scale are unchanged. No wider memory access is introduced. Only `dequantize_V_q8_0<half, ...>` compiled for `__CUDA_ARCH__ == 750` changes. F32 conversion and other device targets keep the old code. There is no runtime flag for this helper; compare two builds.
+
+F16 V/KV caches do not use this conversion. Q8 storage alone is not proof that this helper runs: another attention route may perform its own conversion. This helper can affect both prefill and decode when that route is used.
+
+## Build and apply
+
+This is a unified diff for `git apply`, not a commit for `git am`. It does not change the branch or publish anything. A local branch or uncommitted edits may differ from the public base; check before applying. Do not use `--reject` or force a partially applicable patch.
+
+```sh
+git rev-parse HEAD
+git apply --check /path/to/sm75-ninfer-prefill-af084752.patch && \
+    git apply /path/to/sm75-ninfer-prefill-af084752.patch
+```
+
+Reconfigure and rebuild the CUDA backend. Preserve the existing AVX1, NUMA, compiler, and build settings. For an existing build directory, these are the relevant additions:
+
+```sh
+cmake -S . -B build -DGGML_CUDA=ON -DCMAKE_CUDA_ARCHITECTURES=75 \
+    -DGGML_CUDA_SM75_Q8_0_FAST_HALF=OFF
+cmake --build build --parallel 4
+```
+
+Build a second directory with the same options and `-DGGML_CUDA_SM75_Q8_0_FAST_HALF=ON` for the Q8 experiment. Do not compare two builds with different compiler versions or unrelated CPU/CUDA flags.
+
+## Measurement plan
+
+Use the same model file, context, prompt, microbatch, GPU split, NUMA policy, clocks, power limit, MTP, KV types, and CUDA Graph settings in every A/B pair. Do not enable all experiments at once.
+
+1. Compare the unpatched build with patched defaults: `auto`, J=0, Q8 helper OFF. This is the baseline-regression check.
+2. Compare `auto`, `mmq`, and `cublas` with all other controls unchanged. Use N=512/1024/2048; N=512 requires setting the cublas minimum to 512.
+3. Set mode `mmq` and compare J=0/32/64/96/128 on Q4_K/IQ4_XS. This avoids the old FFN override hiding the MMQ experiment.
+4. Separately compare Q8-helper OFF/ON with Q8 V and the same dispatcher. Also test F16 KV as an unaffected control.
+
+Example environment settings (then run the existing full benchmark/server command):
+
+```sh
+# Baseline
+unset GGML_CUDA_SM75_PREFILL_MODE GGML_CUDA_SM75_PREFILL_MIN_BATCH
+unset GGML_CUDA_SM75_PREFILL_MAX_SCRATCH_MIB GGML_CUDA_SM75_MMQ_MAX_J
+unset GGML_CUDA_SM75_PREFILL_TRACE
+
+# cuBLAS experiment, default 512 MiB estimate cap
+export GGML_CUDA_SM75_PREFILL_MODE=cublas
+export GGML_CUDA_SM75_PREFILL_MIN_BATCH=1024
+
+# Separate MMQ experiment
+export GGML_CUDA_SM75_PREFILL_MODE=mmq
+export GGML_CUDA_SM75_MMQ_MAX_J=64
+```
+
+Trace is only for a diagnostic run. It logs every relevant dispatch and can distort timing. A reused CUDA Graph may not repeat host-side dispatch logging.
+
+Check real backend numerical results, logits/perplexity, short decode, MTP acceptance, non-aligned tails, CUDA Graph replay, and each GPU separately. Then test the two-GPU configuration. A different MMQ reduction order or the existing cuBLAS compute policy can change floating-point results; require an agreed accuracy tolerance rather than identical generated text.
+
+Record individual operation timings, full prefill timings, decode timings, and peak VRAM per GPU. Warm up each case, then repeat A/B/B/A. Keep raw commands and logs. There are no GPU measurements in this patch.
+
+## Validation status at delivery
+
+Performed: C++ host tests with GCC and Clang, GCC AddressSanitizer/UndefinedBehaviorSanitizer, strict environment parsing checks, exact old-policy comparisons over a shape grid, all 65,536 INT8 pairs, and a CPU IEEE binary16 reference check for all 256 codes multiplied by all 63,488 finite FP16 scales.
+
+Not performed: a full repository build, nvcc compilation, actual CUDA execution, backend accuracy tests, long-context or two-GPU tests, performance or VRAM measurements. No speedup is claimed.
+
+The separate delivery archive contains standalone host/CUDA tests and the exact validation report. These tests are not added to the repository's main test suite or CI. The CUDA test must still be compiled and run on SM75; a skipped GPU test is not a pass.

ggml/src/ggml-cuda/CMakeLists.txt

@@ -1,5 +1,7 @@
 cmake_minimum_required(VERSION 3.18)  # for CMAKE_CUDA_ARCHITECTURES
 
+option(GGML_CUDA_SM75_Q8_0_FAST_HALF "SM75: experimental packed Q8_0 to FP16 conversion in FlashAttention" OFF)
+
 find_package(CUDAToolkit)
 
 if (CUDAToolkit_FOUND)
@@ -120,6 +122,10 @@ if (CUDAToolkit_FOUND)
                              ${GGML_SOURCES_CUDA}
                             )
 
+    if (GGML_CUDA_SM75_Q8_0_FAST_HALF)
+        target_compile_definitions(ggml-cuda PRIVATE GGML_CUDA_SM75_Q8_0_FAST_HALF)
+    endif()
+
     if (GGML_CUDA_GRAPHS)
         add_compile_definitions(GGML_CUDA_USE_GRAPHS)
     endif()

ggml/src/ggml-cuda/fattn-common.cuh

@@ -3,6 +3,9 @@
 #include "common.cuh"
 #include "convert.cuh"
 #include "vecdotq.cuh"
+#ifdef GGML_CUDA_SM75_Q8_0_FAST_HALF
+#include "sm75-q8.cuh"
+#endif
 
 #include <cstdint>
 
@@ -601,7 +604,11 @@ static __device__ __forceinline__ void dequantize_V_q8_0(const void * __restrict
 
 #pragma unroll
         for (int l0 = 0; l0 < ne; l0 += 2) {
+#if defined(GGML_CUDA_SM75_Q8_0_FAST_HALF) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ == 750 && !defined(GGML_USE_HIP)
+            ((half2 *) dst)[l0/2] = d * ggml_cuda_sm75_q8_pair_to_half2(qs[l0 + 0], qs[l0 + 1]);
+#else
             ((half2 *) dst)[l0/2] = d * make_half2(qs[l0 + 0], qs[l0 + 1]);
+#endif
         }
     } else
 #endif // FP16_AVAILABLE

ggml/src/ggml-cuda/ggml-cuda.cu

@@ -31,6 +31,7 @@
 #include "ggml-cuda/im2col.cuh"
 #include "ggml-cuda/mmf.cuh"
 #include "ggml-cuda/mmq.cuh"
+#include "ggml-cuda/sm75-tuning.cuh"
 #include "ggml-cuda/mmvf.cuh"
 #include "ggml-cuda/mmvq.cuh"
 #include "ggml-cuda/moe-weighted-reduction.cuh"
@@ -1931,16 +1932,25 @@ static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor
         return;
     }
 #ifndef GGML_CUDA_FORCE_MMQ
-    // Bound FP16 weight workspace and keep small batches on MMQ.
-    const bool sm75_ffn = (ne00 == 4096 && ne01 == 12288) || (ne00 == 12288 && ne01 == 4096)
-                      || (ne00 == 5120 && ne01 == 17408) || (ne00 == 17408 && ne01 == 5120);
-    if (cc == GGML_CUDA_CC_TURING && sm75_ffn && ne11 >= 1024 && ne11 <= 2048
-            && (src0->type == GGML_TYPE_Q4_K || src0->type == GGML_TYPE_IQ4_XS)
+    // Preserve the existing FFN override in auto mode. Experimental modes are
+    // restricted to dense prefill; MMVQ/decode and MUL_MAT_ID are untouched.
+    const bool sm75_legacy_type = src0->type == GGML_TYPE_Q4_K || src0->type == GGML_TYPE_IQ4_XS;
+    const bool sm75_test_type = sm75_legacy_type || src0->type == GGML_TYPE_Q5_K || src0->type == GGML_TYPE_Q6_K;
+    if (cc == GGML_CUDA_CC_TURING && sm75_test_type && ne11 >= 512
             && ne02 == 1 && ne03 == 1 && ne12 == 1 && ne13 == 1
             && ggml_is_contiguous(src0) && ggml_is_contiguous(src1)
             && !ggml_backend_buffer_is_host(src0->buffer)) {
-        ggml_cuda_mul_mat_cublas(ctx, src0, src1, dst);
-        return;
+        const auto & tuning = ggml_cuda_sm75::get_options();
+        const bool use_cublas = ggml_cuda_sm75::use_cublas(tuning, sm75_legacy_type, sm75_test_type, ne00, ne01, ne11);
+        if (tuning.trace) {
+            GGML_LOG_INFO("SM75 prefill: device=%d type=%s M=%lld K=%lld N=%lld mode=%s override=%s\n",
+                    ctx.device, ggml_type_name(src0->type), (long long) ne01, (long long) ne00, (long long) ne11,
+                    ggml_cuda_sm75::mode_name(tuning.mode), use_cublas ? "cublas" : "none");
+        }
+        if (use_cublas) {
+            ggml_cuda_mul_mat_cublas(ctx, src0, src1, dst);
+            return;
+        }
     }
 #endif
     if (ggml_cuda_should_use_mmq(src0->type, cc, ne11, /*n_experts =*/ 0)) {

ggml/src/ggml-cuda/mmq.cuh

@@ -1,6 +1,7 @@
 #pragma once
 
 #include "common.cuh"
+#include "sm75-tuning.cuh"
 
 #include <climits>
 #include <cstdint>
@@ -1480,25 +1481,42 @@ void mul_mat_q_switch_J(ggml_backend_cuda_context & ctx, const mmq_args & args,
     const int    cc    = ggml_cuda_info().devices[id].cc;
     const size_t smpbo = ggml_cuda_info().devices[id].smpbo;
 
-    int J_best        = 0;
-    int ntiles_J_best = INT_MAX;
+    const bool sm75_dense = cc == GGML_CUDA_CC_TURING && !fallback &&
+        (type == GGML_TYPE_Q4_K || type == GGML_TYPE_IQ4_XS) &&
+        !args.ids_dst && !args.expert_bounds && args.nchannels_x == 1 && args.nchannels_y == 1 &&
+        args.nsamples_x == 1 && args.nsamples_y == 1 && args.ncols_y >= 512 && args.ncols_y % 128 == 0;
+    const auto * tuning = sm75_dense ? &ggml_cuda_sm75::get_options() : nullptr;
+    const int J_limit = tuning ? ggml_cuda_sm75::mmq_j_limit(*tuning, true, args.ncols_y) : 128;
 
-    for (int J = 8; J <= 128 && ntiles_J_best > 1; J += 8) {
-        const ggml_cuda_mmq_config config = ggml_cuda_mmq_get_config(type, J, fallback, cc);
-        if (config.type == GGML_TYPE_COUNT) {
-            continue;
+    const auto select_J = [&](const int limit) {
+        int best = 0;
+        int ntiles_best = INT_MAX;
+        for (int J = 8; J <= limit && ntiles_best > 1; J += 8) {
+            const ggml_cuda_mmq_config config = ggml_cuda_mmq_get_config(type, J, fallback, cc);
+            if (config.type == GGML_TYPE_COUNT || mmq_get_nbytes_shared(config, cc) > smpbo) {
+                continue;
+            }
+            const int ntiles_x = (args.ncols_opt + config.J - 1) / config.J;
+            if (ntiles_x < ntiles_best) {
+                best = J;
+                ntiles_best = ntiles_x;
+            }
         }
+        return best;
+    };
 
-        if (mmq_get_nbytes_shared(config, cc) > smpbo) {
-            continue;
-        }
-
-        const int ntiles_x = (args.ncols_opt + config.J - 1) / config.J;
-
-        if (ntiles_x < ntiles_J_best) {
-            J_best = J;
-            ntiles_J_best = ntiles_x;
-        }
+    int J_best = select_J(J_limit);
+    if (J_best == 0 && J_limit < 128) {
+        // A cap is a tuning hint, not permission to launch an invalid kernel.
+        J_best = select_J(128);
+    }
+    // Keep ggml_cuda_mmq_get_J_max and all padding allocations unchanged: a
+    // narrower launch must not shrink buffers needed by another dispatch path.
+    if (tuning && tuning->trace && J_best != 0) {
+        GGML_LOG_INFO("SM75 MMQ: device=%d type=%s M=%lld K=%lld N=%lld J=%d cap=%d shared=%zu\n",
+                id, ggml_type_name(type), (long long) args.nrows_x, (long long) args.ncols_x,
+                (long long) args.ncols_y, J_best, J_limit,
+                mmq_get_nbytes_shared(ggml_cuda_mmq_get_config(type, J_best, fallback, cc), cc));
     }
 
     switch (J_best) {

ggml/src/ggml-cuda/sm75-q8.cuh

@@ -0,0 +1,30 @@
+#pragma once
+
+#include <cstdint>
+
+#if defined(__CUDACC__) && !defined(GGML_USE_HIP)
+#include <cuda_fp16.h>
+#define GGML_SM75_Q8_HD __host__ __device__
+#else
+#define GGML_SM75_Q8_HD
+#endif
+
+// An unsigned byte u in the mantissa of 0x6400 encodes 1024 + u exactly.
+// Flipping the INT8 sign bit gives u = q + 128. Subtracting 1152 then recovers
+// every q in [-128, 127], including -128. No float conversion or wide load.
+GGML_SM75_Q8_HD static constexpr uint32_t ggml_cuda_sm75_q8_pair_bits(uint8_t a, uint8_t b) {
+    return 0x64006400u | uint32_t(a ^ 0x80u) | (uint32_t(b ^ 0x80u) << 16);
+}
+
+#undef GGML_SM75_Q8_HD
+
+#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ == 750 && !defined(GGML_USE_HIP)
+static __device__ __forceinline__ half2 ggml_cuda_sm75_q8_pair_to_half2(int8_t a, int8_t b) {
+    const uint32_t bits = ggml_cuda_sm75_q8_pair_bits(uint8_t(a), uint8_t(b));
+    const half2 encoded = __halves2half2(__ushort_as_half(uint16_t(bits)),
+                                        __ushort_as_half(uint16_t(bits >> 16)));
+    // The explicit rounding intrinsic prevents a following scale multiply from
+    // contracting with the subtraction. Keep the same multiply as the old path.
+    return __hsub2_rn(encoded, __float2half2_rn(1152.0f));
+}
+#endif

ggml/src/ggml-cuda/sm75-tuning.cuh

@@ -0,0 +1,144 @@
+#pragma once
+
+// Opt-in host policy. Keep environment variables fixed while the process runs.
+// CUDA Graphs can retain the selected kernels.
+#include <cstdint>
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+
+namespace ggml_cuda_sm75 {
+
+enum class prefill_mode { automatic, mmq, cublas };
+
+struct options {
+    prefill_mode mode = prefill_mode::automatic;
+    int64_t min_batch = 1024;
+    uint64_t max_scratch_mib = 512;
+    int mmq_max_j = 0;
+    bool trace = false;
+};
+
+// Reject overflow, signs, whitespace, and trailing text.
+inline bool parse_uint(const char * text, uint64_t lo, uint64_t hi, uint64_t & value) {
+    if (!text || !*text || lo > hi) {
+        return false;
+    }
+    uint64_t n = 0;
+    for (const char * p = text; *p; ++p) {
+        if (*p < '0' || *p > '9') {
+            return false;
+        }
+        const uint64_t digit = uint64_t(*p - '0');
+        if (digit > hi || n > (hi - digit) / 10) {
+            return false;
+        }
+        n = 10*n + digit;
+    }
+    if (n < lo) {
+        return false;
+    }
+    value = n;
+    return true;
+}
+
+inline void warn_env(const char * name, const char * value) {
+    std::fprintf(stderr, "ggml_cuda_sm75: invalid %s='%s'; using its default\n", name, value ? value : "(unset)");
+}
+
+inline uint64_t env_uint(const char * name, uint64_t fallback, uint64_t lo, uint64_t hi) {
+    const char * text = std::getenv(name);
+    if (!text) {
+        return fallback;
+    }
+    uint64_t value;
+    if (!parse_uint(text, lo, hi, value)) {
+        warn_env(name, text);
+        return fallback;
+    }
+    return value;
+}
+
+inline options read_options() {
+    options o;
+    const char * mode = std::getenv("GGML_CUDA_SM75_PREFILL_MODE");
+    if (mode) {
+        if (std::strcmp(mode, "mmq") == 0) {
+            o.mode = prefill_mode::mmq;
+        } else if (std::strcmp(mode, "cublas") == 0) {
+            o.mode = prefill_mode::cublas;
+        } else if (std::strcmp(mode, "auto") != 0) {
+            warn_env("GGML_CUDA_SM75_PREFILL_MODE", mode);
+        }
+    }
+    o.min_batch = int64_t(env_uint("GGML_CUDA_SM75_PREFILL_MIN_BATCH", 1024, 512, 1048576));
+    o.max_scratch_mib = env_uint("GGML_CUDA_SM75_PREFILL_MAX_SCRATCH_MIB", 512, 1, 65536);
+    o.mmq_max_j = int(env_uint("GGML_CUDA_SM75_MMQ_MAX_J", 0, 0, 128));
+    if (o.mmq_max_j != 0 && o.mmq_max_j != 32 && o.mmq_max_j != 64 &&
+        o.mmq_max_j != 96 && o.mmq_max_j != 128) {
+        warn_env("GGML_CUDA_SM75_MMQ_MAX_J", std::getenv("GGML_CUDA_SM75_MMQ_MAX_J"));
+        o.mmq_max_j = 0;
+    }
+    o.trace = env_uint("GGML_CUDA_SM75_PREFILL_TRACE", 0, 0, 1) != 0;
+    return o;
+}
+
+inline const options & get_options() {
+    static const options o = read_options();
+    return o;
+}
+
+inline const char * mode_name(prefill_mode mode) {
+    switch (mode) {
+        case prefill_mode::automatic: return "auto";
+        case prefill_mode::mmq:       return "mmq";
+        case prefill_mode::cublas:    return "cublas";
+    }
+    return "auto";
+}
+
+inline bool legacy_ffn_shape(int64_t k, int64_t m) {
+    return (k == 4096 && m == 12288) || (k == 12288 && m == 4096) ||
+           (k == 5120 && m == 17408) || (k == 17408 && m == 5120);
+}
+
+// Bound conversion scratch by 4*(M*K + K*N + M*N), also covering FP32 overrides.
+// Excludes cuBLAS workspace and cached pool memory. Check products before multiplication.
+inline bool scratch_fits(int64_t k, int64_t m, int64_t n, uint64_t limit_mib) {
+    if (k <= 0 || m <= 0 || n <= 0 || limit_mib == 0 || limit_mib > 65536) {
+        return false;
+    }
+    uint64_t budget = limit_mib * (1024u * 1024u / 4u);
+    const uint64_t a[] = { uint64_t(m), uint64_t(k), uint64_t(m) };
+    const uint64_t b[] = { uint64_t(k), uint64_t(n), uint64_t(n) };
+    for (int i = 0; i < 3; ++i) {
+        if (a[i] > budget / b[i]) {
+            return false;
+        }
+        budget -= a[i]*b[i];
+    }
+    return true;
+}
+
+// Caller checks SM75, dense shapes, F32 I/O, contiguous inputs, and device weights.
+// legacy_type: Q4_K/IQ4_XS; experiment_type: also Q5_K/Q6_K.
+inline bool use_cublas(const options & o, bool legacy_type, bool experiment_type,
+                       int64_t k, int64_t m, int64_t n) {
+    if (o.mode == prefill_mode::automatic) {
+        // Exactly the pre-existing af084752 policy, including its tail behavior.
+        return legacy_type && legacy_ffn_shape(k, m) && n >= 1024 && n <= 2048;
+    }
+    if (o.mode == prefill_mode::mmq) {
+        return false;
+    }
+    return experiment_type && n >= o.min_batch && n >= 512 && n % 128 == 0 &&
+           k >= 256 && m >= 128 && k % 256 == 0 && m % 128 == 0 &&
+           scratch_fits(k, m, n, o.max_scratch_mib);
+}
+
+// Select existing specializations only; do not change layouts or buffer sizes.
+inline int mmq_j_limit(const options & o, bool eligible, int64_t n) {
+    return eligible && n >= 512 && n % 128 == 0 && o.mmq_max_j != 0 ? o.mmq_max_j : 128;
+}
+
+} // namespace ggml_cuda_sm75