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/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index c79e87864f26..6e072c4f1ea0 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,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)) { diff --git a/src/models/deepseek4.cpp b/src/models/deepseek4.cpp index c89171043b2c..c45f9a07c6c4 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,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) {