delta-net: fix np>1 hybrid recurrent-state corruption (batched multi-seq)

ggml/src/ggml-cuda/ssm-conv.cu

@@ -574,10 +574,15 @@ void ggml_cuda_op_ssm_conv(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
 
         // Fast path for multi-sequence decode-like batches:
         // one token per unique sequence, no copy-to-multiple-sequences routing.
-        ggml_cuda_pool_alloc<int32_t> seq_ids(ctx.pool(), n_t);
-        ggml_cuda_pool_alloc<int32_t> seq_seen(ctx.pool(), n_kv);
+        // PXA_LLAMA_FIX_v4: the VMM cuda pool requires deallocations in REVERSE order of allocation. fast_path_ok_d is
+        // declared at function scope (freed LAST), so it must also be ALLOCATED FIRST among these three, otherwise when
+        // seq_ids/seq_seen (block-scoped) are freed at the end of this block while fast_path_ok_d is still live, the free
+        // order violates LIFO -> GGML_ASSERT(ptr == pool_addr + pool_used). Previously latent: the n_kv>1 path was never
+        // hit (decode batches were single-token-chunked); the batched mixed-seq delta-net now exercises it.
         int32_t fast_path_ok = 1;
         fast_path_ok_d.alloc(1);
+        ggml_cuda_pool_alloc<int32_t> seq_ids(ctx.pool(), n_t);
+        ggml_cuda_pool_alloc<int32_t> seq_seen(ctx.pool(), n_kv);
 
         CUDA_CHECK(cudaMemsetAsync(seq_seen.get(), 0, n_kv * sizeof(int32_t), ctx.stream()));
         CUDA_CHECK(cudaMemcpyAsync(fast_path_ok_d.get(), &fast_path_ok, sizeof(int32_t), cudaMemcpyHostToDevice, ctx.stream()));

ggml/src/ggml.c

@@ -19012,8 +19012,16 @@ static void ggml_compute_forward_set_rows_f32(
 
                     GGML_ASSERT(i1 >= 0 && i1 < ne1);
 
-                    from_float((const float *) ((char *) src0->data +  i*nb01 + i02*nb02 + i03*nb03),
-                            ((char *)  dst->data + i1*nb1  + i02*nb2  + i03*nb3), nc);
+                    if (dst->type == GGML_TYPE_F32) {
+                        // PXA: CPU set_rows into an F32 dst -- type_traits[F32].from_float is NULL,
+                        // so copy floats directly (mirrors the CUDA set_rows F32 path). This is what
+                        // the batched delta-net recurrent-state scatter needs on the CPU backend.
+                        memcpy(((char *) dst->data + i1*nb1 + i02*nb2 + i03*nb3),
+                               ((const char *) src0->data + i*nb01 + i02*nb02 + i03*nb03), nc*sizeof(float));
+                    } else {
+                        from_float((const float *) ((char *) src0->data +  i*nb01 + i02*nb02 + i03*nb03),
+                                ((char *)  dst->data + i1*nb1  + i02*nb2  + i03*nb3), nc);
+                    }
                 }
             }
         }
@@ -19030,8 +19038,16 @@ static void ggml_compute_forward_set_rows_f32(
 
                     GGML_ASSERT(i1 >= 0 && i1 < ne1);
 
-                    from_float((const float *) ((char *) src0->data +  i*nb01 + i02*nb02 + i03*nb03),
-                            ((char *)  dst->data + i1*nb1  + i02*nb2  + i03*nb3), nc);
+                    if (dst->type == GGML_TYPE_F32) {
+                        // PXA: CPU set_rows into an F32 dst -- type_traits[F32].from_float is NULL,
+                        // so copy floats directly (mirrors the CUDA set_rows F32 path). This is what
+                        // the batched delta-net recurrent-state scatter needs on the CPU backend.
+                        memcpy(((char *) dst->data + i1*nb1 + i02*nb2 + i03*nb3),
+                               ((const char *) src0->data + i*nb01 + i02*nb02 + i03*nb03), nc*sizeof(float));
+                    } else {
+                        from_float((const float *) ((char *) src0->data +  i*nb01 + i02*nb02 + i03*nb03),
+                                ((char *)  dst->data + i1*nb1  + i02*nb2  + i03*nb3), nc);
+                    }
                 }
             }
         }

src/graphs/build_qwen35.cpp

@@ -40,6 +40,16 @@ ggml_cgraph * llm_build_context::build_qwen35moe() {
         cb(lctx.inp_s_seq_qnext, "inp_s_seq_qnext", -1);
         ggml_set_input(lctx.inp_s_seq_qnext);
 
+        // PXA_LLAMA_FIX_v4: conv seq-map [n_kv=n_tokens, n_tokens] for the ONE-batched mixed-seq delta-net path.
+        lctx.inp_conv_seq_map = ggml_new_tensor_2d(ctx0, GGML_TYPE_I32, n_tokens, n_tokens);
+        cb(lctx.inp_conv_seq_map, "inp_conv_seq_map", -1);
+        ggml_set_input(lctx.inp_conv_seq_map);
+
+        // PXA_LLAMA_FIX_v4: per-seq recurrent-state reset mask (0=reset to zero, 1=keep).
+        lctx.inp_qnext_state_mask = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, 1, n_tokens);
+        cb(lctx.inp_qnext_state_mask, "inp_qnext_state_mask", -1);
+        ggml_set_input(lctx.inp_qnext_state_mask);
+
         float KQ_scale = hparams.f_attention_scale == 0.0f ? 1.0f / sqrtf(float(n_embd_head)) : hparams.f_attention_scale;
 
         const int n_transformer_layers = n_layer - hparams.nextn_predict_layers;
@@ -124,6 +134,16 @@ ggml_cgraph * llm_build_context::build_qwen35() {
         cb(lctx.inp_s_seq_qnext, "inp_s_seq_qnext", -1);
         ggml_set_input(lctx.inp_s_seq_qnext);
 
+        // PXA_LLAMA_FIX_v4: conv seq-map [n_kv=n_tokens, n_tokens] for the ONE-batched mixed-seq delta-net path.
+        lctx.inp_conv_seq_map = ggml_new_tensor_2d(ctx0, GGML_TYPE_I32, n_tokens, n_tokens);
+        cb(lctx.inp_conv_seq_map, "inp_conv_seq_map", -1);
+        ggml_set_input(lctx.inp_conv_seq_map);
+
+        // PXA_LLAMA_FIX_v4: per-seq recurrent-state reset mask (0=reset to zero, 1=keep).
+        lctx.inp_qnext_state_mask = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, 1, n_tokens);
+        cb(lctx.inp_qnext_state_mask, "inp_qnext_state_mask", -1);
+        ggml_set_input(lctx.inp_qnext_state_mask);
+
         float KQ_scale = hparams.f_attention_scale == 0.0f ? 1.0f / sqrtf(float(n_embd_head)) : hparams.f_attention_scale;
 
         cur = nullptr;

src/graphs/build_qwen3next.cpp

@@ -21,6 +21,17 @@ ggml_cgraph * llm_build_context::build_qwen3next() {
     cb(lctx.inp_s_seq_qnext, "inp_s_seq_qnext", -1);
     ggml_set_input(lctx.inp_s_seq_qnext);
 
+    // PXA_LLAMA_FIX_v4: conv seq-map [n_kv=n_tokens, n_tokens] for the ONE-batched mixed-seq delta-net path.
+    // Row 0 of column t = state column for token t (identity, since states are gathered in token order); rows >=1 = -1.
+    lctx.inp_conv_seq_map = ggml_new_tensor_2d(ctx0, GGML_TYPE_I32, n_tokens, n_tokens);
+    cb(lctx.inp_conv_seq_map, "inp_conv_seq_map", -1);
+    ggml_set_input(lctx.inp_conv_seq_map);
+
+    // PXA_LLAMA_FIX_v4: per-seq recurrent-state reset mask (0=reset to zero, 1=keep). One column per token/seq.
+    lctx.inp_qnext_state_mask = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, 1, n_tokens);
+    cb(lctx.inp_qnext_state_mask, "inp_qnext_state_mask", -1);
+    ggml_set_input(lctx.inp_qnext_state_mask);
+
     float KQ_scale = hparams.f_attention_scale == 0.0f ? 1.0f / sqrtf(float(n_embd_head)) : hparams.f_attention_scale;
 
     ggml_tensor * cur = nullptr;

src/llama-build-context.cpp

@@ -109,6 +109,8 @@ void llm_build_context::init() {
     lctx.inp_s_mask      = nullptr;
     lctx.inp_s_seq       = nullptr;
     lctx.inp_s_seq_qnext = nullptr;
+    lctx.inp_conv_seq_map  = nullptr; // PXA_LLAMA_FIX_v4
+    lctx.inp_qnext_state_mask = nullptr; // PXA_LLAMA_FIX_v4
     lctx.inp_pos_bucket    = nullptr;
     lctx.inp_embd_enc      = nullptr;
     lctx.inp_KQ_mask_cross = nullptr;

src/llama-context.h

@@ -292,6 +292,8 @@ struct llama_context {
     struct ggml_tensor * inp_s_mask;      // F32 [1, n_kv]
     struct ggml_tensor * inp_s_seq;       // I32 [n_kv, n_batch]
     struct ggml_tensor * inp_s_seq_qnext; // I32 [1, n_batch]
+    struct ggml_tensor * inp_conv_seq_map;  // PXA_LLAMA_FIX_v4: I32 [n_batch, n_batch] conv seq-map for batched mixed-seq delta-net
+    struct ggml_tensor * inp_qnext_state_mask; // PXA_LLAMA_FIX_v4: F32 [1, n_batch] per-seq recurrent-state reset mask (0=reset)
     struct ggml_tensor * inp_pos_bucket;    // I32 [n_batch|n_kv, n_batch]
     struct ggml_tensor * inp_embd_enc;      // F32 [n_embd, n_outputs_enc]
     struct ggml_tensor * inp_KQ_mask_cross; // F32 [n_outputs_enc, n_batch]