From 13df7dfe3d7d2662d0a1c2480e4a92104ce69e02 Mon Sep 17 00:00:00 2001 From: Chris Chuter Date: Wed, 13 May 2026 18:20:36 -0500 Subject: [PATCH 1/7] cuda: allow V4 DSV4 ops to run with split-buffer sources Multi-GPU scheduler was rejecting the 5 V4 custom ops (rope_tail, hc_split_sinkhorn, hc_weighted_sum, hc_expand, fp8_kv_quantize) at ggml_backend_cuda_device_supports_op because their weight tensors land in cuda_split buffers when V4 is split across multiple devices. The op then fell back to CPU, corrupting data via mismatched host<->device transfers and crashing in cudaMemcpyPeerAsync during prompt processing. Add the 5 DSV4 ops to the same exemption list that GGML_OP_MUL_MAT already uses. Single-GPU is unaffected (split-buffer check only fires when ggml_backend_buft_is_cuda_split() returns true, which requires multi-device tensor split). Reported, root-caused, and fixed independently by @DenisVASI9 on an 8x A100 40GB rig (his parallel feat/v4-port branch, commit 7a6a7a29d). This commit ports the same fix to feat/v4-port-cuda. Validates on his rig; needs cross-validation: - single-GPU: should be no-op (no behavior change at -ngl 999 on 1 GPU) - multi-GPU: expected to fix the cudaMemcpyPeerAsync crash on V4 inference Co-Authored-By: Denis Vasilyev Co-Authored-By: Claude Opus 4.7 (1M context) --- ggml/src/ggml-cuda/ggml-cuda.cu | 13 +++++++++++-- 1 file changed, 11 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index c79e87864f26..33fcc0c68856 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -5058,8 +5058,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; From 5db52f6d98f6a7cb06685dd99c16f5fdab420e6e Mon Sep 17 00:00:00 2001 From: Chris Chuter Date: Wed, 13 May 2026 18:34:58 -0500 Subject: [PATCH 2/7] cuda: env-gated DSV4 debug logging for multi-GPU diagnosis Adds two log sites, both gated by GGML_DSV4_DEBUG=1 environment variable (no overhead and zero log noise when unset): 1. At entry of ggml_cuda_compute_forward: when the op is one of the 5 DSV4 custom ops, prints the op type, target device, dst tensor name + shape, and for every source: its buffer type, whether it's a cuda_split buffer, the raw data pointer, and the extra pointer. 2. Just before every cudaMemcpyPeerAsync call: prints src/dst devices, bytes, both tensor names + types + originating ops, and both raw data pointers. The line printed immediately before the crash identifies the exact tensor and src->dst device pair that failed. Usage on the multi-GPU host where the cudaMemcpyPeerAsync crash repros: GGML_DSV4_DEBUG=1 ./build/bin/llama-server 2>&1 | tee crash.log Then trigger one request, wait for the crash, share the tail of crash.log. The diagnostic answers two questions directly: (a) are any DSV4 op sources landing in cuda_split buffers? (b) which peer copy was in flight when the illegal-memory-access fired? Co-Authored-By: Claude Opus 4.7 (1M context) --- ggml/src/ggml-cuda/ggml-cuda.cu | 54 +++++++++++++++++++++++++++++++++ 1 file changed, 54 insertions(+) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 33fcc0c68856..e49e9b74742c 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/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,6 +3277,15 @@ 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) dst=%s(%s,op=%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), + dst->name, ggml_type_name(dst->type), ggml_op_name(dst->op), + src->data, dst->data); + fflush(stderr); + } CUDA_CHECK(cudaMemcpyPeerAsync(dst->data, cuda_ctx_dst->device, src->data, cuda_ctx_src->device, ggml_nbytes(dst), cuda_ctx_src->stream())); #endif // GGML_CUDA_NO_PEER_COPY } From 1aed9b5e9213ff0d4e23e6ce8e9d18a978f8c73a Mon Sep 17 00:00:00 2001 From: Chris Chuter Date: Wed, 13 May 2026 18:39:44 -0500 Subject: [PATCH 3/7] cuda: tighten DSV4 peer-copy diagnostics (codex review nit) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Two improvements per codex review of the prior diagnostic commit: 1. After cudaMemcpyPeerAsync issuance, when GGML_DSV4_DEBUG=1, force a cudaStreamSynchronize on the src stream. CUDA async errors otherwise surface on a later, unrelated API call — meaning the last log line before the abort might NOT be the actual offending copy. The synchronize forces the error to surface at the offending peer copy so the diagnostic is accurate. Heavy perturbation (every peer copy becomes synchronous), only with debug enabled. 2. Before issuance, when debug enabled, check cudaGetLastError() and log any pre-existing stale error. Same goal: pin the failure to the copy that actually caused it. 3. Peer-copy log line now also includes src->buffer->buft and dst->buffer->buft names, so the cuda_split vs regular distinction shows up directly in the log line. Co-Authored-By: Claude Opus 4.7 (1M context) --- ggml/src/ggml-cuda/ggml-cuda.cu | 16 +++++++++++++++- 1 file changed, 15 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index e49e9b74742c..e80d44267877 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -3279,14 +3279,28 @@ static bool ggml_backend_cuda_cpy_tensor_async(ggml_backend_t backend_src, ggml_ #else if (dsv4_debug_enabled()) { fprintf(stderr, "[DSV4_DEBUG] peer-copy: src_dev=%d dst_dev=%d bytes=%zu " - "src=%s(%s,op=%s) dst=%s(%s,op=%s) src_ptr=%p dst_ptr=%p\n", + "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 } From ca8734ab6999bf4190160e61884739944cd231f0 Mon Sep 17 00:00:00 2001 From: Chris Chuter Date: Wed, 13 May 2026 19:18:07 -0500 Subject: [PATCH 4/7] cuda: refuse dispatch when destination buffer is on a different device MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The existing supports_op check validated that all SOURCE tensors live on ctx.device, but did NOT check the destination buffer. Most ops have a sched-allocated dst (always on ctx.device), so the gap is invisible — except for ops like GGML_OP_SET_ROWS that write into a pre-allocated buffer (e.g. a KV cache). For those, dst lives wherever the cache was allocated, which may differ from ctx.device. The kernel then writes through dst->data (a foreign-device pointer) → cudaErrorIllegalAddress. Diagnosed via CUDA_LAUNCH_BLOCKING=1 + GGML_DSV4_DEBUG=1 on @DenisVASI9's 8x A100 40GB rig: V4 Flash's dsv4_store_cache_rows emits SET_ROWS into the K-cache at layer 7 (cache buffer on CUDA1, V4 ops dispatched on CUDA1). Sched then dispatched SET_ROWS on CUDA0 because its newly-peer- copied src tensors were there. The kernel wrote to a CUDA1 pointer from a CUDA0 stream → illegal address. Without LAUNCH_BLOCKING the error surfaced later at the next sync API call (peer-copy or flash-attn), and we chased the wrong site for three rounds of diagnostics. Adding the dst-device check forces sched to only consider devices where dst lives. Sched will then schedule the necessary peer-copies for srcs. Co-Authored-By: Claude Opus 4.7 (1M context) --- ggml/src/ggml-cuda/ggml-cuda.cu | 15 +++++++++++++++ 1 file changed, 15 insertions(+) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index e80d44267877..7b88289beb55 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -5154,6 +5154,21 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g } } + // Some ops write through a pre-allocated destination buffer that wasn't + // sched-allocated for this dispatch (e.g. GGML_OP_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. + // 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. + if (op->buffer && ggml_backend_buft_is_cuda(op->buffer->buft)) { + ggml_backend_cuda_buffer_type_context * buft_ctx = (ggml_backend_cuda_buffer_type_context *)op->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)) { From 19c3f8fd9f2ca14314224ceb5ca89a1ca23b7aa4 Mon Sep 17 00:00:00 2001 From: Chris Chuter Date: Wed, 13 May 2026 20:08:08 -0500 Subject: [PATCH 5/7] cuda: walk view_src chain to find dst buffer in supports_op ggml_set_rows returns a view tensor (ggml_view_tensor(ctx, a)), so op->buffer is nullptr at supports_op query time. The dst-device guard added in ca8734ab6 was guarded by `if (op->buffer)` and silently skipped, leaving the original multi-GPU SET_ROWS crash in place. Walk view_src to the root tensor (matching the pattern used in ggml_backend_cuda_cpy_tensor_async) and check the real buffer's device. The walk is a loop rather than a one-hop because V4's dsv4_store_cache_rows builds a multi-level view chain (cache -> set_rows view of reshape of cont of source). Verified in https://github.com/cchuter/llama.cpp on fix/v4-cuda-multigpu-supports-op with V4 Q2_K-XL Flash on 2x RTX 6000 Ada. Co-Authored-By: Claude Opus 4.7 (1M context) --- ggml/src/ggml-cuda/ggml-cuda.cu | 27 ++++++++++++++++++--------- 1 file changed, 18 insertions(+), 9 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 7b88289beb55..6e072c4f1ea0 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -5154,18 +5154,27 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g } } - // Some ops write through a pre-allocated destination buffer that wasn't - // sched-allocated for this dispatch (e.g. GGML_OP_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. + // 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. - if (op->buffer && ggml_backend_buft_is_cuda(op->buffer->buft)) { - ggml_backend_cuda_buffer_type_context * buft_ctx = (ggml_backend_cuda_buffer_type_context *)op->buffer->buft->context; - if (buft_ctx->device != dev_ctx->device) { - return false; + { + 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; + } } } From a4bb644d6ec1f33db0de739e879e1a308220823a Mon Sep 17 00:00:00 2001 From: Chris Chuter Date: Wed, 13 May 2026 21:20:24 -0500 Subject: [PATCH 6/7] v4: replace ggml_set_rows with ggml_cpy-into-view in cache writes MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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 cache destination is anchored to a different device — the kernel then writes through a foreign-device pointer (dst->data), surfacing as cudaErrorIllegalAddress in ggml_cuda_compute_forward / SET_ROWS. ggml_cpy into a contiguous view of the destination routes by dst-buffer affinity and works in production multi-GPU today (cf. dsv4_store_state_segment at lines 375-389 of this file, which uses this exact pattern and has never been implicated in any multi-GPU crash log). 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. The cpy result tensor has op=GGML_OP_CPY and src[0]=src, so downstream consumers of the returned kv_state/score_state get a proper data dependency on the cpy through normal ggml graph traversal — no explicit ggml_build_forward_expand needed at site 2 (gf isn't plumbed there anyway). No ggml/src/ changes; this is a graph-construction fix only. Co-Authored-By: Claude Opus 4.7 (1M context) --- docs/plans/v4-multigpu-set-rows-fix-v2.md | 142 ++++++++++++++++++++++ src/models/deepseek4.cpp | 30 ++++- 2 files changed, 167 insertions(+), 5 deletions(-) create mode 100644 docs/plans/v4-multigpu-set-rows-fix-v2.md diff --git a/docs/plans/v4-multigpu-set-rows-fix-v2.md b/docs/plans/v4-multigpu-set-rows-fix-v2.md new file mode 100644 index 000000000000..a281f294498f --- /dev/null +++ b/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 +Co-Authored-By: Claude Opus 4.7 (1M context) +``` diff --git a/src/models/deepseek4.cpp b/src/models/deepseek4.cpp index c89171043b2c..b2ccaa881e65 100644 --- a/src/models/deepseek4.cpp +++ b/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,21 @@ 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, dst is on a different device); + // see dsv4_store_cache_rows for the same problem and fix. + // ggml_cpy returns a view-of-dst with op=GGML_OP_CPY and src[0]=src, so + // downstream consumers of kv_state/score_state get a proper data + // dependency on the cpy without needing ggml_build_forward_expand here. + auto cpy_into_row = [&](ggml_tensor * dst, ggml_tensor * row_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, row_src, view); + }; + 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) { From 72e6de0acf711bf114be1d02bfcbfb809e095625 Mon Sep 17 00:00:00 2001 From: Chris Chuter Date: Wed, 13 May 2026 21:26:11 -0500 Subject: [PATCH 7/7] =?UTF-8?q?v4:=20site=202=20=E2=80=94=20return=20full-?= =?UTF-8?q?shape=20view=20with=20cpy=20as=20src[0]=20dependency?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Codex code-review caught that returning the cpy result (which is a [width, 1] row-shape view) as kv_state breaks downstream readers: dsv4_view_cols at lines 869-894 slices the full [width, rows] state by columns AND rows. With the row-shape cpy result, offsets become relative to a single row rather than the full state origin. Mimic ggml_set_rows's internal construction: create a view_tensor of prev_kv_state (full shape, view_src = prev_kv_state) and manually set src[0] to the cpy node. Sched orders cpy before any consumer of this view, and the view reports prev_kv_state's full shape so dsv4_view_cols slicing produces the expected coordinate system. This is the same pattern ggml_set_rows uses internally (view-of-dst + op-tagged + src[] populated), just split across two ops we control at the graph-builder level so we get cpy's dst-affinity routing instead of set_rows's src-affinity routing. Co-Authored-By: Claude Opus 4.7 (1M context) --- src/models/deepseek4.cpp | 21 ++++++++++++++------- 1 file changed, 14 insertions(+), 7 deletions(-) diff --git a/src/models/deepseek4.cpp b/src/models/deepseek4.cpp index b2ccaa881e65..c45f9a07c6c4 100644 --- a/src/models/deepseek4.cpp +++ b/src/models/deepseek4.cpp @@ -870,17 +870,24 @@ static dsv4_decode_compressor dsv4_build_compressor_decode_projected( const bool should_compress = (pos + 1) % compress_ratio == 0; // Single-row write via cpy-into-view. ggml_set_rows would crash on - // multi-GPU (sched routes by src device, dst is on a different device); - // see dsv4_store_cache_rows for the same problem and fix. - // ggml_cpy returns a view-of-dst with op=GGML_OP_CPY and src[0]=src, so - // downstream consumers of kv_state/score_state get a proper data - // dependency on the cpy without needing ggml_build_forward_expand here. + // 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 * view = ggml_view_2d(ctx, dst, + ggml_tensor * row_view = ggml_view_2d(ctx, dst, dst->ne[0], 1, dst->nb[1], row * dst->nb[1]); - return ggml_cpy(ctx, row_src, view); + 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);