From e3d9df5bba5648a2da61b18c93a10f400be0a29a Mon Sep 17 00:00:00 2001 From: sb32445 Date: Sat, 3 Oct 2026 11:09:51 +0200 Subject: [PATCH 1/3] cuda: prefetch the next PTQ1_0 mat-vec's weights into L2 from the last CTAs Every PTQ1_0 mat-vec in the decode graph pays about 2.5 us of ramp-up and ramp-down in which the DRAM is not busy. The node loop now looks ahead for the next PTQ1_0 MUL_MAT (a gate/up pair feeding one GLU counts as one op) and hands its weight pointer to the launcher. The last 46 CTAs of the running kernel issue prefetch.global.L2 for the first 50 % (at most 16 MiB) of those weights right after their own loads. More CTAs or more bytes compete with the running kernel and get slower (184 CTAs: -2.3 %). Values are not changed. GGML_CUDA_L2_PREFETCH_PCT=0 turns it off; _CTAS and _MAX_KB tune it. Decode with MTP n-max 2, same binary, outputs identical: +2.00 % (greedy benchmark), +2.10 % (Hermes setup, ctx 114688), +2.2 / +2.1 / +1.7 % at depth 0 / 16k / 65k. Co-Authored-By: Claude Sonnet 5.5 --- ggml/src/ggml-cuda/ggml-cuda.cu | 34 ++++++++++++++++++++++++++++++ ggml/src/ggml-cuda/l2-hint.cuh | 11 ++++++++++ ggml/src/ggml-cuda/mmvq-ptq1_0.cuh | 33 +++++++++++++++++++++++++++-- ggml/src/ggml-cuda/mmvq.cu | 3 +++ 4 files changed, 79 insertions(+), 2 deletions(-) create mode 100644 ggml/src/ggml-cuda/l2-hint.cuh diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 41b3be6317d8..8886c824a0d0 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -32,6 +32,7 @@ #include "ggml-cuda/mmq.cuh" #include "ggml-cuda/mmvf.cuh" #include "ggml-cuda/mmvq.cuh" +#include "ggml-cuda/l2-hint.cuh" #include "ggml-cuda/norm.cuh" #include "ggml-cuda/opt-step-adamw.cuh" #include "ggml-cuda/opt-step-sgd.cuh" @@ -3200,6 +3201,38 @@ static bool ggml_cuda_topk_moe_fusion(const struct ggml_cgraph * cgraph, int nod return true; } +// L2 prefetch hint (see mmvq-ptq1_0.cuh): the weights of the next PTQ1_0 mat-vec after node i. Gate/up pairs feeding one GLU +// count as one op, so the partner of a fused gate/up kernel is skipped. +static ggml_cuda_l2_hint_t ggml_cuda_l2_hint_for_node(const ggml_cgraph * cgraph, const int i) { + static const bool enabled = [] { const char * e = getenv("GGML_CUDA_L2_PREFETCH_PCT"); return !e || atoi(e) > 0; }(); + ggml_cuda_l2_hint_t hint; + const ggml_tensor * cur = cgraph->nodes[i]; + if (!enabled || cur->op != GGML_OP_MUL_MAT || cur->src[0]->type != GGML_TYPE_PTQ1_0) { + return hint; + } + const int last = std::min(i + 64, cgraph->n_nodes - 1); + for (int j = i + 1; j <= last; j++) { + const ggml_tensor * nj = cgraph->nodes[j]; + if (nj->op != GGML_OP_MUL_MAT || nj->src[0]->type != GGML_TYPE_PTQ1_0 || nj->src[0]->data == cur->src[0]->data) { + continue; + } + bool partner = false; + if (nj->src[1] == cur->src[1]) { + for (int k = i + 1; k <= std::min(j + 8, cgraph->n_nodes - 1) && !partner; k++) { + const ggml_tensor * g = cgraph->nodes[k]; + partner = g->op == GGML_OP_GLU && ((g->src[0] == cur && g->src[1] == nj) || (g->src[0] == nj && g->src[1] == cur)); + } + } + if (partner) { + continue; + } + hint.ptr = (const char *) nj->src[0]->data; + hint.bytes = ggml_nbytes(nj->src[0]); + break; + } + return hint; +} + // returns whether the write (out) nodes overwrite the read nodes in operation static bool ggml_cuda_check_fusion_memory_ranges(const ggml_cgraph * cgraph, const int node_idx, @@ -4552,6 +4585,7 @@ static void ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context * cud } prev_i = i; + g_ggml_cuda_l2_hint = ggml_cuda_l2_hint_for_node(cgraph, i); if (ggml_cuda_is_view_or_noop(node)) { continue; diff --git a/ggml/src/ggml-cuda/l2-hint.cuh b/ggml/src/ggml-cuda/l2-hint.cuh new file mode 100644 index 000000000000..d1b99a641509 --- /dev/null +++ b/ggml/src/ggml-cuda/l2-hint.cuh @@ -0,0 +1,11 @@ +#pragma once + +#include + +// L2 prefetch hint: weights of the next PTQ1_0 mat-vec in the graph, set by the node loop before it dispatches a node, +// read by the PTQ1_0 mat-vec launcher (same thread). nullptr = no hint. +struct ggml_cuda_l2_hint_t { + const char * ptr = nullptr; + size_t bytes = 0; +}; +extern thread_local ggml_cuda_l2_hint_t g_ggml_cuda_l2_hint; diff --git a/ggml/src/ggml-cuda/mmvq-ptq1_0.cuh b/ggml/src/ggml-cuda/mmvq-ptq1_0.cuh index 4755d7a86d00..4cf399db9156 100644 --- a/ggml/src/ggml-cuda/mmvq-ptq1_0.cuh +++ b/ggml/src/ggml-cuda/mmvq-ptq1_0.cuh @@ -25,9 +25,13 @@ #pragma once #include "common.cuh" +#include "l2-hint.cuh" #include "unary.cuh" #include "vecdotq.cuh" +#include +#include + #define PTQ1_0_PT_PLANES 9 // dedicated 2D kernel geometry, see mul_mat_vec_ptq1_0_pt below @@ -298,7 +302,8 @@ static __global__ void mul_mat_vec_ptq1_0_pt( const void * vx_, const void * vy_, const ggml_cuda_mm_fusion_args_device fusion, float * dst_, const int ncols_x, const int nrows_x, const int stride_row_x, const int stride_col_y, const int stride_col_dst, - const int rows_per_cta, const uint3 bpr_fd, const uint3 rpc_fd, const bool invariant) { + const int rows_per_cta, const uint3 bpr_fd, const uint3 rpc_fd, const bool invariant, + const char * pf_ptr, const int pf_lines_per_cta, const int pf_ctas) { // GGML_CUDA_RESTRICT stays off the formal parameters: cudafe's host stub drops __restrict // from the explicit specialization and MSVC/GCC then reject it (C2912 / "does not match // any template declaration") when compiling sm_90/sm_120. Same pattern as mul_mat_vec_q. @@ -366,6 +371,15 @@ static __global__ void mul_mat_vec_ptq1_0_pt( } } + // the last CTAs have finished their loads and the DRAM is about to go idle: pull the head of the next + // mat-vec's weights into L2 (prefetch.global.L2 keeps the loads off the critical path of this kernel) + if (pf_ptr != nullptr && (int) blockIdx.x + pf_ctas >= (int) gridDim.x) { + const char * pf = pf_ptr + (size_t) ((int) blockIdx.x - ((int) gridDim.x - pf_ctas)) * pf_lines_per_cta * 128; + for (int i = tid; i < pf_lines_per_cta; i += PTQ1_0_PT_THREADS) { + asm volatile("prefetch.global.L2 [%0];" :: "l"(pf + (size_t) i * 128)); + } + } + __syncthreads(); if (invariant) { @@ -504,10 +518,25 @@ static void mul_mat_vec_ptq1_0_pt_launch( const size_t smem = ptq1_0_pt_smem_bytes(bpr, ncols, nrows_x, has_gate); const ggml_cuda_kernel_launch_params lp = ggml_cuda_kernel_launch_params(block_nums, block_dims, smem, stream); + // L2 prefetch of the next mat-vec's weights (hint from the node loop; GGML_CUDA_L2_PREFETCH_PCT/_MAX_KB/_CTAS, PCT=0 turns it off; defaults 50 / 16384 / 46) + static const int pf_pct = [] { const char * e = getenv("GGML_CUDA_L2_PREFETCH_PCT"); return e ? atoi(e) : 50; }(); + static const int pf_max_kb = [] { const char * e = getenv("GGML_CUDA_L2_PREFETCH_MAX_KB"); return e ? atoi(e) : 16384; }(); + static const int pf_cta_cfg = [] { const char * e = getenv("GGML_CUDA_L2_PREFETCH_CTAS"); return e ? atoi(e) : 46; }(); + const char * pf_ptr = nullptr; + int pf_lines_per_cta = 0; + const int pf_ctas = std::min(pf_cta_cfg, (int) block_nums.x); + if (pf_pct > 0 && g_ggml_cuda_l2_hint.ptr != nullptr && pf_ctas > 0) { + const size_t pf_bytes = std::min(g_ggml_cuda_l2_hint.bytes / 100 * pf_pct, (size_t) pf_max_kb * 1024); + pf_lines_per_cta = (int) (pf_bytes / 128 / pf_ctas); + if (pf_lines_per_cta > 0) { + pf_ptr = g_ggml_cuda_l2_hint.ptr; + } + } + #define PTQ1_0_PT_LAUNCH(FUS, GATE) \ ggml_cuda_kernel_launch(mul_mat_vec_ptq1_0_pt, lp, \ vx, vy, fusion, dst, ncols_x, nrows_x, stride_row_x, stride_col_y, stride_col_dst, rows_per_cta, bpr_fd, rpc_fd, \ - ggml_cuda_batch_invariant()) + ggml_cuda_batch_invariant(), pf_ptr, pf_lines_per_cta, pf_ctas) if (has_fusion) { GGML_ASSERT(ncols == 1 && "fusion only supported for ncols_dst=1"); diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 0008fdefc4b5..48ace2b92041 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -1,4 +1,5 @@ #include "mmvq.cuh" +#include "l2-hint.cuh" #include "mmvq-ptq1_0.cuh" #include "quantize.cuh" #include "unary.cuh" @@ -7,6 +8,8 @@ #include #include +thread_local ggml_cuda_l2_hint_t g_ggml_cuda_l2_hint; + typedef float (*vec_dot_q_cuda_t)(const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs); static constexpr __device__ vec_dot_q_cuda_t get_vec_dot_q_cuda(ggml_type type) { From a4249f5d57a76429cf43fb030d8b529661911f0c Mon Sep 17 00:00:00 2001 From: sb32445 Date: Sat, 3 Oct 2026 14:58:26 +0200 Subject: [PATCH 2/3] cuda: fix the L2 prefetch parameters and remove its switches GGML_CUDA_L2_PREFETCH_PCT, _MAX_KB and _CTAS of the previous commit were only there to tune and measure the change. The measured values (50 % of the next tensor, at most 16 MiB, 46 CTAs) are now constants. Co-Authored-By: Claude Sonnet 5.5 --- ggml/src/ggml-cuda/ggml-cuda.cu | 3 +-- ggml/src/ggml-cuda/mmvq-ptq1_0.cuh | 19 +++++++++++-------- 2 files changed, 12 insertions(+), 10 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 8886c824a0d0..ba65aa9a8a98 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -3204,10 +3204,9 @@ static bool ggml_cuda_topk_moe_fusion(const struct ggml_cgraph * cgraph, int nod // L2 prefetch hint (see mmvq-ptq1_0.cuh): the weights of the next PTQ1_0 mat-vec after node i. Gate/up pairs feeding one GLU // count as one op, so the partner of a fused gate/up kernel is skipped. static ggml_cuda_l2_hint_t ggml_cuda_l2_hint_for_node(const ggml_cgraph * cgraph, const int i) { - static const bool enabled = [] { const char * e = getenv("GGML_CUDA_L2_PREFETCH_PCT"); return !e || atoi(e) > 0; }(); ggml_cuda_l2_hint_t hint; const ggml_tensor * cur = cgraph->nodes[i]; - if (!enabled || cur->op != GGML_OP_MUL_MAT || cur->src[0]->type != GGML_TYPE_PTQ1_0) { + if (cur->op != GGML_OP_MUL_MAT || cur->src[0]->type != GGML_TYPE_PTQ1_0) { return hint; } const int last = std::min(i + 64, cgraph->n_nodes - 1); diff --git a/ggml/src/ggml-cuda/mmvq-ptq1_0.cuh b/ggml/src/ggml-cuda/mmvq-ptq1_0.cuh index 4cf399db9156..a8677dc3ffef 100644 --- a/ggml/src/ggml-cuda/mmvq-ptq1_0.cuh +++ b/ggml/src/ggml-cuda/mmvq-ptq1_0.cuh @@ -30,7 +30,6 @@ #include "vecdotq.cuh" #include -#include #define PTQ1_0_PT_PLANES 9 @@ -40,6 +39,13 @@ #define PTQ1_0_PT_MAX_COLS 8 // equals MMVQ_MAX_BATCH_SIZE, checked in mmvq.cu #define PTQ1_0_PT_SMEM_FLOATS 4096 // 16 KiB target when choosing rows per CTA; the launch may request more for one item +// L2 prefetch of the next mat-vec's weights by the last CTAs of the running kernel (see mul_mat_vec_ptq1_0_pt and +// ggml_cuda_l2_hint_for_node): percent of the next tensor, upper bound in bytes, number of prefetching CTAs. +// Measured on an RTX 4070: more CTAs or more bytes take DRAM bandwidth from the running kernel and are slower. +#define PTQ1_0_L2_PREFETCH_PCT 50 +#define PTQ1_0_L2_PREFETCH_BYTES (16u << 20) +#define PTQ1_0_L2_PREFETCH_CTAS 46 + // the PT path is CUDA only; HIP keeps the block_q8_1 layout and the old vec_dot static constexpr __host__ __device__ bool ptq1_0_pt_enabled() { #if defined(GGML_USE_HIP) @@ -518,15 +524,12 @@ static void mul_mat_vec_ptq1_0_pt_launch( const size_t smem = ptq1_0_pt_smem_bytes(bpr, ncols, nrows_x, has_gate); const ggml_cuda_kernel_launch_params lp = ggml_cuda_kernel_launch_params(block_nums, block_dims, smem, stream); - // L2 prefetch of the next mat-vec's weights (hint from the node loop; GGML_CUDA_L2_PREFETCH_PCT/_MAX_KB/_CTAS, PCT=0 turns it off; defaults 50 / 16384 / 46) - static const int pf_pct = [] { const char * e = getenv("GGML_CUDA_L2_PREFETCH_PCT"); return e ? atoi(e) : 50; }(); - static const int pf_max_kb = [] { const char * e = getenv("GGML_CUDA_L2_PREFETCH_MAX_KB"); return e ? atoi(e) : 16384; }(); - static const int pf_cta_cfg = [] { const char * e = getenv("GGML_CUDA_L2_PREFETCH_CTAS"); return e ? atoi(e) : 46; }(); + // L2 prefetch of the next mat-vec's weights (hint from the node loop) const char * pf_ptr = nullptr; int pf_lines_per_cta = 0; - const int pf_ctas = std::min(pf_cta_cfg, (int) block_nums.x); - if (pf_pct > 0 && g_ggml_cuda_l2_hint.ptr != nullptr && pf_ctas > 0) { - const size_t pf_bytes = std::min(g_ggml_cuda_l2_hint.bytes / 100 * pf_pct, (size_t) pf_max_kb * 1024); + const int pf_ctas = std::min(PTQ1_0_L2_PREFETCH_CTAS, (int) block_nums.x); + if (g_ggml_cuda_l2_hint.ptr != nullptr) { + const size_t pf_bytes = std::min(g_ggml_cuda_l2_hint.bytes / 100 * PTQ1_0_L2_PREFETCH_PCT, (size_t) PTQ1_0_L2_PREFETCH_BYTES); pf_lines_per_cta = (int) (pf_bytes / 128 / pf_ctas); if (pf_lines_per_cta > 0) { pf_ptr = g_ggml_cuda_l2_hint.ptr; From e697877d9c5e4db5f312f41625b082e3374b1015 Mon Sep 17 00:00:00 2001 From: sb32445 Date: Wed, 7 Oct 2026 22:53:12 +0200 Subject: [PATCH 3/3] cuda : cap the PTQ1_0 L2 prefetch at 8 MiB on Ada and 2 MiB elsewhere The cap of 16 MiB and the fixed 46 CTAs were tuned on an RTX 4070 only. On an RTX 2060 SUPER (4 MiB L2, 34 SMs) that setting is 2.4 to 2.7 % slower than no prefetch (measured by a reviewer, see the PR). Prefetch with one CTA per SM, and cap the bytes at 8 MiB on Ada (cc 8.9) and at 2 MiB on every other GPU. RTX 4070 (36 MiB L2), MTP n-max 2, same outputs in all runs, 12 pairs against the previous setting: 8 MiB +0.02 % (n.s.), 16 MiB -0.03 % (n.s.), 4 MiB -0.17 %, 2 MiB -0.40 %. The 2 MiB default is not measured on other GPUs. The 4070 has 46 SMs, so one CTA per SM is the old count. --- ggml/src/ggml-cuda/mmvq-ptq1_0.cuh | 16 +++++++++------- 1 file changed, 9 insertions(+), 7 deletions(-) diff --git a/ggml/src/ggml-cuda/mmvq-ptq1_0.cuh b/ggml/src/ggml-cuda/mmvq-ptq1_0.cuh index a8677dc3ffef..7d55bb655eda 100644 --- a/ggml/src/ggml-cuda/mmvq-ptq1_0.cuh +++ b/ggml/src/ggml-cuda/mmvq-ptq1_0.cuh @@ -40,11 +40,11 @@ #define PTQ1_0_PT_SMEM_FLOATS 4096 // 16 KiB target when choosing rows per CTA; the launch may request more for one item // L2 prefetch of the next mat-vec's weights by the last CTAs of the running kernel (see mul_mat_vec_ptq1_0_pt and -// ggml_cuda_l2_hint_for_node): percent of the next tensor, upper bound in bytes, number of prefetching CTAs. -// Measured on an RTX 4070: more CTAs or more bytes take DRAM bandwidth from the running kernel and are slower. -#define PTQ1_0_L2_PREFETCH_PCT 50 -#define PTQ1_0_L2_PREFETCH_BYTES (16u << 20) -#define PTQ1_0_L2_PREFETCH_CTAS 46 +// ggml_cuda_l2_hint_for_node): percent of the next tensor and upper bound in bytes. One CTA per SM prefetches. +// 8 MiB costs nothing on an RTX 4070 (36 MiB L2, measured). Other GPUs get 2 MiB, not measured there. +#define PTQ1_0_L2_PREFETCH_PCT 50 +#define PTQ1_0_L2_PREFETCH_BYTES_ADA (8u << 20) +#define PTQ1_0_L2_PREFETCH_BYTES (2u << 20) // the PT path is CUDA only; HIP keeps the block_q8_1 layout and the old vec_dot static constexpr __host__ __device__ bool ptq1_0_pt_enabled() { @@ -527,9 +527,11 @@ static void mul_mat_vec_ptq1_0_pt_launch( // L2 prefetch of the next mat-vec's weights (hint from the node loop) const char * pf_ptr = nullptr; int pf_lines_per_cta = 0; - const int pf_ctas = std::min(PTQ1_0_L2_PREFETCH_CTAS, (int) block_nums.x); + const auto & pf_dev = ggml_cuda_info().devices[ggml_cuda_get_device()]; + const int pf_ctas = std::min(pf_dev.nsm, (int) block_nums.x); if (g_ggml_cuda_l2_hint.ptr != nullptr) { - const size_t pf_bytes = std::min(g_ggml_cuda_l2_hint.bytes / 100 * PTQ1_0_L2_PREFETCH_PCT, (size_t) PTQ1_0_L2_PREFETCH_BYTES); + const size_t pf_cap = pf_dev.cc == GGML_CUDA_CC_ADA_LOVELACE ? PTQ1_0_L2_PREFETCH_BYTES_ADA : PTQ1_0_L2_PREFETCH_BYTES; + const size_t pf_bytes = std::min(g_ggml_cuda_l2_hint.bytes / 100 * PTQ1_0_L2_PREFETCH_PCT, pf_cap); pf_lines_per_cta = (int) (pf_bytes / 128 / pf_ctas); if (pf_lines_per_cta > 0) { pf_ptr = g_ggml_cuda_l2_hint.ptr;