v4: fix multi-GPU SET_ROWS crash via cpy-into-view substitution

docs/plans/v4-multigpu-set-rows-fix-v2.md

@@ -0,0 +1,142 @@
+# V4 multi-GPU SET_ROWS fix — Plan v2 (second attempt, V4-graph-only)
+
+Spec: `.claude/agents/v4-multigpu-set-rows-fix.md` (rewritten 2026-05-13)
+Branch: `fix/v4-cuda-multigpu-supports-op` (baseline commit `19c3f8fd9`, kept in place)
+Hard constraint: changes confined to `src/models/deepseek4.cpp`. No edits under `ggml/src/`.
+
+## Why this attempt is different
+
+Attempt 1 tried to patch ggml-cuda's `supports_op` to walk the view-src chain and read the destination buffer. That was a no-op because the scheduler invokes `supports_op` *before* it has allocated buffers for the intermediate K-cache state tensors, so `op->view_src->buffer` (and its chain root) is `nullptr` at query time. The commit (`19c3f8fd9`) is harmless and stays as defensive logic for other failure modes.
+
+This attempt avoids the scheduler entirely by sidestepping the problematic op at the V4-graph level. `ggml_set_rows` is routed by source-device affinity; `ggml_cpy`-into-view is routed by destination-buffer affinity. The latter already works correctly in production via the `dsv4_store_state_segment` precedent (which has never been implicated in any multi-GPU crash log). We substitute `cpy`-into-view for `set_rows` at the two V4 sites that touch recurrent state / K-cache.
+
+## Site inventory
+
+`grep -n ggml_set_rows src/models/deepseek4.cpp` will turn up three call sites. Only two are in scope:
+
+| # | Function                                       | Line(s) | In scope? |
+|---|------------------------------------------------|---------|-----------|
+| 1 | `dsv4_store_cache_rows`                        | 405-406 | YES       |
+| 2 | `dsv4_build_compressor_decode_projected`       | 864-866 | YES       |
+| 3 | Lightning indexer top-k write (`mask = ggml_set_rows(...)`) | ~1053 | NO — different pattern (sparse), not implicated in crash, spec explicitly excludes |
+
+## Change 1 — `dsv4_store_cache_rows` (lines 391-407)
+
+Replace the trailing two statements:
+```cpp
+ggml_tensor * rows = dsv4_arange_i32(ctx, row_start, row_start + n_rows);
+ggml_build_forward_expand(gf, ggml_set_rows(ctx, cache, src, rows));
+```
+
+with a single contiguous CPY-into-view, modeled on `dsv4_store_state_segment` (lines 375-389):
+```cpp
+// Avoid ggml_set_rows here: on multi-GPU, sched routes set_rows by SOURCE
+// device, but the cache destination has its own device affinity → illegal
+// memory access when those differ. ggml_cpy into a contiguous view of
+// cache routes correctly by dst affinity (same pattern as
+// dsv4_store_state_segment, which works in production multi-GPU).
+ggml_tensor * cache_view = ggml_view_2d(ctx, cache,
+        cache->ne[0], n_rows,
+        cache->nb[1],
+        row_start * cache->nb[1]);
+ggml_build_forward_expand(gf, ggml_cpy(ctx, src, cache_view));
+```
+
+Semantics: `src` has already been made contiguous and reshaped to `[cache->ne[0], n_rows]` two lines above. The view spans rows `[row_start, row_start + n_rows)` of `cache` with stride `cache->nb[1]`. CPY of the matching-shape tensor into that view writes exactly the same bytes that `set_rows` would have, into the same destination addresses. Because we're writing CONTIGUOUS rows here (the `rows` argument was just `arange(row_start, row_start + n_rows)`), the row-stride view is equivalent — no indirection needed.
+
+Drop the now-unused `rows` local. The `dsv4_arange_i32` import call disappears from this call site (still used elsewhere — site 2, indexer).
+
+## Change 2 — `dsv4_build_compressor_decode_projected` (lines 864-866)
+
+Single-row write at row index `row` (a scalar known at graph-build time). Replace:
+```cpp
+ggml_tensor * row_idx = dsv4_arange_i32(ctx, row, row + 1);
+ggml_tensor * kv_state    = ggml_set_rows(ctx, prev_kv_state,    kv_cur, row_idx);
+ggml_tensor * score_state = ggml_set_rows(ctx, prev_score_state, sc_cur, row_idx);
+```
+
+with:
+```cpp
+// Single-row write via cpy-into-view. set_rows would crash on multi-GPU
+// (see dsv4_store_cache_rows for the same problem and fix).
+auto cpy_into_row = [&](ggml_tensor * dst, ggml_tensor * src) -> ggml_tensor * {
+    ggml_tensor * view = ggml_view_2d(ctx, dst,
+            dst->ne[0], 1,
+            dst->nb[1],
+            row * dst->nb[1]);
+    return ggml_cpy(ctx, src, view);  // returns a view-of-dst with op=CPY, src[0]=src
+};
+
+ggml_tensor * kv_state    = cpy_into_row(prev_kv_state,    kv_cur);
+ggml_tensor * score_state = cpy_into_row(prev_score_state, sc_cur);
+```
+
+### Dependency-chain correctness
+
+`ggml_cpy(ctx, src, dst)` returns a tensor whose `op = GGML_OP_CPY`, `src[0] = src`, `view_src = dst`. Downstream code at lines 869-894 reads from `kv_state` / `score_state` via `dsv4_view_cols(...)` (`ggml_view_2d`). Those new views have `view_src = kv_state` (the cpy result), and their own descendants ultimately bottom out in the cpy node when ggml walks the dataflow during scheduling.
+
+Because `kv_state` IS the cpy result (not a fresh view of `prev_kv_state`), any consumer of `kv_state` in the forward graph creates a transitive dependency on the cpy node. The scheduler will therefore order the cpy before the consumer. This matches the dependency semantics of the original `ggml_set_rows` (which also returned a view-tagged-as-set_rows-op).
+
+NOTE: The orchestrator's earlier draft attempted to wrap the cpy in `ggml_build_forward_expand(gf, ...)` and return a fresh view of `prev_kv_state`. That approach is broken on TWO counts:
+1. The cpy result tensor would have no in-graph consumer, so the dependency would rely on the explicit forward_expand call.
+2. `gf` is NOT plumbed into `dsv4_build_compressor_decode_projected` — it's not in the function signature. Trying to use it would not compile.
+
+Returning the cpy result directly resolves both issues. The cpy will be visited via the normal graph traversal that the caller's `ggml_build_forward_expand(gf, ...)` already does on whatever consumes `kv_state` downstream.
+
+### Edge cases checked
+
+- `row_start = 0`: view offset is 0, valid for `ggml_view_2d`.
+- `n_rows = 1`: same shape as site 2 case, valid.
+- `src not yet contiguous`: site 1 already calls `ggml_cont` + `ggml_reshape_2d` before the write — no change. Site 2's `kv_cur`/`sc_cur` are mul_mat outputs, which are F32 contiguous `[width, 1]` tensors and match the view shape exactly. No additional `ggml_cont` needed.
+- Non-contiguous row stride: cache->nb[1] is the natural row stride; the view we create has the same nb[1]. CPY preserves the dst layout, so this is fine for any cache->nb[1] (contiguous or padded).
+
+## What is NOT changing
+
+- Commit `19c3f8fd9` (view-chain walk in supports_op) — kept in place as harmless defense.
+- Commits `13df7dfe3`, `5db52f6d9`, `1aed9b5e9`, `ca8734ab6` — kept (split-buffer source allowance + diagnostics + earlier dst-device check). They cover different failure modes from this bug.
+- Line 1053 lightning-indexer set_rows — different pattern (sparse top-k mask write), not in the crash logs.
+- `dsv4_store_state_segment` — already uses the correct cpy-into-view pattern.
+- Any file under `ggml/src/`.
+
+## Build & verify locally (Mac/Metal)
+
+```
+cmake --build build --target llama-server -j
+```
+
+Mac build only validates that the C++ compiles. CUDA-path correctness needs gpudual.
+
+## Validation on gpudual (matches spec §"Validation plan")
+
+1. `git fetch && git reset --hard origin/fix/v4-cuda-multigpu-supports-op` then `cmake --build build-cuda -j --target llama-cli test-backend-ops`. Expect a clean build.
+2. `./build-cuda/bin/test-backend-ops -o DSV4_ROPE_TAIL,DSV4_HC_SPLIT_SINKHORN,DSV4_HC_WEIGHTED_SUM,DSV4_HC_EXPAND,DSV4_FP8_KV_QUANTIZE`. Expect 19/19 pass.
+3. The repro script with `--model …Q2_K-XL…00001-of-00003.gguf`, `-ngl 999 -cmoe -ub 128 --ctx-size 8192 -p "hi" -n 20`, kill-timer `perl -e "alarm 600; exec @ARGV"`. Expect 20 generated tokens, no `SET_ROWS failed` / `CUDA error` / `Aborted` in output.
+
+## Failure handling
+
+If check 3 fails with a CPY-related error (e.g., `CPY failed` + illegal memory access), per the spec we DON'T retry V4-graph variations. We update the JSON state to `blocked-on-ggml-core-change` with the empirical finding "cpy also misrouted by sched" and stop. The user must then decide whether to relax the no-ggml-core constraint.
+
+If check 3 fails with the same `SET_ROWS failed` somewhere we didn't expect, that means there's a third site we missed. Grep for `ggml_set_rows` in deepseek4.cpp again, double-check, and if it really is one of these two sites, mark blocked and report.
+
+If check 1 or 2 fails (build broken or per-op regression), this is a coding error — fix the code, recommit, rerun. Don't escalate to ggml-core constraint.
+
+## Commit message draft
+
+```
+v4: replace ggml_set_rows with ggml_cpy-into-view in compressor decode/store paths
+
+On multi-GPU CUDA, ggml_set_rows is routed by the device affinity of its
+SOURCES (kv_cur, row indices) after peer-copy, while the K-cache destination
+is anchored to a different device, producing an illegal memory access when
+they differ. ggml_cpy into a contiguous view of the destination routes by
+dst affinity and works in production today (cf. dsv4_store_state_segment).
+
+Substitute the working pattern at the two implicated V4 sites:
+- dsv4_store_cache_rows: contiguous N-row write via ggml_view_2d + ggml_cpy.
+- dsv4_build_compressor_decode_projected: single-row write via the same.
+
+No ggml/src/ changes; this is a graph-construction fix only.
+
+Co-Authored-By: cchuter <cchuter@yahoo.com>
+Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
+```

ggml/src/ggml-cuda/ggml-cuda.cu

@@ -2783,7 +2783,52 @@ static void ggml_cuda_mul_mat_id(ggml_backend_cuda_context & ctx, ggml_tensor *
         nb1, nb2, nb3, stream);
 }
 
+// ---------- DSV4 debug logging (env-gated, set GGML_DSV4_DEBUG=1 to enable) ----------
+static bool dsv4_debug_enabled() {
+    static const bool enabled = (getenv("GGML_DSV4_DEBUG") != nullptr);
+    return enabled;
+}
+
+static const char * dsv4_op_short(enum ggml_op op) {
+    switch (op) {
+        case GGML_OP_DSV4_ROPE_TAIL:         return "DSV4_ROPE_TAIL";
+        case GGML_OP_DSV4_HC_SPLIT_SINKHORN: return "DSV4_HC_SPLIT_SINKHORN";
+        case GGML_OP_DSV4_HC_WEIGHTED_SUM:   return "DSV4_HC_WEIGHTED_SUM";
+        case GGML_OP_DSV4_HC_EXPAND:         return "DSV4_HC_EXPAND";
+        case GGML_OP_DSV4_FP8_KV_QUANTIZE:   return "DSV4_FP8_KV_QUANTIZE";
+        default:                              return nullptr;
+    }
+}
+
+static bool dsv4_op_is_v4(enum ggml_op op) {
+    return dsv4_op_short(op) != nullptr;
+}
+
+static void dsv4_log_op_entry(int device, const struct ggml_tensor * dst) {
+    if (!dsv4_debug_enabled() || !dsv4_op_is_v4(dst->op)) return;
+    fprintf(stderr, "[DSV4_DEBUG] dev=%d op=%s dst=%s(%s) shape=[%lld,%lld,%lld,%lld]\n",
+            device, dsv4_op_short(dst->op),
+            dst->name, ggml_type_name(dst->type),
+            (long long) dst->ne[0], (long long) dst->ne[1],
+            (long long) dst->ne[2], (long long) dst->ne[3]);
+    for (int i = 0; i < GGML_MAX_SRC; i++) {
+        if (!dst->src[i]) continue;
+        const char * buft_name = "(null-buf)";
+        int is_split = 0;
+        if (dst->src[i]->buffer) {
+            buft_name = ggml_backend_buft_name(dst->src[i]->buffer->buft);
+            is_split  = ggml_backend_buft_is_cuda_split(dst->src[i]->buffer->buft) ? 1 : 0;
+        }
+        fprintf(stderr, "[DSV4_DEBUG]   src[%d]=%s(%s) buft=%s split=%d data=%p extra=%p\n",
+                i, dst->src[i]->name, ggml_type_name(dst->src[i]->type),
+                buft_name, is_split, dst->src[i]->data, (void *) dst->src[i]->extra);
+    }
+    fflush(stderr);
+}
+// ---------- end DSV4 debug ----------
+
 static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct ggml_tensor * dst) {
+    dsv4_log_op_entry(ctx.device, dst);
     switch (dst->op) {
         case GGML_OP_ARGMAX:
             ggml_cuda_argmax(ctx, dst);
@@ -3232,7 +3277,30 @@ static bool ggml_backend_cuda_cpy_tensor_async(ggml_backend_t backend_src, ggml_
 #ifdef GGML_CUDA_NO_PEER_COPY
             return false;
 #else
+            if (dsv4_debug_enabled()) {
+                fprintf(stderr, "[DSV4_DEBUG] peer-copy: src_dev=%d dst_dev=%d bytes=%zu "
+                                "src=%s(%s,op=%s,buft=%s) dst=%s(%s,op=%s,buft=%s) src_ptr=%p dst_ptr=%p\n",
+                        cuda_ctx_src->device, cuda_ctx_dst->device, ggml_nbytes(dst),
+                        src->name, ggml_type_name(src->type), ggml_op_name(src->op),
+                        src->buffer ? ggml_backend_buft_name(src->buffer->buft) : "?",
+                        dst->name, ggml_type_name(dst->type), ggml_op_name(dst->op),
+                        dst->buffer ? ggml_backend_buft_name(dst->buffer->buft) : "?",
+                        src->data, dst->data);
+                fflush(stderr);
+                // Force any deferred CUDA error to surface BEFORE the next op, so the log line
+                // immediately above truly identifies the failing copy (codex review nit #1).
+                cudaError_t pre_err = cudaGetLastError();
+                if (pre_err != cudaSuccess) {
+                    fprintf(stderr, "[DSV4_DEBUG] pre-copy stale error: %s\n", cudaGetErrorString(pre_err));
+                    fflush(stderr);
+                }
+            }
             CUDA_CHECK(cudaMemcpyPeerAsync(dst->data, cuda_ctx_dst->device, src->data, cuda_ctx_src->device, ggml_nbytes(dst), cuda_ctx_src->stream()));
+            if (dsv4_debug_enabled()) {
+                // Synchronous wait to force the async error (if any) to surface at the offending copy,
+                // not at some later API call. Heavy perturbation — only with GGML_DSV4_DEBUG=1.
+                CUDA_CHECK(cudaStreamSynchronize(cuda_ctx_src->stream()));
+            }
 #endif // GGML_CUDA_NO_PEER_COPY
         }
 
@@ -5058,8 +5126,17 @@ static ggml_backend_buffer_type_t ggml_backend_cuda_device_get_host_buffer_type(
 static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const ggml_tensor * op) {
     ggml_backend_cuda_device_context * dev_ctx = (ggml_backend_cuda_device_context *) dev->context;
 
-    // split buffers can only be used with GGML_OP_MUL_MAT
-    if (op->op != GGML_OP_MUL_MAT) {
+    // split buffers can only be used with GGML_OP_MUL_MAT and DeepSeek V4 custom ops.
+    // Without the DSV4 exception, multi-GPU scheduler rejects the V4 ops once their
+    // weight tensors land in cuda_split buffers and falls back to CPU — which then
+    // corrupts data via host<->device transfer mismatches and crashes during decode.
+    // Reported and root-caused by @DenisVASI9 on an 8x A100 40GB rig.
+    if (op->op != GGML_OP_MUL_MAT &&
+        op->op != GGML_OP_DSV4_HC_SPLIT_SINKHORN &&
+        op->op != GGML_OP_DSV4_HC_WEIGHTED_SUM &&
+        op->op != GGML_OP_DSV4_HC_EXPAND &&
+        op->op != GGML_OP_DSV4_FP8_KV_QUANTIZE &&
+        op->op != GGML_OP_DSV4_ROPE_TAIL) {
         for (int i = 0; i < GGML_MAX_SRC; i++) {
             if (op->src[i] && op->src[i]->buffer && ggml_backend_buft_is_cuda_split(op->src[i]->buffer->buft)) {
                 return false;
@@ -5077,6 +5154,30 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
         }
     }
 
+    // Some ops write through a pre-allocated destination buffer (e.g. SET_ROWS
+    // into a KV cache). For those, the dst lives on a specific device — dispatching
+    // the op on a different device causes the CUDA kernel to write through a
+    // foreign-device pointer (dst->data), surfacing as cudaErrorIllegalAddress.
+    //
+    // SET_ROWS returns a view tensor (ggml_view_tensor(ctx, a)) so op->buffer is
+    // nullptr. We must walk the view chain to find the real buffer.
+    // Diagnosed via CUDA_LAUNCH_BLOCKING=1 + GGML_DSV4_DEBUG=1 on @DenisVASI9's
+    // 8x A100 rig: V4's dsv4_store_cache_rows emits SET_ROWS at layer-7 K-cache
+    // (on CUDA1) while sched dispatched on CUDA0 → illegal access.
+    {
+        const ggml_tensor * t = op;
+        while (t->view_src) {
+            t = t->view_src;
+        }
+        if (t->buffer && ggml_backend_buft_is_cuda(t->buffer->buft)) {
+            ggml_backend_cuda_buffer_type_context * buft_ctx =
+                (ggml_backend_cuda_buffer_type_context *) t->buffer->buft->context;
+            if (buft_ctx->device != dev_ctx->device) {
+                return false;
+            }
+        }
+    }
+
     switch (op->op) {
         case GGML_OP_UNARY:
             switch (ggml_get_unary_op(op)) {

src/models/deepseek4.cpp

@@ -402,8 +402,16 @@ static void dsv4_store_cache_rows(
     src = ggml_cont(ctx, src);
     src = ggml_reshape_2d(ctx, src, cache->ne[0], n_rows);
 
-    ggml_tensor * rows = dsv4_arange_i32(ctx, row_start, row_start + n_rows);
-    ggml_build_forward_expand(gf, ggml_set_rows(ctx, cache, src, rows));
+    // Avoid ggml_set_rows here: on multi-GPU, sched routes set_rows by SOURCE
+    // device, but the cache destination has its own device affinity → illegal
+    // memory access when those differ. ggml_cpy into a contiguous view of
+    // cache routes correctly by dst affinity (same pattern as
+    // dsv4_store_state_segment, which works in production multi-GPU).
+    ggml_tensor * cache_view = ggml_view_2d(ctx, cache,
+            cache->ne[0], n_rows,
+            cache->nb[1],
+            row_start * cache->nb[1]);
+    ggml_build_forward_expand(gf, ggml_cpy(ctx, src, cache_view));
 }
 
 static dsv4_rope_cfg dsv4_make_rope_cfg(
@@ -861,9 +869,28 @@ static dsv4_decode_compressor dsv4_build_compressor_decode_projected(
     const int64_t row = compress_ratio == 4 ? compress_ratio + pos_mod : pos_mod;
     const bool should_compress = (pos + 1) % compress_ratio == 0;
 
-    ggml_tensor * row_idx = dsv4_arange_i32(ctx, row, row + 1);
-    ggml_tensor * kv_state    = ggml_set_rows(ctx, prev_kv_state,    kv_cur, row_idx);
-    ggml_tensor * score_state = ggml_set_rows(ctx, prev_score_state, sc_cur, row_idx);
+    // Single-row write via cpy-into-view. ggml_set_rows would crash on
+    // multi-GPU (sched routes by src device while dst is on a different
+    // device; see dsv4_store_cache_rows for the same problem and fix).
+    //
+    // We need to return a FULL-shape view of dst (downstream code at
+    // dsv4_view_cols slices the full state by columns/rows) AND establish
+    // a data dependency on the cpy. We mimic ggml_set_rows's internal
+    // construction: create a view_tensor of dst (which inherits dst's full
+    // shape), then manually set src[0] to the cpy result so sched orders
+    // the cpy before any consumer reading from this view.
+    auto cpy_into_row = [&](ggml_tensor * dst, ggml_tensor * row_src) -> ggml_tensor * {
+        ggml_tensor * row_view = ggml_view_2d(ctx, dst,
+                dst->ne[0], 1,
+                dst->nb[1],
+                row * dst->nb[1]);
+        ggml_tensor * cpy = ggml_cpy(ctx, row_src, row_view);
+        ggml_tensor * full_state = ggml_view_tensor(ctx, dst);
+        full_state->src[0] = cpy;  // dependency: full_state's consumers wait for cpy
+        return full_state;
+    };
+    ggml_tensor * kv_state    = cpy_into_row(prev_kv_state,    kv_cur);
+    ggml_tensor * score_state = cpy_into_row(prev_score_state, sc_cur);
     ggml_tensor * kv_comp = nullptr;
 
     if (should_compress) {