cuda: AMD RDNA4 Q1_0/Q2_0 — HIP-path vec_dots (+37%/+16% decode), opt-in quant dedup, opt-in hipBLASLt prefill routes

ggml/src/ggml-cuda/common.cuh

@@ -868,6 +868,62 @@ static __device__ __forceinline__ uint8_t ggml_cuda_fp32_to_ue4m3(float x) {
 #endif // defined(BLACKWELL_MMA_AVAILABLE)
 }
 
+// Signed E4M3 (OCP e4m3fn) encoder. Portable closed-form bit manipulation, no
+// hardware fp8 instruction dependency -- usable on any device (unlike the
+// unsigned Blackwell-only scale encoder above). Used by the RDNA4 Q1_0/Q2_0
+// hipBLASLt prefill routes to requantize weights to e4m3 for the fp8 GEMM
+// variant (mul_mat_q{1,2}_0_hipblaslt.cu).
+static __device__ __forceinline__ uint8_t ggml_cuda_fp32_to_e4m3(float x) {
+    uint32_t bits;
+    memcpy(&bits, &x, 4);
+    const int sign = (bits >> 31) & 1;
+
+    if (x != x) { // NaN in -> NaN out
+        return (uint8_t) ((sign << 7) | 0x7F);
+    }
+
+    float ax = fabsf(x);
+    if (ax > 448.0f) {
+        ax = 448.0f; // clamp, e4m3fn has no infinity
+    }
+    if (!(ax > 0.0f)) {
+        return (uint8_t) (sign << 7); // +-0
+    }
+
+    memcpy(&bits, &ax, 4);
+    int fp32_exp = ((bits >> 23) & 0xFF) - 127;
+    int fp32_man = (bits >> 20) & 0x7;
+    int e4_exp   = fp32_exp + 7;
+
+    if (e4_exp <= 0) {
+        // subnormal: value = man * 2^-9, man = round(ax * 2^9)
+        int man = (int) (ax * 512.0f + 0.5f);
+        if (man > 7) {
+            man = 7;
+        }
+        if (man < 1) {
+            return (uint8_t) (sign << 7);
+        }
+        return (uint8_t) ((sign << 7) | man);
+    }
+
+    const int round_bit = (bits >> 19) & 1;
+    int e4_man = fp32_man + round_bit;
+    if (e4_man > 7) {
+        e4_man = 0;
+        e4_exp++;
+    }
+    if (e4_exp >= 15 && e4_man == 7) {
+        // never emit the NaN pattern (S.1111.111) from rounding: clamp to max finite
+        e4_exp = 15;
+        e4_man = 6;
+    } else if (e4_exp > 15) {
+        e4_exp = 15;
+        e4_man = 6;
+    }
+    return (uint8_t) ((sign << 7) | (e4_exp << 3) | e4_man);
+}
+
 __device__ __forceinline__ uint8_t ggml_cuda_float_to_fp4_e2m1(float x, float e) {
     const uint8_t sign_bit = (x < 0.0f) << 3;
     float         ax       = fabsf(x) * e;
@@ -1405,6 +1461,15 @@ struct ggml_backend_cuda_context {
 
     int curr_stream_no = 0;
 
+    // Opt-in (GGML_HIP_DEDUP_MMVQ_QUANT) sibling-matmul activation-quant dedup
+    // cache: the last src1 tensor whose q8_1 quantization was computed, and the
+    // buffer holding it. Keyed by tensor pointer, valid only within one graph
+    // evaluation -- reset at every graph_compute (ggml-cuda.cu) so an entry can
+    // never survive into a graph where the pointer identifies a different
+    // tensor. See mmvq.cu for the full rationale.
+    const ggml_tensor * mmvq_quant_cache_tensor = nullptr;
+    std::unique_ptr<ggml_cuda_pool_alloc<char>> mmvq_quant_cache_buf;
+
 #ifdef USE_CUDA_GRAPH
     // Map from first_node_ptr to cuda_graph - allows multiple graphs per context
     // when the computation is split across CPU/GPU (e.g., with --n-cpu-moe)

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

@@ -29,6 +29,9 @@
 #include "ggml-cuda/im2col.cuh"
 #include "ggml-cuda/mmf.cuh"
 #include "ggml-cuda/mmq.cuh"
+#include "ggml-cuda/hipblaslt_wcache.cuh"
+#include "ggml-cuda/mul_mat_q2_0_hipblaslt.cuh"
+#include "ggml-cuda/mul_mat_q1_0_hipblaslt.cuh"
 #include "ggml-cuda/mmvf.cuh"
 #include "ggml-cuda/mmvq.cuh"
 #include "ggml-cuda/norm.cuh"
@@ -639,6 +642,10 @@ struct ggml_backend_cuda_buffer_context {
 
 static void ggml_backend_cuda_buffer_free_buffer(ggml_backend_buffer_t buffer) {
     ggml_backend_cuda_buffer_context * ctx = (ggml_backend_cuda_buffer_context *)buffer->context;
+    // The hipBLASLt prefill routes key their converted-weight caches on device
+    // addresses inside this buffer; drop those entries before the address range
+    // can be reused by a later allocation (see hipblaslt_wcache.cuh).
+    ggml_hipblaslt_wcache_invalidate(ctx->dev_ptr, buffer->size);
     delete ctx;
 }
 
@@ -2587,6 +2594,27 @@ static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor
         any_gpus_with_slow_fp16 = any_gpus_with_slow_fp16   || !fast_fp16_hardware_available(cc);
     }
 
+    // Opt-in RDNA4 prefill levers (see mul_mat_q2_0_hipblaslt.cuh): route Q2_0 /
+    // Q1_0 large-M (prefill) matmuls through AMD's tuned hipBLASLt int8 GEMM
+    // instead of the dp4a/mmq path. W8A8, per-channel weight scale. Both are
+    // env-gated OFF by default; a soft-fail (false) falls through to the
+    // unmodified paths below, and non-HIP builds compile the routes to stubs
+    // that always return false.
+    if (!split) {
+        static const bool q2_0_hipblaslt_prefill_enabled = (getenv("GGML_HIP_Q2_0_HIPBLASLT_PREFILL") != nullptr);
+        if (q2_0_hipblaslt_prefill_enabled && ggml_cuda_q2_0_hipblaslt_prefill_supports(src0, src1, dst)) {
+            if (ggml_cuda_op_mul_mat_q2_0_hipblaslt(ctx, src0, src1, dst)) {
+                return;
+            }
+        }
+        static const bool q1_0_hipblaslt_prefill_enabled = (getenv("GGML_HIP_Q1_0_HIPBLASLT_PREFILL") != nullptr);
+        if (q1_0_hipblaslt_prefill_enabled && ggml_cuda_q1_0_hipblaslt_prefill_supports(src0, src1, dst)) {
+            if (ggml_cuda_op_mul_mat_q1_0_hipblaslt(ctx, src0, src1, dst)) {
+                return;
+            }
+        }
+    }
+
     // debug helpers
     //printf("src0: %8d %8d %8d %8d\n", src0->ne[0], src0->ne[1], src0->ne[2], src0->ne[3]);
     //printf("      %8d %8d %8d %8d\n", src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3]);
@@ -4471,6 +4499,12 @@ static enum ggml_status ggml_backend_cuda_graph_compute(ggml_backend_t backend,
 
     ggml_cuda_set_device(cuda_ctx->device);
 
+    // Reset the opt-in mmvq activation-quant dedup cache at every graph build:
+    // it is keyed by tensor pointer, which is only meaningful within one graph
+    // evaluation (see common.cuh / mmvq.cu).
+    cuda_ctx->mmvq_quant_cache_tensor = nullptr;
+    cuda_ctx->mmvq_quant_cache_buf.reset();
+
     bool use_cuda_graph             = false;
     bool cuda_graph_update_required = false;
     const void * graph_key = nullptr;

ggml/src/ggml-cuda/hipblaslt_wcache.cu

@@ -0,0 +1,23 @@
+#include "hipblaslt_wcache.cuh"
+#include <mutex>
+#include <vector>
+
+namespace {
+std::vector<ggml_hipblaslt_wcache_invalidator> & registry() {
+    static std::vector<ggml_hipblaslt_wcache_invalidator> v;   // function-local: no static-init order issue
+    return v;
+}
+std::mutex & registry_mtx() { static std::mutex m; return m; }
+}
+
+void ggml_hipblaslt_wcache_register(ggml_hipblaslt_wcache_invalidator fn) {
+    if (!fn) return;
+    std::lock_guard<std::mutex> lk(registry_mtx());
+    registry().push_back(fn);
+}
+
+void ggml_hipblaslt_wcache_invalidate(const void * base, size_t size) {
+    if (!base || size == 0) return;
+    std::lock_guard<std::mutex> lk(registry_mtx());
+    for (auto fn : registry()) fn(base, size);
+}

ggml/src/ggml-cuda/hipblaslt_wcache.cuh

@@ -0,0 +1,22 @@
+#pragma once
+#include <cstddef>
+
+// Invalidation registry for the hipBLASLt per-route converted-weight caches.
+//
+// Each mul_mat_*_hipblaslt.cu keeps a cache of weights it has already converted
+// (requantised to int8 or e4m3), keyed on the weight tensor's device address.
+// That key is only unique while the underlying buffer is alive: once a buffer is
+// freed, a later allocation can land on the same address and would silently hit
+// a stale entry. That is not hypothetical -- it produces garbage output on
+// multi-model runs and on llama-server model swaps.
+//
+// So every route registers an invalidator here, and the CUDA backend calls
+// ggml_hipblaslt_wcache_invalidate() when it frees a device buffer. Cost is
+// zero on the hot path; the work happens only at teardown.
+
+typedef void (*ggml_hipblaslt_wcache_invalidator)(const void * base, size_t size);
+
+void ggml_hipblaslt_wcache_register(ggml_hipblaslt_wcache_invalidator fn);
+
+// Drop every cached entry whose weight pointer lies in [base, base+size).
+void ggml_hipblaslt_wcache_invalidate(const void * base, size_t size);

ggml/src/ggml-cuda/mmvq.cu

@@ -4,6 +4,36 @@
 #include "vecdotq.cuh"
 
 #include <cstdint>
+#include <memory>
+
+// GGML_HIP_DEDUP_MMVQ_QUANT (default: unset/OFF): dedup the quantize_row_q8_1
+// dispatch across sibling mmvq matmuls that share one input tensor. Sibling
+// matmuls reading the exact same activation tensor (wq/wk/wv, ffn_gate/ffn_up,
+// wqkv/wqkv_gate within one decoder layer) each launch a quantize that
+// recomputes byte-identical output; this lever keeps quantize as a separate
+// launch (the packed dp4a decode path is untouched) and only skips launches
+// whose result is already cached for the same src1 tensor. The cache lives on
+// the backend context and is reset at every graph_compute, so entries can
+// never survive a graph rebuild (see ggml-cuda.cu). Measured on RDNA4
+// (2x R9700): Bonsai-27B Q2_0 MTP decode 51 -> 67 t/s; lossless, the cached
+// bytes are the same bytes the skipped launch would have produced.
+static bool ggml_cuda_dedup_mmvq_quant_enabled() {
+    static const bool enabled = getenv("GGML_HIP_DEDUP_MMVQ_QUANT") != nullptr;
+    return enabled;
+}
+
+// GGML_HIP_DEDUP_MMVQ_QUANT_BATCH (default OFF, REQUIRES the base flag too):
+// widen the sibling quant-dedup above to ncols_dst>1 (ne11>1) -- i.e. a
+// spec-decode verify pass batching n_draft+1 candidate tokens through the
+// target in one forward, not just the ne11==1 single-token decode path the
+// base flag is scoped to. The underlying redundancy is identical at any batch
+// size: sibling matmuls still read the exact same activation tensor whether
+// it holds 1 token or N. Kept as a separate flag so the single-stream win and
+// the verify-pass extension can be A/B'd independently.
+static bool ggml_cuda_dedup_mmvq_quant_batch_enabled() {
+    static const bool enabled = getenv("GGML_HIP_DEDUP_MMVQ_QUANT_BATCH") != nullptr;
+    return enabled;
+}
 
 typedef float (*vec_dot_q_cuda_t)(const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs);
 
@@ -1192,12 +1222,32 @@ void ggml_cuda_mul_mat_vec_q(
     }
 
     const int64_t ne10_padded = GGML_PAD(ne10, MATRIX_ROW_PADDING);
-    ggml_cuda_pool_alloc<char> src1_q8_1(ctx.pool(), ne13*ne12 * ne11*ne10_padded * sizeof(block_q8_1)/QK8_1);
-    {
+
+    // Opt-in sibling-matmul activation-quant dedup (see the flag doc at the top
+    // of this file). MUL_MAT_ID is excluded: the expert path reorders rows.
+    const bool dedup_quant_batch_ok = ne11 == 1 || (ne11 > 1 && ggml_cuda_dedup_mmvq_quant_batch_enabled());
+    const bool dedup_quant = ggml_cuda_dedup_mmvq_quant_enabled() && !ids && dedup_quant_batch_ok;
+    const bool dedup_hit   = dedup_quant &&
+        ctx.mmvq_quant_cache_tensor == src1 && ctx.mmvq_quant_cache_buf;
+
+    // dedup_quant (hit OR miss-that-populates) ALWAYS routes data through
+    // ctx.mmvq_quant_cache_buf, never through the local src1_q8_1 -- so the
+    // local buffer must be skipped (size 0) whenever dedup_quant is true, not
+    // just on a hit: gating on dedup_hit alone would leave the miss
+    // occurrence's vy pointer aimed at an allocated-but-never-written local
+    // buffer while the real quantized data went into the cache.
+    ggml_cuda_pool_alloc<char> src1_q8_1(ctx.pool(), dedup_quant ? 0 : ne13*ne12 * ne11*ne10_padded * sizeof(block_q8_1)/QK8_1);
+    if (!dedup_hit) {
         const int64_t s11 = src1->nb[1] / ts_src1;
         const int64_t s12 = src1->nb[2] / ts_src1;
         const int64_t s13 = src1->nb[3] / ts_src1;
-        quantize_row_q8_1_cuda(src1_d, nullptr, src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, ne11, ne12, ne13, stream);
+        if (dedup_quant) {
+            ctx.mmvq_quant_cache_buf = std::make_unique<ggml_cuda_pool_alloc<char>>(ctx.pool(), ne13*ne12 * ne11*ne10_padded * sizeof(block_q8_1)/QK8_1);
+            ctx.mmvq_quant_cache_tensor = src1;
+            quantize_row_q8_1_cuda(src1_d, nullptr, ctx.mmvq_quant_cache_buf->get(), src0->type, ne10, s11, s12, s13, ne10_padded, ne11, ne12, ne13, stream);
+        } else {
+            quantize_row_q8_1_cuda(src1_d, nullptr, src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, ne11, ne12, ne13, stream);
+        }
     }
 
     const int64_t s01 = src0->nb[1] / ts_src0;
@@ -1222,8 +1272,11 @@ void ggml_cuda_mul_mat_vec_q(
 
     const int64_t ids_stride = ids ? ids->nb[1] / ggml_type_size(ids->type) : 0;
 
+    const void * vy_ptr = dedup_quant ? (const void *) ctx.mmvq_quant_cache_buf->get()
+                                      : (const void *) src1_q8_1.get();
+
     mul_mat_vec_q_switch_type(
-        src0->data, src0->type, src1_q8_1.get(), ids_d, fusion_local, dst_d, ne00,
+        src0->data, src0->type, vy_ptr, ids_d, fusion_local, dst_d, ne00,
         ne01,              ncols_dst,     s01, stride_col_y,     stride_col_dst,
         ne02, nchannels_y, nchannels_dst, s02, stride_channel_y, stride_channel_dst,
         ne03,              ne3,           s03, s13,              s3,               ids_stride, stream);