perf(qwen35): add SpecLA factor-capture speculative decoding

server/CMakeLists.txt

@@ -401,6 +401,7 @@ add_library(dflash_common STATIC
     src/kv_cache.cpp
     src/kv_quant.cpp
     src/delta_net_chunked.cpp
+    src/delta_net_specla.cpp
     # Laguna-XS.2 (Poolside) target arch
     src/laguna/laguna_target_loader.cpp
     src/laguna/laguna_target_graph.cpp
@@ -563,6 +564,10 @@ if(DFLASH27B_GPU_BACKEND STREQUAL "hip")
     target_sources(dflash_common PRIVATE src/common/geometric_draft_topk_cuda.cu)
     set_source_files_properties(src/common/geometric_draft_topk_cuda.cu
         PROPERTIES LANGUAGE HIP)
+    # Fused SpecLA accepted-state commit (one launch for all delta layers).
+    target_sources(dflash_common PRIVATE src/common/specla_commit_cuda.cu)
+    set_source_files_properties(src/common/specla_commit_cuda.cu
+        PROPERTIES LANGUAGE HIP)
     # PUBLIC so test consumers (test_dflash / test_draft_topk_cuda) also take the
     # GPU draft top-K path instead of the CPU fallback.
     target_compile_definitions(dflash_common PUBLIC DFLASH27B_HAVE_DRAFT_TOPK=1)
@@ -597,7 +602,8 @@ elseif(DFLASH27B_GPU_BACKEND STREQUAL "cuda")
     target_sources(dflash_common PRIVATE
         src/flashprefill_select.cpp
         src/flashprefill.cpp
-        src/common/geometric_draft_topk_cuda.cu)
+        src/common/geometric_draft_topk_cuda.cu
+        src/common/specla_commit_cuda.cu)
     # PUBLIC so consumers (e.g. the test_dflash executable) also see the macro
     # and take the GPU draft top-K path instead of the CPU fallback. Same macro
     # name as the HIP branch above (backend-neutral).
@@ -1459,6 +1465,18 @@ if(DFLASH27B_TESTS)
         list(APPEND _raw_unit_test_targets test_recurrent_snapshot)
     endif()
 
+    if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/test/test_delta_net_specla.cpp")
+        # GPU parity test: SpecLA topology-masked verify + factor-based state
+        # reconstruction vs the fused sequential gated-delta-net kernel.
+        # Exits 77 (ctest SKIP) when no GPU is present.
+        add_executable(test_delta_net_specla test/test_delta_net_specla.cpp)
+        target_include_directories(test_delta_net_specla PRIVATE
+            ${DFLASH27B_SRC_INCLUDE_DIRS})
+        target_link_libraries(test_delta_net_specla PRIVATE
+            dflash_common ggml ${DFLASH27B_GGML_BACKEND_TARGET})
+        list(APPEND _raw_unit_test_targets test_delta_net_specla)
+    endif()
+
     if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/test/test_server_unit.cpp")
         set(_server_unit_sources
             test/test_unit_main.cpp
@@ -1471,6 +1489,7 @@ if(DFLASH27B_TESTS)
             test/test_admission.cpp
             test/test_restore_delta.cpp
             test/test_chain_rollback_policy.cpp
+            test/test_ddtree_tau.cpp
             test/test_anchor_transitive.cpp
             test/test_drafter_early_exit_score_range.cpp
             test/test_drafter_tail_capture_guard.cpp
@@ -1612,7 +1631,8 @@ if(DFLASH27B_TESTS)
                     set(_unit_ctest_name paged_kv_pool)
                 endif()
                 add_test(NAME "${_unit_ctest_name}" COMMAND ${_unit_target})
-                if(_unit_target STREQUAL "test_deepseek4_mmid_grouped_cuda")
+                if(_unit_target STREQUAL "test_deepseek4_mmid_grouped_cuda" OR
+                   _unit_target STREQUAL "test_delta_net_specla")
                     set_tests_properties("${_unit_ctest_name}" PROPERTIES SKIP_RETURN_CODE 77)
                 endif()
                 if(_unit_target STREQUAL "test_rocmfp3_mix_registry")

server/deps/llama.cpp/ggml/include/ggml.h

@@ -2696,6 +2696,26 @@ extern "C" {
             struct ggml_tensor  * c,
             struct ggml_tensor  * parent_ids);
 
+    // SpecLA heavy-light convolution. Applies compact accepted inputs to the
+    // durable conv state, then verifies the current tree without committing
+    // speculative inputs. Current input factors are written directly to the
+    // persistent double-buffer selected by factor_ptrs/layer/bank; the result
+    // packs [conv output | boundary windows] and already includes SiLU.
+    GGML_API struct ggml_tensor * ggml_ssm_conv_specla(
+            struct ggml_context * ctx,
+            struct ggml_tensor  * x,
+            struct ggml_tensor  * c,
+            struct ggml_tensor  * state,
+            struct ggml_tensor  * hld,
+            struct ggml_tensor  * factor_ptrs,
+            int                   n_layers,
+            int                   layer,
+            int                   pending_bank,
+            int                   n_boundaries,
+            int                   n_chains,
+            int                   n_waves,
+            int                   max_parallel_chains);
+
     GGML_API struct ggml_tensor * ggml_ssm_scan(
             struct ggml_context * ctx,
             struct ggml_tensor  * s,
@@ -2866,6 +2886,28 @@ extern "C" {
             struct ggml_tensor  * parent_ids,
             struct ggml_tensor  * persist_inter);
 
+    // SpecLA state-resident heavy-light verify. The kernel applies the compact
+    // factors accepted in the preceding step, writes only that committed base
+    // state, then verifies the current HLD chains while writing raw
+    // (k, delta, log-gate) factors directly to the opposite persistent bank.
+    GGML_API struct ggml_tensor * ggml_gated_delta_net_specla(
+            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  * hld,
+            struct ggml_tensor  * factor_ptrs,
+            int                   n_layers,
+            int                   layer,
+            int                   pending_bank,
+            int                   n_boundaries,
+            int                   n_chains,
+            int                   n_waves,
+            int                   max_parallel_chains);
+
     // custom operators
 
     typedef void (*ggml_custom1_op_t)(struct ggml_tensor * dst , const struct ggml_tensor * a, int ith, int nth, void * userdata);

server/deps/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.cpp

@@ -472,6 +472,15 @@ static bool ggml_backend_cpu_device_supports_op(ggml_backend_dev_t dev, const st
             return !ggml_flash_attn_ext_is_ds4(op);
         case GGML_OP_PAGED_ATTN:
             return false;
+        case GGML_OP_SSM_CONV:
+            // The Specla layout (op param 0 == 1) needs the packed HLD state and
+            // is only supported by the CUDA kernel; the generic CPU kernel would
+            // silently compute garbage.
+            return ggml_get_op_params_i32(op, 0) != 1;
+        case GGML_OP_GATED_DELTA_NET:
+            // The Specla GDN variant (op param 2 == 1) is stateful via HLD and is
+            // only supported by the CUDA kernel.
+            return ggml_get_op_params_i32(op, 2) != 1;
         case GGML_OP_OUT_PROD:
             return (src0->type == GGML_TYPE_F32 || (ggml_is_quantized(src0->type) && src0->ne[2] == src1->ne[2] && src0->ne[3] == src1->ne[3])) &&
                 src1->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32;

server/deps/llama.cpp/ggml/src/ggml-cuda/gated_delta_net.cu

@@ -545,7 +545,226 @@ static void launch_gated_delta_net(
     }
 }
 
+template <int S_v>
+__global__ void gated_delta_net_specla_hld_cuda(
+        const float * __restrict__ q,
+        const float * __restrict__ k,
+        const float * __restrict__ v,
+        const float * __restrict__ g,
+        const float * __restrict__ beta,
+        float * __restrict__ durable_state,
+        const int * __restrict__ meta,
+        const int64_t * __restrict__ factor_ptrs,
+        float * __restrict__ packed,
+        int64_t H,
+        int64_t n_tokens,
+        int n_layers,
+        int layer,
+        int pending_bank,
+        int64_t sq1,
+        int64_t sq2,
+        int64_t sv1,
+        int64_t sv2,
+        int64_t sb1,
+        int64_t sb2,
+        int n_chains,
+        int wave,
+        float scale) {
+    const int h_idx = blockIdx.x;
+    const int wave_chain = blockIdx.y;
+    constexpr int warp_size =
+        ggml_cuda_get_physical_warp_size() < S_v ?
+        ggml_cuda_get_physical_warp_size() : S_v;
+    constexpr int rows_per_lane = (S_v + warp_size - 1) / warp_size;
+    const int lane = threadIdx.x;
+    const int col = blockIdx.z * blockDim.y + threadIdx.y;
+    if (col >= S_v) return;
+
+    const int order_off    = meta[6];
+    const int offsets_off  = meta[7];
+    const int parent_off   = meta[8];
+    const int boundary_off = meta[9];
+    const int wave_off     = meta[10];
+    int chain = 0;
+    while (chain < n_chains && meta[wave_off + chain] < wave) ++chain;
+    chain += wave_chain;
+    if (chain >= n_chains || meta[wave_off + chain] != wave) return;
+
+    const int64_t plane_offset = (int64_t)h_idx*S_v*S_v;
+    float * state_plane = durable_state + plane_offset;
+    const int pbase = pending_bank*4;
+    const int cbase = (1 - pending_bank)*4;
+    const float * pending_k = (const float *)(uintptr_t)factor_ptrs[pbase + 0];
+    const float * pending_v = (const float *)(uintptr_t)factor_ptrs[pbase + 1];
+    const float * pending_g = (const float *)(uintptr_t)factor_ptrs[pbase + 2];
+    float * current_k = (float *)(uintptr_t)factor_ptrs[cbase + 0];
+    float * current_v = (float *)(uintptr_t)factor_ptrs[cbase + 1];
+    float * current_g = (float *)(uintptr_t)factor_ptrs[cbase + 2];
+    float state_shard[rows_per_lane];
+    const int parent_boundary = meta[parent_off + chain];
+    const int64_t attn_elems = (int64_t)S_v*H*n_tokens;
+    const int64_t boundary_base = attn_elems;
+
+    if (parent_boundary < 0) {
+#pragma unroll
+        for (int r = 0; r < rows_per_lane; ++r) {
+            const int row = r*warp_size + lane;
+            state_shard[r] = state_plane[(int64_t)col*S_v + row];
+        }
+
+        // Delayed commit of the preceding accepted path. Factors have already
+        // been compacted into path order, so this is the exact serial
+        // recurrence and does not touch any rejected branch.
+        const int pending_count = meta[5];
+        for (int t = 0; t < pending_count; ++t) {
+            const int64_t th = ((int64_t)t*n_layers + layer)*H + h_idx;
+            const float g_val = expf(pending_g[th]);
+            const float delta = pending_v[th*S_v + col];
+#pragma unroll
+            for (int r = 0; r < rows_per_lane; ++r) {
+                const int row = r*warp_size + lane;
+                const float k_val = pending_k[th*S_v + row];
+                state_shard[r] = fmaf(k_val, delta, g_val*state_shard[r]);
+            }
+        }
+#pragma unroll
+        for (int r = 0; r < rows_per_lane; ++r) {
+            const int row = r*warp_size + lane;
+            state_plane[(int64_t)col*S_v + row] = state_shard[r];
+        }
+    } else {
+        const float * boundary = packed + boundary_base +
+            ((int64_t)parent_boundary*H + h_idx)*S_v*S_v +
+            (int64_t)col*S_v;
+#pragma unroll
+        for (int r = 0; r < rows_per_lane; ++r) {
+            const int row = r*warp_size + lane;
+            state_shard[r] = boundary[row];
+        }
+    }
+
+    const int begin = meta[offsets_off + chain];
+    const int end   = meta[offsets_off + chain + 1];
+    for (int p = begin; p < end; ++p) {
+        const int node = meta[order_off + p];
+        const float * q_t = q + (int64_t)node*sq2 + (int64_t)h_idx*sq1;
+        const float * k_t = k + (int64_t)node*sq2 + (int64_t)h_idx*sq1;
+        const float * v_t = v + (int64_t)node*sv2 + (int64_t)h_idx*sv1;
+        const int64_t gb = (int64_t)node*sb2 + (int64_t)h_idx*sb1;
+        const float g_log = g[gb];
+        const float g_val = expf(g_log);
+        const float beta_val = beta[gb];
+
+        float kv_partial = 0.0f;
+        float k_reg[rows_per_lane];
+        float q_reg[rows_per_lane];
+#pragma unroll
+        for (int r = 0; r < rows_per_lane; ++r) {
+            const int row = r*warp_size + lane;
+            k_reg[r] = k_t[row];
+            q_reg[r] = q_t[row];
+            kv_partial += state_shard[r]*k_reg[r];
+        }
+        const float kv_col = warp_reduce_sum<warp_size>(kv_partial);
+        const float delta = (v_t[col] - g_val*kv_col)*beta_val;
+
+        float attn_partial = 0.0f;
+#pragma unroll
+        for (int r = 0; r < rows_per_lane; ++r) {
+            state_shard[r] = fmaf(k_reg[r], delta, g_val*state_shard[r]);
+            attn_partial += state_shard[r]*q_reg[r];
+        }
+        const float attn_col = warp_reduce_sum<warp_size>(attn_partial);
+        if (lane == 0) {
+            packed[((int64_t)node*H + h_idx)*S_v + col] =
+                attn_col*scale;
+            const int64_t nh = ((int64_t)node*n_layers + layer)*H + h_idx;
+            current_v[nh*S_v + col] = delta;
+            if (col == 0) current_g[nh] = g_log;
+        }
+        if (col == 0) {
+            const int64_t nh = ((int64_t)node*n_layers + layer)*H + h_idx;
+#pragma unroll
+            for (int r = 0; r < rows_per_lane; ++r) {
+                const int row = r*warp_size + lane;
+                current_k[nh*S_v + row] = k_reg[r];
+            }
+        }
+
+        const int boundary_slot = meta[boundary_off + node];
+        if (boundary_slot >= 0) {
+            float * boundary = packed + boundary_base +
+                ((int64_t)boundary_slot*H + h_idx)*S_v*S_v +
+                (int64_t)col*S_v;
+#pragma unroll
+            for (int r = 0; r < rows_per_lane; ++r) {
+                const int row = r*warp_size + lane;
+                boundary[row] = state_shard[r];
+            }
+        }
+    }
+}
+
+static void launch_gated_delta_net_specla(
+        ggml_backend_cuda_context & ctx,
+        ggml_tensor * dst) {
+    ggml_tensor * q = dst->src[0];
+    ggml_tensor * k = dst->src[1];
+    ggml_tensor * v = dst->src[2];
+    ggml_tensor * g = dst->src[3];
+    ggml_tensor * beta = dst->src[4];
+    ggml_tensor * state = dst->src[5];
+    ggml_tensor * hld = dst->src[6];
+    ggml_tensor * factor_ptrs = dst->src[7];
+    const int S_v = (int)v->ne[0];
+    const int H = (int)v->ne[1];
+    const int n_tokens = (int)v->ne[2];
+    const int n_chains = ggml_get_op_params_i32(dst, 4);
+    const int n_waves = ggml_get_op_params_i32(dst, 5);
+    const int n_layers = ggml_get_op_params_i32(dst, 6);
+    const int layer = ggml_get_op_params_i32(dst, 7);
+    const int pending_bank = ggml_get_op_params_i32(dst, 8);
+    const int max_parallel_chains = ggml_get_op_params_i32(dst, 9);
+    GGML_ASSERT(v->ne[3] == 1 && g->ne[0] == 1);
+    GGML_ASSERT(hld->type == GGML_TYPE_I32 && n_chains > 0 && n_waves > 0);
+    GGML_ASSERT(ggml_is_contiguous(state));
+
+    const int warp_size = ggml_cuda_info().devices[ggml_cuda_get_device()].warp_size;
+    constexpr int num_warps = 4;
+    const dim3 block(warp_size <= S_v ? warp_size : S_v, num_warps, 1);
+    const dim3 grid((unsigned)H, (unsigned)max_parallel_chains,
+                    (unsigned)((S_v + num_warps - 1)/num_warps));
+    const float scale = 1.0f/sqrtf((float)S_v);
+    auto launch = [&](auto SV, int wave) {
+        constexpr int kSV = decltype(SV)::value;
+        gated_delta_net_specla_hld_cuda<kSV><<<grid, block, 0, ctx.stream()>>>(
+            (const float *)q->data, (const float *)k->data,
+            (const float *)v->data, (const float *)g->data,
+            (const float *)beta->data, (float *)state->data,
+            (const int *)hld->data,
+            (const int64_t *)factor_ptrs->data,
+            (float *)dst->data, H, n_tokens, n_layers, layer, pending_bank,
+            q->nb[1]/sizeof(float), q->nb[2]/sizeof(float),
+            v->nb[1]/sizeof(float), v->nb[2]/sizeof(float),
+            beta->nb[1]/sizeof(float), beta->nb[2]/sizeof(float),
+            n_chains, wave, scale);
+    };
+    for (int wave = 0; wave < n_waves; ++wave) {
+        switch (S_v) {
+            case 16:  launch(std::integral_constant<int, 16>{}, wave); break;
+            case 32:  launch(std::integral_constant<int, 32>{}, wave); break;
+            case 64:  launch(std::integral_constant<int, 64>{}, wave); break;
+            case 128: launch(std::integral_constant<int, 128>{}, wave); break;
+            default: GGML_ABORT("Unsupported SpecLA GDN state size");
+        }
+    }
+}
+
 void ggml_cuda_op_gated_delta_net(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
+    if (ggml_get_op_params_i32(dst, 2) == 1) {
+        launch_gated_delta_net_specla(ctx, dst);
+        return;
+    }
     ggml_tensor * src_q     = dst->src[0];
     ggml_tensor * src_k     = dst->src[1];
     ggml_tensor * src_v     = dst->src[2];

server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu

@@ -4359,6 +4359,11 @@ static bool ggml_cuda_can_fuse(const struct ggml_cgraph *                cgraph,
         const ggml_tensor * ssm_conv = cgraph->nodes[node_idx];
         const ggml_tensor * silu     = cgraph->nodes[node_idx+1];
 
+        if (ggml_get_op_params_i32(ssm_conv, 0) == 1) {
+            // the Specla ssm_conv kernel applies SiLU itself
+            return false;
+        }
+
         if (ssm_conv->type != GGML_TYPE_F32 || silu->type != GGML_TYPE_F32) {
             return false;
         }
@@ -6148,6 +6153,11 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
             }
         }
         case GGML_OP_SSM_CONV: {
+            if (ggml_get_op_params_i32(op, 0) == 1) {
+                // Specla layout: x is [d_inner, n_tokens], so d_inner = ne[0]
+                // and the kernel requires d_inner % threads == 0
+                return op->src[0]->ne[0] % 128 == 0;
+            }
             // assumes d_inner % threads == 0
             return op->src[0]->ne[1] % 128 == 0;
         }