From 4452cdbae810a5bdfd11dd5c5683a3206d469543 Mon Sep 17 00:00:00 2001 From: Justin Monk Date: Tue, 11 Aug 2026 00:48:11 -0400 Subject: [PATCH 1/3] =?UTF-8?q?cuda:=20AMD=20Q1=5F0/Q2=5F0=20=E2=80=94=20H?= =?UTF-8?q?IP-path=20Q2=5F0=20vec=5Fdot,=20opt-in=20quant=20dedup,=20opt-i?= =?UTF-8?q?n=20hipBLASLt=20prefill=20routes?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit HIP-guarded vec_dot_q2_0_q8_1 (dynamic-selector __byte_perm lowers poorly on AMD): Bonsai-27B Q2_0 decode 34.52 -> 47.20 t/s on gfx1201, PPL byte-identical. Opt-in GGML_HIP_DEDUP_MMVQ_QUANT skips redundant sibling activation quantize launches (+2-3% decode here, more under spec-decode verify). Opt-in GGML_HIP_{Q1_0,Q2_0}_HIPBLASLT_PREFILL routes large-M matmuls to hipBLASLt int8 GEMM with a self-tuning algo cache and buffer-free weight-cache invalidation: pp1024 +37.5% (Q2_0) / +17.8% (Q1_0) warm. hipBLASLt is an optional dependency; without it (Windows HIP) and on CUDA builds the routes compile to inert stubs. Co-Authored-By: Claude Fable 5 --- ggml/src/ggml-cuda/common.cuh | 65 ++ ggml/src/ggml-cuda/ggml-cuda.cu | 34 + ggml/src/ggml-cuda/hipblaslt_wcache.cu | 23 + ggml/src/ggml-cuda/hipblaslt_wcache.cuh | 22 + ggml/src/ggml-cuda/mmvq.cu | 61 +- ggml/src/ggml-cuda/mul_mat_q1_0_hipblaslt.cu | 621 ++++++++++++++++++ ggml/src/ggml-cuda/mul_mat_q1_0_hipblaslt.cuh | 17 + ggml/src/ggml-cuda/mul_mat_q2_0_hipblaslt.cu | 586 +++++++++++++++++ ggml/src/ggml-cuda/mul_mat_q2_0_hipblaslt.cuh | 30 + ggml/src/ggml-cuda/vecdotq.cuh | 36 +- ggml/src/ggml-hip/CMakeLists.txt | 13 + 11 files changed, 1501 insertions(+), 7 deletions(-) create mode 100644 ggml/src/ggml-cuda/hipblaslt_wcache.cu create mode 100644 ggml/src/ggml-cuda/hipblaslt_wcache.cuh create mode 100644 ggml/src/ggml-cuda/mul_mat_q1_0_hipblaslt.cu create mode 100644 ggml/src/ggml-cuda/mul_mat_q1_0_hipblaslt.cuh create mode 100644 ggml/src/ggml-cuda/mul_mat_q2_0_hipblaslt.cu create mode 100644 ggml/src/ggml-cuda/mul_mat_q2_0_hipblaslt.cuh diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 36f1d3cdabf..83ae0b5d76e 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/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> 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) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 13e1b8a2e73..401b2903658 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/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; diff --git a/ggml/src/ggml-cuda/hipblaslt_wcache.cu b/ggml/src/ggml-cuda/hipblaslt_wcache.cu new file mode 100644 index 00000000000..0424351f177 --- /dev/null +++ b/ggml/src/ggml-cuda/hipblaslt_wcache.cu @@ -0,0 +1,23 @@ +#include "hipblaslt_wcache.cuh" +#include +#include + +namespace { +std::vector & registry() { + static std::vector 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 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 lk(registry_mtx()); + for (auto fn : registry()) fn(base, size); +} diff --git a/ggml/src/ggml-cuda/hipblaslt_wcache.cuh b/ggml/src/ggml-cuda/hipblaslt_wcache.cuh new file mode 100644 index 00000000000..88c89630e65 --- /dev/null +++ b/ggml/src/ggml-cuda/hipblaslt_wcache.cuh @@ -0,0 +1,22 @@ +#pragma once +#include + +// 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); diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index d67abeac3c0..9c7a9628c70 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -4,6 +4,36 @@ #include "vecdotq.cuh" #include +#include + +// 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 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 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>(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); diff --git a/ggml/src/ggml-cuda/mul_mat_q1_0_hipblaslt.cu b/ggml/src/ggml-cuda/mul_mat_q1_0_hipblaslt.cu new file mode 100644 index 00000000000..f543d9f9d64 --- /dev/null +++ b/ggml/src/ggml-cuda/mul_mat_q1_0_hipblaslt.cu @@ -0,0 +1,621 @@ +// See mul_mat_q1_0_hipblaslt.cuh for the full rationale. +// Stage 1 of Q1_0 hipBLASLt route: wire the tuned hipBLASLt int8/fp8 GEMM +// as an opt-in prefill route for Q1_0. Correctness-first; algo/handle caching +// and the fp8 front-end are part of this implementation (both modes built, +// measurement decides which is best). Q1_0 block: {ggml_half d; uint8_t qs[QK1_0/8]}, +// QK1_0=128, 1 bit/weight. Dequant: bit=1 -> +d, bit=0 -> -d. + +#include "mul_mat_q1_0_hipblaslt.cuh" +#include "hipblaslt_wcache.cuh" + +// Real implementation only on the AMD/HIP build (hipBLASLt is a ROCm lib). +// On a CUDA build these become stubs so the globbed source is a no-op. +#if defined(__HIP_PLATFORM_AMD__) && !defined(GGML_HIP_NO_HIPBLASLT) + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace { + +constexpr int Q1K = QK1_0; // 128 elements per Q1_0 block +constexpr float I8MAX = 127.0f; + +#define LT_OK(x) do { hipblasStatus_t s_ = (x); if (s_ != HIPBLAS_STATUS_SUCCESS) { \ + GGML_LOG_ERROR("%s: hipBLASLt error %d at %s:%d\n", __func__, (int)s_, __FILE__, __LINE__); \ + return false; } } while(0) + +// ---- device helpers -------------------------------------------------------- + +__device__ __forceinline__ float q1_half2float(ggml_half h) { + return __half2float(*reinterpret_cast(&h)); +} + +// Requant a Q1_0 weight row -> int8 with ONE symmetric scale per output channel. +// One block per output row n. Weight row layout: n_blocks contiguous block_q1_0. +// Output int8 is row-major [N x K] (== col-major [K x N], the TN A-operand). +__global__ void k_requant_q1_0_to_int8_perchannel( + const char * __restrict__ wdata, int64_t nb01, + int8_t * __restrict__ q8, float * __restrict__ wscale, + int64_t K, int64_t n_blocks) { + const int64_t n = blockIdx.x; // output channel + const block_q1_0 * row = (const block_q1_0 *)(wdata + n * nb01); + + // pass 1: block-wide amax over the K weights of this row + float amax = 0.0f; + for (int64_t l = threadIdx.x; l < K; l += blockDim.x) { + const int64_t b = l / Q1K; + const int t = (int)(l % Q1K); + const int byte_index = t / 8; + const int bit_offset = t % 8; + const uint8_t bit = (row[b].qs[byte_index] >> bit_offset) & 1; + const float d = q1_half2float(row[b].d); + const float w = bit ? d : -d; + amax = fmaxf(amax, fabsf(w)); + } + __shared__ float sred[1024]; + sred[threadIdx.x] = amax; + __syncthreads(); + for (int s = blockDim.x >> 1; s > 0; s >>= 1) { + if (threadIdx.x < s) sred[threadIdx.x] = fmaxf(sred[threadIdx.x], sred[threadIdx.x + s]); + __syncthreads(); + } + const float scale = (sred[0] > 0.0f) ? sred[0] / I8MAX : 1.0f; + if (threadIdx.x == 0) wscale[n] = scale; + const float inv = 1.0f / scale; + + // pass 2: quantize + int8_t * out = q8 + n * K; + for (int64_t l = threadIdx.x; l < K; l += blockDim.x) { + const int64_t b = l / Q1K; + const int t = (int)(l % Q1K); + const int byte_index = t / 8; + const int bit_offset = t % 8; + const uint8_t bit = (row[b].qs[byte_index] >> bit_offset) & 1; + const float d = q1_half2float(row[b].d); + const float w = bit ? d : -d; + int v = __float2int_rn(w * inv); + v = max(-127, min(127, v)); + out[l] = (int8_t)v; + } +} + +// Quantize activations -> int8 with ONE symmetric scale per token (column). +// src1 fp32, column j at (src1 + j*nb1), K contiguous floats. Output col-major +// [K x M] (ld=K) + per-token scale. +__global__ void k_quantize_act_int8_percol( + const char * __restrict__ src1, int64_t nb1, + int8_t * __restrict__ x8, float * __restrict__ ascale, int64_t K) { + const int64_t j = blockIdx.x; // token / column + const float * col = (const float *)(src1 + j * nb1); + + float amax = 0.0f; + for (int64_t l = threadIdx.x; l < K; l += blockDim.x) amax = fmaxf(amax, fabsf(col[l])); + __shared__ float sred[1024]; + sred[threadIdx.x] = amax; + __syncthreads(); + for (int s = blockDim.x >> 1; s > 0; s >>= 1) { + if (threadIdx.x < s) sred[threadIdx.x] = fmaxf(sred[threadIdx.x], sred[threadIdx.x + s]); + __syncthreads(); + } + const float scale = (sred[0] > 0.0f) ? sred[0] / I8MAX : 1.0f; + if (threadIdx.x == 0) ascale[j] = scale; + const float inv = 1.0f / scale; + + int8_t * out = x8 + j * K; + for (int64_t l = threadIdx.x; l < K; l += blockDim.x) { + int v = __float2int_rn(col[l] * inv); + v = max(-127, min(127, v)); + out[l] = (int8_t)v; + } +} + +// Dequant the int32 GEMM result: dst[n,j] = i32[n,j] * wscale[n] * ascale[j]. +// Both dst and i32 are col-major [N x M] (dst ld from nb1, i32 ld=N). +__global__ void k_apply_scales( + const int32_t * __restrict__ i32, char * __restrict__ dst, int64_t nb1, + const float * __restrict__ wscale, const float * __restrict__ ascale, + int64_t N, int64_t M) { + const int64_t idx = blockIdx.x * (int64_t)blockDim.x + threadIdx.x; + if (idx >= N * M) return; + const int64_t n = idx % N; + const int64_t j = idx / N; + float * out = (float *)(dst + j * nb1); + out[n] = (float)i32[j * N + n] * wscale[n] * ascale[j]; +} + +// ============================ fp8 (E4M3) variant ============================ +// Same v1 scheme, but encode to OCP e4m3 and use hipBLASLt's fp8 GEMM (202 +// tuned gfx1201 kernels vs 16 int8). e4m3 max = 448; scale to fill the range. +constexpr float F8MAX = 448.0f; + +__global__ void k_requant_q1_0_to_e4m3_perchannel( + const char * __restrict__ wdata, int64_t nb01, + uint8_t * __restrict__ q8, float * __restrict__ wscale, + int64_t K, int64_t n_blocks) { + const int64_t n = blockIdx.x; + const block_q1_0 * row = (const block_q1_0 *)(wdata + n * nb01); + float amax = 0.0f; + for (int64_t l = threadIdx.x; l < K; l += blockDim.x) { + const int64_t b = l / Q1K; + const int t = (int)(l % Q1K); + const int byte_index = t / 8; + const int bit_offset = t % 8; + const uint8_t bit = (row[b].qs[byte_index] >> bit_offset) & 1; + const float d = q1_half2float(row[b].d); + amax = fmaxf(amax, fabsf(bit ? d : -d)); + } + __shared__ float sred[1024]; + sred[threadIdx.x] = amax; __syncthreads(); + for (int s = blockDim.x >> 1; s > 0; s >>= 1) { + if (threadIdx.x < s) sred[threadIdx.x] = fmaxf(sred[threadIdx.x], sred[threadIdx.x + s]); + __syncthreads(); + } + const float scale = (sred[0] > 0.0f) ? sred[0] / F8MAX : 1.0f; + if (threadIdx.x == 0) wscale[n] = scale; + const float inv = 1.0f / scale; + uint8_t * out = q8 + n * K; + for (int64_t l = threadIdx.x; l < K; l += blockDim.x) { + const int64_t b = l / Q1K; + const int t = (int)(l % Q1K); + const int byte_index = t / 8; + const int bit_offset = t % 8; + const uint8_t bit = (row[b].qs[byte_index] >> bit_offset) & 1; + const float d = q1_half2float(row[b].d); + out[l] = ggml_cuda_fp32_to_e4m3((bit ? d : -d) * inv); + } +} + +__global__ void k_quantize_act_e4m3_percol( + const char * __restrict__ src1, int64_t nb1, + uint8_t * __restrict__ x8, float * __restrict__ ascale, int64_t K) { + const int64_t j = blockIdx.x; + const float * col = (const float *)(src1 + j * nb1); + float amax = 0.0f; + for (int64_t l = threadIdx.x; l < K; l += blockDim.x) amax = fmaxf(amax, fabsf(col[l])); + __shared__ float sred[1024]; + sred[threadIdx.x] = amax; __syncthreads(); + for (int s = blockDim.x >> 1; s > 0; s >>= 1) { + if (threadIdx.x < s) sred[threadIdx.x] = fmaxf(sred[threadIdx.x], sred[threadIdx.x + s]); + __syncthreads(); + } + const float scale = (sred[0] > 0.0f) ? sred[0] / F8MAX : 1.0f; + if (threadIdx.x == 0) ascale[j] = scale; + const float inv = 1.0f / scale; + uint8_t * out = x8 + j * K; + for (int64_t l = threadIdx.x; l < K; l += blockDim.x) out[l] = ggml_cuda_fp32_to_e4m3(col[l] * inv); +} + +// fp8 GEMM emits f32 accumulate directly -> apply scales from a float buffer. +__global__ void k_apply_scales_f32( + const float * __restrict__ acc, char * __restrict__ dst, int64_t nb1, + const float * __restrict__ wscale, const float * __restrict__ ascale, + int64_t N, int64_t M) { + const int64_t idx = blockIdx.x * (int64_t)blockDim.x + threadIdx.x; + if (idx >= N * M) return; + const int64_t n = idx % N, j = idx / N; + float * out = (float *)(dst + j * nb1); + out[n] = acc[j * N + n] * wscale[n] * ascale[j]; +} + +// One hipBLASLt handle per process (single-GPU prefill use). Thread-safe lazy init. +hipblasLtHandle_t get_lt_handle() { + static hipblasLtHandle_t h = [](){ hipblasLtHandle_t t; hipblasLtCreate(&t); return t; }(); + return h; +} + +// ---- per-shape plan cache ------------------------------------------------- +// The dominant integration cost is the per-call heuristic search + descriptor +// churn (Stage 4: isolated GEMM = 190-338 TOPS but integrated = 55 because this +// ran EVERY matmul). A prefill touches only a handful of distinct (M,N,K) +// shapes (ubatch M x each weight's N,K), so build the descriptors + run the +// heuristic ONCE per shape and reuse the plan. hipblasLtMatmul takes fresh data +// pointers each call; the desc/layouts/algo are shape-only and safely shared. +constexpr size_t LT_WS_BYTES = 32ull << 20; // fixed workspace budget for algo selection + +struct lt_plan { + hipblasLtMatmulDesc_t op = nullptr; + hipblasLtMatrixLayout_t lA = nullptr, lB = nullptr, lD = nullptr; + hipblasLtMatmulHeuristicResult_t heur{}; + bool ok = false; +}; + +enum gemm_mode { MODE_I8 = 0, MODE_F8 = 1 }; // int8 (i8->i32) or fp8 (e4m3->f32) + +std::map, lt_plan> g_plan_cache; +std::mutex g_plan_mtx; + +// ---- persistent tuned-algo cache ------------------------------------------ +// The autotune benchmark (measure best algo per shape) costs ~10s of warmup per +// process. hipblasLtMatmulAlgo_t is a trivially-serializable POD (docs: "can be +// trivially serialized and later restored for use with the same version of the +// library"), so persist the winners to disk keyed by (N,M,K,mode). A version +// tag (hipblaslt major.minor.patch-githash) invalidates the file if the library +// changes. On a hit we skip the benchmark (one validation matmul instead). +// Persist the winning heuristic-candidate INDEX (not the opaque algo blob -- +// restoring a serialized hipblasLtMatmulAlgo_t segfaults on use, and there is no +// AlgoCheck API to validate it safely). The heuristic candidate list is +// deterministic per (shape, library version), so re-requesting it and picking +// the saved index reproduces the tuned kernel; a stale index just falls back to +// a valid heuristic candidate (never a crash). +#define TUNE_S2(x) #x +#define TUNE_S(x) TUNE_S2(x) +constexpr char TUNE_MAGIC[8] = {'R','D','N','4','G','T','3','\0'}; + +std::map, int32_t> g_disk_best; +bool g_disk_loaded = false; + +const char * tune_cache_path() { + static std::string path = [](){ + if (const char * e = getenv("GGML_HIP_Q1_0_HIPBLASLT_TUNE_CACHE")) return std::string(e); + const char * home = getenv("HOME"); + return std::string(home ? home : "/tmp") + "/.cache/ggml-rdna4-gemm-tune.bin"; + }(); + return path.c_str(); +} +std::string tune_version_tag() { + return std::string("hipblaslt-") + TUNE_S(HIPBLASLT_VERSION_MAJOR) "." TUNE_S(HIPBLASLT_VERSION_MINOR) + "." TUNE_S(HIPBLASLT_VERSION_PATCH) "-" TUNE_S(HIPBLASLT_VERSION_TWEAK); +} +struct TuneRec { int64_t N, M, K; int32_t mode; int32_t best_index; }; + +void load_disk_algos() { + if (g_disk_loaded) return; + g_disk_loaded = true; + FILE * f = fopen(tune_cache_path(), "rb"); + if (!f) return; + char magic[8] = {0}; + uint32_t vlen = 0; + std::string want = tune_version_tag(); + std::string ver; + if (fread(magic, 1, 8, f) == 8 && memcmp(magic, TUNE_MAGIC, 8) == 0 && + fread(&vlen, 4, 1, f) == 1 && vlen <= 256) { + ver.resize(vlen); + if (fread(&ver[0], 1, vlen, f) == vlen && ver == want) { + TuneRec r; + while (fread(&r, sizeof(TuneRec), 1, f) == 1) { + g_disk_best[std::make_tuple(r.N, r.M, r.K, (int)r.mode)] = r.best_index; + } + } + } + fclose(f); +} +void save_disk_algos() { // rewrite whole file (few records); caller holds g_plan_mtx + std::string p = tune_cache_path(); + auto slash = p.find_last_of('/'); + if (slash != std::string::npos) mkdir(p.substr(0, slash).c_str(), 0755); // best-effort + FILE * f = fopen(p.c_str(), "wb"); + if (!f) return; + std::string ver = tune_version_tag(); + uint32_t vlen = (uint32_t)ver.size(); + fwrite(TUNE_MAGIC, 1, 8, f); + fwrite(&vlen, 4, 1, f); + fwrite(ver.data(), 1, vlen, f); + for (auto & kv : g_disk_best) { + TuneRec r{ std::get<0>(kv.first), std::get<1>(kv.first), std::get<2>(kv.first), + (int32_t)std::get<3>(kv.first), kv.second }; + fwrite(&r, sizeof(TuneRec), 1, f); + } + fclose(f); +} + +// Returns a cached (or freshly built) plan for D(NxM)=op(A=W)[NxK]*B(X)[KxM], TN. +// plan.ok == false means the heuristic found no algo for this shape/mode. +const lt_plan & get_plan(int64_t N, int64_t M, int64_t K, int mode) { + std::lock_guard lk(g_plan_mtx); + auto key = std::make_tuple(N, M, K, mode); + auto it = g_plan_cache.find(key); + if (it != g_plan_cache.end()) { + return it->second; + } + const hipblasComputeType_t compute = (mode == MODE_F8) ? HIPBLAS_COMPUTE_32F : HIPBLAS_COMPUTE_32I; + const hipDataType scaleT = (mode == MODE_F8) ? HIP_R_32F : HIP_R_32I; + const hipDataType abT = (mode == MODE_F8) ? HIP_R_8F_E4M3 : HIP_R_8I; + const hipDataType dT = (mode == MODE_F8) ? HIP_R_32F : HIP_R_32I; + lt_plan p; + hipblasLtHandle_t h = get_lt_handle(); + hipblasOperation_t opT = HIPBLAS_OP_T, opN = HIPBLAS_OP_N; + bool built = + hipblasLtMatmulDescCreate(&p.op, compute, scaleT) == HIPBLAS_STATUS_SUCCESS && + hipblasLtMatmulDescSetAttribute(p.op, HIPBLASLT_MATMUL_DESC_TRANSA, &opT, sizeof(opT)) == HIPBLAS_STATUS_SUCCESS && + hipblasLtMatmulDescSetAttribute(p.op, HIPBLASLT_MATMUL_DESC_TRANSB, &opN, sizeof(opN)) == HIPBLAS_STATUS_SUCCESS && + hipblasLtMatrixLayoutCreate(&p.lA, abT, K, N, K) == HIPBLAS_STATUS_SUCCESS && // stored KxN, op=T -> NxK + hipblasLtMatrixLayoutCreate(&p.lB, abT, K, M, K) == HIPBLAS_STATUS_SUCCESS && + hipblasLtMatrixLayoutCreate(&p.lD, dT, N, M, N) == HIPBLAS_STATUS_SUCCESS; + if (built) { + load_disk_algos(); + const auto dkey = std::make_tuple(N, M, K, mode); + + hipblasLtMatmulPreference_t pref = nullptr; + size_t ws = LT_WS_BYTES; int nAlgo = 0; + constexpr int REQ = 64; + std::vector cand(REQ); + // ALWAYS fetch candidates (fast, deterministic order). The benchmark loop + // (the ~5s/shape cost) is what a persisted index lets us skip. + if (hipblasLtMatmulPreferenceCreate(&pref) == HIPBLAS_STATUS_SUCCESS && + hipblasLtMatmulPreferenceSetAttribute(pref, HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES, &ws, sizeof(ws)) == HIPBLAS_STATUS_SUCCESS && + hipblasLtMatmulAlgoGetHeuristic(h, p.op, p.lA, p.lB, p.lD, p.lD, pref, REQ, cand.data(), &nAlgo) == HIPBLAS_STATUS_SUCCESS && + nAlgo > 0) { + static const bool notune = (getenv("GGML_HIP_Q1_0_HIPBLASLT_NOTUNE") != nullptr); + int best = 0; + auto di = g_disk_best.find(dkey); + if (di != g_disk_best.end() && di->second >= 0 && di->second < nAlgo) { + best = di->second; // persisted winner -> skip benchmark + } else if (!notune && nAlgo > 1) { + // benchmark all candidates with scratch buffers, persist the winning index + void *sA=nullptr,*sB=nullptr,*sD=nullptr,*sW=nullptr; + if (hipMalloc(&sA,(size_t)K*N)==hipSuccess && hipMalloc(&sB,(size_t)K*M)==hipSuccess && + hipMalloc(&sD,(size_t)N*M*4)==hipSuccess && hipMalloc(&sW,LT_WS_BYTES)==hipSuccess) { + hipMemset(sA,1,(size_t)K*N); hipMemset(sB,1,(size_t)K*M); + const int32_t ai=1,bi=0; const float af=1.f,bf=0.f; + const void *alpha=(mode==MODE_F8)?(const void*)&af:(const void*)&ai; + const void *beta =(mode==MODE_F8)?(const void*)&bf:(const void*)&bi; + auto run=[&](hipblasLtMatmulAlgo_t &a){ return hipblasLtMatmul(h,p.op,alpha,sA,p.lA,sB,p.lB,beta,sD,p.lD,sD,p.lD,&a,sW,LT_WS_BYTES,0); }; + hipEvent_t e0,e1; hipEventCreate(&e0); hipEventCreate(&e1); + double bestMs = 1e30; + for (int i=0;i0 && mssecond; +} + +// ---- bounded int8 weight cache --------------------------------------------- +// The residual prefill loss after the plan cache is the per-call requant pass +// (Q1_0 -> int8/e4m3, a full memory pass that dp4a fuses into its kernel). Weights +// are constant, so cache the int8/e4m3 copy + per-channel scale keyed by the weight +// pointer. A full-model int8/e4m3 copy won't fit beside the resident +// Q1_0 on a 32 GB card, so cap the cache at a VRAM budget: weights that fit are +// cached (requant paid once), the rest fall back to on-the-fly pool requant. +// hipMalloc failure also falls back -- never OOM-crash. +struct cached_w { int8_t * q8 = nullptr; float * wscale = nullptr; size_t bytes = 0; }; +struct cached_w_f8 { uint8_t * q8f8 = nullptr; float * wscale = nullptr; size_t bytes = 0; }; +std::map g_wcache_i8; +std::map g_wcache_f8; +size_t g_wcache_bytes = 0; +std::mutex g_wcache_mtx; + +size_t wcache_budget_bytes() { + static size_t b = [](){ + const char * e = getenv("GGML_HIP_Q1_0_HIPBLASLT_WCACHE_MB"); + size_t mb = e ? (size_t)atoll(e) : (size_t)12000; // ~12 GB default: leaves headroom + // for model + llama.cpp compute bufs + return mb << 20; + }(); + return b; +} + +// Returns cached int8 weight (building it on first miss if within budget), or +// nullptr -> caller must requant on-the-fly. Build + all uses share the stream, +// so the one-time requant is correctly ordered before any GEMM that reads it. +const cached_w * try_cache_weight_i8(const void * key, const char * wdata, int64_t nb01, + int64_t N, int64_t K, int64_t n_blocks, cudaStream_t stream) { + std::lock_guard lk(g_wcache_mtx); + auto it = g_wcache_i8.find(key); + if (it != g_wcache_i8.end()) return &it->second; + + const size_t need = (size_t)N * K + (size_t)N * sizeof(float); + if (g_wcache_bytes + need > wcache_budget_bytes()) return nullptr; // budget hit + + cached_w c; + if (hipMalloc(&c.q8, (size_t)N * K) != hipSuccess) return nullptr; + if (hipMalloc(&c.wscale, (size_t)N * sizeof(float)) != hipSuccess) { hipFree(c.q8); return nullptr; } + c.bytes = need; + const dim3 grid((unsigned)N), block(256); + k_requant_q1_0_to_int8_perchannel<<>>(wdata, nb01, c.q8, c.wscale, K, n_blocks); + g_wcache_bytes += need; + auto res = g_wcache_i8.emplace(key, c); + return &res.first->second; +} + +const cached_w_f8 * try_cache_weight_f8(const void * key, const char * wdata, int64_t nb01, + int64_t N, int64_t K, int64_t n_blocks, cudaStream_t stream) { + std::lock_guard lk(g_wcache_mtx); + auto it = g_wcache_f8.find(key); + if (it != g_wcache_f8.end()) return &it->second; + + const size_t need = (size_t)N * K + (size_t)N * sizeof(float); + if (g_wcache_bytes + need > wcache_budget_bytes()) return nullptr; // budget hit + + cached_w_f8 c; + if (hipMalloc(&c.q8f8, (size_t)N * K) != hipSuccess) return nullptr; + if (hipMalloc(&c.wscale, (size_t)N * sizeof(float)) != hipSuccess) { hipFree(c.q8f8); return nullptr; } + c.bytes = need; + const dim3 grid((unsigned)N), block(256); + k_requant_q1_0_to_e4m3_perchannel<<>>(wdata, nb01, c.q8f8, c.wscale, K, n_blocks); + g_wcache_bytes += need; + auto res = g_wcache_f8.emplace(key, c); + return &res.first->second; +} + +// Drop cached conversions whose weight pointer falls inside a buffer being freed. +// Q1_0 keeps two caches (int8 and fp8 modes); both key on the raw device address +// and both must be purged, or a reused allocation returns stale weights. +void wcache_invalidate_range(const void * base, size_t size) { + std::lock_guard lk(g_wcache_mtx); + const char * b = (const char *) base; + for (auto it = g_wcache_i8.begin(); it != g_wcache_i8.end(); ) { + const char * k = (const char *) it->first; + if (k >= b && k < b + size) { + if (it->second.q8) hipFree(it->second.q8); + if (it->second.wscale) hipFree(it->second.wscale); + g_wcache_bytes -= it->second.bytes; + it = g_wcache_i8.erase(it); + } else { ++it; } + } + for (auto it = g_wcache_f8.begin(); it != g_wcache_f8.end(); ) { + const char * k = (const char *) it->first; + if (k >= b && k < b + size) { + if (it->second.q8f8) hipFree(it->second.q8f8); + if (it->second.wscale) hipFree(it->second.wscale); + g_wcache_bytes -= it->second.bytes; + it = g_wcache_f8.erase(it); + } else { ++it; } + } +} +struct wcache_registrar { wcache_registrar() { ggml_hipblaslt_wcache_register(wcache_invalidate_range); } }; +wcache_registrar g_wcache_registrar; + +} // namespace + +bool ggml_cuda_q1_0_hipblaslt_prefill_supports(const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * dst) { + if (src0->type != GGML_TYPE_Q1_0) return false; + if (src1->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_F32) return false; + if (src0->ne[2] != 1 || src0->ne[3] != 1) return false; + if (src1->ne[2] != 1 || src1->ne[3] != 1) return false; + if (src0->ne[0] != src1->ne[0] || src0->ne[0] % Q1K != 0) return false; + + // Prefill only: M must clear the threshold (decode stays on dp4a). Tunable. + static const int64_t M_THRESH = [](){ + const char * e = getenv("GGML_HIP_Q1_0_HIPBLASLT_MTHRESH"); + return e ? (int64_t)atoll(e) : (int64_t)384; // was 32: M-sweep showed int8 route regresses M<256 (fable review); engage only in the win regime + }(); + if (src1->ne[1] <= M_THRESH) return false; + + const int device = ggml_cuda_get_device(); + const int cc = ggml_cuda_info().devices[device].cc; + return GGML_CUDA_CC_IS_RDNA4(cc); +} + +bool ggml_cuda_op_mul_mat_q1_0_hipblaslt(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { + GGML_ASSERT(src0->type == GGML_TYPE_Q1_0); + GGML_ASSERT(src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32); + + const int64_t K = src0->ne[0]; + const int64_t N = src0->ne[1]; + const int64_t M = src1->ne[1]; + const int64_t n_blocks = K / Q1K; + cudaStream_t stream = ctx.stream(); + + // int8 (default) or fp8/e4m3 GEMM -- env-checked once. fp8 rides the 202 + // tuned gfx1201 fp8 kernels (vs 16 int8); see the fp8 kernels above. + static const int mode = (getenv("GGML_HIP_Q1_0_HIPBLASLT_FP8") != nullptr) ? MODE_F8 : MODE_I8; + + // ---- weight -> int8/e4m3 (per-output-channel): bounded cache, pool fallback ---- + int8_t * wq8_ptr = nullptr; + float * wsc_ptr = nullptr; + uint8_t * wq8f8_ptr = nullptr; + ggml_cuda_pool_alloc wq8_pool; // lazily allocated only on cache miss + ggml_cuda_pool_alloc wsc_pool; + ggml_cuda_pool_alloc wq8f8_pool; + + const cached_w * cw = nullptr; + const cached_w_f8 * cw_f8 = nullptr; + + if (mode == MODE_I8) { + cw = try_cache_weight_i8(src0->data, (const char *)src0->data, src0->nb[1], + N, K, n_blocks, stream); + if (cw) { + wq8_ptr = cw->q8; + wsc_ptr = cw->wscale; + } + } else { + cw_f8 = try_cache_weight_f8(src0->data, (const char *)src0->data, src0->nb[1], + N, K, n_blocks, stream); + if (cw_f8) { + wq8f8_ptr = cw_f8->q8f8; + wsc_ptr = cw_f8->wscale; + } + } + + if (mode == MODE_I8 && !wq8_ptr) { + wq8_ptr = wq8_pool.alloc(ctx.pool(), (size_t)N * K); + wsc_ptr = wsc_pool.alloc(ctx.pool(), (size_t)N); + const dim3 grid((unsigned)N), block(256); + k_requant_q1_0_to_int8_perchannel<<>>( + (const char *)src0->data, src0->nb[1], wq8_ptr, wsc_ptr, K, n_blocks); + } else if (mode == MODE_F8 && !wq8f8_ptr) { + wq8f8_ptr = wq8f8_pool.alloc(ctx.pool(), (size_t)N * K); + wsc_ptr = wsc_pool.alloc(ctx.pool(), (size_t)N); + const dim3 grid((unsigned)N), block(256); + k_requant_q1_0_to_e4m3_perchannel<<>>( + (const char *)src0->data, src0->nb[1], wq8f8_ptr, wsc_ptr, K, n_blocks); + } + + // ---- activation int8/e4m3 (per-token) + 4-byte accumulator, from the pool ---- + ggml_cuda_pool_alloc x8 (ctx.pool(), (size_t)K * M); + ggml_cuda_pool_alloc asc (ctx.pool(), (size_t)M); + ggml_cuda_pool_alloc acc (ctx.pool(), (size_t)N * M); // i32 (int8) or reinterpreted f32 (fp8) + { + const dim3 grid((unsigned)M), block(256); + if (mode == MODE_F8) { + k_quantize_act_e4m3_percol<<>>( + (const char *)src1->data, src1->nb[1], (uint8_t *)x8.get(), asc.get(), K); + } else { + k_quantize_act_int8_percol<<>>( + (const char *)src1->data, src1->nb[1], x8.get(), asc.get(), K); + } + } + + // ---- hipBLASLt GEMM: D(NxM) = op(A=W)[NxK] * B(X)[KxM], TN. Plan cached per (N,M,K,mode). ---- + hipblasLtHandle_t h = get_lt_handle(); + const lt_plan & plan = get_plan(N, M, K, mode); + if (!plan.ok) { + GGML_LOG_ERROR("%s: no hipBLASLt algo for %ldx%ldx%ld mode=%d\n", __func__, N, M, K, mode); + return false; + } + + ggml_cuda_pool_alloc ws(ctx.pool(), LT_WS_BYTES); + + if (mode == MODE_F8) { + const float alpha = 1.0f, beta = 0.0f; + LT_OK(hipblasLtMatmul(h, plan.op, &alpha, wq8f8_ptr, plan.lA, x8.get(), plan.lB, &beta, + acc.get(), plan.lD, acc.get(), plan.lD, + &plan.heur.algo, ws.get(), LT_WS_BYTES, stream)); + } else { + const int32_t alpha = 1, beta = 0; + LT_OK(hipblasLtMatmul(h, plan.op, &alpha, wq8_ptr, plan.lA, x8.get(), plan.lB, &beta, + acc.get(), plan.lD, acc.get(), plan.lD, + &plan.heur.algo, ws.get(), LT_WS_BYTES, stream)); + } + + // ---- dequant: dst = acc * wscale[row] * ascale[col] (acc is i32 or f32) ---- + { + const int64_t total = N * M; + const dim3 block(256), grid((unsigned)((total + 255) / 256)); + if (mode == MODE_F8) { + k_apply_scales_f32<<>>( + (const float *)acc.get(), (char *)dst->data, dst->nb[1], wsc_ptr, asc.get(), N, M); + } else { + k_apply_scales<<>>( + acc.get(), (char *)dst->data, dst->nb[1], wsc_ptr, asc.get(), N, M); + } + } + return true; +} + +#else // ---- non-HIP / disabled: stubs so the globbed source is a no-op ---- + +bool ggml_cuda_q1_0_hipblaslt_prefill_supports(const ggml_tensor *, const ggml_tensor *, const ggml_tensor *) { + return false; +} +bool ggml_cuda_op_mul_mat_q1_0_hipblaslt(ggml_backend_cuda_context &, const ggml_tensor *, const ggml_tensor *, ggml_tensor *) { + return false; +} + +#endif diff --git a/ggml/src/ggml-cuda/mul_mat_q1_0_hipblaslt.cuh b/ggml/src/ggml-cuda/mul_mat_q1_0_hipblaslt.cuh new file mode 100644 index 00000000000..3d39c66e577 --- /dev/null +++ b/ggml/src/ggml-cuda/mul_mat_q1_0_hipblaslt.cuh @@ -0,0 +1,17 @@ +// mul_mat_q1_0_hipblaslt.cuh - Q1_0 hipBLASLt prefill GEMM route header +// Route Q1_0 (binary {-1,+1} weights, QK1_0=128) to hipBLASLt int8 or fp8 (e4m3) GEMM +// for prefill (M > M_THRESH). Per-channel weight scale, per-token activation scale. + +#pragma once + +#include "common.cuh" + +// Whether this call site is in-scope for the hipBLASLt prefill path (Q1_0 +// weight, F32 acts/dst, 2D, RDNA4, and M > threshold). Checked by the +// ggml_cuda_mul_mat() intercept before dispatch. Soft/opt-in: a false return +// falls straight through to the unmodified mmq/dp4a path. +bool ggml_cuda_q1_0_hipblaslt_prefill_supports(const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * dst); + +// Run the Q1_0 prefill matmul through hipBLASLt int8/fp8. Returns false if the +// build has no hipBLASLt (non-HIP / disabled) so the caller can fall back. +bool ggml_cuda_op_mul_mat_q1_0_hipblaslt(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst); diff --git a/ggml/src/ggml-cuda/mul_mat_q2_0_hipblaslt.cu b/ggml/src/ggml-cuda/mul_mat_q2_0_hipblaslt.cu new file mode 100644 index 00000000000..c5c9d725c71 --- /dev/null +++ b/ggml/src/ggml-cuda/mul_mat_q2_0_hipblaslt.cu @@ -0,0 +1,586 @@ +// See mul_mat_q2_0_hipblaslt.cuh for the full rationale. +// Wire the tuned hipBLASLt int8 GEMM +// as an opt-in prefill route for Q2_0. Correctness-first; algo/handle caching +// and the fp8 front-end are Stage 4. + +#include "mul_mat_q2_0_hipblaslt.cuh" +#include "hipblaslt_wcache.cuh" + +// Real implementation only on the AMD/HIP build (hipBLASLt is a ROCm lib). +// On a CUDA build these become stubs so the globbed source is a no-op. +#if defined(__HIP_PLATFORM_AMD__) && !defined(GGML_HIP_NO_HIPBLASLT) + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace { + +constexpr int Q2K = QK2_0; // 128 elements per Q2_0 block +constexpr float I8MAX = 127.0f; + +#define LT_OK(x) do { hipblasStatus_t s_ = (x); if (s_ != HIPBLAS_STATUS_SUCCESS) { \ + GGML_LOG_ERROR("%s: hipBLASLt error %d at %s:%d\n", __func__, (int)s_, __FILE__, __LINE__); \ + return false; } } while(0) + +// ---- device helpers -------------------------------------------------------- + +__device__ __forceinline__ float q2_half2float(ggml_half h) { + return __half2float(*reinterpret_cast(&h)); +} + +// Requant a Q2_0 weight row -> int8 with ONE symmetric scale per output channel. +// One block per output row n. Weight row layout: n_blocks contiguous block_q2_0. +// Output int8 is row-major [N x K] (== col-major [K x N], the TN A-operand). +__global__ void k_requant_q2_0_to_int8_perchannel( + const char * __restrict__ wdata, int64_t nb01, + int8_t * __restrict__ q8, float * __restrict__ wscale, + int64_t K, int64_t n_blocks) { + const int64_t n = blockIdx.x; // output channel + const block_q2_0 * row = (const block_q2_0 *)(wdata + n * nb01); + + // pass 1: block-wide amax over the K weights of this row + float amax = 0.0f; + for (int64_t l = threadIdx.x; l < K; l += blockDim.x) { + const int64_t b = l / Q2K; + const int t = (int)(l % Q2K); + const float d = q2_half2float(row[b].d); + const uint8_t q = (row[b].qs[t >> 2] >> ((t & 3) * 2)) & 0x3; + const float w = ((int)q - 1) * d; // {-d,0,+d,+2d} + amax = fmaxf(amax, fabsf(w)); + } + __shared__ float sred[1024]; + sred[threadIdx.x] = amax; + __syncthreads(); + for (int s = blockDim.x >> 1; s > 0; s >>= 1) { + if (threadIdx.x < s) sred[threadIdx.x] = fmaxf(sred[threadIdx.x], sred[threadIdx.x + s]); + __syncthreads(); + } + const float scale = (sred[0] > 0.0f) ? sred[0] / I8MAX : 1.0f; + if (threadIdx.x == 0) wscale[n] = scale; + const float inv = 1.0f / scale; + + // pass 2: quantize + int8_t * out = q8 + n * K; + for (int64_t l = threadIdx.x; l < K; l += blockDim.x) { + const int64_t b = l / Q2K; + const int t = (int)(l % Q2K); + const float d = q2_half2float(row[b].d); + const uint8_t q = (row[b].qs[t >> 2] >> ((t & 3) * 2)) & 0x3; + const float w = ((int)q - 1) * d; + int v = __float2int_rn(w * inv); + v = max(-127, min(127, v)); + out[l] = (int8_t)v; + } +} + +// Quantize activations -> int8 with ONE symmetric scale per token (column). +// src1 fp32, column j at (src1 + j*nb1), K contiguous floats. Output col-major +// [K x M] (ld=K) + per-token scale. +__global__ void k_quantize_act_int8_percol( + const char * __restrict__ src1, int64_t nb1, + int8_t * __restrict__ x8, float * __restrict__ ascale, int64_t K) { + const int64_t j = blockIdx.x; // token / column + const float * col = (const float *)(src1 + j * nb1); + + float amax = 0.0f; + for (int64_t l = threadIdx.x; l < K; l += blockDim.x) amax = fmaxf(amax, fabsf(col[l])); + __shared__ float sred[1024]; + sred[threadIdx.x] = amax; + __syncthreads(); + for (int s = blockDim.x >> 1; s > 0; s >>= 1) { + if (threadIdx.x < s) sred[threadIdx.x] = fmaxf(sred[threadIdx.x], sred[threadIdx.x + s]); + __syncthreads(); + } + const float scale = (sred[0] > 0.0f) ? sred[0] / I8MAX : 1.0f; + if (threadIdx.x == 0) ascale[j] = scale; + const float inv = 1.0f / scale; + + int8_t * out = x8 + j * K; + for (int64_t l = threadIdx.x; l < K; l += blockDim.x) { + int v = __float2int_rn(col[l] * inv); + v = max(-127, min(127, v)); + out[l] = (int8_t)v; + } +} + +// Dequant the int32 GEMM result: dst[n,j] = i32[n,j] * wscale[n] * ascale[j]. +// Both dst and i32 are col-major [N x M] (dst ld from nb1, i32 ld=N). +__global__ void k_apply_scales( + const int32_t * __restrict__ i32, char * __restrict__ dst, int64_t nb1, + const float * __restrict__ wscale, const float * __restrict__ ascale, + int64_t N, int64_t M) { + const int64_t idx = blockIdx.x * (int64_t)blockDim.x + threadIdx.x; + if (idx >= N * M) return; + const int64_t n = idx % N; + const int64_t j = idx / N; + float * out = (float *)(dst + j * nb1); + out[n] = (float)i32[j * N + n] * wscale[n] * ascale[j]; +} + +// ============================ fp8 (E4M3) variant ============================ +// Same v1 scheme, but encode to OCP e4m3 and use hipBLASLt's fp8 GEMM (202 +// tuned gfx1201 kernels vs 16 int8). e4m3 max = 448; scale to fill the range. +constexpr float F8MAX = 448.0f; + +__global__ void k_requant_q2_0_to_e4m3_perchannel( + const char * __restrict__ wdata, int64_t nb01, + uint8_t * __restrict__ q8, float * __restrict__ wscale, + int64_t K, int64_t n_blocks) { + const int64_t n = blockIdx.x; + const block_q2_0 * row = (const block_q2_0 *)(wdata + n * nb01); + float amax = 0.0f; + for (int64_t l = threadIdx.x; l < K; l += blockDim.x) { + const int64_t b = l / Q2K; const int t = (int)(l % Q2K); + const float d = q2_half2float(row[b].d); + const uint8_t q = (row[b].qs[t >> 2] >> ((t & 3) * 2)) & 0x3; + amax = fmaxf(amax, fabsf(((int)q - 1) * d)); + } + __shared__ float sred[1024]; + sred[threadIdx.x] = amax; __syncthreads(); + for (int s = blockDim.x >> 1; s > 0; s >>= 1) { + if (threadIdx.x < s) sred[threadIdx.x] = fmaxf(sred[threadIdx.x], sred[threadIdx.x + s]); + __syncthreads(); + } + const float scale = (sred[0] > 0.0f) ? sred[0] / F8MAX : 1.0f; + if (threadIdx.x == 0) wscale[n] = scale; + const float inv = 1.0f / scale; + uint8_t * out = q8 + n * K; + for (int64_t l = threadIdx.x; l < K; l += blockDim.x) { + const int64_t b = l / Q2K; const int t = (int)(l % Q2K); + const float d = q2_half2float(row[b].d); + const uint8_t q = (row[b].qs[t >> 2] >> ((t & 3) * 2)) & 0x3; + out[l] = ggml_cuda_fp32_to_e4m3(((int)q - 1) * d * inv); + } +} + +__global__ void k_quantize_act_e4m3_percol( + const char * __restrict__ src1, int64_t nb1, + uint8_t * __restrict__ x8, float * __restrict__ ascale, int64_t K) { + const int64_t j = blockIdx.x; + const float * col = (const float *)(src1 + j * nb1); + float amax = 0.0f; + for (int64_t l = threadIdx.x; l < K; l += blockDim.x) amax = fmaxf(amax, fabsf(col[l])); + __shared__ float sred[1024]; + sred[threadIdx.x] = amax; __syncthreads(); + for (int s = blockDim.x >> 1; s > 0; s >>= 1) { + if (threadIdx.x < s) sred[threadIdx.x] = fmaxf(sred[threadIdx.x], sred[threadIdx.x + s]); + __syncthreads(); + } + const float scale = (sred[0] > 0.0f) ? sred[0] / F8MAX : 1.0f; + if (threadIdx.x == 0) ascale[j] = scale; + const float inv = 1.0f / scale; + uint8_t * out = x8 + j * K; + for (int64_t l = threadIdx.x; l < K; l += blockDim.x) out[l] = ggml_cuda_fp32_to_e4m3(col[l] * inv); +} + +// fp8 GEMM emits f32 accumulate directly -> apply scales from a float buffer. +__global__ void k_apply_scales_f32( + const float * __restrict__ acc, char * __restrict__ dst, int64_t nb1, + const float * __restrict__ wscale, const float * __restrict__ ascale, + int64_t N, int64_t M) { + const int64_t idx = blockIdx.x * (int64_t)blockDim.x + threadIdx.x; + if (idx >= N * M) return; + const int64_t n = idx % N, j = idx / N; + float * out = (float *)(dst + j * nb1); + out[n] = acc[j * N + n] * wscale[n] * ascale[j]; +} + +// One hipBLASLt handle per process (single-GPU prefill use). Thread-safe lazy init. +hipblasLtHandle_t get_lt_handle() { + static hipblasLtHandle_t h = [](){ hipblasLtHandle_t t; hipblasLtCreate(&t); return t; }(); + return h; +} + +// ---- per-shape plan cache ------------------------------------------------- +// The dominant integration cost is the per-call heuristic search + descriptor +// churn (Stage 4: isolated GEMM = 190-338 TOPS but integrated = 55 because this +// ran EVERY matmul). A prefill touches only a handful of distinct (M,N,K) +// shapes (ubatch M x each weight's N,K), so build the descriptors + run the +// heuristic ONCE per shape and reuse the plan. hipblasLtMatmul takes fresh data +// pointers each call; the desc/layouts/algo are shape-only and safely shared. +constexpr size_t LT_WS_BYTES = 32ull << 20; // fixed workspace budget for algo selection + +struct lt_plan { + hipblasLtMatmulDesc_t op = nullptr; + hipblasLtMatrixLayout_t lA = nullptr, lB = nullptr, lD = nullptr; + hipblasLtMatmulHeuristicResult_t heur{}; + bool ok = false; +}; + +enum gemm_mode { MODE_I8 = 0, MODE_F8 = 1 }; // int8 (i8->i32) or fp8 (e4m3->f32) + +std::map, lt_plan> g_plan_cache; +std::mutex g_plan_mtx; + +// ---- persistent tuned-algo cache ------------------------------------------ +// The autotune benchmark (measure best algo per shape) costs ~10s of warmup per +// process. hipblasLtMatmulAlgo_t is a trivially-serializable POD (docs: "can be +// trivially serialized and later restored for use with the same version of the +// library"), so persist the winners to disk keyed by (N,M,K,mode). A version +// tag (hipblaslt major.minor.patch-githash) invalidates the file if the library +// changes. On a hit we skip the benchmark (one validation matmul instead). +// Persist the winning heuristic-candidate INDEX (not the opaque algo blob -- +// restoring a serialized hipblasLtMatmulAlgo_t segfaults on use, and there is no +// AlgoCheck API to validate it safely). The heuristic candidate list is +// deterministic per (shape, library version), so re-requesting it and picking +// the saved index reproduces the tuned kernel; a stale index just falls back to +// a valid heuristic candidate (never a crash). +#define TUNE_S2(x) #x +#define TUNE_S(x) TUNE_S2(x) +constexpr char TUNE_MAGIC[8] = {'R','D','N','4','G','T','3','\0'}; + +std::map, int32_t> g_disk_best; +bool g_disk_loaded = false; + +const char * tune_cache_path() { + static std::string path = [](){ + if (const char * e = getenv("GGML_HIP_Q2_0_HIPBLASLT_TUNE_CACHE")) return std::string(e); + const char * home = getenv("HOME"); + return std::string(home ? home : "/tmp") + "/.cache/ggml-rdna4-gemm-tune.bin"; + }(); + return path.c_str(); +} +std::string tune_version_tag() { + return std::string("hipblaslt-") + TUNE_S(HIPBLASLT_VERSION_MAJOR) "." TUNE_S(HIPBLASLT_VERSION_MINOR) + "." TUNE_S(HIPBLASLT_VERSION_PATCH) "-" TUNE_S(HIPBLASLT_VERSION_TWEAK); +} +struct TuneRec { int64_t N, M, K; int32_t mode; int32_t best_index; }; + +void load_disk_algos() { + if (g_disk_loaded) return; + g_disk_loaded = true; + FILE * f = fopen(tune_cache_path(), "rb"); + if (!f) return; + char magic[8] = {0}; + uint32_t vlen = 0; + std::string want = tune_version_tag(); + std::string ver; + if (fread(magic, 1, 8, f) == 8 && memcmp(magic, TUNE_MAGIC, 8) == 0 && + fread(&vlen, 4, 1, f) == 1 && vlen <= 256) { + ver.resize(vlen); + if (fread(&ver[0], 1, vlen, f) == vlen && ver == want) { + TuneRec r; + while (fread(&r, sizeof(TuneRec), 1, f) == 1) { + g_disk_best[std::make_tuple(r.N, r.M, r.K, (int)r.mode)] = r.best_index; + } + } + } + fclose(f); +} +void save_disk_algos() { // rewrite whole file (few records); caller holds g_plan_mtx + std::string p = tune_cache_path(); + auto slash = p.find_last_of('/'); + if (slash != std::string::npos) mkdir(p.substr(0, slash).c_str(), 0755); // best-effort + FILE * f = fopen(p.c_str(), "wb"); + if (!f) return; + std::string ver = tune_version_tag(); + uint32_t vlen = (uint32_t)ver.size(); + fwrite(TUNE_MAGIC, 1, 8, f); + fwrite(&vlen, 4, 1, f); + fwrite(ver.data(), 1, vlen, f); + for (auto & kv : g_disk_best) { + TuneRec r{ std::get<0>(kv.first), std::get<1>(kv.first), std::get<2>(kv.first), + (int32_t)std::get<3>(kv.first), kv.second }; + fwrite(&r, sizeof(TuneRec), 1, f); + } + fclose(f); +} + +// Returns a cached (or freshly built) plan for D(NxM)=op(A=W)[NxK]*B(X)[KxM], TN. +// plan.ok == false means the heuristic found no algo for this shape/mode. +const lt_plan & get_plan(int64_t N, int64_t M, int64_t K, int mode) { + std::lock_guard lk(g_plan_mtx); + auto key = std::make_tuple(N, M, K, mode); + auto it = g_plan_cache.find(key); + if (it != g_plan_cache.end()) { + return it->second; + } + const hipblasComputeType_t compute = (mode == MODE_F8) ? HIPBLAS_COMPUTE_32F : HIPBLAS_COMPUTE_32I; + const hipDataType scaleT = (mode == MODE_F8) ? HIP_R_32F : HIP_R_32I; + const hipDataType abT = (mode == MODE_F8) ? HIP_R_8F_E4M3 : HIP_R_8I; + const hipDataType dT = (mode == MODE_F8) ? HIP_R_32F : HIP_R_32I; + lt_plan p; + hipblasLtHandle_t h = get_lt_handle(); + hipblasOperation_t opT = HIPBLAS_OP_T, opN = HIPBLAS_OP_N; + bool built = + hipblasLtMatmulDescCreate(&p.op, compute, scaleT) == HIPBLAS_STATUS_SUCCESS && + hipblasLtMatmulDescSetAttribute(p.op, HIPBLASLT_MATMUL_DESC_TRANSA, &opT, sizeof(opT)) == HIPBLAS_STATUS_SUCCESS && + hipblasLtMatmulDescSetAttribute(p.op, HIPBLASLT_MATMUL_DESC_TRANSB, &opN, sizeof(opN)) == HIPBLAS_STATUS_SUCCESS && + hipblasLtMatrixLayoutCreate(&p.lA, abT, K, N, K) == HIPBLAS_STATUS_SUCCESS && // stored KxN, op=T -> NxK + hipblasLtMatrixLayoutCreate(&p.lB, abT, K, M, K) == HIPBLAS_STATUS_SUCCESS && + hipblasLtMatrixLayoutCreate(&p.lD, dT, N, M, N) == HIPBLAS_STATUS_SUCCESS; + if (built) { + load_disk_algos(); + const auto dkey = std::make_tuple(N, M, K, mode); + + hipblasLtMatmulPreference_t pref = nullptr; + size_t ws = LT_WS_BYTES; int nAlgo = 0; + constexpr int REQ = 64; + std::vector cand(REQ); + // ALWAYS fetch candidates (fast, deterministic order). The benchmark loop + // (the ~5s/shape cost) is what a persisted index lets us skip. + if (hipblasLtMatmulPreferenceCreate(&pref) == HIPBLAS_STATUS_SUCCESS && + hipblasLtMatmulPreferenceSetAttribute(pref, HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES, &ws, sizeof(ws)) == HIPBLAS_STATUS_SUCCESS && + hipblasLtMatmulAlgoGetHeuristic(h, p.op, p.lA, p.lB, p.lD, p.lD, pref, REQ, cand.data(), &nAlgo) == HIPBLAS_STATUS_SUCCESS && + nAlgo > 0) { + static const bool notune = (getenv("GGML_HIP_Q2_0_HIPBLASLT_NOTUNE") != nullptr); + int best = 0; + auto di = g_disk_best.find(dkey); + if (di != g_disk_best.end() && di->second >= 0 && di->second < nAlgo) { + best = di->second; // persisted winner -> skip benchmark + } else if (!notune && nAlgo > 1) { + // benchmark all candidates with scratch buffers, persist the winning index + void *sA=nullptr,*sB=nullptr,*sD=nullptr,*sW=nullptr; + if (hipMalloc(&sA,(size_t)K*N)==hipSuccess && hipMalloc(&sB,(size_t)K*M)==hipSuccess && + hipMalloc(&sD,(size_t)N*M*4)==hipSuccess && hipMalloc(&sW,LT_WS_BYTES)==hipSuccess) { + hipMemset(sA,1,(size_t)K*N); hipMemset(sB,1,(size_t)K*M); + const int32_t ai=1,bi=0; const float af=1.f,bf=0.f; + const void *alpha=(mode==MODE_F8)?(const void*)&af:(const void*)&ai; + const void *beta =(mode==MODE_F8)?(const void*)&bf:(const void*)&bi; + auto run=[&](hipblasLtMatmulAlgo_t &a){ return hipblasLtMatmul(h,p.op,alpha,sA,p.lA,sB,p.lB,beta,sD,p.lD,sD,p.lD,&a,sW,LT_WS_BYTES,0); }; + hipEvent_t e0,e1; hipEventCreate(&e0); hipEventCreate(&e1); + double bestMs = 1e30; + for (int i=0;i0 && mssecond; +} + +// ---- bounded int8 weight cache --------------------------------------------- +// The residual prefill loss after the plan cache is the per-call requant pass +// (Q2_0 -> int8, a full memory pass that dp4a fuses into its kernel). Weights +// are constant, so cache the int8 copy + per-channel scale keyed by the weight +// pointer. A full-model int8 copy (~27 GB on 27B) won't fit beside the resident +// Q2_0 on a 32 GB card, so cap the cache at a VRAM budget: weights that fit are +// cached (requant paid once), the rest fall back to on-the-fly pool requant. +// hipMalloc failure also falls back -- never OOM-crash (the Stage-3 lesson). +struct cached_w { int8_t * q8 = nullptr; float * wscale = nullptr; size_t bytes = 0; }; +std::map g_wcache; +size_t g_wcache_bytes = 0; +std::mutex g_wcache_mtx; + +size_t wcache_budget_bytes() { + static size_t b = [](){ + const char * e = getenv("GGML_HIP_Q2_0_HIPBLASLT_WCACHE_MB"); + size_t mb = e ? (size_t)atoll(e) : (size_t)12000; // ~12 GB default: leaves headroom + // for model + llama.cpp compute bufs + return mb << 20; + }(); + return b; +} + +// Returns cached int8 weight (building it on first miss if within budget), or +// nullptr -> caller must requant on-the-fly. Build + all uses share the stream, +// so the one-time requant is correctly ordered before any GEMM that reads it. +const cached_w * try_cache_weight(const void * key, const char * wdata, int64_t nb01, + int64_t N, int64_t K, int64_t n_blocks, int mode, cudaStream_t stream) { + std::lock_guard lk(g_wcache_mtx); + auto it = g_wcache.find(key); + if (it != g_wcache.end()) return &it->second; // mode is fixed per process (env-checked once) + + const size_t need = (size_t)N * K + (size_t)N * sizeof(float); + if (g_wcache_bytes + need > wcache_budget_bytes()) return nullptr; // budget hit + { size_t freeb = 0, totb = 0; // VRAM-adaptive: keep headroom for the transient pool bufs + if (hipMemGetInfo(&freeb, &totb) == hipSuccess && freeb < need + (size_t)(2ull << 30)) return nullptr; } + + cached_w c; + if (hipMalloc(&c.q8, (size_t)N * K) != hipSuccess) return nullptr; + if (hipMalloc(&c.wscale, (size_t)N * sizeof(float)) != hipSuccess) { hipFree(c.q8); return nullptr; } + c.bytes = need; + const dim3 grid((unsigned)N), block(256); + if (mode == 1 /*MODE_F8*/) { + k_requant_q2_0_to_e4m3_perchannel<<>>(wdata, nb01, (uint8_t *)c.q8, c.wscale, K, n_blocks); + } else { + k_requant_q2_0_to_int8_perchannel<<>>(wdata, nb01, c.q8, c.wscale, K, n_blocks); + } + g_wcache_bytes += need; + auto res = g_wcache.emplace(key, c); + return &res.first->second; +} + +// Drop cached conversions whose weight pointer falls inside a buffer being freed. +// Without this the cache key (a raw device address) can be reused by a later +// allocation and silently return the previous model's converted weights. +void wcache_invalidate_range(const void * base, size_t size) { + std::lock_guard lk(g_wcache_mtx); + const char * b = (const char *) base; + for (auto it = g_wcache.begin(); it != g_wcache.end(); ) { + const char * k = (const char *) it->first; + if (k >= b && k < b + size) { + if (it->second.q8) hipFree(it->second.q8); + if (it->second.wscale) hipFree(it->second.wscale); + g_wcache_bytes -= it->second.bytes; + it = g_wcache.erase(it); + } else { + ++it; + } + } +} +struct wcache_registrar { wcache_registrar() { ggml_hipblaslt_wcache_register(wcache_invalidate_range); } }; +wcache_registrar g_wcache_registrar; + +} // namespace + +bool ggml_cuda_q2_0_hipblaslt_prefill_supports(const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * dst) { + if (src0->type != GGML_TYPE_Q2_0) return false; + if (src1->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_F32) return false; + if (src0->ne[2] != 1 || src0->ne[3] != 1) return false; + if (src1->ne[2] != 1 || src1->ne[3] != 1) return false; + if (src0->ne[0] != src1->ne[0] || src0->ne[0] % Q2K != 0) return false; + + // Prefill only: M must clear the threshold (decode stays on dp4a). Tunable. + static const int64_t M_THRESH = [](){ + // THRESHOLD NOTE (2026-08-05). This constant is deliberately conservative and + // it is model-dependent -- two measurements disagree in SIGN at M=128: + // Q2_0 campaign model -23% at M=32, -9.5% at M=64 (why it was raised 32 -> 384) + // sweep model +6.0% at M=128, +40.8% at M=256 + // Both are real. A single global constant cannot satisfy both, so 384 is chosen + // to never regress, at the cost of forgoing wins on models like the second. + // The principled fix is to measure it per model at load time, the way + // scripts/auto-batch-serve.sh already does for continuous-batching -np, rather + // than ship a constant. Until then, override per model with the env var below. + const char * e = getenv("GGML_HIP_Q2_0_HIPBLASLT_MTHRESH"); + return e ? (int64_t)atoll(e) : (int64_t)384; // was 32: M-sweep showed int8 route regresses M<256 (fable review); engage only in the win regime + }(); + if (src1->ne[1] <= M_THRESH) return false; + + const int device = ggml_cuda_get_device(); + const int cc = ggml_cuda_info().devices[device].cc; + return GGML_CUDA_CC_IS_RDNA4(cc); +} + +bool ggml_cuda_op_mul_mat_q2_0_hipblaslt(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { + GGML_ASSERT(src0->type == GGML_TYPE_Q2_0); + GGML_ASSERT(src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32); + + const int64_t K = src0->ne[0]; + const int64_t N = src0->ne[1]; + const int64_t M = src1->ne[1]; + const int64_t n_blocks = K / Q2K; + cudaStream_t stream = ctx.stream(); + + // int8 (default) or fp8/e4m3 GEMM -- env-checked once. fp8 rides the 202 + // tuned gfx1201 fp8 kernels (vs 16 int8); see the fp8 kernels above. + static const int mode = (getenv("GGML_HIP_Q2_0_HIPBLASLT_FP8") != nullptr) ? MODE_F8 : MODE_I8; + + // ---- weight -> int8/e4m3 (per-output-channel): bounded cache, pool fallback ---- + int8_t * wq8_ptr = nullptr; + float * wsc_ptr = nullptr; + ggml_cuda_pool_alloc wq8_pool; // lazily allocated only on cache miss + ggml_cuda_pool_alloc wsc_pool; + const cached_w * cw = try_cache_weight(src0->data, (const char *)src0->data, src0->nb[1], + N, K, n_blocks, mode, stream); + // Pre-flight VRAM guard: fall back to mmq (return false) instead of letting a pool + // alloc OOM-abort when a near-full model leaves too little free VRAM for the transient + // int8/fp8 + accumulator + GEMM-workspace buffers. + { + size_t freeb = 0, totb = 0; + const size_t wq8_need = cw ? 0 : (size_t)N * K; // pool alloc for weight only on cache miss + const size_t transient = wq8_need + (size_t)K*M + (size_t)N*M*4 + (size_t)M*4 + + LT_WS_BYTES + (size_t)(64ull << 20); + if (hipMemGetInfo(&freeb, &totb) == hipSuccess && freeb < transient) return false; + } + if (cw) { + wq8_ptr = cw->q8; + wsc_ptr = cw->wscale; + } else { + wq8_ptr = wq8_pool.alloc(ctx.pool(), (size_t)N * K); + wsc_ptr = wsc_pool.alloc(ctx.pool(), (size_t)N); + const dim3 grid((unsigned)N), block(256); + if (mode == MODE_F8) { + k_requant_q2_0_to_e4m3_perchannel<<>>( + (const char *)src0->data, src0->nb[1], (uint8_t *)wq8_ptr, wsc_ptr, K, n_blocks); + } else { + k_requant_q2_0_to_int8_perchannel<<>>( + (const char *)src0->data, src0->nb[1], wq8_ptr, wsc_ptr, K, n_blocks); + } + } + + // ---- activation int8/e4m3 (per-token) + 4-byte accumulator, from the pool ---- + ggml_cuda_pool_alloc x8 (ctx.pool(), (size_t)K * M); + ggml_cuda_pool_alloc asc (ctx.pool(), (size_t)M); + ggml_cuda_pool_alloc acc (ctx.pool(), (size_t)N * M); // i32 (int8) or reinterpreted f32 (fp8) + { + const dim3 grid((unsigned)M), block(256); + if (mode == MODE_F8) { + k_quantize_act_e4m3_percol<<>>( + (const char *)src1->data, src1->nb[1], (uint8_t *)x8.get(), asc.get(), K); + } else { + k_quantize_act_int8_percol<<>>( + (const char *)src1->data, src1->nb[1], x8.get(), asc.get(), K); + } + } + + // ---- hipBLASLt GEMM: D(NxM) = op(A=W)[NxK] * B(X)[KxM], TN. Plan cached per (N,M,K,mode). ---- + hipblasLtHandle_t h = get_lt_handle(); + const lt_plan & plan = get_plan(N, M, K, mode); + if (!plan.ok) { + GGML_LOG_ERROR("%s: no hipBLASLt algo for %ldx%ldx%ld mode=%d\n", __func__, N, M, K, mode); + return false; + } + + ggml_cuda_pool_alloc ws(ctx.pool(), LT_WS_BYTES); + if (mode == MODE_F8) { + const float alpha = 1.0f, beta = 0.0f; + LT_OK(hipblasLtMatmul(h, plan.op, &alpha, wq8_ptr, plan.lA, x8.get(), plan.lB, &beta, + acc.get(), plan.lD, acc.get(), plan.lD, + &plan.heur.algo, ws.get(), LT_WS_BYTES, stream)); + } else { + const int32_t alpha = 1, beta = 0; + LT_OK(hipblasLtMatmul(h, plan.op, &alpha, wq8_ptr, plan.lA, x8.get(), plan.lB, &beta, + acc.get(), plan.lD, acc.get(), plan.lD, + &plan.heur.algo, ws.get(), LT_WS_BYTES, stream)); + } + + // ---- dequant: dst = acc * wscale[row] * ascale[col] (acc is i32 or f32) ---- + { + const int64_t total = N * M; + const dim3 block(256), grid((unsigned)((total + 255) / 256)); + if (mode == MODE_F8) { + k_apply_scales_f32<<>>( + (const float *)acc.get(), (char *)dst->data, dst->nb[1], wsc_ptr, asc.get(), N, M); + } else { + k_apply_scales<<>>( + acc.get(), (char *)dst->data, dst->nb[1], wsc_ptr, asc.get(), N, M); + } + } + return true; +} + +#else // ---- non-HIP / disabled: stubs so the globbed source is a no-op ---- + +bool ggml_cuda_q2_0_hipblaslt_prefill_supports(const ggml_tensor *, const ggml_tensor *, const ggml_tensor *) { + return false; +} +bool ggml_cuda_op_mul_mat_q2_0_hipblaslt(ggml_backend_cuda_context &, const ggml_tensor *, const ggml_tensor *, ggml_tensor *) { + return false; +} + +#endif diff --git a/ggml/src/ggml-cuda/mul_mat_q2_0_hipblaslt.cuh b/ggml/src/ggml-cuda/mul_mat_q2_0_hipblaslt.cuh new file mode 100644 index 00000000000..cdbb4de4fa7 --- /dev/null +++ b/ggml/src/ggml-cuda/mul_mat_q2_0_hipblaslt.cuh @@ -0,0 +1,30 @@ +// Experimental PREFILL lever: route Q2_0 (ternary, g128) large-M matmuls through +// AMD's tuned hipBLASLt int8 GEMM instead of llama.cpp's ~5%-efficient mmq/dp4a +// path. Measured hipBLASLt int8 on gfx1201 = 156-266 TOPS (20-35% of peak) vs +// dp4a's ~64 TOPS -> 2.5-4x on the GEMM at the SAME W8A8 precision. +// +// Data flow (per prefill matmul, M > threshold): +// Q2_0 weight ----requant (once, cached)----> int8 + per-output-channel scale +// activations ---quantize (per call)--------> int8 + per-token scale +// hipBLASLt int8 GEMM (i8 x i8 -> i32, TN) +// i32 * wscale[row] * ascale[col] -> f32 dst +// +// This is v1 (per-output-channel weight scale): +// Stage 2 measured the added accuracy cost over the int8-activation floor at +// only ~0.23% rel-RMS. Decode (M<=threshold) is UNTOUCHED -- dp4a stays the +// GEMV path; hipBLASLt is a prefill/large-M play (matrix engine needs a full +// tile). Gated behind env var GGML_HIP_Q2_0_HIPBLASLT_PREFILL, opt-in, dormant +// by default. Single GPU, 2D weights only (ne2==ne3==1), MUL_MAT only. +#pragma once + +#include "common.cuh" + +// Whether this call site is in-scope for the hipBLASLt prefill path (Q2_0 +// weight, F32 acts/dst, 2D, RDNA4, and M > threshold). Checked by the +// ggml_cuda_mul_mat() intercept before dispatch. Soft/opt-in: a false return +// falls straight through to the unmodified mmq/dp4a path. +bool ggml_cuda_q2_0_hipblaslt_prefill_supports(const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * dst); + +// Run the Q2_0 prefill matmul through hipBLASLt int8. Returns false if the +// build has no hipBLASLt (non-HIP / disabled) so the caller can fall back. +bool ggml_cuda_op_mul_mat_q2_0_hipblaslt(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst); diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh index 5cd3cb1b0ce..68711df93d1 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -747,12 +747,41 @@ static __device__ __forceinline__ float vec_dot_q2_0_q8_1( // iqs selects which of the 4 chunks of 32 elements to process (0-3) const float d2 = bq2_0->d; - // each 32-element chunk occupies 8 bytes of qs (4 int16), regardless of group size - const int16_t * qs = (const int16_t *) bq2_0->qs + iqs * 4; - // Process only the chunk specified by iqs const block_q8_1 * bq8_1_chunk = bq8_1 + iqs; +#if defined(GGML_USE_HIP) || defined(GGML_USE_MUSA) + // AMD path (same split as unpack_q1_0_bytes above): the dynamic-selector + // __byte_perm chain below is built around NVIDIA's PRMT; on HIP it lowers + // poorly and the decode GEMV loses ~30% end-to-end. Instead, bit-spread + // the raw codes c in {0,1,2,3} into bytes and use the identity + // dot(s, u) = dot(c, u) - sum(u), s = c - 1 + // applying the -sum(u) offset once at the end via the q8_1 stored sum + // (ds.y = d8 * sum(u)) -- no per-code subtract, plain shift/mask + dp4a. + const int offset = iqs * 8; + const int qs0 = bq2_0->qs[offset + 0] | (bq2_0->qs[offset + 1] << 8) | + (bq2_0->qs[offset + 2] << 16) | (bq2_0->qs[offset + 3] << 24); + const int qs1 = bq2_0->qs[offset + 4] | (bq2_0->qs[offset + 5] << 8) | + (bq2_0->qs[offset + 6] << 16) | (bq2_0->qs[offset + 7] << 24); + + int sumi = 0; // = dot(c, u), c in {0,1,2,3} +#pragma unroll + for (int j = 0; j < 4; ++j) { + const int b0 = (qs0 >> (j*8)) & 0xFF; + const int s0 = (b0 | (b0 << 6) | (b0 << 12) | (b0 << 18)) & 0x03030303; // 4 codes -> 4 bytes + sumi = ggml_cuda_dp4a(s0, get_int_b4(bq8_1_chunk->qs, j), sumi); + const int b1 = (qs1 >> (j*8)) & 0xFF; + const int s1 = (b1 | (b1 << 6) | (b1 << 12) | (b1 << 18)) & 0x03030303; + sumi = ggml_cuda_dp4a(s1, get_int_b4(bq8_1_chunk->qs, 4 + j), sumi); + } + + const float d8 = __low2float(bq8_1_chunk->ds); + const float s8 = __high2float(bq8_1_chunk->ds); // = d8 * sum(u) + return d2 * (d8 * sumi - s8); +#else + // each 32-element chunk occupies 8 bytes of qs (4 int16), regardless of group size + const int16_t * qs = (const int16_t *) bq2_0->qs + iqs * 4; + int sumi = 0; #pragma unroll for (int j = 0; j < 4; ++j) { @@ -775,6 +804,7 @@ static __device__ __forceinline__ float vec_dot_q2_0_q8_1( // symbols are already signed, so no deferred sum(act) correction is needed const float d8 = __low2float(bq8_1_chunk->ds); return d2 * d8 * sumi; +#endif // defined(GGML_USE_HIP) || defined(GGML_USE_MUSA) } static __device__ __forceinline__ float vec_dot_q4_0_q8_1( diff --git a/ggml/src/ggml-hip/CMakeLists.txt b/ggml/src/ggml-hip/CMakeLists.txt index a7d4e0ea2b5..113bedfeee5 100644 --- a/ggml/src/ggml-hip/CMakeLists.txt +++ b/ggml/src/ggml-hip/CMakeLists.txt @@ -47,6 +47,11 @@ find_package(hip REQUIRED) find_package(hipblas REQUIRED) find_package(rocblas REQUIRED) +# Optional: the RDNA4 Q1_0/Q2_0 prefill routes (mul_mat_q{1,2}_0_hipblaslt.cu) +# use hipBLASLt. It is Linux-only in current ROCm distributions, so it must not +# be a hard requirement -- without it the routes compile to inert stubs. +find_package(hipblaslt QUIET) + if (GGML_HIP_RCCL) find_package(rccl REQUIRED) endif() @@ -155,3 +160,11 @@ if (GGML_HIP_RCCL) endif() target_link_libraries(ggml-hip PRIVATE ggml-base hip::host roc::rocblas roc::hipblas) + +if (hipblaslt_FOUND) + message(STATUS "hipBLASLt found: RDNA4 Q1_0/Q2_0 hipBLASLt prefill routes enabled") + target_link_libraries(ggml-hip PRIVATE roc::hipblaslt) +else() + message(STATUS "hipBLASLt not found: RDNA4 Q1_0/Q2_0 hipBLASLt prefill routes disabled") + target_compile_definitions(ggml-hip PRIVATE GGML_HIP_NO_HIPBLASLT) +endif() From ce5a32be370f4b6744cf9f699dac67d4041d5acb Mon Sep 17 00:00:00 2001 From: Justin Monk Date: Tue, 11 Aug 2026 02:09:37 -0400 Subject: [PATCH 2/3] cuda: HIP-path vec_dot_q1_0_q8_1 via the 2c-1 identity (+16% Q1_0 decode) Same treatment as the Q2_0 HIP path: bit-spread raw code bits c in {0,1} into {0,1}-byte dp4a operands and fold dot(s,u) = 2*dot(c,u) - sum(u) through the q8_1 stored sum, replacing the select-chain fallback in unpack_q1_0_bytes. Bonsai-27B Q1_0 tg128 on 1x R9700: 54.53 -> 63.03 t/s (64.93 with GGML_HIP_DEDUP_MMVQ_QUANT=1). test-backend-ops MUL_MAT: all pass. CUDA path untouched. Co-Authored-By: Claude Fable 5 --- ggml/src/ggml-cuda/vecdotq.cuh | 26 +++++++++++++++++++++++++- 1 file changed, 25 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh index 68711df93d1..7887f4cdc37 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -712,11 +712,34 @@ static __device__ __forceinline__ float vec_dot_q1_0_q8_1( // iqs selects which of the 4 chunks of 32 elements to process (0-3) const float d1 = bq1_0->d; - const uint16_t * qs = (const uint16_t *) bq1_0->qs + iqs * 2; // Process only the chunk specified by iqs const block_q8_1 * bq8_1_chunk = bq8_1 + iqs; +#if defined(GGML_USE_HIP) || defined(GGML_USE_MUSA) + // AMD path, same identity as the Q2_0 HIP path above: with raw code bits + // c in {0,1}, s = 2c - 1, so dot(s,u) = 2*dot(c,u) - sum(u). Bit-spread + // each qs byte into two {0,1}-byte dp4a operands (4 fused shift-or ops + // each) and apply -sum(u) once at the end via the q8_1 stored sum + // (ds.y = d8*sum(u)). This replaces the select-chain fallback in + // unpack_q1_0_bytes on a VALU-bound kernel (~2.5x fewer VALU ops). + const int offset = iqs * 4; + int sumi = 0; // = dot(c, u), c in {0,1} +#pragma unroll + for (int j2 = 0; j2 < 4; ++j2) { + const int b = bq1_0->qs[offset + j2]; + const int lo = ( b | (b << 7) | (b << 14) | (b << 21)) & 0x01010101; // bits 0..3 -> bytes + const int hi = ((b >> 4) | (b << 3) | (b << 10) | (b << 17)) & 0x01010101; // bits 4..7 -> bytes + sumi = ggml_cuda_dp4a(lo, get_int_b4(bq8_1_chunk->qs, 2*j2 + 0), sumi); + sumi = ggml_cuda_dp4a(hi, get_int_b4(bq8_1_chunk->qs, 2*j2 + 1), sumi); + } + + const float d8 = __low2float(bq8_1_chunk->ds); + const float s8 = __high2float(bq8_1_chunk->ds); // = d8 * sum(u) + return d1 * (2.0f * d8 * (float) sumi - s8); +#else + const uint16_t * qs = (const uint16_t *) bq1_0->qs + iqs * 2; + int sumi = 0; #pragma unroll for (int j = 0; j < 2; ++j) { @@ -735,6 +758,7 @@ static __device__ __forceinline__ float vec_dot_q1_0_q8_1( const float d8 = __low2float(bq8_1_chunk->ds); return d1 * d8 * sumi; +#endif // defined(GGML_USE_HIP) || defined(GGML_USE_MUSA) } static __device__ __forceinline__ float vec_dot_q2_0_q8_1( From 65cd5bd2945387b4cce602dbd265687ef438a97c Mon Sep 17 00:00:00 2001 From: jmonk Date: Tue, 11 Aug 2026 14:09:58 -0400 Subject: [PATCH 3/3] hipblaslt wcache: enforce cross-stream ordering of cached weight builds Review hardening (Copilot flag on the converted-weight caches): entries are published under the mutex right after the requant kernel is launched, so ordering vs the GEMM that reads them was guaranteed only for same-stream consumers. Record a build_done event on the build stream and make any consumer on a different stream hipStreamWaitEvent on it before using the entry (falling back to on-the-fly requant if the wait cannot be issued). Events are destroyed with their entries in the invalidators. No behavior change on the current single-compute-stream-per-device backend (the wait never fires); the previously documented invariant is now enforced. Verified: build clean, both routes smoke-tested on 1x R9700 (Q1_0 pp1024 1500 t/s, Q2_0 1487 t/s, build+hit paths exercised). CCA (Claude Code Augmented) --- ggml/src/ggml-cuda/mul_mat_q1_0_hipblaslt.cu | 49 ++++++++++++++++---- ggml/src/ggml-cuda/mul_mat_q2_0_hipblaslt.cu | 28 +++++++++-- 2 files changed, 65 insertions(+), 12 deletions(-) diff --git a/ggml/src/ggml-cuda/mul_mat_q1_0_hipblaslt.cu b/ggml/src/ggml-cuda/mul_mat_q1_0_hipblaslt.cu index f543d9f9d64..c9bd51aa0b3 100644 --- a/ggml/src/ggml-cuda/mul_mat_q1_0_hipblaslt.cu +++ b/ggml/src/ggml-cuda/mul_mat_q1_0_hipblaslt.cu @@ -393,8 +393,17 @@ const lt_plan & get_plan(int64_t N, int64_t M, int64_t K, int mode) { // Q1_0 on a 32 GB card, so cap the cache at a VRAM budget: weights that fit are // cached (requant paid once), the rest fall back to on-the-fly pool requant. // hipMalloc failure also falls back -- never OOM-crash. -struct cached_w { int8_t * q8 = nullptr; float * wscale = nullptr; size_t bytes = 0; }; -struct cached_w_f8 { uint8_t * q8f8 = nullptr; float * wscale = nullptr; size_t bytes = 0; }; +// build_stream/build_done: entries are published under the mutex right after +// the requant kernel is LAUNCHED, not after it completes. Same-stream +// consumers are ordered by the stream itself; a consumer on any other stream +// must wait on build_done before its GEMM reads the converted weights +// (review hardening -- today's backend uses one compute stream per device, +// so the wait never fires, but the invariant is now enforced rather than +// assumed). +struct cached_w { int8_t * q8 = nullptr; float * wscale = nullptr; size_t bytes = 0; + hipStream_t build_stream = nullptr; hipEvent_t build_done = nullptr; }; +struct cached_w_f8 { uint8_t * q8f8 = nullptr; float * wscale = nullptr; size_t bytes = 0; + hipStream_t build_stream = nullptr; hipEvent_t build_done = nullptr; }; std::map g_wcache_i8; std::map g_wcache_f8; size_t g_wcache_bytes = 0; @@ -417,7 +426,13 @@ const cached_w * try_cache_weight_i8(const void * key, const char * wdata, int64 int64_t N, int64_t K, int64_t n_blocks, cudaStream_t stream) { std::lock_guard lk(g_wcache_mtx); auto it = g_wcache_i8.find(key); - if (it != g_wcache_i8.end()) return &it->second; + if (it != g_wcache_i8.end()) { + if (it->second.build_stream != stream && + hipStreamWaitEvent(stream, it->second.build_done, 0) != hipSuccess) { + return nullptr; // can't prove ordering -> caller requants on the fly + } + return &it->second; + } const size_t need = (size_t)N * K + (size_t)N * sizeof(float); if (g_wcache_bytes + need > wcache_budget_bytes()) return nullptr; // budget hit @@ -425,9 +440,14 @@ const cached_w * try_cache_weight_i8(const void * key, const char * wdata, int64 cached_w c; if (hipMalloc(&c.q8, (size_t)N * K) != hipSuccess) return nullptr; if (hipMalloc(&c.wscale, (size_t)N * sizeof(float)) != hipSuccess) { hipFree(c.q8); return nullptr; } + if (hipEventCreateWithFlags(&c.build_done, hipEventDisableTiming) != hipSuccess) { + hipFree(c.q8); hipFree(c.wscale); return nullptr; + } c.bytes = need; + c.build_stream = stream; const dim3 grid((unsigned)N), block(256); k_requant_q1_0_to_int8_perchannel<<>>(wdata, nb01, c.q8, c.wscale, K, n_blocks); + hipEventRecord(c.build_done, stream); g_wcache_bytes += need; auto res = g_wcache_i8.emplace(key, c); return &res.first->second; @@ -437,7 +457,13 @@ const cached_w_f8 * try_cache_weight_f8(const void * key, const char * wdata, in int64_t N, int64_t K, int64_t n_blocks, cudaStream_t stream) { std::lock_guard lk(g_wcache_mtx); auto it = g_wcache_f8.find(key); - if (it != g_wcache_f8.end()) return &it->second; + if (it != g_wcache_f8.end()) { + if (it->second.build_stream != stream && + hipStreamWaitEvent(stream, it->second.build_done, 0) != hipSuccess) { + return nullptr; // can't prove ordering -> caller requants on the fly + } + return &it->second; + } const size_t need = (size_t)N * K + (size_t)N * sizeof(float); if (g_wcache_bytes + need > wcache_budget_bytes()) return nullptr; // budget hit @@ -445,9 +471,14 @@ const cached_w_f8 * try_cache_weight_f8(const void * key, const char * wdata, in cached_w_f8 c; if (hipMalloc(&c.q8f8, (size_t)N * K) != hipSuccess) return nullptr; if (hipMalloc(&c.wscale, (size_t)N * sizeof(float)) != hipSuccess) { hipFree(c.q8f8); return nullptr; } + if (hipEventCreateWithFlags(&c.build_done, hipEventDisableTiming) != hipSuccess) { + hipFree(c.q8f8); hipFree(c.wscale); return nullptr; + } c.bytes = need; + c.build_stream = stream; const dim3 grid((unsigned)N), block(256); k_requant_q1_0_to_e4m3_perchannel<<>>(wdata, nb01, c.q8f8, c.wscale, K, n_blocks); + hipEventRecord(c.build_done, stream); g_wcache_bytes += need; auto res = g_wcache_f8.emplace(key, c); return &res.first->second; @@ -462,8 +493,9 @@ void wcache_invalidate_range(const void * base, size_t size) { for (auto it = g_wcache_i8.begin(); it != g_wcache_i8.end(); ) { const char * k = (const char *) it->first; if (k >= b && k < b + size) { - if (it->second.q8) hipFree(it->second.q8); - if (it->second.wscale) hipFree(it->second.wscale); + if (it->second.q8) hipFree(it->second.q8); + if (it->second.wscale) hipFree(it->second.wscale); + if (it->second.build_done) hipEventDestroy(it->second.build_done); g_wcache_bytes -= it->second.bytes; it = g_wcache_i8.erase(it); } else { ++it; } @@ -471,8 +503,9 @@ void wcache_invalidate_range(const void * base, size_t size) { for (auto it = g_wcache_f8.begin(); it != g_wcache_f8.end(); ) { const char * k = (const char *) it->first; if (k >= b && k < b + size) { - if (it->second.q8f8) hipFree(it->second.q8f8); - if (it->second.wscale) hipFree(it->second.wscale); + if (it->second.q8f8) hipFree(it->second.q8f8); + if (it->second.wscale) hipFree(it->second.wscale); + if (it->second.build_done) hipEventDestroy(it->second.build_done); g_wcache_bytes -= it->second.bytes; it = g_wcache_f8.erase(it); } else { ++it; } diff --git a/ggml/src/ggml-cuda/mul_mat_q2_0_hipblaslt.cu b/ggml/src/ggml-cuda/mul_mat_q2_0_hipblaslt.cu index c5c9d725c71..b4541990533 100644 --- a/ggml/src/ggml-cuda/mul_mat_q2_0_hipblaslt.cu +++ b/ggml/src/ggml-cuda/mul_mat_q2_0_hipblaslt.cu @@ -381,7 +381,15 @@ const lt_plan & get_plan(int64_t N, int64_t M, int64_t K, int mode) { // Q2_0 on a 32 GB card, so cap the cache at a VRAM budget: weights that fit are // cached (requant paid once), the rest fall back to on-the-fly pool requant. // hipMalloc failure also falls back -- never OOM-crash (the Stage-3 lesson). -struct cached_w { int8_t * q8 = nullptr; float * wscale = nullptr; size_t bytes = 0; }; +// build_stream/build_done: entries are published under the mutex right after +// the requant kernel is LAUNCHED, not after it completes. Same-stream +// consumers are ordered by the stream itself; a consumer on any other stream +// must wait on build_done before its GEMM reads the converted weights +// (review hardening -- today's backend uses one compute stream per device, +// so the wait never fires, but the invariant is now enforced rather than +// assumed). +struct cached_w { int8_t * q8 = nullptr; float * wscale = nullptr; size_t bytes = 0; + hipStream_t build_stream = nullptr; hipEvent_t build_done = nullptr; }; std::map g_wcache; size_t g_wcache_bytes = 0; std::mutex g_wcache_mtx; @@ -403,7 +411,13 @@ const cached_w * try_cache_weight(const void * key, const char * wdata, int64_t int64_t N, int64_t K, int64_t n_blocks, int mode, cudaStream_t stream) { std::lock_guard lk(g_wcache_mtx); auto it = g_wcache.find(key); - if (it != g_wcache.end()) return &it->second; // mode is fixed per process (env-checked once) + if (it != g_wcache.end()) { // mode is fixed per process (env-checked once) + if (it->second.build_stream != stream && + hipStreamWaitEvent(stream, it->second.build_done, 0) != hipSuccess) { + return nullptr; // can't prove ordering -> caller requants on the fly + } + return &it->second; + } const size_t need = (size_t)N * K + (size_t)N * sizeof(float); if (g_wcache_bytes + need > wcache_budget_bytes()) return nullptr; // budget hit @@ -413,13 +427,18 @@ const cached_w * try_cache_weight(const void * key, const char * wdata, int64_t cached_w c; if (hipMalloc(&c.q8, (size_t)N * K) != hipSuccess) return nullptr; if (hipMalloc(&c.wscale, (size_t)N * sizeof(float)) != hipSuccess) { hipFree(c.q8); return nullptr; } + if (hipEventCreateWithFlags(&c.build_done, hipEventDisableTiming) != hipSuccess) { + hipFree(c.q8); hipFree(c.wscale); return nullptr; + } c.bytes = need; + c.build_stream = stream; const dim3 grid((unsigned)N), block(256); if (mode == 1 /*MODE_F8*/) { k_requant_q2_0_to_e4m3_perchannel<<>>(wdata, nb01, (uint8_t *)c.q8, c.wscale, K, n_blocks); } else { k_requant_q2_0_to_int8_perchannel<<>>(wdata, nb01, c.q8, c.wscale, K, n_blocks); } + hipEventRecord(c.build_done, stream); g_wcache_bytes += need; auto res = g_wcache.emplace(key, c); return &res.first->second; @@ -434,8 +453,9 @@ void wcache_invalidate_range(const void * base, size_t size) { for (auto it = g_wcache.begin(); it != g_wcache.end(); ) { const char * k = (const char *) it->first; if (k >= b && k < b + size) { - if (it->second.q8) hipFree(it->second.q8); - if (it->second.wscale) hipFree(it->second.wscale); + if (it->second.q8) hipFree(it->second.q8); + if (it->second.wscale) hipFree(it->second.wscale); + if (it->second.build_done) hipEventDestroy(it->second.build_done); g_wcache_bytes -= it->second.bytes; it = g_wcache.erase(it); } else {