Skip to content

Enable GLM 5.3 tensor parallelism on the ROCm backend - #1024

Open
davidcanar wants to merge 34 commits into
antirez:mainfrom
davidcanar:glm-tp-rocm
Open

Enable GLM 5.3 tensor parallelism on the ROCm backend#1024
davidcanar wants to merge 34 commits into
antirez:mainfrom
davidcanar:glm-tp-rocm

Conversation

@davidcanar

@davidcanar davidcanar commented Sep 10, 2026

Copy link
Copy Markdown

Enables GLM 5.3 (GLM_DSA family) on the two-node ROCm tensor-parallel backend from #813. Before this, GLM TP never armed on ROCm: the GLM graph's TP initialization is guarded #if defined(__APPLE__), so tp_world stayed 0 and every rank silently computed the full expert set against its half-resident weights — word-salad output, zero gate exchanges, no error reported. With this PR, GLM-5.3-Flash Q4_K runs across two 128 GB Strix Halo boxes at 177.76 GiB model size (unrunnable single-box), including 47k-token contexts and ds4-server tool calls.

Opt-in, same as the backend itself: set DS4_ROCM_GLM_TP=1 on both ranks.

Branch layout (stacked)

This branch is the integration lineage of rocm-tensor-parallel (#813); it also merges current upstream main into it. Commits new to this PR:

  • ea7057d fix up the upstream-main merge for the ROCm TP build
  • 4dd3f4d enable GLM 5.3 tensor parallelism on the ROCm backend (the core change)
  • 848fb67 ROCm streaming: lock static weight spans and keep them after upload
  • 1a646f6 ROCm GLM TP: avoid the faulting Q4_K sorted down kernels
  • 682cdb4 restore the CPU build after the upstream merge

Added since the description was first written:

  • 0e02528, 4402b61, cb7cb70 verbs transport fixes from this lineage: reset the RDMA bulk window on drain timeout and name DS4_TP_GATE_TIMEOUT_MS, negotiate RC queue pairs on devices that support them, report the real ibv_post_* return codes and log the QP caps
  • bd177c9 + da15223 merge upstream main (DeepSeek v4.1 Flash, bd66c40) into the branch, plus the merge-resolution fixups
  • 5272bda, 5b31f6d, b6f965b, ebfe289 ROCm decode performance work, described in its own section below
  • 7598395 OdinLink loss recovery, the retransmission commit reviewed in Add OdinLink (odl_tb5) TP transport: --transport odl #1018, described in the transport reliability section below

Also present in the lineage: the OdinLink transport commit (reviewed separately in #1018) and the pipeline speculative work (reviewed separately in #835), plus a few small GLM ROCm fixes from this lineage (matvec 128-bit vectorization, sorted-path env kill switches, pipeline slice boundary fix). Suggested review order: #813#835#1018 → this PR.

Changes

  • Extend the GLM graph TP init guard to DS4_ROCM_BUILD (the ownership-aware Q4_K routed kernels and the slab gate machinery are shared between the Metal and ROCm backends).
  • Gate the old refusal behind DS4_ROCM_GLM_TP=1 (experimental opt-in).
  • Exempt GLM from the Q8_0 output-head requirement: GLM keeps its output head replicated/unsplit under TP (DeepSeek splits it), so the ownership check does not apply.
  • Fix an operator-precedence slip in the ROCm GLM streaming gate so single-box --ssd-streaming runs work again.
  • Streaming hardening: extend the SSD static-span mlock to ROCm builds and make the post-upload DONTNEED discards skip locked spans — on Linux the GPU cannot fault file-backed pages back in, so an evicted locked-weights page is a hard fault.
  • Q4_K sorted-pairs down GEMM: route the sorted/tile down-projection through the proven per-pair kernel and keep the moe_sum epilogue for q4k (details under notable failures).
  • Decode performance on ROCm (four commits, see the section below): pre-quantized lane-strided k-slice Q8_0 matvec, gate arrive/wait split with the shared expert overlapping the FFN exchange, vectorized small-output BF16 row matvec, and no periodic decode-stream flush on ROCm.

Test evidence (per CONTRIBUTING.md)

Machine: 2x AMD Ryzen AI Max+ 395 (Radeon 8060S, gfx1151, 128 GB unified), Fedora 44, ROCm, Thunderbolt 5 direct link. Model: GLM-5.3-Flash Q4_K (glm53-q4 GGUF, 177.76 GiB, sha256 c7a0d950...). Build on both boxes:

make strix-halo -j32 CFLAGS="-O3 -ffast-math -g -march=native -Wall -Wextra -std=c99 -fPIC -D_GNU_SOURCE -fno-finite-math-only" \
  ROCM_LDLIBS="-lm -pthread -lamdhip64 -lhipblas -lhipblaslt -lrocblas"
make cpu   # also verified clean (this PR restores it)

Launch (worker first, then coordinator; transport odl from #1018, rdma/tcp also work):

# box2
env DS4_ROCM_GLM_TP=1 ~/ds4/ds4 -m ~/models/GLM-5.3-Flash-Q4_K.gguf \
  --tensor-parallel --role worker --coordinator 10.0.2.1 7300 --transport odl --ctx 8192
# box1
env DS4_ROCM_GLM_TP=1 ~/ds4/ds4 -m ~/models/GLM-5.3-Flash-Q4_K.gguf \
  --tensor-parallel --role coordinator --listen 10.0.2.1 7300 --transport odl --ctx 8192 \
  -p "<prompt>"            # or --prompt-file for the 47k run

Correctness and speed, all with no workarounds or extra env beyond the opt-in above (numbers as of 682cdb4; the performance section below has the current ones):

  • Short exact-prefill prompts (<= 64 tokens): correct output ("two plus two" -> "Four"), 8.7 t/s generation.
  • 227-token prompt (sorted-pairs batch path): prefill 53.7 t/s, generation 8.65 t/s, accurate two-sentence summary.
  • 47,288-token prompt at ctx 49152: prefill 39.5 t/s (~20 min), generation 6.98 t/s, accurate summary of the fed source file, zero faults, zero lockstep divergence across both ranks.
  • ds4-server with --batched-session 4: two-turn weather-tool conversation (163-token prompt, correct get_weather call, KV cache reuse on the tool-result turn); the same conversation while a concurrent 450-token generation runs; and a 6,181-token multi-chunk tool prompt whose tool-result turn reuses 6,192 cached tokens. All correct, zero faults on both ranks.
  • DS4_TP_TIMING=1: row-gate exchange 172 us avg, 10k+ gates per run.

Decode performance (added after the initial submission)

Profiled a steady-state decode token on both ranks with a roctracer kernel timeline (official Q4_K GGUF, 45-token prompt, greedy, ctx 4096). Decode is neither compute- nor network-bound; per token (114 ms) the coordinator GPU spent:

  • 56.7 ms in the BF16 matvec streaming the KDA attention projections (q/k/v/output, 4 x 67 MB per KDA layer; BF16 in the official GGUF and replicated on both ranks) at ~160 GB/s, against ~240 GB/s achievable on gfx1151 from either device or host-mapped memory;
  • 13 ms in the routed experts (Q4_K kernels at 175-210 GB/s; the 50/50 expert ownership puts E[max] ≈ 5.1 of the 8 selected experts on the slower rank per layer);
  • 8.7 ms parked in the gate spin kernels: 53 gates per token (42 FFN + 11 DSA attention), while the odl exchange itself is 80-175 us per gate;
  • 4.9 ms in the k-sliced Q8_0 attention-output kernel, which ran at 80 GB/s;
  • ~17 ms in ~1,500 small kernels and launch gaps.

(DS4_TP_TIMING's arrival is the local GPU time of the layer up to the partial, not the wait for the peer.) Four commits take the code-side part of that:

  • 5272bda k-sliced Q8_0 matvec on pre-quantized activations: the owned activation slice is quantized to Q8 like the dense decode path and the lanes stride over the row's 34-byte blocks (one wave instruction per 32 blocks, dp4a). 444 -> 176 us per DSA layer (~200 GB/s), about 3 ms per token. Same entry point as the DeepSeek TP attention-output / shared-expert expand paths, which were not exercised here.
  • 5b31f6d ds4_gpu_tp_gate_encode split into ds4_gpu_tp_gate_arrive (arrival event + exchange request) and ds4_gpu_tp_gate_wait (release spin kernel). The GLM sparse FFN encodes the replicated shared expert between the two on ROCm, so the GPU computes it while the CPUs swap the routed halves. The exchanged partial is unchanged, so both ranks still rebuild the routed output bit-identically; Metal keeps the single-call gate. About 2.5 ms per token.
  • b6f965b the small-output BF16 row matvec (hyper-connection mixers, KDA f_a/g_a/beta, indexer projections; ~214 launches per token) issues eight independent 128-bit weight loads per thread instead of one dependent 2-byte load per iteration: 16 us -> a few us per launch.
  • ebfe289 ROCm no longer flushes the decode stream every 4 layers (indexed attention) / 32 layers. On the eager in-order stream that flush is a full device sync that only drains the CPU run-ahead; it cost ~2 ms per token under TP. DS4_GLM_DECODE_FLUSH_INTERVAL still forces a cadence; Metal and the CPU build are unchanged. Exercised under TP only, not on single-box resident decode.

Result on the official GGUF, same 45-token prompt, greedy, 96 tokens, 2x gfx1151 over odl: 8.93 -> 9.65 t/s (per-gate arrival 1.81 -> 1.69 ms). 323-token prompt: 57 t/s prefill, 9.62 t/s decode, correct two-sentence summary. Zero faults and zero desyncs across all runs. Greedy output matched the previous binary for ~50 tokens and then diverged at a near-tie token: the k-slice path now quantizes its activation slice to Q8 like every other Q8 decode matmul in the ROCm backend, so this is the expected int8-activation drift rather than a lockstep problem.

For the remaining gap, out of scope here: with the KDA projections and the output head requantized to Q8_0 (a data-only change to the GGUF, output equivalent on the prompts tested) the same code reaches 15.2 t/s, i.e. the BF16 KDA tensors of the glm53-q4 recipe are now the dominant cost of GLM TP decode on this hardware. Code-side, the next levers are the efficiency of the big BF16 matvec (~160 vs ~230 GB/s for the Q8 kernels), head-split TP for the KDA layers, and the routed-expert ownership imbalance.

Transport reliability (added, 7598395)

The OdinLink transport of #1018 lost whole 1 MiB gate messages whenever a Thunderbolt frame failed CRC (odl_tb5: RX CRC error ... dropping message in dmesg), which parked both ranks on their spin kernels for the 300 s bulk timeout and left the TP session dead; it happened 4096 tokens into an 11k-token ds4-server prefill on this pair. 7598395 is the loss-recovery commit reviewed in #1018 (per-message headers, receiver-driven retransmission from retained frames, idle ranks keep serving requests from their control-channel waits), cherry-picked here with this branch's tp_read_full_deadline gate-header reads polling the same way.

Validated here with DS4_TP_ODL_DROP_TEST: 96-token greedy decode 15.35 t/s without and 15.23 t/s with every 400th gate message dropped (12 per rank recovered, output byte-identical); the 323-token prompt prefilled at 69.1 vs 67.5 t/s with every 50th message dropped (75 per rank recovered, output identical). The server script now defaults to odl again.

Speculative decoding (MTP)

Validated working with two-rank TP (CLI and ds4-server, contexts from 226 tokens to 47k, including tool-call conversations) with no code changes beyond this PR: the GLM graph TP arming above is what lets the MTP path engage at all (the worker mirrors each cycle through the GLM_MTP command). Requirements are runtime settings: --mtp on both ranks, DS4_GLM_MTP_DRAFTS=1 (the fork deliberately keeps GLM drafting off by default on this hardware; TP cycles cap at 2 rows, so 1 draft is the maximum over TP), and --batched-session 1 on ds4-server (its spec gate refuses to run with more than one resident slot). Cycles are observable via --mtp-timing (CLI) or DS4_MTP_SPEC_LOG=1 (server).

Speed note, for honesty: on the 2x gfx1151 cluster the 2-row verify span costs roughly double a plain decode (~215 ms at small context, ~270 ms at 47k) while observed draft agreement stays below the ~50% break-even, so MTP is a net slowdown at every context size tested (6.4 vs 8.7 t/s at 226 tokens; 4.7 vs 7.0 t/s at 47k; 4.2-5.5 t/s on the server). Outputs remain correct through accepts, rejects and rollbacks. Default-off is the right setting for this hardware.

Notable failures root-caused on the way

  1. Q4_K sorted-pairs down-projection fault. Any batch prefill >= 32 tokens aborted both ranks with HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION (surfacing through ds4-server as "GLM TP batch gate failed"). Stage-skipping isolated it to the sorted/tile down kernels (moe_down_q4K_expert_tile{4,8}_row32, moe_down_q4K_sorted_qwarp32): the sort kernels, the gate/up tile kernels, and every input tensor validate clean, and the per-pair decode down kernel reads the same weight images without faulting. 1a646f6 routes the sorted-path down GEMM through that per-pair kernel and excludes q4k from use_atomic_down so the moe_sum epilogue reduces its down->ptr writes (with the epilogue skipped, batches >= 128 tokens silently dropped the MoE output). Why the tile down kernels fault is still open — all bounds check out statically and no VM fault is even logged in dmesg; they remain unused for q4k until that is understood. Possibly related to GLM-5.3-Flash Metal: prefill fails for prompts > 4096 tokens (compact indexed path) — HTTP 500 'metal GLM-5.3 prefill failed at token 0' #890 / GLM-5.2 Q4: prefill fails at ≥64K context on M3 Ultra 512GB (KV payload staging failed) #998 / GLM-5.2 (glm-dsa) prefill fails on ROCm backend initialized on AMD Radeon Graphics (sm_115): "rocm GLM prefill failed" with --ssd-streaming #608 / GLM-5.2 (glm-dsa) prefill fails on CUDA / DGX Spark GB10 (sm_121): "cuda GLM prefill failed" with --ssd-streaming #595.
  2. Static-weight pageout under streaming — fixed by the mlock extension + discard guards in 848fb67.

Not run: the full make test suite (needs the CPU test binary model fixtures on hand); behavior was validated end-to-end on real hardware as above on both ranks.

🤖 Generated with Claude Code

davidcanar and others added 24 commits August 16, 2026 13:21
ds4's two-machine tensor parallelism (--tensor-parallel with
--role coordinator/worker) previously required the Metal backend and
Apple's Thunderbolt RDMA stack.  This change ports it to the ROCm
(HIP) backend so pairs of AMD machines can split a DeepSeek model
50/50 and exchange gate partials over any verbs device.

- ds4_server.c: parse TP options, validate them, bind the leader and
  run the mirrored worker loop (previously ds4 CLI only).
- ds4.c: enable ds4_engine_tp_bind on the ROCm build, validate
  Q4_K/Q8_0 tensor layouts (the ownership-aware kernels cover those),
  and allocate the transport slab as GPU/CPU-shared memory.
- rocm/ds4_rocm_tp.cuh (new): TP gate machinery for the eager default
  stream — arrival/release flag words in the slab, a service thread,
  spin-wait kernels, and big-gate staging through pinned memory.
  Plus the TP compute kernels: k-sliced Q8 matvec, attention output
  projection for owned head groups, fused HC-expand combine, and the
  prefill row-range attention variants.
- rocm/ds4_rocm_moe*: expert ownership (tp_rank/tp_world/rebased
  weights) in the Q4_K kernels and sort filtering so each rank
  computes only its resident expert half, with the shared-expert
  partial folded into the down projection.
- ds4_tp.c: Linux verbs support via dlopen (Apple's librdma.dylib on
  macOS, libibverbs on Linux), a dedicated bulk QP for prefill row
  swaps with per-round recv barriers, a decode lookahead window armed
  at bring-up, and error-code logging instead of strerror(errno)
  after verbs calls (the inline API returns the code directly).

Tested on two Strix Halo boxes (gfx1151, ROCm 7.14) linked by the
thunderbolt-ibverbs module: 50/50 expert split of DeepSeek-V4-Flash,
~14 tok/s decode over RDMA, lockstep hidden-state hashes passing,
long prefills and concurrent sessions stable.
Found while running a resumed session with the disk KV cache: the
leader's restore short-circuits its prefill without mirroring the
skipped prefix to the worker, so the ranks ran different prefill
graphs and the gate sequences diverged.  The divergence surfaced as
a 300s bulk-round timeout, and the abandoned round left stale
receives posted on the bulk QP, poisoning every later request.

- ds4_server.c: skip kv-disk restores on the TP leader (the worker
  mirrors only the tokens the leader actually evaluates; the leader
  now re-prefills the full transcript and mirrors it exactly).
- ds4_tp.c: carry the payload byte count in the gate header so a
  size mismatch aborts at the header exchange instead of timing out
  mid-round, and rebuild the bulk QP in lockstep (data-socket
  rendezvous, QPN/PSN re-exchange, RTR/RTS) after any bulk failure
  so stale receives can never misalign later gates.
- ds4.h/ds4.c: export ds4_session_tp_leader for the server layer.

Retested: session resume (the failing scenario), 12.5k-token tools
prompt, concurrent sessions, multi-turn — all clean over RDMA.
The eager stream already orders everything after the partial-producing
kernels, so the GPU-side arrival signal can be a recorded stream event
waited by the transport thread — eliminating one tiny kernel launch per
gate (86 per decode token on DS4).  Events share a single monotonic
counter across row/batch/big gates, and the ring is sized to the
request queue so a slot is always consumed before it wraps.

Measured on the two-box ROCm pair (DeepSeek-V4-Flash, RDMA transport,
ctx 393216): decode 13.8-14.0 -> 14.3-14.4 tok/s (~+4%).
Each rank scores only its parity rows of the compressed index, packs its
local top-k, and exchanges (id, score) candidate lists over a dedicated
third UC queue pair; both ranks then merge into the identical global
selection on the GPU. Payloads are small fixed slots staged through new
split slab regions, so the split channel never misorders the row or bulk
gates. Opt-in via DS4_TP_KV_SPLIT=1, set on both ranks.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Cancelling a request mid-prefill deadlocked the pair: the leader mirrors
the whole prompt in one sync frame, and its cooperative-cancel check
stopped only the leader at a chunk boundary while the worker sailed into
the next chunk and blocked forever inside a gate exchange the leader no
longer serves (leader stuck in the ack read, worker parked in the release
spin kernel).

Make every cancellation check inside a mirrored sync a lockstep barrier:
the leader publishes go/stop on the otherwise idle control socket and the
worker blocks on it and adopts the verdict, so both ranks take the same
INTERRUPTED unwind at the same chunk boundary and keep identical live
prefixes (still valid for KV reuse). The worker acks an interrupted sync
and keeps serving instead of exiting. Hard mid-sync errors poison the
socket pair so blocked peers unblock with an error instead of hanging,
and gate-header reads honor the pair timeout. Protocol version bumped to
8 for the new frame.

The barrier lives in backend-agnostic code, so this also fixes the same
latent deadlock in Metal TP. Verified on a Strix Halo pair: a 66k-token
prefill aborted mid-chunk left both ranks healthy and the next request
was answered ~20s later.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Measured on the Strix Halo pair (greedy 400-token code generation,
DS4_DSPARK_STATS): verify drops 333 -> 279 ms per <=6-row block and
DSpark TP decode rises 9.9 -> 10.8 t/s, with identical acceptance
(80%, same histograms) and valid output.

- Batch the TP verify routed MoE: one grouped dispatch into the slab
  batch rows instead of one single-token dispatch per draft row (the
  ROCm launcher already restricts experts to the owned half).  ROCm
  only; Metal keeps the per-row path.
- Size the shared-x batch matmul token tile to the batch: a 32-token
  tile reserves 64 KiB LDS (one workgroup per CU) while a verify block
  fills 6 lanes.  Applies to the generic Q8 matmul and the attention
  output A projection; >=32-token prefill shapes are unchanged.
- One-round-trip batch gates: post the bulk recvs before the header
  handshake so the header read doubles as the posted-recv barrier,
  replacing the header + one-byte barrier pair.
- DS4_TP_TIMING now buckets service-thread gate timing by kind
  (row/batch/split/big).

Profiling notes: per-layer verify compute is ~4.6 ms (dense stages and
owned-expert streaming), batch-gate wire is ~1.45 ms of which ~1.2 ms
is the 16 KiB UC per-message cost (6 messages per 6-row gate); routed
expert bytes scale with rows, so speculation cannot amortize them.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
The worker owns the DSpark capture target layers (40-42) and the output
head in the pipeline split, so it runs the whole propose pipeline (stage
chain, confidence gate, markov bias) and returns drafts through the
existing F_OUTPUT_DRAFTS route; the coordinator drives the same
verify/rollback cycle as legacy MTP.  Specifically:

- ds4_engine_mtp_draft_tokens: drop the distributed-role exclusion, so
  the coordinator cycle, the server gate, and the worker draft sizing
  all see the DSpark block under the pipeline split.
- ds4_session_dist_dspark_draft + ds4_session_dist_support_draft: the
  worker-side drafter and the support-kind dispatch.  Zero drafts on a
  scheduler or confidence skip is a normal outcome and arms the same
  no-draft backoff as the local path.
- Feed the DSpark capture ring from the layer-slice paths: per-layer
  capture in the single-token decode loop, and in multi-token spans a
  verified-suffix continuation when the span extends a captured
  checkpoint (the verify span) with a from-scratch batch window
  otherwise (prefill chunks, rollback re-evals).  No-ops on slices
  without capture targets.
- Batched output head for multi-row slice spans (one pass over the head
  weights into spec_logits instead of one full stream per draft row).
- Fix the LOGITS_DRAFTS payload for zero-draft cycles: always write the
  draft count (previously an uninitialized count under a drafts_max
  sized payload made the coordinator reject the result and kill the
  request; latent in the MTP path, hit constantly by DSpark skips).
- Keep the exact two-row verifier for legacy MTP only (it skips capture,
  which would starve the next DSpark propose).
- DS4_DIST_SPEC_MIN_DRAFT (default 3): shallow blocks are not worth a
  verify span round trip; emit the base token and keep its logits.
- Log the failing span error under DS4_MTP_SPEC_LOG before rebuilds.

Measured on the Strix Halo pair (greedy code + prose, pipeline split
0:22/23:output, support gguf on both sides): drafting and verification
work end to end with valid output; 4-5 row verify spans cost ~300 ms
(two serialized full-expert slices), acceptance is healthy (13/25
verifies accepted 4, 6/25 accepted 5 at DS4_DIST_SPEC_MIN_DRAFT=4), and
throughput is ~11.5 t/s vs ~13 t/s plain pipeline.  The remaining gap
is the fixed span cost plus rollback re-evals on partial accepts;
prefix-snapshot commits over the dist route are the next lever.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Partial accepts previously paid a full rollback re-eval span (frontier
restore plus re-decoding the accepted prefix through both slices, about
as expensive as the verify itself).  Port the single-node prefix-snapshot
mechanism to the distributed route instead:

- Layer-slice spans capture per-prefix compressor/indexer state with the
  same gating as the single-node verifier (1 < n <= DS4_SPEC_PREFIX_SLOTS
  + 1); a spec_prefix_rows counter tracks freshness and is zeroed by any
  other eval on the session.
- ds4_session_dist_spec_commit_prefix rewinds the timeline, restores the
  captured state for the accepted row (owned layers only on slices), and
  re-appends the accepted tokens -- the same recipe as the single-node
  direct-partial commit; rejected append-only rows stay invisible.
- New DS4_DIST_WORK_F_SPEC_COMMIT work item: tokens carry the accepted
  ids, no hidden payload, the reply is a bare ack carrying the new token
  hash (the incoming prefix-hash check is skipped since a commit rewinds
  the timeline by design; the next span check covers divergence).
- The coordinator commits locally, sends the commit item, and uses the
  accepted row logits it already holds from the verify; on any failure it
  falls back to the old rollback re-eval (the local rewind is overridden
  by the frontier restore there).

Measured on the Strix Halo pair (greedy code + prose, default
DS4_DIST_SPEC_MIN_DRAFT=3): all partial accepts commit without fallback
or hash errors, output stays valid, and pipeline DSpark decode improves
to ~11.7-11.9 t/s (from ~10.8-11.3; plain pipeline is ~13.0).  The
remaining gap is the ~268 ms fixed verify span (two serialized
full-expert slices); per-row streamed verify handoffs are the next lever.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Streaming the verify span row-by-row would lose to physics here: the
span cost is ~265 ms of weight streaming shared across rows plus ~8 ms
per row, so sub-batching re-streams weights faster than overlap can
recover.  Attack the serialization from the other side instead: the
worker already computes every verified row's logits, so it now proposes
continuation drafts from the last row (its greedy token is the next base
token whenever the span is fully accepted) and appends them to the
verify reply.  On a full accept the next cycle skips the separate base
decode entirely and verifies [next_token, drafts...] in one fused span:
one wire round trip and one span instead of two of each, chaining
span-to-span while acceptance holds.

- Verify replies (LOGITS_NROWS) may carry trailing continuation drafts;
  OUTPUT_DRAFTS combined with OUTPUT_ALL_LOGITS proposes from the last
  row at pos0 + n_tokens.
- The coordinator caches the continuation block (token- and
  position-keyed, consumed or invalidated at every cycle entry) and runs
  the fused span when it matches; every other outcome falls back to the
  classic two-item cycle.
- Fused spans are capped at 4 drafts so the 5-row span keeps prefix
  capture and partial accepts stay commit-only (prefix lengths reach
  1 + accept_n <= DS4_SPEC_PREFIX_SLOTS).
- decode2 stays legacy-MTP only and never combines with drafts.

Measured (greedy code + prose, default --dspark-confidence 0.7 and
DS4_DIST_SPEC_MIN_DRAFT=3): 37 fused spans with 27 accepting the full
4-draft cap, zero fallbacks or hash errors, valid output.  Code
generation now beats plain pipeline decode (36.6 s vs 40.3 s wall for
400 tokens, ~13.4 vs ~13.0 t/s); prose remains slightly under (~11.6).
A looser confidence gate (0.45) was measured strictly worse on both.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…he ROCm verify floor

- DS4_SPEC_PREFIX_SLOTS 4 -> 5 so fused verify spans carry a whole DSpark
  block (drafted=5 chains confirmed end to end: 18 of 28 fused spans at 5
  drafts, 16 full accepts, zero fallbacks; prose ~3% faster, codegen
  within run variance).
- Add a barrier-free direct Q8_0 batch matmul kernel for 2..8-row
  speculative batches (bit-identical reduction order; the shared-x tile
  kernel keeps larger batches). On the resident-weight gfx1151 config the
  dense stages ride hipBLASLt f16 GEMMs instead, so this engages only
  where the Q8 batch path runs.
- Wire the DSpark selected-expert profiler into the distributed verify
  span (DS4_DSPARK_VERIFY_SELECTED_PROFILE): 6-row spans hit 14-28 of 256
  unique experts (5.5-11% of expert bytes) on the worker slice.
- Add measured diagnostic rollbacks for the verify span:
  DS4_ROCM_ATTN_OUTPUT_CUBLAS=0 (-14% on the output_proj stage, e2e
  neutral), DS4_ROCM_SHARED_DOWN_CUBLAS=0 (no effect; that stage does not
  route through the gated entry), DS4_ROCM_Q4K_SORTED_MIN (sorted and
  sorted+tiled MoE both measured 4.0 ms vs 2.4 ms per layer at 6 rows -
  the per-pair path already runs at DRAM bandwidth, so the 32-token
  threshold stands). Expert tiles now decouple from pair sorting below
  32 tokens so the two effects can be measured separately.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…e DSpark

Conflicts resolved:
- ds4_rocm.cu: keep TP host_alias + ds4_rocm_tp.cuh, add upstream
  rocblas include, MMQ decls, MXFP4 tables, GLM53 vision include
- ds4_tp.h/c: renumber SYNC_GO=18/SYNC_MULTIMODAL=19, keep kv-split
  helpers + upstream gate-mask helpers, take upstream chained RDMA
  send/recv with rc logging, drop orphaned recv-window drain fn,
  wrap multimodal sync dispatch in TP lockstep
- ds4_server.c: mtp spec gate allows single-slot batched (ours) and
  non-greedy temperature (upstream)
- rocm kernels: upstream antirez#887 gfx1151 tuning + MXFP4 path, keep TP
  expert rebasing and q_row0 attention row ranges, keep both new
  q8 verify kernels, keep DS4_ROCM_Q4K_SORTED_MIN env override
GLM 5.3 layer-slice boundaries carry the hyper-connection-expanded
state (DS4_N_HC * DS4_N_EMBD floats per token) that the graph reads
and writes, but the distributed path sized the boundary with plain
DS4_N_EMBD. The hc-expanded readback then overran the coordinator's
hidden-state buffer and aborted with heap corruption on the first
pipeline request.

Size every boundary through ds4_engine_hidden_f32_values(), which now
returns the hc-expanded size for GLM 5.3 and keeps the plain hidden
size for GLM 5.2, and step the layer-slice chunk window with the same
helper instead of a hard-coded DS4_N_EMBD.
DS4_ROCM_DISABLE_SORTED_PAIRS=1 forces the per-pair routed-expert path
and DS4_ROCM_DISABLE_EXPERT_TILES=1 disables only the expert-tile
variant. These are needed as a workaround for GLM 5.3 Flash on Strix
Halo: the sorted-pair batch path faults with an HSA aperture violation
in the first sparse layer whenever more than ~24 tokens are batched,
while the per-pair path runs clean. Pipeline prefill then runs at
~26 tok/s for a 44k-token prompt.

Suspected upstream kernel bug in the GLM resident batch MoE on ROCm;
kept behind an env flag so upstream defaults are untouched.
The decode-shape BF16 matvec (KDA q/k/v/output projections, DSA o_proj,
tied embedding head) loaded two bytes per lane per iteration, moving only
64 bytes per warp instruction. Load eight BF16 weights per lane with one
uint4 and pair them with two float4 activation loads; keep the scalar
path for unaligned dimensions. Pipeline decode goes 5.0 -> 6.3 tok/s on
GLM-5.3-Flash Q4_K.

Also carries an env-gated KDA kernel timing probe (DS4_GLM53DBG).
Implements speculative decoding for GLM 5.3 Flash across the two-box
pipeline, reusing the DeepSeek dist MTP machinery plus GLM-specific
state handling:

- ds4_session_glm_dist_draft(): chained nextn drafting on the worker
  (the slice that owns blk.45 and the output head). glm_graph_mtp_step
  gains a chain_hidden override so the nextn block can run multiple
  steps, staging each step's hidden through g->hc_output.
- Mid-span KDA snapshots: spec spans split each KDA layer's recurrence
  into base-row + speculative-rows calls and snapshot (conv, recurrent)
  state right after the committed base row, giving a cheap post-base
  rollback point on both machines (env kill switch
  DS4_GLM_SPEC_SPLIT_DISABLE).
- Worker slices map token_embd + the nextn block when GLM MTP is on
  (weights_bind include_nextn, span-builder g_nextn_slice_include).
- ds4_session_eval_layer_slice_logits_all() GLM path: run the shared
  head per row from the batch hc states with inter-row syncs.
- Coordinator: fused verify path for GLM (min_fused=1), stash of
  continuation drafts (drafts[1..], since drafts[0] is the pending
  token), all-or-nothing acceptance with post-base KDA rollback and
  a cheap single-row SPEC_ROLLBACK re-eval on fused-span failure
  (no full-timeline rebuild).
- ds4_session_dist_frontier_snapshot/restore GLM branches +
  ds4_session_glm_dist_spec_span_begin() arming/implicit rollback.

Measured on the 2x Strix Halo ROCm cluster: the machinery is correct
(verify targets match the base model, cheap rollback works, no
rebuilds), but with ~40-60% first-draft agreement, ~500ms fused spans,
and chain decay beyond the first step, speculation is currently
net-negative vs 6.3 tok/s ordinary decode. Disabled by default:
DS4_GLM_MTP_DRAFTS (1..5) enables it explicitly.

Also carries env-gated diagnostics (DS4_MTP_SPEC_LOG, DS4_GLM53DBG)
and the KDA kernel timing probe in ds4_rocm_glm.cuh.
- ds4_server.c: ds4_engine_tp_gate_schedule call takes the gate-slot
  mask argument added upstream (required to link; the merge commit
  carried the pre-fix version).
- ds4_rocm_q8.cuh: indentation of the merged exact8 kernel tail.
Native OdinLink stream API over Thunderbolt, dlopen-loaded like the verbs
stack (no link-time dependency). Each rank receives on its own stream,
ODL_READY barrier in attach, non-blocking recv poll with control-socket
peer-death checks and the gate deadline, ODL branches in the gate/batch/
big/kv-split exchanges, bulk path with a 4 MiB send-ahead window.

Transport A/B on 2x Strix Halo (tp_odl_test.c): 24KB gate vectors
odl p50 30us / p99 52-63us vs tcp p50 32-262us / p99 ~320us; 50k-gate
soak clean; bulk ~1.1 GB/s per direction (stream-API kernel ceiling).

Also: exempt TP from the full-model GLM ROCm ssd-streaming gate (TP
holds half the routed experts per rank), and label odl in bind logs.
# Conflicts:
#	ds4.c
#	ds4_gpu.h
#	ds4_server.c
#	ds4_tp.c
#	ds4_tp.h
- g_tp_block_ctx guard now includes DS4_ROCM_BUILD (upstream was Apple-only)
- restore the fork KDA chained-MTP state backup alongside upstream spec
  backup (mtp_kda_backup field, kda_state_bytes, alloc in mtp_ensure)
- ds4_gpu_tp_init: accept upstream out_off/vec_bytes params in the ROCm TP
  implementation (unused there; the slab layout call carries the offsets)
- ds4-server: keep fork worker-mode support, declare block_start for
  upstream speculative rewind boundary
- drop stray static on ds4_session_tp_leader prototype
- frame enum: SYNC_GO=22, ODL_READY=23 past upstream GLM_MTP=21;
  SPLIT_WR_TAG moved off the warm-up tag bit
- extend the GLM graph TP initialization guard to DS4_ROCM_BUILD: the
  50/50 expert split, slab views and gate schedule were never armed on
  ROCm, so every rank silently computed the full expert set against its
  half-resident weights (word-salad output, zero gate exchanges)
- gate the GLM refusal behind DS4_ROCM_GLM_TP=1 (opt-in, experimental)
- exempt GLM from the Q8_0 output-head requirement: GLM keeps its head
  replicated and unsplit under TP
- fix the upstream-merge operator precedence in the ROCm streaming gate
  so single-box runs with --ssd-streaming work again

Validated 2x Strix Halo (gfx1151) over OdinLink, GLM-5.3-Flash Q4_K:
exact prefill (<=64 tokens) + decode produce correct output at
8.7 tok/s generation; prefill 25.8 t/s odl vs 18.4 t/s tcp.
Batch prefill (prompts >64 tokens, incl. server chat) faults in the GPU
on GLM+ROCm independent of TP (single-box repro) - documented in notes.
Extend the SSD static-span mlock to DS4_ROCM_BUILD (it was Apple-only):
on Linux the GPU cannot fault in file-backed pages, so any static
weight left pageable becomes a memory-access fault once the page cache
drops it under streaming churn. Register every locked span with the GPU
runtime and make cuda_model_drop_file_pages/cuda_model_discard_source_pages
skip them: fadvise/madvise DONTNEED evicts even mlocked file pages, and
the post-upload discard was dropping the very spans just locked.

Evidence (2x Strix Halo, GLM-5.3-Flash Q4_K, --ssd-streaming): with the
lock, statics stay resident (14.49 GiB locked, 0 pageable); batch-prefill
no longer faults on static weights. Note: GLM Q4_K batch prefill still
needs DS4_ROCM_DISABLE_SORTED_PAIRS=1 (separate sorted-pairs kernel bug,
traced to a GPU access in the model mmap past the mapped tables).
The Q4_K sorted-pairs prefill path (engages at >=32 tokens) reliably kills
both ranks with HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION inside the
sorted/tile down-projection kernels (moe_down_q4K_expert_tile{4,8}_row32
and moe_down_q4K_sorted_qwarp32).  The sort machinery, the gate/up tile
kernels, and all input tensors validate clean; reads of the same weight
images through the per-pair decode kernel work.  Route the sorted-path
down GEMM through moe_down_q4K_qwarp32_kernel instead and exclude q4k
from use_atomic_down so the moe_sum epilogue reduces the per-pair
down->ptr writes (the skipped epilogue silently dropped the MoE output
for batches >=128 tokens).

Validated on 2x Strix Halo (gfx1151) over OdinLink, GLM-5.3-Flash Q4_K,
no workarounds: 62- and 227-token prompts coherent; 47,288-token prefill
39.5 t/s, generation 6.98 t/s (36.6/6.95 previously with sorted pairs
disabled); ds4-server tool calls (sequential, concurrent with another
generation, and with a 6.2k-token multi-chunk prompt) clean with KV
cache reuse.
Two fork-side entry points broke make cpu after the upstream-main
merge: ds4_session_eval_layer_slice_span() used GPU-only types and
ds4_gpu_synchronize() without a DS4_NO_GPU stub, and ds4_distributed.c
still references ds4_session_glm_dist_spec_span_begin(), whose
definition lives inside a #ifndef DS4_NO_GPU block.  Give the span
eval the same error-stub its sibling entry point uses and add a CPU
stub for the spec span begin.  make cpu links ds4, ds4-eval and
ds4-agent again; GPU builds are unchanged (guards only).
@MerlijnWajer

MerlijnWajer commented Sep 11, 2026

Copy link
Copy Markdown

Thanks for working on this, I'm excited to try to get this working locally (using docker.io/kyuz0/strix-halo-ds4-toolbox:rocm-10.0 as docker base). I gave this a try with GLM 5.3 Flash (GLM-5.3-Flash-Q4_K.gguf) and I get this error on the coordinator:

0911 11:07:21 ds4-server: KV disk cache /tmp/ds4-kv (budget=100000 MiB, cross-quant=accept, min=512, cold_max=100000, continued=10000, trim=32, align=2048, hit_half_life=21600s)
0911 11:07:21 ds4-server: listening on http://0.0.0.0:8000
0911 11:07:53 ds4-server: chat ctx=0..1322:1322 TOOLS prompt start
ds4: ROCm GLM causal indexed prefill using fp16 hipBLAS attention GEMMs (tokens=1318 rows=1318 cache=f16)
ds4-tp: big gate window barrier failed
ds4: GLM TP batch gate failed (layer 3)

And this error on the worker:

ds4: memory: KV 0.56 GiB (raw 0.00 + compressed 0.56) + buffers 3.01 GiB + resident model 96.13 GiB = 99.69 GiB planned
ds4: memory detail: ctx=50000 prefill_cap=2048 raw_kv_rows=0 compressed_kv_rows=50000 backend=rocm
ds4-tp: rdma device rocep3s0f0 (port state 4)
ds4-tp: rdma device limits: max_qp_wr 8192, max_sge 30, max_cqe 4194303, max_mr_size 18446744073709551615
ds4-tp: leader connected, transport=rdma gate-timeout=750ms
ds4-tp: rdma warm-up ok (1 attempt), queue depth recv 1024 send 1024
tensor parallelism bound: rank 1, 50/50 expert split, rdma transporttp worker ready for mirrored sessionsds4: GLM session ctx=50000 (model max=1048576); prefill/work cap=4096; dense attention limit=2051
ds4: GLM graph using compact DSA KV only; expanded full-attention KV cache is skipped
ds4: GLM graph allocating compact DSA cache: rows=50000 logical_ctx=50000 kv_layers=45 indexer_layers=11 f16 0.56 GiB
ds4: GLM compact indexed prefill chunk=2048 score_rows=2048 score_scratch=97.66 MiB
ds4: ROCm GLM causal indexed prefill using fp16 hipBLAS attention GEMMs (tokens=1318 rows=1318 cache=f16)
ds4-tp: timeout in big gate window (237/256 recvs, 4/4 sends, 256 sent)
ds4: GLM TP batch gate failed (layer 3)

Commands that I use:

merlijn@strixhalo:~$ cat ./run-ds4-tp.sh
#!/bin/sh
sudo docker run --rm -it --net host --ulimit memlock=-1 -e DS4_ROCM_GLM_TP=1   --device /dev/kfd --device /dev/dri --device /dev/infiniband   --group-add video --group-add render   --ipc=host --cap-add=SYS_PTRACE --security-opt seccomp=unconfined  -v ~/ds4:/ds4  -v ~/ds4/gguf:/models -v ~/ds4/ds4-kv:/tmp/ds4-kv  ds4-glm   /ds4/ds4-server -m /models/GLM-5.3-Flash-Q4_K.gguf --kv-disk-dir /tmp/ds4-kv --kv-disk-space-mb 100000 -t 32 --ctx 50000 --role coordinator --tensor-parallel --transport rdma --listen 0.0.0.0 1234  --host 0.0.0.0  --kv-cache-cold-max-tokens 100000
merlijn@strix-halo-2:~$ cat ./run-ds4-worker-tp.sh
#!/bin/sh
sudo docker run --rm -it --net host --ulimit memlock=-1 -e DS4_ROCM_GLM_TP=1  --device /dev/kfd --device /dev/dri --device /dev/infiniband   --group-add video --group-add render   --ipc=host --cap-add=SYS_PTRACE --security-opt seccomp=unconfined   -v ~/ds4/gguf:/models -v ~/ds4:/ds4   ds4-glm   /ds4/ds4-server -m /models/GLM-5.3-Flash-Q4_K.gguf -t 32 --ctx 50000 --kv-disk-dir /tmp/ds4-kv --kv-disk-space-mb 8192   --role worker --tensor-parallel --coordinator 192.168.100.1 1234  --kv-cache-cold-max-tokens 100000

I see a different error when using the 0731 model (DeepSeek-V4-Flash-Q4KExperts-F16HC-F16Compressor-F16Indexer-Q8Attn-Q8Shared-Q8Out-chat-v2-imatrix-0731.gguf):

coordinator:

ds4: sharded warm done in 26.1s (checksum=2515874312)
ds4: ROCm preparing model tensor mappings: 8.20 GiB
ds4: ROCm startup model preparation covered 8.20 GiB of tensor spans in 0.001s
ds4: rocm backend initialized for graph diagnostics
ds4: memory: KV 1.00 GiB (raw 0.36 + compressed 0.64) + buffers 0.38 GiB + resident model 80.76 GiB = 82.14 GiB planned
ds4: memory detail: ctx=50000 prefill_cap=4096 raw_kv_rows=4352 compressed_kv_rows=12502 backend=rocm
ds4-tp: waiting for worker on 0.0.0.0:1234 ...
ds4-tp: rdma device rocep3s0f0 (port state 4)
ds4-tp: rdma device limits: max_qp_wr 8192, max_sge 30, max_cqe 4194303, max_mr_size 18446744073709551615
ds4-tp: worker connected, transport=rdma gate-timeout=750ms
ds4-tp: rdma warm-up ok (1 attempt), queue depth recv 1024 send 1024
tensor parallelism bound: rank 0, 50/50 expert split, rdma transport0911 11:14:00 ds4-server: context buffers 1414.48 MiB (ctx=50000, backend=rocm, prefill_chunk=4096, raw_kv_rows=4352, compressed_kv_rows=12502)
0911 11:14:01 ds4-server: KV disk cache /tmp/ds4-kv (budget=100000 MiB, cross-quant=accept, min=512, cold_max=100000, continued=10000, trim=32, align=2048, hit_half_life=21600s)
0911 11:14:01 ds4-server: listening on http://0.0.0.0:8000
0911 11:14:39 ds4-server: chat ctx=0..1534:1534 TOOLS prompt start
ds4-tp: timeout in big gate window (151/256 recvs, 4/4 sends, 256 sent)
ds4-tp: timeout in big gate window (124/256 recvs, 4/4 sends, 256 sent)
ds4: TP prefill attention row gate failed (layer 1)
ds4: gpu layer 1 attention batch encode failed
ds4: gpu layer-major prefill layer 1 encode failed

worker:

ds4: warming sharded tensor pages (rank 1): 80.76 of 153.32 GiB
ds4: sharded warm done in 19.5s (checksum=2515656570)
ds4: ROCm preparing model tensor mappings: 8.20 GiB
ds4: ROCm startup model preparation covered 8.20 GiB of tensor spans in 0.001s
ds4: rocm backend initialized for graph diagnostics
ds4: memory: KV 1.00 GiB (raw 0.36 + compressed 0.64) + buffers 0.38 GiB + resident model 80.76 GiB = 82.14 GiB planned
ds4: memory detail: ctx=50000 prefill_cap=4096 raw_kv_rows=4352 compressed_kv_rows=12502 backend=rocm
ds4-tp: connecting to 192.168.100.1:1234 ... (Connection refused)
ds4-tp: rdma device rocep3s0f0 (port state 4)
ds4-tp: rdma device limits: max_qp_wr 8192, max_sge 30, max_cqe 4194303, max_mr_size 18446744073709551615
ds4-tp: leader connected, transport=rdma gate-timeout=750ms
ds4-tp: rdma warm-up ok (1 attempt), queue depth recv 1024 send 1024
tensor parallelism bound: rank 1, 50/50 expert split, rdma transporttp worker ready for mirrored sessionsds4-tp: timeout in big gate window (145/256 recvs, 4/4 sends, 256 sent)
ds4: TP prefill attention row gate failed (layer 1)
ds4: gpu layer 1 attention batch encode failed
ds4: gpu whole-prefill layer 1 encode failed
ds4-tp: timeout in big gate window (140/256 recvs, 4/4 sends, 256 sent)

EDIT: I'll try to use an agent to see if I can track this down to an issue on my side or not.

…TIMEOUT_MS

A timed-out big-gate window left its chunks orphaned in the QP: the next
exchange inherited the stale state and failed with even fewer completions
(observed on a RoCE setup whose per-layer arrival skew exceeded the 750 ms
default gate deadline; the retry degradation 151 -> 124 recvs in the
reporter log matches).  Reset the bulk state on the timeout path like the
completion-error path already does, and mention DS4_TP_GATE_TIMEOUT_MS in
the barrier / window / odl timeout messages so the knob is discoverable
from the log.
@davidcanar

Copy link
Copy Markdown
Author

Thanks for the detailed report — it's diagnosable from your logs, and the good news is this doesn't look like the GLM code in this PR.

What your logs show. All three lines are one event. Your coordinator arrived at a TP bulk-gate window barrier first and waited 750 ms for the worker: the barrier is a one-byte handshake on the control socket, and that socket's receive timeout is DS4_TP_GATE_TIMEOUT_MS (750 ms by default). The worker arrived late — its handshake byte was already buffered, so it proceeded — and streamed its half of the exchange into a partner that had already bailed. That's the worker's timeout in big gate window (237/256 recvs, ...): that timeout only fires when the RDMA completion queue is completely idle, i.e. the coordinator had genuinely stopped sending. GLM TP batch gate failed (layer 3) is just the graceful fallout on the coordinator. In short: per-layer arrival skew between the two ranks exceeded the 750 ms gate deadline.

Why I don't think it's the GLM changes: your DeepSeek-V4-Flash run fails identically (TP prefill attention row gate failed (layer 1), same window timeouts) — and that code path predates this PR entirely. When both models fail the same way on the first prefill chunk, the cause is transport/environment/timing rather than the GLM enablement. We see the same failure family on our own boxes at very large contexts: at 128k the per-layer GPU work between gates grows past 750 ms and the pair dies exactly like yours.

Things to try, in order:

  1. Add -e DS4_TP_GATE_TIMEOUT_MS=10000 to both docker runs. This is the knob for exactly this signature; at 128k context it took us from hard failure to fully clean runs.
  2. Run the worker as the CLI binary instead of ds4-server: /ds4/ds4 -m /models/GLM-5.3-Flash-Q4_K.gguf --role worker --tensor-parallel --coordinator 192.168.100.1 1234 --ctx 50000 (plus the same env). That is the configuration validated in this PR; ds4-server --role worker is permitted but was not part of our validation matrix. Also drop --kv-disk-dir and -t 32 from the worker — it needs neither.
  3. A/B with --transport tcp on both sides. If TCP works, the RDMA path is implicated — RoCE without priority flow control is notoriously stall-prone under bursty 16 KiB message streams, and Docker's /dev/infiniband plumbing can contribute.
  4. If it still fails at a modest --ctx (say 8192) with the larger timeout, please share rdma link show and ibv_devinfo -d rocep3s0f0 output (hardware NIC vs rxe/soft-RoCE, MTU, direct-attached or via a switch).

I've also pushed a small hardening commit to this branch (0e02528): the window-timeout path now resets the RDMA bulk state — previously only the completion-error path did, so after one timeout every later window failed worse, which matches the degrading recv counts in your DeepSeek log (151 → 124) — and the barrier/window timeout messages now name DS4_TP_GATE_TIMEOUT_MS so the knob is discoverable from the log itself.

@MerlijnWajer

Copy link
Copy Markdown

Thanks for the reply. I did try the higher timeout (in both the coordinator and the worker) and that alone didn't help.

Please find below the output that you asked for from both machines. I have a local deepseek v4 flash trying to see if it can figure out the problem as well (but I don't know if it will).

merlijn@strixhalo:~$ rdma link show
link rocep3s0f0/1 state ACTIVE physical_state LINK_UP netdev enp3s0f0np0 
link rocep3s0f1/1 state DOWN physical_state DISABLED netdev enp3s0f1np1 
merlijn@strixhalo:~$ ibv_devinfo -d rocep3s0f0
hca_id: rocep3s0f0
        transport:                      InfiniBand (0)
        fw_ver:                         14.32.1912
        node_guid:                      9803:9b03:0001:9952
        sys_image_guid:                 9803:9b03:0001:9952
        vendor_id:                      0x02c9
        vendor_part_id:                 4117
        hw_ver:                         0x0
        board_id:                       MT_2470112034
        phys_port_cnt:                  1
                port:   1
                        state:                  PORT_ACTIVE (4)
                        max_mtu:                4096 (5)
                        active_mtu:             1024 (3)
                        sm_lid:                 0
                        port_lid:               0
                        port_lmc:               0x00
                        link_layer:             Ethernet
merlijn@strix-halo-2:~$ rdma link show
link rocep3s0f0/1 state ACTIVE physical_state LINK_UP netdev enp3s0f0np0 
link rocep3s0f1/1 state DOWN physical_state DISABLED netdev enp3s0f1np1 
merlijn@strix-halo-2:~$ ibv_devinfo -d rocep3s0f0
hca_id: rocep3s0f0
        transport:                      InfiniBand (0)
        fw_ver:                         14.32.1912
        node_guid:                      9803:9b03:0096:49f0
        sys_image_guid:                 9803:9b03:0096:49f0
        vendor_id:                      0x02c9
        vendor_part_id:                 4117
        hw_ver:                         0x0
        board_id:                       MT_2470112034
        phys_port_cnt:                  1
                port:   1
                        state:                  PORT_ACTIVE (4)
                        max_mtu:                4096 (5)
                        active_mtu:             1024 (3)
                        sm_lid:                 0
                        port_lid:               0
                        port_lmc:               0x00
                        link_layer:             Ethernet

@davidcanar

Copy link
Copy Markdown
Author

Thanks — that devinfo helps, and the negative result on the timeout is itself informative: the drain loop only declares a timeout when the RDMA completion queue is completely idle, so a larger DS4_TP_GATE_TIMEOUT_MS can't fix chunks that are never going to arrive. Combined with your devinfo, my suspicion moves to the RDMA link/Docker path itself.

Two things stand out from the output:

  1. active_mtu: 1024 (while max_mtu: 4096). The ds4 RDMA path streams 16 KiB SEND messages in 256-chunk windows, so at MTU 1024 each window is ~4096 wire packets. MTU 1024 usually means the module/DAC or port config negotiated down; ConnectX-4 Lx is also picky about cables. Worth trying: set ip link set enp3s0f0np0 mtu 4096 on both boxes and confirm active_mtu rises to 4096, then retry.
  2. ConnectX-4 Lx on fw 14.32 is an old-ish stack; RC retransmission under a lossy or flow-control-less RoCE link can stall a QP for a long time, which matches a silent CQ.

On our side, we just ran the TCP A/B I mentioned, end-to-end on our own two Strix Halo boxes with the same model and your context scale: --transport tcp completes the full 47,288-token prefill + decode cleanly (prefill 38.5 t/s, generation 6.5 t/s, zero gate failures), and a short 227-token prompt also passes (51.3 / 8.1 t/s). So the TCP transport is fully functional at your context size — if it also works for you, the RoCE/Docker path is confirmed as the culprit and you can either stay on TCP (25G is more than enough; our TCP numbers match our OdinLink TB5 numbers within noise) or debug the RDMA path with:

# one box:
ib_send_bw -d rocep3s0f0 -R
# other box:
ib_send_bw -d rocep3s0f0 -R <ip-of-first-box>

ib_send_bw is the standard Mellanox send-burst benchmark; if a plain 16 KiB send stream can't sustain ~line rate with zero errors there, the link (cable/MTU/switch) is the problem and ds4 is just the messenger. Also useful after a failed ds4 run: ethtool -S enp3s0f0np0 | grep -iE "crc|drop|error|discard" on both boxes — rising counters would directly confirm link-level drops.

Two more things that would help narrow it down:

  • The failure logs from the raised-timeout run — did the failing line move (e.g. the barrier timeout now at 10 s instead of 750 ms), and at which layer?
  • The CLI-worker and bare-metal (outside Docker) variants if you get a chance: the PR's validation used the CLI worker binary, and bare metal rules out the container's /dev/infiniband plumbing.

For transparency: our boxes have no verbs devices at all (we run the OdinLink TB5 transport from #1018), so we can't reproduce your RoCE path directly — the tests above are the fastest way to pin it down from your side.

@MerlijnWajer

MerlijnWajer commented Sep 11, 2026

Copy link
Copy Markdown

Thank you for the detailed reply, I will report back later this weekend. I do have this fork (https://github.com/wkljohn/ds4-strix-halo-tp-odinlink) working with RDMA on the same hardware at 20t/s tg, so I think the hardware stack is capable/functional, but you mention some good points and I'll make sure to check them and report back.

@MerlijnWajer

Copy link
Copy Markdown

Okay, so I got things to work with this commit created by deepseek-v4-flash-0731 using pi:
MerlijnWajer@6d2dc68

The tcp transport already worked as you mentioned, but the above commit makes the rdma transport work for me as well.

GLM Flash 5.3 runs at similar speed: ~50 t/s prefill, 8t/s gen. Deepseek v4 flash runs at 14 t/s gen (slower than the 20t/s from the repo I mentioned earlier).

As you can see the link itself is stable, although I think I might be able to do some work in configuring my switch so that it can handle larger MTUs.

I don't know if the above fix belongs in this branch or not, but it made things work for me. The commit also adds some diagnostics code which may or may not be useful.

In any case, thanks for all the help, this is really great work, as it also makes GLM Flash 5.3 work in a distributed manner (it doesn't on the main branch). I might ask GLM Flash 5.3 if it can figure out why the https://github.com/wkljohn/ds4-strix-halo-tp-odinlink fork runs at 20 t/s gen as opposed to the 15 t/s gen on deepseek v4 flash, but I don't want to hijack the PR any further. Please let me know if I can help with anything or what other tests you'd like me to run if any.

The rdma transport hardcodes UC because it was built for the Apple
thunderbolt-ibverbs driver, where only UC queue pairs exist.  On real
NICs UC silently drops a SEND that lands before the peer recv WR is
armed (no RNR retry): a big-gate window then stalls with silent
completion queues (e.g. 237/256 recvs, no error completion) and no
DS4_TP_GATE_TIMEOUT_MS setting can help.  Probe RC first and fall back
to UC for the thunderbolt driver; on RC set the RTR/RTS attributes UC
does not consume (max_dest_rd_atomic/min_rnr_timer at RTR, timeout/
retry_cnt/rnr_retry=7/max_rd_atomic at RTS).

Also drop the bulk_reset call the window-timeout path made: it rebuilt
the bulk QP while the orphaned window lives on the latency QP, and its
peer-coordinated handshake cannot complete on an asymmetric failure.
Recovery from a window timeout needs a symmetric latency-QP rebuild and
is future work; under RC the transport-level retries make reaching the
timeout a true dead-peer signal.  Add a DS4_TP_CQ_DEBUG completion-queue
dump to the timeout path so a future occurrence reports queue state.

Based on the working fix by MerlijnWajer
(github.com/MerlijnWajer/ds4/commit/6d2dc68), verified by him on two
Strix Halo hosts over ConnectX-4 Lx RoCE at MTU 1024: GLM-5.3-Flash
50.1 t/s prefill and 8.1 t/s generation with zero gate failures; the
21.6 MB big-gate exchange drops from a ~12 s stall to ~18 ms.  The link
itself measures clean (ib_send_bw 2.5 GiB/s, 164K msg/s).
@davidcanar

Copy link
Copy Markdown
Author

This is a great catch, and your diagnosis is exactly right — thank you. The UC in our transport is a fossil from its origin: it was written for the Apple thunderbolt-ibverbs driver, where only UC queue pairs exist. On a real NIC, UC's "drop a SEND that lands before the peer's recv WR is armed" behavior explains everything about your failure signature, including why no timeout setting could help (the chunks were genuinely never going to arrive, and the completion queue stays silent because UC reports nothing).

You were also right about our 0e02528 hardening being wrong: tp_rdma_bulk_reset() rebuilds the bulk QP while the orphaned window lives on the latency QP, and its peer-coordinated handshake can't complete on an asymmetric failure anyway. I've dropped that call — recovery from a window timeout needs a symmetric latency-QP rebuild and is noted as future work; under RC the transport-level retries make the timeout a true dead-peer signal.

I've integrated your fix into this branch as 4402b61, with one change: instead of switching to RC unconditionally, the transport now probes RC and falls back to UC — the Apple thunderbolt driver is UC-only, and the rdma path is shared code, so a hard switch would have broken macOS tensor parallelism. On RC it sets the RTR/RTS attributes UC doesn't consume (same values as your commit: min_rnr_timer=12, timeout=14, retry_cnt=7, rnr_retry=7, max_rd_atomic=1), and it includes your DS4_TP_CQ_DEBUG queue dump on the timeout path.

Since you have the only verbs hardware in the loop: could you pull davidcanar:glm-tp-rocm and re-run your RDMA setup once? The probe should log ds4-tp: rdma using RC queue pairs at startup, and everything else should behave like your commit.

And nicely done getting DeepSeek-V4 to produce a working verbs fix — the commit message reads like it genuinely understood the failure mode. Thanks for the kind words too: GLM-5.3 Flash distributed across two Strix Halos is exactly what this PR is for, so reports like yours are the payoff.

@MerlijnWajer

Copy link
Copy Markdown

I tried the latest commit that you pushed and it doesn't quite work for me yet. I get this on the coordinator:

ds4: warming sharded tensor pages (rank 0): 96.13 of 177.76 GiB
ds4: sharded warm done in 102.6s (checksum=3149873359)
ds4: ROCm preparing model tensor mappings: 14.49 GiB
ds4: ROCm startup model preparation covered 14.49 GiB of tensor spans in 0.002s
ds4: ROCm mapped 1.05 GiB auxiliary model
ds4: rocm backend initialized for graph diagnostics
ds4: memory: KV 2.79 GiB (raw 0.00 + compressed 2.79) + buffers 3.16 GiB + resident model 96.13 GiB = 1
02.08 GiB planned
ds4: memory detail: ctx=250000 prefill_cap=2048 raw_kv_rows=0 compressed_kv_rows=250000 backend=rocm
ds4-tp: waiting for worker on 0.0.0.0:8000 ...
ds4-tp: rdma device rocep3s0f0 (port state 4)
ds4-tp: rdma device limits: max_qp_wr 8192, max_sge 30, max_cqe 4194303, max_mr_size 184467440737095516
15
ds4-tp: rdma using RC queue pairs
ds4-tp: worker connected, transport=rdma gate-timeout=10000ms
0912 21:42:26 ds4-server: tp rdma: warm-up post_send: No space left on device

and this on the worker:

ds4: warming sharded tensor pages (rank 1): 96.13 of 177.76 GiB
ds4: sharded warm done in 61.6s (checksum=3148094405)
ds4: ROCm preparing model tensor mappings: 14.49 GiB
ds4: ROCm startup model preparation covered 14.49 GiB of tensor spans in 0.000s
ds4: ROCm mapped 1.05 GiB auxiliary model
ds4: rocm backend initialized for graph diagnostics
ds4: memory: KV 2.79 GiB (raw 0.00 + compressed 2.79) + buffers 3.16 GiB + resident model 96.13 GiB = 102.08 GiB planned
ds4: memory detail: ctx=250000 prefill_cap=2048 raw_kv_rows=0 compressed_kv_rows=250000 backend=rocm
ds4-tp: connecting to 192.168.100.1:8000 ... (Connection refused)
ds4-tp: connecting to 192.168.100.1:8000 ... (Connection refused)
ds4-tp: rdma device rocep3s0f0 (port state 4)
ds4-tp: rdma device limits: max_qp_wr 8192, max_sge 30, max_cqe 4194303, max_mr_size 18446744073709551615
ds4-tp: rdma using RC queue pairs
ds4-tp: leader connected, transport=rdma gate-timeout=10000ms
tp worker: tp rdma: warm-up post_send: No space left on device

I am not quite sure how to interpret 'No space left on device' here.

ibv_post_send/ibv_post_recv return the errno-style code themselves and
do not reliably set errno, so the warm-up and big-gate failure prints
could show a stale value (a reporter saw "No space left on device" on
the first warm-up send of an RC pair, which may not be the real
reason).  Capture the return code and print strerror() on it at the
warm-up and big-gate window posts, and log the driver-adjusted
max_send_wr/max_recv_wr the QP was actually created with.
@davidcanar

Copy link
Copy Markdown
Author

Progress — the probe works (rdma using RC queue pairs), so RC creation and the modified QP transitions are fine on your ConnectX-4 Lx. The new failure is at the first warm-up post_send, but I don't yet trust the message itself: ibv_post_send returns the error code directly and does not reliably set errno, so our print was showing a possibly-stale errno — "No space left on device" (ENOMEM) may not be the real reason.

I pushed cb7cb70, which:

  • captures the real return code from ibv_post_send/ibv_post_recv at the warm-up and big-gate sites and prints strerror() on that;
  • logs the driver-adjusted QP caps after creation (ds4-tp: rdma RC QP caps: max_send_wr=... max_recv_wr=...) — for RC the driver can adjust the requested counts, and I want to see what it actually gave us vs. what your own working commit used.

Could you pull and re-run once (both ranks, and add DS4_TP_CQ_DEBUG=1 to both so the hello exchange details are logged too)? With the real return code and the caps we should be able to pin this quickly. Two early hypotheses, depending on what the re-run says: if it turns out to be EINVAL, something about the RC transition differs from your commit; if it is genuinely ENOMEM, the RC send queue is being sized smaller than requested and I'll size the warm-up/window posts to the negotiated depth instead.

Frustrating that this is now a three-way handshake between your hardware, my code, and both of our assumptions — but we're close: the transport already negotiates RC, connects, and fails at a well-defined point with (soon) honest error codes.

davidcanar and others added 7 commits September 12, 2026 21:25
…TP branch

Upstream independently added a mirrored-sync checkpoint protocol that
collided with this fork's lockstep frames at wire value 22.  Resolved by
renumbering (SYNC_GO 22, SYNC_CHECKPOINT 23, ODL_READY 24, protocol v13)
and keeping both mechanisms: the fork lockstep (SYNC_GO +
wait_command_ack_status + poison) keeps driving GLM TP, while upstream
checkpoint/sync_lockstep and its frame stay for upstream flows.

- ds4_engine_tp_bind: keep the fork kv-split registration; adopt the
  upstream vocab-split rule (only DeepSeek V4 splits its output head,
  so GLM and V4.1 stay replicated).
- ds4_session_sync: keep the fork lockstep flow (validated for GLM TP
  through 47k contexts) instead of upstream sync_lockstep; both remain
  in the tree.
- ack waits: keep the fork bounded/poisoning wait_command_ack_status and
  add upstream reserved-field validation; ds4_tp_wait_command_status
  becomes a wrapper over it (upstream auto-merged call sites unchanged).
- big gate header: keep the fork bytes field (desync check + odl path),
  adopt the upstream cold-prefill grace window around the header
  exchange; adapt the timeout print to upstream reworked drain
  accounting (bitmap dedup replaced the signaled counter).

Both cpu and strix-halo targets build clean.
The post-merge stitching fixes (adapting the window-timeout print to
upstream reworked drain accounting, dropping the duplicated
wait_command_ack definition, and removing a message that referenced a
variable the resolution dropped) were applied after the merge was
staged and did not make it into the merge commit; this restores the
state that was built and validated.
The k-slice kernel behind the GLM DSA attention output projection (and
the DeepSeek TP attention-output / shared-expert expand paths) read one
byte per lane per 34-byte block with a serial loop over the owned K
range, which ran at ~80 GB/s on gfx1151: 444 us for the 35.6 MB half of
a 16384x4096 Q8_0 matrix.  Quantize the owned activation slice to Q8
like the dense decode path and let the lanes stride over the row's
blocks, so one wave instruction covers 32 consecutive blocks and the
int8 dot uses dp4a: 176 us, ~200 GB/s, about 3 ms per decoded token
over the 11 DSA layers of GLM-5.3-Flash.

Measured on 2x Strix Halo (gfx1151) over OdinLink; the DeepSeek callers
of the same entry point were not exercised.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Split ds4_gpu_tp_gate_encode into ds4_gpu_tp_gate_arrive (arrival event
plus exchange request) and ds4_gpu_tp_gate_wait (release spin kernel).
On the eager ROCm stream the GLM sparse FFN now encodes the replicated
shared expert between the two, so the GPU computes it while the CPUs
swap the routed halves instead of parking on the spin kernel; the
combine and the residual run after both.  The exchanged partial is
unchanged, so both ranks still rebuild the routed output bit-identically.
Metal keeps the single-call gate; ds4_gpu_tp_gate_encode is now arrive
followed by wait.

Worth ~2.5 ms per decoded token on 2x Strix Halo (53 gates per token,
~130 us of shared-expert compute available to hide each FFN exchange).

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
The row kernel (out_dim <= 128: hyper-connection mixers, KDA
f_a/g_a/beta, indexer projections) streamed a whole 8-32 KB row per
block with one dependent 2-byte load per iteration, so each of its
~214 launches per decoded token was latency bound (16 us for ~1 MB on
gfx1151).  Issue eight independent 128-bit weight loads per thread
before consuming them; the scalar loop remains for in_dim not divisible
by 8.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
glm_graph_forward_token flushed the command stream every 4 layers under
indexed attention (32 otherwise), a cadence inherited from the Metal
command-buffer budget.  On ROCm the flush is a full device sync that
only drains the CPU run-ahead and idles the GPU until the encoder
catches up: under two-node TP it cost ~2 ms per decoded token
(14.9 -> 15.4 t/s on GLM-5.3-Flash with Q8 attention weights, 45-token
prompt, greedy).  Keep it off on ROCm; DS4_GLM_DECODE_FLUSH_INTERVAL
still forces a cadence when wanted.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
…timeout

odl_tb5 delivers whole messages in order, but when one Thunderbolt frame of
a message fails CRC the driver drops the entire message ("fragment gap ...
dropping message" in dmesg) and the sender is never told.  A gate that
waited for such a message parked both ranks on their spin kernels until the
timeout and took the TP session down: seen twice on a 2x Strix Halo pair
in one afternoon, once in an MTP verify block and once 4096 tokens into an
11k-token ds4-server prefill (the link negotiates 10 Gb/s per lane and logs
CRC errors every few GB).

Every stream message now carries a 16-byte header (magic, kind, chunk
index, exchange sequence).  A receiver asks for the first chunk it misses
either as soon as a later index reveals the gap or after a retry interval
without progress (2 ms for gate vectors, 50 ms for bulk), and a sender
keeps the framed copies of its current and previous exchange: lockstep
bounds the peer's lag to one exchange (it cannot finish N+1 without our
N+1 data, which we only send after finishing N), so those two buffers
always cover a request.  Chunks of the next exchange that arrive while
this one still misses data are parked and replayed.  Requests are served
from the exchange loop and, for a rank that is idle or waiting on the
control channel (worker between commands, leader waiting for an ack,
either side in the big-gate header handshake), from the polling wait that
now fronts those reads.  No extra round trip on the fast path: a decode
gate costs one 16 KB frame copy on each side.

Measured with tp_odl_test over the real link, 3000 gates of 16 KB plus the
batch and 2 MiB big exchanges: p50 29.0 us (29.3 us before), 2 MiB swap
2.2 ms (1.8 ms before, the header/frame copies).  With
DS4_TP_ODL_DROP_TEST=5 on both ranks (every fifth data send dropped) all
exchanges complete with verified payloads and one retransmission per
drop; p90 gate latency becomes the 2 ms detection interval.  Dropping 1 in
3 sends on one rank only also completes cleanly.  The knob is a diagnostic
that exercises this path; it is off unless set.
(cherry-picked from 73d9909 on odinlink-transport; this branch's
tp_read_full_deadline gate-header reads poll the same way.)

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants