From 2576c152ff87737a9294fdd33958b0104c24d30b Mon Sep 17 00:00:00 2001 From: dusterbloom <32869278+dusterbloom@users.noreply.github.com> Date: Tue, 30 Jun 2026 21:11:48 +0200 Subject: [PATCH 1/2] perf(cuda): add in-place GDN state write --- server/deps/llama.cpp/ggml/include/ggml.h | 9 +++++ .../ggml/src/ggml-cuda/gated_delta_net.cu | 34 +++++++++++-------- server/deps/llama.cpp/ggml/src/ggml.c | 14 ++++++++ 3 files changed, 43 insertions(+), 14 deletions(-) diff --git a/server/deps/llama.cpp/ggml/include/ggml.h b/server/deps/llama.cpp/ggml/include/ggml.h index df742fbc5..d76166f16 100644 --- a/server/deps/llama.cpp/ggml/include/ggml.h +++ b/server/deps/llama.cpp/ggml/include/ggml.h @@ -2547,6 +2547,15 @@ extern "C" { struct ggml_tensor * beta, struct ggml_tensor * state); + GGML_API struct ggml_tensor * ggml_gated_delta_net_inplace( + struct ggml_context * ctx, + struct ggml_tensor * q, + struct ggml_tensor * k, + struct ggml_tensor * v, + struct ggml_tensor * g, + struct ggml_tensor * beta, + struct ggml_tensor * state); + GGML_API void ggml_gated_delta_net_set_skip_intermediate( struct ggml_tensor * tensor, bool skip_intermediate); diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/gated_delta_net.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/gated_delta_net.cu index 8cb989b1a..75fbbab1b 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/gated_delta_net.cu +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/gated_delta_net.cu @@ -51,6 +51,7 @@ gated_delta_net_cuda(const float * q, const float * beta, const float * curr_state, float * dst, + float * state_out, const int * parent_ids, // TREE_MODE only; else ignored InterT * persist_inter, // optional external buffer for per-token intermediates bool skip_intermediate, @@ -81,7 +82,7 @@ gated_delta_net_cuda(const float * q, const int64_t attn_score_elems = S_v * H * n_tokens * n_seqs; const int64_t final_state_elems = S_v * S_v * H * n_seqs; float * attn_data = dst; - float * state = dst + attn_score_elems; + float * state = state_out; // intermediate_states region: one S_v*S_v*H*n_seqs state per token. Written // inside the token loop below (one state per `t`) to enable spec-decode // rollback without a replay forward pass. See ggml.c::ggml_gated_delta_net. @@ -275,6 +276,7 @@ gated_delta_net_cuda_grouped_cols(const float * q, const float * beta, const float * curr_state, float * dst, + float * state_out, InterT * persist_inter, bool skip_intermediate, int64_t H, @@ -318,7 +320,7 @@ gated_delta_net_cuda_grouped_cols(const float * q, const int64_t attn_score_elems = S_v * H * n_tokens * n_seqs; const int64_t final_state_elems = S_v * S_v * H * n_seqs; float * attn_data = dst; - float * state = dst + attn_score_elems; + float * state = state_out; InterT * inter_states = persist_inter ? persist_inter : (InterT *)(dst + attn_score_elems + final_state_elems); @@ -452,6 +454,7 @@ 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 * s_d, float * dst_d, + float * state_out_d, const int * parent_ids_d, InterT * persist_inter_d, int64_t S_v, int64_t H, int64_t n_tokens, int64_t n_seqs, @@ -475,19 +478,19 @@ static void launch_gated_delta_net( switch (S_v) { case 16: gated_delta_net_cuda<16, KDA, TREE_MODE, InterT><<>>( - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, parent_ids_d, persist_inter_d, skip_intermediate, H, + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, parent_ids_d, persist_inter_d, skip_intermediate, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2, sb3, neqk1_magic, rq3_magic, scale); break; case 32: gated_delta_net_cuda<32, KDA, TREE_MODE, InterT><<>>( - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, parent_ids_d, persist_inter_d, skip_intermediate, H, + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, parent_ids_d, persist_inter_d, skip_intermediate, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2, sb3, neqk1_magic, rq3_magic, scale); break; case 64: { gated_delta_net_cuda<64, KDA, TREE_MODE, InterT><<>>( - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, parent_ids_d, persist_inter_d, skip_intermediate, H, + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, parent_ids_d, persist_inter_d, skip_intermediate, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2, sb3, neqk1_magic, rq3_magic, scale); break; @@ -505,7 +508,7 @@ static void launch_gated_delta_net( dim3 grouped_grid_dims(H, n_seqs, (groups + column_groups_per_block * groups_per_warp - 1) / (column_groups_per_block * groups_per_warp)); dim3 grouped_block_dims(32, column_groups_per_block, 1); gated_delta_net_cuda_grouped_cols<128, cols, width, 32, InterT><<>>( - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, persist_inter_d, skip_intermediate, H, + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, persist_inter_d, skip_intermediate, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2, sb3, neqk1_magic, rq3_magic, scale); } else if (warp_size == 64) { @@ -513,24 +516,24 @@ static void launch_gated_delta_net( dim3 grouped_grid_dims(H, n_seqs, (groups + column_groups_per_block * groups_per_warp - 1) / (column_groups_per_block * groups_per_warp)); dim3 grouped_block_dims(64, column_groups_per_block, 1); gated_delta_net_cuda_grouped_cols<128, cols, width, 64, InterT><<>>( - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, persist_inter_d, skip_intermediate, H, + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, persist_inter_d, skip_intermediate, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2, sb3, neqk1_magic, rq3_magic, scale); } else { gated_delta_net_cuda<128, KDA, TREE_MODE, InterT><<>>( - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, parent_ids_d, persist_inter_d, skip_intermediate, H, + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, parent_ids_d, persist_inter_d, skip_intermediate, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2, sb3, neqk1_magic, rq3_magic, scale); } } else { gated_delta_net_cuda<128, KDA, TREE_MODE, InterT><<>>( - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, parent_ids_d, persist_inter_d, skip_intermediate, H, + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, parent_ids_d, persist_inter_d, skip_intermediate, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2, sb3, neqk1_magic, rq3_magic, scale); } } else { gated_delta_net_cuda<128, KDA, TREE_MODE, InterT><<>>( - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, parent_ids_d, persist_inter_d, skip_intermediate, H, + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, parent_ids_d, persist_inter_d, skip_intermediate, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2, sb3, neqk1_magic, rq3_magic, scale); } @@ -586,6 +589,9 @@ void ggml_cuda_op_gated_delta_net(ggml_backend_cuda_context & ctx, ggml_tensor * const float * s_d = (const float *) src_state->data; float * dst_d = (float *) dst->data; + const int64_t attn_score_elems = S_v * H * n_tokens * n_seqs; + const bool inplace_state = ggml_get_op_params_i32(dst, 1) != 0; + float * state_out_d = inplace_state ? (float *) src_state->data : dst_d + attn_score_elems; const int * parent_ids_d = src_parent ? (const int *) src_parent->data : nullptr; @@ -640,24 +646,24 @@ void ggml_cuda_op_gated_delta_net(ggml_backend_cuda_context & ctx, ggml_tensor * if (kda) { \ if (tree_mode) { \ launch_gated_delta_net( \ - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, parent_ids_d, persist_typed, \ + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, parent_ids_d, persist_typed, \ S_v, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, \ sb1, sb2, sb3, neqk1, rq3, skip_intermediate, scale, stream); \ } else { \ launch_gated_delta_net( \ - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, nullptr, persist_typed, \ + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, nullptr, persist_typed, \ S_v, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, \ sb1, sb2, sb3, neqk1, rq3, skip_intermediate, scale, stream); \ } \ } else { \ if (tree_mode) { \ launch_gated_delta_net( \ - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, parent_ids_d, persist_typed, \ + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, parent_ids_d, persist_typed, \ S_v, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, \ sb1, sb2, sb3, neqk1, rq3, skip_intermediate, scale, stream); \ } else { \ launch_gated_delta_net( \ - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, nullptr, persist_typed, \ + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, nullptr, persist_typed, \ S_v, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, \ sb1, sb2, sb3, neqk1, rq3, skip_intermediate, scale, stream); \ } \ diff --git a/server/deps/llama.cpp/ggml/src/ggml.c b/server/deps/llama.cpp/ggml/src/ggml.c index c1ffccafb..fda42ef76 100644 --- a/server/deps/llama.cpp/ggml/src/ggml.c +++ b/server/deps/llama.cpp/ggml/src/ggml.c @@ -6262,6 +6262,7 @@ struct ggml_tensor * ggml_gated_delta_net( // roll back SSM state to the accepted prefix without a full replay forward pass. const int64_t ne[4] = { S_v * H, n_tokens * n_seqs + S_v * n_seqs + S_v * n_tokens * n_seqs, 1, 1 }; struct ggml_tensor * result = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, ne); + ggml_set_op_params_i32(result, 1, 0); result->op = GGML_OP_GATED_DELTA_NET; result->src[0] = q; @@ -6274,6 +6275,19 @@ struct ggml_tensor * ggml_gated_delta_net( return result; } +struct ggml_tensor * ggml_gated_delta_net_inplace( + struct ggml_context * ctx, + struct ggml_tensor * q, + struct ggml_tensor * k, + struct ggml_tensor * v, + struct ggml_tensor * g, + struct ggml_tensor * beta, + struct ggml_tensor * state) { + struct ggml_tensor * result = ggml_gated_delta_net(ctx, q, k, v, g, beta, state); + ggml_set_op_params_i32(result, 1, 1); + return result; +} + void ggml_gated_delta_net_set_skip_intermediate( struct ggml_tensor * tensor, bool skip_intermediate) { From c3d44ed490a1d623c02595b887632bf50c094426 Mon Sep 17 00:00:00 2001 From: dusterbloom <32869278+dusterbloom@users.noreply.github.com> Date: Wed, 8 Jul 2026 10:37:11 +0200 Subject: [PATCH 2/2] perf(cuda): recover qwen dense decode path --- .../ggml/src/ggml-cuda/gated_delta_net.cu | 135 ++++++++++-------- .../llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu | 15 +- 2 files changed, 87 insertions(+), 63 deletions(-) diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/gated_delta_net.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/gated_delta_net.cu index 75fbbab1b..c561ace62 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/gated_delta_net.cu +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/gated_delta_net.cu @@ -2,6 +2,7 @@ #ifndef GGML_USE_HIP #include #endif +#include #include // Tree-mode parent index sentinel: a node whose parent is the pre-block state @@ -42,7 +43,7 @@ __device__ __forceinline__ float gdn_subgroup_broadcast_lane0(float value, int w return __shfl_sync(0xffffffffU, value, 0, width); } -template +template __global__ void __launch_bounds__((ggml_cuda_get_physical_warp_size() < S_v ? ggml_cuda_get_physical_warp_size() : S_v) * 4, 2) gated_delta_net_cuda(const float * q, const float * k, @@ -54,7 +55,6 @@ gated_delta_net_cuda(const float * q, float * state_out, const int * parent_ids, // TREE_MODE only; else ignored InterT * persist_inter, // optional external buffer for per-token intermediates - bool skip_intermediate, int64_t H, int64_t n_tokens, int64_t n_seqs, @@ -83,30 +83,22 @@ gated_delta_net_cuda(const float * q, const int64_t final_state_elems = S_v * S_v * H * n_seqs; float * attn_data = dst; float * state = state_out; - // intermediate_states region: one S_v*S_v*H*n_seqs state per token. Written - // inside the token loop below (one state per `t`) to enable spec-decode - // rollback without a replay forward pass. See ggml.c::ggml_gated_delta_net. - // - // dflash27b_ggml: if persist_inter != nullptr, the kernel writes the - // intermediate states DIRECTLY to that external buffer instead of the - // embedded region inside dst. InterT selects the storage precision (float - // or __half). f16 halves the memory footprint — enough to fit larger - // DDtree budgets on the 24 GB 3090. - // When persist_inter is null, InterT MUST be float (the embedded region - // inside dst is f32). - InterT * inter_states = persist_inter - ? persist_inter - : (InterT *)(dst + attn_score_elems + final_state_elems); - const bool write_intermediate = !skip_intermediate || TREE_MODE || persist_inter != nullptr; + InterT * inter_states = nullptr; + InterT * inter_base = nullptr; + if constexpr (WRITE_INTER || TREE_MODE) { + // One S_v*S_v*H*n_seqs state per token for rollback/tree paths. + // Pure AR instantiates WRITE_INTER=false, so this address arithmetic + // and the store loop below compile out completely. + inter_states = persist_inter + ? persist_inter + : (InterT *)(dst + attn_score_elems + final_state_elems); + inter_base = inter_states + (sequence * n_tokens * H + h_idx) * S_v * S_v; + } const int64_t state_offset = (sequence * H + h_idx) * S_v * S_v; state += state_offset; curr_state += state_offset + col * S_v; attn_data += (sequence * n_tokens * H + h_idx) * S_v; - // Per-sequence per-head base for this block's intermediates, token t=0. - // Advance by (H * S_v * S_v) each iteration. - InterT * inter_base = inter_states + (sequence * n_tokens * H + h_idx) * S_v * S_v; - constexpr int warp_size = ggml_cuda_get_physical_warp_size() < S_v ? ggml_cuda_get_physical_warp_size() : S_v; static_assert(S_v % warp_size == 0, "S_v must be a multiple of warp_size"); constexpr int rows_per_lane = (S_v + warp_size - 1) / warp_size; @@ -247,14 +239,14 @@ gated_delta_net_cuda(const float * q, // final-state write below). Used by dflash27b_ggml spec-decode rollback. // Plain chain prefill does not consume it, so qwen35 can opt out to // avoid large transient global writes. - if (write_intermediate) { + if constexpr (WRITE_INTER) { #pragma unroll for (int r = 0; r < rows_per_lane; r++) { const int i = r * warp_size + lane; store_inter_state(inter_base, col * S_v + i, s_shard[r]); } + inter_base += S_v * S_v * H; } - inter_base += S_v * S_v * H; attn_data += S_v * H; } @@ -267,7 +259,7 @@ gated_delta_net_cuda(const float * q, } } -template +template __global__ void __launch_bounds__(WARP_THREADS * 8, 2) gated_delta_net_cuda_grouped_cols(const float * q, const float * k, @@ -278,7 +270,6 @@ gated_delta_net_cuda_grouped_cols(const float * q, float * dst, float * state_out, InterT * persist_inter, - bool skip_intermediate, int64_t H, int64_t n_tokens, int64_t n_seqs, @@ -321,17 +312,19 @@ gated_delta_net_cuda_grouped_cols(const float * q, const int64_t final_state_elems = S_v * S_v * H * n_seqs; float * attn_data = dst; float * state = state_out; - InterT * inter_states = persist_inter - ? persist_inter - : (InterT *)(dst + attn_score_elems + final_state_elems); - const bool write_intermediate = !skip_intermediate || persist_inter != nullptr; + InterT * inter_states = nullptr; + InterT * inter_base = nullptr; + if constexpr (WRITE_INTER) { + inter_states = persist_inter + ? persist_inter + : (InterT *)(dst + attn_score_elems + final_state_elems); + inter_base = inter_states + (sequence * n_tokens * H + h_idx) * S_v * S_v; + } const int64_t state_offset = (sequence * H + h_idx) * S_v * S_v; state += state_offset; curr_state += state_offset; attn_data += (sequence * n_tokens * H + h_idx) * S_v; - InterT * inter_base = inter_states + (sequence * n_tokens * H + h_idx) * S_v * S_v; - float state_shard[COLS][rows_per_lane]; #pragma unroll @@ -422,7 +415,7 @@ gated_delta_net_cuda_grouped_cols(const float * q, } } - if (write_intermediate) { + if constexpr (WRITE_INTER) { #pragma unroll for (int c = 0; c < COLS; ++c) { const int col = col_base + c; @@ -432,10 +425,10 @@ gated_delta_net_cuda_grouped_cols(const float * q, store_inter_state(inter_base, col * S_v + row, state_shard[c][r]); } } + inter_base += S_v * S_v * H; } attn_data += S_v * H; - inter_base += S_v * S_v * H; } #pragma unroll @@ -449,7 +442,7 @@ gated_delta_net_cuda_grouped_cols(const float * q, } } -template +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 * s_d, @@ -462,7 +455,6 @@ static void launch_gated_delta_net( int64_t sv1, int64_t sv2, int64_t sv3, int64_t sb1, int64_t sb2, int64_t sb3, int64_t neqk1, int64_t rq3, - bool skip_intermediate, float scale, cudaStream_t stream) { //TODO: Add chunked kernel for even faster pre-fill const int warp_size = ggml_cuda_info().devices[ggml_cuda_get_device()].warp_size; @@ -474,31 +466,39 @@ static void launch_gated_delta_net( const uint3 rq3_magic = init_fastdiv_values(rq3); int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc; + const bool ampere_nvidia = GGML_CUDA_CC_IS_NVIDIA(cc) + && cc >= GGML_CUDA_CC_AMPERE + && cc < GGML_CUDA_CC_ADA_LOVELACE; + const bool force_grouped_cols = getenv("DFLASH_GDN_FORCE_GROUPED_COLS") != nullptr; + const bool disable_grouped_cols = getenv("DFLASH_GDN_NO_GROUPED_COLS") != nullptr; + const bool use_grouped_cols = force_grouped_cols || + (!disable_grouped_cols && !ampere_nvidia); switch (S_v) { case 16: - gated_delta_net_cuda<16, KDA, TREE_MODE, InterT><<>>( - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, parent_ids_d, persist_inter_d, skip_intermediate, H, + gated_delta_net_cuda<16, KDA, TREE_MODE, WRITE_INTER, InterT><<>>( + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, parent_ids_d, persist_inter_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2, sb3, neqk1_magic, rq3_magic, scale); break; case 32: - gated_delta_net_cuda<32, KDA, TREE_MODE, InterT><<>>( - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, parent_ids_d, persist_inter_d, skip_intermediate, H, + gated_delta_net_cuda<32, KDA, TREE_MODE, WRITE_INTER, InterT><<>>( + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, parent_ids_d, persist_inter_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2, sb3, neqk1_magic, rq3_magic, scale); break; case 64: { - gated_delta_net_cuda<64, KDA, TREE_MODE, InterT><<>>( - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, parent_ids_d, persist_inter_d, skip_intermediate, H, + gated_delta_net_cuda<64, KDA, TREE_MODE, WRITE_INTER, InterT><<>>( + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, parent_ids_d, persist_inter_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2, sb3, neqk1_magic, rq3_magic, scale); break; } case 128: { if constexpr (!KDA && !TREE_MODE) { - if ((GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_AMPERE) || - GGML_CUDA_CC_IS_AMD(cc)) { + if (use_grouped_cols && + ((GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_AMPERE) || + GGML_CUDA_CC_IS_AMD(cc))) { constexpr int cols = 4; constexpr int width = 16; constexpr int column_groups_per_block = 8; @@ -507,33 +507,33 @@ static void launch_gated_delta_net( constexpr int groups_per_warp = 32 / width; dim3 grouped_grid_dims(H, n_seqs, (groups + column_groups_per_block * groups_per_warp - 1) / (column_groups_per_block * groups_per_warp)); dim3 grouped_block_dims(32, column_groups_per_block, 1); - gated_delta_net_cuda_grouped_cols<128, cols, width, 32, InterT><<>>( - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, persist_inter_d, skip_intermediate, H, + gated_delta_net_cuda_grouped_cols<128, cols, width, 32, WRITE_INTER, InterT><<>>( + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, persist_inter_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2, sb3, neqk1_magic, rq3_magic, scale); } else if (warp_size == 64) { constexpr int groups_per_warp = 64 / width; dim3 grouped_grid_dims(H, n_seqs, (groups + column_groups_per_block * groups_per_warp - 1) / (column_groups_per_block * groups_per_warp)); dim3 grouped_block_dims(64, column_groups_per_block, 1); - gated_delta_net_cuda_grouped_cols<128, cols, width, 64, InterT><<>>( - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, persist_inter_d, skip_intermediate, H, + gated_delta_net_cuda_grouped_cols<128, cols, width, 64, WRITE_INTER, InterT><<>>( + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, persist_inter_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2, sb3, neqk1_magic, rq3_magic, scale); } else { - gated_delta_net_cuda<128, KDA, TREE_MODE, InterT><<>>( - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, parent_ids_d, persist_inter_d, skip_intermediate, H, + gated_delta_net_cuda<128, KDA, TREE_MODE, WRITE_INTER, InterT><<>>( + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, parent_ids_d, persist_inter_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2, sb3, neqk1_magic, rq3_magic, scale); } } else { - gated_delta_net_cuda<128, KDA, TREE_MODE, InterT><<>>( - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, parent_ids_d, persist_inter_d, skip_intermediate, H, + gated_delta_net_cuda<128, KDA, TREE_MODE, WRITE_INTER, InterT><<>>( + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, parent_ids_d, persist_inter_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2, sb3, neqk1_magic, rq3_magic, scale); } } else { - gated_delta_net_cuda<128, KDA, TREE_MODE, InterT><<>>( - q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, parent_ids_d, persist_inter_d, skip_intermediate, H, + gated_delta_net_cuda<128, KDA, TREE_MODE, WRITE_INTER, InterT><<>>( + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, parent_ids_d, persist_inter_d, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2, sb3, neqk1_magic, rq3_magic, scale); } @@ -637,35 +637,46 @@ void ggml_cuda_op_gated_delta_net(ggml_backend_cuda_context & ctx, ggml_tensor * const bool tree_mode = (parent_ids_d != nullptr); const bool skip_intermediate = ggml_get_op_params_i32(dst, 0) != 0; + const bool write_intermediate = tree_mode || !skip_intermediate || persist_inter_d != nullptr; - // Macro to expand the 4 (KDA × TREE_MODE) cases for a given InterT. + // Macro to expand KDA × TREE_MODE × WRITE_INTER for a given InterT. // The persist_is_f16 branch picks between __half and float instantiations. #define GDN_LAUNCH(INTER_T) \ do { \ INTER_T * persist_typed = (INTER_T *)persist_inter_d; \ if (kda) { \ if (tree_mode) { \ - launch_gated_delta_net( \ + launch_gated_delta_net( \ q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, parent_ids_d, persist_typed, \ S_v, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, \ - sb1, sb2, sb3, neqk1, rq3, skip_intermediate, scale, stream); \ + sb1, sb2, sb3, neqk1, rq3, scale, stream); \ + } else if (write_intermediate) { \ + launch_gated_delta_net( \ + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, nullptr, persist_typed, \ + S_v, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, \ + sb1, sb2, sb3, neqk1, rq3, scale, stream); \ } else { \ - launch_gated_delta_net( \ + launch_gated_delta_net( \ q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, nullptr, persist_typed, \ S_v, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, \ - sb1, sb2, sb3, neqk1, rq3, skip_intermediate, scale, stream); \ + sb1, sb2, sb3, neqk1, rq3, scale, stream); \ } \ } else { \ if (tree_mode) { \ - launch_gated_delta_net( \ + launch_gated_delta_net( \ q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, parent_ids_d, persist_typed, \ S_v, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, \ - sb1, sb2, sb3, neqk1, rq3, skip_intermediate, scale, stream); \ + sb1, sb2, sb3, neqk1, rq3, scale, stream); \ + } else if (write_intermediate) { \ + launch_gated_delta_net( \ + q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, nullptr, persist_typed, \ + S_v, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, \ + sb1, sb2, sb3, neqk1, rq3, scale, stream); \ } else { \ - launch_gated_delta_net( \ + launch_gated_delta_net( \ q_d, k_d, v_d, g_d, b_d, s_d, dst_d, state_out_d, nullptr, persist_typed, \ S_v, H, n_tokens, n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, \ - sb1, sb2, sb3, neqk1, rq3, skip_intermediate, scale, stream); \ + sb1, sb2, sb3, neqk1, rq3, scale, stream); \ } \ } \ } while (0) diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu index 29e77ae72..feb9104a9 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu @@ -4309,6 +4309,14 @@ static bool ggml_cuda_graph_set_enabled(ggml_backend_cuda_context * cuda_ctx, co } #endif // USE_CUDA_GRAPH +static thread_local bool ggml_cuda_skip_props_check = false; + +extern "C" void ggml_cuda_set_skip_props_check(bool skip); + +extern "C" void ggml_cuda_set_skip_props_check(bool skip) { + ggml_cuda_skip_props_check = skip; +} + static enum ggml_status ggml_backend_cuda_graph_compute(ggml_backend_t backend, ggml_cgraph * cgraph) { ggml_backend_cuda_context * cuda_ctx = (ggml_backend_cuda_context *) backend->context; @@ -4332,7 +4340,12 @@ static enum ggml_status ggml_backend_cuda_graph_compute(ggml_backend_t backend, if (graph->is_enabled()) { const bool graph_compatible = ggml_cuda_graph_check_compability(cgraph); if (graph_compatible) { - const bool properties_changed = ggml_cuda_graph_update_required(cuda_ctx, cgraph); + const bool can_skip_props_check = ggml_cuda_skip_props_check + && graph->warmup_complete + && graph->instance != nullptr; + const bool properties_changed = can_skip_props_check + ? false + : ggml_cuda_graph_update_required(cuda_ctx, cgraph); if (!graph->warmup_complete) { // Warmup: need at least 2 calls with no property change on the 2nd call