From 8fd3a28d9456ac25250d795c1662fffe296a812f Mon Sep 17 00:00:00 2001 From: Marcos Damasceno Date: Thu, 24 Sep 2026 23:45:29 -0500 Subject: [PATCH 1/3] perf(cuda): the recurrent-state GET_ROWS folds into GATED_DELTA_NET's load, f32 or q8_0 cache (#220 plus -cts q8_0) --- ggml/src/ggml-cuda/common.cuh | 33 +++++++++++ ggml/src/ggml-cuda/gated_delta_net.cu | 58 +++++++++++++++----- ggml/src/ggml-cuda/ggml-cuda.cu | 79 +++++++++++++++++++++++++++ tests/test-backend-ops.cpp | 49 ++++++++++++++--- 4 files changed, 197 insertions(+), 22 deletions(-) diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 79b44b10a6b9..ce539ed687f1 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -1486,6 +1486,36 @@ struct ggml_cuda_stream_context { } }; +// Fused recurrent-state gather for GATED_DELTA_NET (PrismML-Eng/llama.cpp#220, plus a q8_0 cache): build_rs +// materialises GET_ROWS(cache, s_copy) into a temp that only the GDN kernel reads; the graph evaluator skips that +// GET_ROWS (ggml_cuda_try_gdn_gather_skip) and records the gather here, so the kernel reads each sequence's state +// row ids[seq] out of the cache itself. +struct ggml_cuda_gated_delta_net_gather { + const void * base = nullptr; // cache rows, f32 or q8_0 + const int32_t * ids = nullptr; // per-seq row index + int64_t row_stride = 0; // between rows, in elements + bool q8_0 = false; // a q8_0 cache (-cts q8_0): the kernel dequantizes as it loads +}; + +// Owned by the backend context that evaluates the graph: registrations are keyed by node pointer, so they only mean +// something for the evaluation that made them. Cleared at the start of every graph evaluation/capture. +struct ggml_cuda_gdn_gather_context { + std::unordered_map gathers; + + void reset() { + gathers.clear(); + } + + void set(const ggml_tensor * gdn, const ggml_cuda_gated_delta_net_gather & gather) { + gathers[gdn] = gather; + } + + const ggml_cuda_gated_delta_net_gather * find(const ggml_tensor * gdn) const { + const auto it = gathers.find(gdn); + return it == gathers.end() ? nullptr : &it->second; + } +}; + // Names one ggml graph across computes, to find its CUDA graph. The first node's address alone names the memory a // graph is built in, and graphs of different shapes built there in turn (a speculative draft context's catch-up batch // and its one-row draft step) would share one CUDA graph, each finding the other's capture stale on every compute and @@ -1589,6 +1619,7 @@ struct ggml_backend_cuda_context { } ggml_cuda_stream_context concurrent_stream_context; + ggml_cuda_gdn_gather_context gdn_gather_context; ~ggml_backend_cuda_context(); @@ -1604,6 +1635,8 @@ struct ggml_backend_cuda_context { ggml_cuda_stream_context & stream_context() { return concurrent_stream_context; } + ggml_cuda_gdn_gather_context & gdn_gathers() { return gdn_gather_context; } + cublasHandle_t cublas_handle() { if (cublas_handles[device][curr_stream_no] == nullptr) { ggml_cuda_set_device(device); diff --git a/ggml/src/ggml-cuda/gated_delta_net.cu b/ggml/src/ggml-cuda/gated_delta_net.cu index 9420d0da3434..42fc91b0a44c 100644 --- a/ggml/src/ggml-cuda/gated_delta_net.cu +++ b/ggml/src/ggml-cuda/gated_delta_net.cu @@ -36,6 +36,19 @@ static __device__ __forceinline__ void gdn_store_state(void * dst, const int64_t } } +// Reads element e of a state source: f32, or a q8_0 cache row (the fused gather, ggml_cuda_try_gdn_gather_skip). q8_0 +// is dequantized with get_rows' formula (dequantize_q8_0), so the kernel sees the values the skipped GET_ROWS wrote. +static __device__ __forceinline__ float gdn_load_state(const void * src, const int64_t e, const bool q8) { + if (q8) { + const block_q8_0 * b = (const block_q8_0 *) src + e / QK8_0; + const float d = b->d; + float v = b->qs[e % QK8_0]; + v *= d; + return v; + } + return ((const float *) src)[e]; +} + // RAW: beta and g arrive pre-activation (ggml_gated_delta_net_set_raw_gates); the kernel applies // sigmoid(beta) and raw_a[h] * softplus(g + raw_dt_bias[h]) with the unary kernels' formulas. // G_PRECOMPUTED: g already holds exp(g) (GB10 long-prompt path); only used with RAW == false. @@ -48,7 +61,7 @@ gated_delta_net_cuda(const float * q, const float * beta, const float * raw_dt_bias, const float * raw_a, - const float * curr_state, + const void * curr_state, float * dst, void * state, int64_t H, @@ -68,7 +81,10 @@ gated_delta_net_cuda(const float * q, float scale, int64_t state_slot_stride, int K, - int64_t attn_seq_stride) { + int64_t attn_seq_stride, + const int32_t * s_ids, + int64_t s_row_stride, + bool s_q8) { const uint32_t h_idx = blockIdx.x; const uint32_t sequence = blockIdx.y; // Each warp owns one or more columns, using warp-level primitives to reduce across rows. @@ -85,11 +101,10 @@ gated_delta_net_cuda(const float * q, float * attn_data = dst; - // input state holds s0 only: [S_v, S_v, H, n_seqs] — seq stride is D = H * S_v * S_v. + // input state holds s0 only: [S_v, S_v, H, n_seqs] — seq stride is D = H * S_v * S_v; with the fused gather + // (s_ids) it is the cache itself, this sequence's live state at row s_ids[sequence] (read after the PDL wait). // output state layout (per-slot D * n_seqs) — same per-(seq,head) offset as before. - const int64_t state_in_offset = sequence * H * S_v * S_v + h_idx * S_v * S_v; const int64_t state_out_offset = (sequence * H + h_idx) * S_v * S_v; // in elements of the destination - curr_state += state_in_offset; // attention rows of one sequence are attn_seq_stride apart: n_tokens * H * S_v, unless this launch // covers only the tail of each sequence (the chunked prefill path hands the last K-1 tokens here) attn_data += sequence * attn_seq_stride + h_idx * S_v; @@ -101,12 +116,14 @@ gated_delta_net_cuda(const float * q, // state is stored transposed: M[col][i] = S[i][col], row col is contiguous ggml_cuda_pdl_sync(); + const int64_t state_in_offset = (s_ids ? (int64_t) s_ids[sequence] * s_row_stride : (int64_t) sequence * H * S_v * S_v) + + h_idx * S_v * S_v; #pragma unroll for (int c = 0; c < cols_per_warp; ++c) { #pragma unroll for (int r = 0; r < rows_per_lane; r++) { const int i = r * warp_size + lane; - s_shard[c][r] = curr_state[(col + c) * S_v + i]; + s_shard[c][r] = gdn_load_state(curr_state, state_in_offset + (col + c) * S_v + i, s_q8); } } @@ -235,14 +252,15 @@ gated_delta_net_cuda(const float * q, template static void launch_gated_delta_net( const float * q_d, const float * k_d, const float * v_d, - const float * g_d, const float * b_d, const float * rb_d, const float * ra_d, const float * s_d, + const float * g_d, const float * b_d, const float * rb_d, const float * ra_d, const void * s_d, float * dst_d, void * state_d, int64_t S_v, int64_t H, int64_t n_tokens, int64_t n_seqs, int64_t sq1, int64_t sq2, int64_t sq3, int64_t sv1, int64_t sv2, int64_t sv3, int64_t sb1, int64_t sb2, int64_t sb3, int64_t neqk1, int64_t rq3, - float scale, int64_t state_slot_stride, int K, int64_t attn_seq_stride, cudaStream_t stream) { + float scale, int64_t state_slot_stride, int K, int64_t attn_seq_stride, + const int32_t * s_ids, int64_t s_row_stride, bool s_q8, cudaStream_t stream) { const int warp_size = ggml_cuda_info().devices[ggml_cuda_get_device()].warp_size; const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc; const int num_warps = 4; @@ -260,7 +278,7 @@ static void launch_gated_delta_net( ggml_cuda_kernel_launch(gated_delta_net_cuda<16, KDA, keep_rs_t, RAW, G_PRECOMPUTED, false>, launch_params, q_d, k_d, v_d, g_d, b_d, rb_d, ra_d, s_d, dst_d, state_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, - sb1, sb2, sb3, neqk1_magic, rq3_magic, scale, state_slot_stride, K, attn_seq_stride); + sb1, sb2, sb3, neqk1_magic, rq3_magic, scale, state_slot_stride, K, attn_seq_stride, s_ids, s_row_stride, s_q8); break; } GGML_ABORT("a q8_0 recurrent state needs S_v >= 32"); @@ -268,20 +286,20 @@ static void launch_gated_delta_net( ggml_cuda_kernel_launch(gated_delta_net_cuda<32, KDA, keep_rs_t, RAW, G_PRECOMPUTED, STATE_Q8>, launch_params, q_d, k_d, v_d, g_d, b_d, rb_d, ra_d, s_d, dst_d, state_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, - sb1, sb2, sb3, neqk1_magic, rq3_magic, scale, state_slot_stride, K, attn_seq_stride); + sb1, sb2, sb3, neqk1_magic, rq3_magic, scale, state_slot_stride, K, attn_seq_stride, s_ids, s_row_stride, s_q8); break; case 64: { ggml_cuda_kernel_launch(gated_delta_net_cuda<64, KDA, keep_rs_t, RAW, G_PRECOMPUTED, STATE_Q8>, launch_params, q_d, k_d, v_d, g_d, b_d, rb_d, ra_d, s_d, dst_d, state_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, - sb1, sb2, sb3, neqk1_magic, rq3_magic, scale, state_slot_stride, K, attn_seq_stride); + sb1, sb2, sb3, neqk1_magic, rq3_magic, scale, state_slot_stride, K, attn_seq_stride, s_ids, s_row_stride, s_q8); break; } case 128: { ggml_cuda_kernel_launch(gated_delta_net_cuda<128, KDA, keep_rs_t, RAW, G_PRECOMPUTED, STATE_Q8>, launch_params, q_d, k_d, v_d, g_d, b_d, rb_d, ra_d, s_d, dst_d, state_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, - sb1, sb2, sb3, neqk1_magic, rq3_magic, scale, state_slot_stride, K, attn_seq_stride); + sb1, sb2, sb3, neqk1_magic, rq3_magic, scale, state_slot_stride, K, attn_seq_stride, s_ids, s_row_stride, s_q8); break; } default: @@ -383,6 +401,14 @@ static void ggml_cuda_op_gated_delta_net_impl( const float * s_d = (const float *) src_state->data; float * dst_d = (float *) dst->data; + // the fused state gather, registered for this node by this context's graph evaluator (ggml_cuda_try_gdn_gather_skip): + // the state comes from the cache rows, not from src_state (the skipped GET_ROWS's temp, never written) + const ggml_cuda_gated_delta_net_gather * gather = ctx.gdn_gathers().find(dst); + const void * s_in = gather ? gather->base : (const void *) s_d; + const int32_t * s_ids = gather ? gather->ids : nullptr; + const int64_t s_row_stride = gather ? gather->row_stride : 0; + const bool s_q8 = gather != nullptr && gather->q8_0; + GGML_ASSERT(ggml_is_contiguous_rows(src_q)); GGML_ASSERT(ggml_is_contiguous_rows(src_k)); GGML_ASSERT(ggml_is_contiguous_rows(src_v)); @@ -434,6 +460,7 @@ static void ggml_cuda_op_gated_delta_net_impl( // the fused cache). if (ggml_cuda_should_use_chunked_gdn(dst)) { GGML_ASSERT(!state_q8); // the chunk pipeline writes f32 (ggml_cuda_try_gdn_cache_fusion keeps q8_0 off it) + GGML_ASSERT(!gather); // and reads a gathered s0 (ggml_cuda_try_gdn_gather_skip keeps the GET_ROWS for it) const int64_t n_tail = K - 1; const int64_t n_chunked = n_tokens - n_tail; float * chunk_state = (float *) state_d + n_tail * state_slot_stride; @@ -469,7 +496,7 @@ static void ggml_cuda_op_gated_delta_net_impl( launch_gated_delta_net(q_d + t0 * sq2, k_d + t0 * sq2, v_d + t0 * sv2, \ g_d + t0 * sb2, b_d + t0 * sb2, rb_d, ra_d, chunk_state, dst_d + t0 * S_v * H, state_d, \ S_v, H, n_tail, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, \ - sb1, sb2, sb3, neqk1, rq3, scale, state_slot_stride, K, S_v * H * n_tokens, stream) + sb1, sb2, sb3, neqk1, rq3, scale, state_slot_stride, K, S_v * H * n_tokens, nullptr, 0, false, stream) if (raw) { GDN_TAIL_LAUNCH(true); } else { GDN_TAIL_LAUNCH(false); } #undef GDN_TAIL_LAUNCH } @@ -492,9 +519,10 @@ static void ggml_cuda_op_gated_delta_net_impl( } #define GDN_LAUNCH(KDA_, KEEP_, RAW_, PRE_, Q8_) \ - launch_gated_delta_net(q_d, k_d, v_d, g_d, b_d, rb_d, ra_d, s_d, dst_d, state_d, \ + launch_gated_delta_net(q_d, k_d, v_d, g_d, b_d, rb_d, ra_d, s_in, dst_d, state_d, \ S_v, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, \ - sb1, sb2, sb3, neqk1, rq3, scale, state_slot_stride, K, S_v * H * n_tokens, stream) + sb1, sb2, sb3, neqk1, rq3, scale, state_slot_stride, K, S_v * H * n_tokens, \ + s_ids, s_row_stride, s_q8, stream) // a q8_0 cache is fused for the scalar gate only (ggml_cuda_try_gdn_cache_fusion) GGML_ASSERT(!(state_q8 && kda)); diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 12fab252a75e..709f2f37138b 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -2814,6 +2814,77 @@ static bool ggml_cuda_should_fuse_rms_norm_mul_rope(const ggml_tensor * rms_norm return true; } +// GET_ROWS(cache, ids) -> [RESHAPE] -> GATED_DELTA_NET src[5], build_rs's recurrent-state gather: skip the GET_ROWS and +// register the gather for the GDN node, whose kernel then reads each sequence's cache row ids[seq] itself (f32, or q8_0 +// dequantized as it loads) instead of the temp, which the allocator still reserved. One sequence only: with several, a +// sequence's source row can be another's destination (a copied sequence), which the GET_ROWS read before any write and +// the kernel's blocks would race on; with one, the block that reads a region is the only one that writes it, after. +// Declined when anything else reads the temp, when a node between the two writes the cache (build_rs's extra-states +// copy relocates rows there; the GET_ROWS read them before it, the kernel would read after), and when the GDN takes +// the chunked prefill path (its pipeline reads a gathered s0). From PrismML-Eng/llama.cpp#220 (f32), plus a q8_0 +// cache (-cts q8_0), the one the served head runs. GGML_CUDA_GDN_STATE_GATHER_LEGACY=1 keeps the GET_ROWS. +static bool ggml_cuda_try_gdn_gather_skip(ggml_backend_cuda_context & ctx, const ggml_cgraph * cgraph, int node_idx) { + static const bool legacy = [] { + const char * e = getenv("GGML_CUDA_GDN_STATE_GATHER_LEGACY"); + return e != nullptr && atoi(e) != 0; + }(); + if (legacy) { + return false; + } + + const ggml_tensor * gr = cgraph->nodes[node_idx]; + if (gr->op != GGML_OP_GET_ROWS || gr->type != GGML_TYPE_F32 || (gr->flags & GGML_TENSOR_FLAG_OUTPUT) || + !ggml_is_contiguous(gr) || gr->ne[1] != 1 || gr->ne[2] != 1 || gr->ne[3] != 1) { + return false; + } + const ggml_tensor * cache = gr->src[0]; + const ggml_tensor * ids = gr->src[1]; + const bool q8 = cache->type == GGML_TYPE_Q8_0; + if ((cache->type != GGML_TYPE_F32 && !q8) || ids->type != GGML_TYPE_I32 || cache->data == nullptr || + ids->data == nullptr || cache->ne[0] != gr->ne[0] || cache->nb[0] != ggml_type_size(cache->type) || + cache->nb[1] % ggml_type_size(cache->type) != 0 || !ggml_is_contiguous(ids) || ids->ne[0] != gr->ne[1] || + ggml_nelements(ids) != gr->ne[1] || ggml_node_get_use_count(cgraph, node_idx) != 1) { + return false; + } + const ggml_tensor * cache_base = cache->view_src ? cache->view_src : cache; + + const ggml_tensor * cur = gr; + for (int j = node_idx + 1; j < cgraph->n_nodes; ++j) { + const ggml_tensor * n = cgraph->nodes[j]; + if (n->op == GGML_OP_GATED_DELTA_NET && n->src[5] == cur) { + const ggml_tensor * v = n->src[2]; + const int64_t D = v->ne[0] * v->ne[0] * v->ne[1]; + if (gr->ne[0] != D || gr->ne[1] != v->ne[3] || ggml_nelements(cur) != D * v->ne[3] || + (q8 && D % QK8_0 != 0) || ggml_cuda_should_use_chunked_gdn(n)) { + return false; + } + ggml_cuda_gated_delta_net_gather gather; + gather.base = cache->data; + gather.ids = (const int32_t *) ids->data; + gather.row_stride = (int64_t) (cache->nb[1] / ggml_type_size(cache->type)) * ggml_blck_size(cache->type); + gather.q8_0 = q8; + ctx.gdn_gathers().set(n, gather); + return true; + } + if (n->op == GGML_OP_RESHAPE && n->src[0] == cur) { + if (ggml_node_get_use_count(cgraph, j) != 1) { + return false; + } + cur = n; + continue; + } + if (!ggml_cuda_is_view_or_noop(n) && ggml_nelements(n) > 0 && (n == cache_base || n->view_src == cache_base)) { + return false; + } + for (int s = 0; s < GGML_MAX_SRC; ++s) { + if (n->src[s] == cur || (n->view_src != nullptr && n->view_src == gr)) { + return false; + } + } + } + return false; +} + // match gated_delta_net + the strided cpy that scatters its state snapshots into the cache // (slot i -> rollback group i, slot 0 newest), so the kernel can write them and skip the cpy. static int ggml_cuda_try_gdn_cache_fusion( @@ -4445,6 +4516,8 @@ static void ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context * cud stream_ctx.concurrent_events.clear(); } + cuda_ctx->gdn_gathers().reset(); + for (int i = 0; i < cgraph->n_nodes; i++) { ggml_tensor * node = cgraph->nodes[i]; if (is_concurrent_event_active) { @@ -4487,6 +4560,12 @@ static void ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context * cud continue; } + // the recurrent-state gather whose only reader is a GATED_DELTA_NET: its kernel reads the cache rows + if (node->op == GGML_OP_GET_ROWS && !is_concurrent_event_active && + ggml_cuda_try_gdn_gather_skip(*cuda_ctx, cgraph, i)) { + continue; + } + // The normalized pre-attention residual is consumed only by a // group of low-bit projections. Preserve residual + one scale per // row and let their shared Q8 quantizer apply the norm weight. diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 6efb7b3085fb..5bd6a3810d5d 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -4516,6 +4516,9 @@ struct test_gated_delta_net_cache_fusion : public test_case { // out ssm_states_all, K slots of mem_size rows, and the view takes n_seqs rows from row kv_head of each slot const int64_t mem_size; const int64_t kv_head; + // >= 0 (mem_size > 0, one sequence): the initial state is build_rs's gather, GET_ROWS(cache, [state_row]) reshaped + // for the op, which CUDA folds into the kernel (ggml_cuda_try_gdn_gather_skip); the cache then starts random + const int64_t state_row; ggml_tensor * cpy_node = nullptr; ggml_tensor * readback_node = nullptr; @@ -4525,15 +4528,19 @@ struct test_gated_delta_net_cache_fusion : public test_case { if (mem_size > 0) { s += "," + VARS_TO_STR2(mem_size, kv_head); } + if (state_row >= 0) { + s += "," + VAR_TO_STR(state_row); + } return s; } test_gated_delta_net_cache_fusion(ggml_type type = GGML_TYPE_F32, int64_t head_count = 4, int64_t head_size = 32, int64_t n_seq_tokens = 2, int64_t n_seqs = 1, int64_t K = 2, int v_repeat = 1, bool raw_gates = false, ggml_type cache_type = GGML_TYPE_F32, - int64_t mem_size = 0, int64_t kv_head = 0) + int64_t mem_size = 0, int64_t kv_head = 0, int64_t state_row = -1) : type(type), head_count(head_count), head_size(head_size), n_seq_tokens(n_seq_tokens), n_seqs(n_seqs), K(K), - v_repeat(v_repeat), raw_gates(raw_gates), cache_type(cache_type), mem_size(mem_size), kv_head(kv_head) {} + v_repeat(v_repeat), raw_gates(raw_gates), cache_type(cache_type), mem_size(mem_size), kv_head(kv_head), + state_row(state_row) {} ggml_tensor * build_graph(ggml_context * ctx) override { const int64_t S_v = head_size; @@ -4550,10 +4557,21 @@ struct test_gated_delta_net_cache_fusion : public test_case { ggml_set_name(v, "v"); ggml_tensor * g = ggml_new_tensor_4d(ctx, type, 1, H_v, n_seq_tokens, n_seqs); ggml_tensor * beta = ggml_new_tensor_4d(ctx, type, 1, H_v, n_seq_tokens, n_seqs); - ggml_tensor * state = ggml_new_tensor_4d(ctx, type, head_size, head_size, H_v, n_seqs); ggml_set_name(g, "g"); ggml_set_name(beta, "beta"); - ggml_set_name(state, "state"); + + ggml_tensor * cache = nullptr; + ggml_tensor * state; + if (state_row >= 0) { + GGML_ASSERT(mem_size > 0 && n_seqs == 1); + cache = ggml_new_tensor_2d(ctx, cache_type, D, mem_size * K); + ggml_tensor * ids = ggml_new_tensor_1d(ctx, GGML_TYPE_I32, n_seqs); + ggml_set_name(ids, "ids"); + state = ggml_reshape_4d(ctx, ggml_get_rows(ctx, cache, ids), head_size, head_size, H_v, n_seqs); + } else { + state = ggml_new_tensor_4d(ctx, type, head_size, head_size, H_v, n_seqs); + ggml_set_name(state, "state"); + } q = ggml_l2_norm(ctx, q, 1e-6f); k = ggml_l2_norm(ctx, k, 1e-6f); @@ -4578,11 +4596,12 @@ struct test_gated_delta_net_cache_fusion : public test_case { // recurrent cache view [D, n_seqs, n_written] const size_t row_size = ggml_row_size(cache_type, D); - ggml_tensor * cache; ggml_tensor * dst; if (mem_size > 0) { // slot stride nb[2] = mem_size rows, not n_seqs, at an offset of kv_head rows - cache = ggml_new_tensor_2d(ctx, cache_type, D, mem_size * K); + if (cache == nullptr) { + cache = ggml_new_tensor_2d(ctx, cache_type, D, mem_size * K); + } dst = ggml_view_3d(ctx, cache, D, n_seqs, n_written, row_size, mem_size * row_size, kv_head * row_size); } else { cache = ggml_new_tensor_3d(ctx, cache_type, D, n_seqs, n_written); @@ -4653,7 +4672,11 @@ struct test_gated_delta_net_cache_fusion : public test_case { } else if (strcmp(t->name, "v") == 0) { init_tensor_uniform(t, -0.3f, 5.0f); } else if (strcmp(t->name, "cache") == 0) { - init_tensor_uniform(t, 0.0f, 0.0f); + // gathered: every row a different state, so reading another row than state_row fails + init_tensor_uniform(t, state_row >= 0 ? -1.0f : 0.0f, state_row >= 0 ? 1.0f : 0.0f); + } else if (strcmp(t->name, "ids") == 0) { + const std::vector rows(n_seqs, (int32_t) state_row); + ggml_backend_tensor_set(t, rows.data(), 0, rows.size() * sizeof(int32_t)); } else { init_tensor_uniform(t); } @@ -10959,6 +10982,18 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_gated_delta_net_cache_fusion(GGML_TYPE_F32, 16, 128, 3, n_seqs, 3, 3, true, GGML_TYPE_Q8_0, n_seqs + 2, 1)); test_cases.emplace_back(new test_gated_delta_net_cache_fusion(GGML_TYPE_F32, 16, 128, 3, n_seqs, 3, 3, true, GGML_TYPE_F32, n_seqs + 2, 1)); } + // build_rs's gather of the initial state from the cache (GET_ROWS -> RESHAPE -> GATED_DELTA_NET), folded into the + // kernel on CUDA: the served layer at decode and at the 3-token verify, the state read from the row the newest + // snapshot overwrites (1, kv_head in slot 0) and from a rollback slot's row the verify also writes (4, slot 1), a + // q8_0 and an f32 cache; and the chunked prefill ubatch, which keeps the GET_ROWS + for (ggml_type cache_type : { GGML_TYPE_Q8_0, GGML_TYPE_F32 }) { + for (int64_t n_tokens : { 1, 3 }) { + for (int64_t state_row : { 1, 4 }) { + test_cases.emplace_back(new test_gated_delta_net_cache_fusion(GGML_TYPE_F32, 16, 128, n_tokens, 1, 3, 3, true, cache_type, 3, 1, state_row)); + } + } + } + test_cases.emplace_back(new test_gated_delta_net_cache_fusion(GGML_TYPE_F32, 16, 128, 512, 1, 3, 3, true, GGML_TYPE_Q8_0, 3, 1, 4)); // a q8_0 cache the fusion refuses keeps its cpy: a 16-wide head, and the chunked prefill ubatch test_cases.emplace_back(new test_gated_delta_net_cache_fusion(GGML_TYPE_F32, 4, 16, 2, 1, 2, 1, false, GGML_TYPE_Q8_0)); test_cases.emplace_back(new test_gated_delta_net_cache_fusion(GGML_TYPE_F32, 16, 128, 512, 1, 3, 3, true, GGML_TYPE_Q8_0)); From b534681b3a78102e88f040269c988073c9fa0049 Mon Sep 17 00:00:00 2001 From: Marcos Damasceno Date: Fri, 25 Sep 2026 00:06:01 -0500 Subject: [PATCH 2/3] fix(cuda): the GDN state gather declines when a node up to the GDN sits on ids' memory The kernel reads ids at the GATED_DELTA_NET, but the GET_ROWS it skips can be ids' last reader in graph order (the last recurrent layer's), after which ggml-alloc frees ids and may place a later node on it. The matcher now declines when any non-view node after the GET_ROWS, the GDN included, overlaps ids' bytes; that layer keeps its GET_ROWS. --- ggml/src/ggml-cuda/ggml-cuda.cu | 16 ++++++++++++++-- 1 file changed, 14 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 709f2f37138b..269b39725466 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -2814,14 +2814,21 @@ static bool ggml_cuda_should_fuse_rms_norm_mul_rope(const ggml_tensor * rms_norm return true; } +static bool ggml_cuda_ranges_overlap(const ggml_tensor * a, const ggml_tensor * b) { + const char * a0 = (const char *) a->data; + const char * b0 = (const char *) b->data; + return a0 < b0 + ggml_nbytes(b) && b0 < a0 + ggml_nbytes(a); +} + // GET_ROWS(cache, ids) -> [RESHAPE] -> GATED_DELTA_NET src[5], build_rs's recurrent-state gather: skip the GET_ROWS and // register the gather for the GDN node, whose kernel then reads each sequence's cache row ids[seq] itself (f32, or q8_0 // dequantized as it loads) instead of the temp, which the allocator still reserved. One sequence only: with several, a // sequence's source row can be another's destination (a copied sequence), which the GET_ROWS read before any write and // the kernel's blocks would race on; with one, the block that reads a region is the only one that writes it, after. // Declined when anything else reads the temp, when a node between the two writes the cache (build_rs's extra-states -// copy relocates rows there; the GET_ROWS read them before it, the kernel would read after), and when the GDN takes -// the chunked prefill path (its pipeline reads a gathered s0). From PrismML-Eng/llama.cpp#220 (f32), plus a q8_0 +// copy relocates rows there; the GET_ROWS read them before it, the kernel would read after), when a node up to the GDN +// was allocated over ids (freed after the GET_ROWS if it was their last reader), and when the GDN takes the chunked +// prefill path (its pipeline reads a gathered s0). From PrismML-Eng/llama.cpp#220 (f32), plus a q8_0 // cache (-cts q8_0), the one the served head runs. GGML_CUDA_GDN_STATE_GATHER_LEGACY=1 keeps the GET_ROWS. static bool ggml_cuda_try_gdn_gather_skip(ggml_backend_cuda_context & ctx, const ggml_cgraph * cgraph, int node_idx) { static const bool legacy = [] { @@ -2851,6 +2858,11 @@ static bool ggml_cuda_try_gdn_gather_skip(ggml_backend_cuda_context & ctx, const const ggml_tensor * cur = gr; for (int j = node_idx + 1; j < cgraph->n_nodes; ++j) { const ggml_tensor * n = cgraph->nodes[j]; + // the kernel reads ids at the GDN, and the GET_ROWS it skips may have been ids' last reader (the last recurrent + // layer's), after which the allocator hands ids' memory on: nothing up to and including the GDN may be on it + if (!ggml_cuda_is_view_or_noop(n) && ggml_cuda_ranges_overlap(n, ids)) { + return false; + } if (n->op == GGML_OP_GATED_DELTA_NET && n->src[5] == cur) { const ggml_tensor * v = n->src[2]; const int64_t D = v->ne[0] * v->ne[0] * v->ne[1]; From f5c16d7069f3cc928bb277de67d65b75ec803603 Mon Sep 17 00:00:00 2001 From: Marcos Damasceno Date: Fri, 25 Sep 2026 01:46:54 -0500 Subject: [PATCH 3/3] docs(torad): the GDN state gather fold row --- TORAD.md | 1 + 1 file changed, 1 insertion(+) diff --git a/TORAD.md b/TORAD.md index 9dec9ab7c6ac..d48025111699 100644 --- a/TORAD.md +++ b/TORAD.md @@ -33,6 +33,7 @@ which pins a commit of this branch as a submodule. | `d79f6e0` | a sequence's first KQ mask row tests the KV cells 16 at a time and masks a group past the sequence's last ubatch position on the positions alone, before any sequence bitset is read (a bit-packed mask is written a word at a time). With `--kv-unified` and another slot holding 118K tokens, building the mask went 230 → 79 µs for a verify and 251 → 73 µs for a draft step (131,328 cells). Served together with the next two rows (RTX 5070 Ti, 2 slots × 131,072, MTP n_max 2, greedy, legs new/legacy/legacy/new): empty pool +2.6 %, a slot decoding beside a resident 118K context +9.5 % (8 of 8 requests +8.9 to +9.7 %), the long slot's own decode at 118K +2.6 %; text identical 13/13 across all legs. Fuzzed against the old builder over 20,000 random caches (streams, multi-sequence cells, 2D positions, SWA, ALiBi, every mask layout): 0 byte mismatches | `LLAMA_KQ_MASK_SCAN_LEGACY=1` | | `6ea8df2` | a speculative rollback's `seq_rm` (every verify, on the target's and the draft's cache) walks back from the head to the last of the 0-2 cells it removes instead of scanning every cell; more than 1,024 cells keep the scan. Measured with the row above; fuzzed against the old walk over 72,709 calls: 0 mismatches | `LLAMA_KV_SEQ_RM_SCAN_LEGACY=1` | | `da382d9` | the batch's coupled-sequence check (n_seq_max² flags per decode, 256² with `--kv-unified`) runs only when a token is in several sequences, the only way a flag is set. Measured with the rows above | `LLAMA_BATCH_CPL_CHECK_LEGACY=1` | +| `8fd3a28`, `b534681` | a Gated DeltaNet layer's recurrent state is read by the GATED_DELTA_NET kernel straight from its cache row (`ids[seq]`), f32 or q8_0 dequantized as it loads with get_rows' formula, and the GET_ROWS that gathered it into a temp is skipped (from PrismML-Eng/llama.cpp#220, plus the q8_0 cache); declined for more than one sequence, another reader of the temp, a node between them that writes the cache or sits on ids' memory, and the chunked prefill. RTX 5070 Ti, served flags (`-cts q8_0`, MTP n_max 2), Nsight per step at 1 / 3 rows: the dequantizing gather's 48 launches (270 µs) gone, the GDN kernel +39 µs, 12.698 → 12.338 ms / 14.774 → 14.427 ms (−2.8 / −2.3 %, A-B-B-A), perplexity identical to the last digit; served with the MTP draft head, 48 greedy eval requests, legs cancelled: +1.69 % tok/s (95 % CI +1.59 to +1.80), 48 of 48 faster, text identical 48/48, same draft acceptance | `GGML_CUDA_GDN_STATE_GATHER_LEGACY=1` | Every switch in the last column is read once per process and parses as an integer: a `*_LEGACY` switch set to `0` is the same as unset (the change stays on), and `=0` turns off `GGML_CUDA_LORA_RANK1_FUSE` and