Repository navigation
Use CUDA P2P / NVLink for multi-GPU layer-split hand-off (dual RTX 2080 Ti) #338
Description
[Feature request] Use CUDA P2P / NVLink for multi-GPU layer-split hand-off (dual RTX 2080 Ti)
Summary
Would you consider adding an optional CUDA P2P / NVLink path for the multi-GPU layer-split hand-off?
Strata currently documents the split as one hand-off per verify window through mapped pinned host memory, and that works well without requiring NVLink. On systems where CUDA peer access is available, especially dual RTX 2080 Ti with NVLink, it may be worth allowing the stage boundary to use direct GPU-to-GPU transfer instead.
A possible opt-in flag could be:
--split-p2p
with the existing mapped pinned-host path kept as the default/fallback.
Test system
- Linux / Ubuntu 24.04
- 2× RTX 2080 Ti 22 GB
- NVLink present (
nvidia-smi topo -mreportsNV2) - Strata engine 0.1.27
- Swift 1.5 / Qwen3.8-Flash-Next
- GSQ-RCO IQ3_XXS
- layer split: auto
- context: 262144
- KV: int8
- KV resident: 32768
- vision: GPU
- prefill: 4096
- VRAM reserve: 1200 MiB
Selected split:
strata serve: layer split: layers 0-25 (CUDA0), 26-47 (CUDA1), one hand-off per window
Runtime:
strata serve: 1475 MiB of VRAM free with everything loaded
strata serve: prompt 201056 tokens = 0 reused + 201056 read in 274704 ms (731.9 tok/s),
362 generated in 8072 ms (44.8 tok/s),
drafts accepted 222 of 337
strata serve: decode expert cache hit rate: 92.8% (212845 hits / 229302 lookups)
strata serve: KV streaming: 93.58% of 1485072 block reads hit VRAM,
384.2 MiB read from RAM
PCIe probes reported by Strata:
CUDA0 host->device: 6.2 GB/s
CUDA1 host->device: 10.9 GB/s
Why this might help
The current implementation intentionally does not require NVLink/P2P and uses a mapped pinned host hand-off, which is great for compatibility.
But on a pair of GPUs with working peer access and NVLink, the boundary could potentially be:
CUDA0 device buffer
|
| cudaMemcpyPeerAsync / peer-visible memory
v
CUDA1 device buffer
instead of:
CUDA0 -> mapped pinned host memory -> CUDA1
Since there is only one hand-off per verify window, I would not expect a huge speedup, but it could potentially:
- reduce stage-boundary latency;
- reduce host involvement;
- avoid PCIe/host-memory jitter;
- improve prefill pipeline smoothness;
- slightly reduce fixed decode-window overhead.
The current verifier API already has a clear stage boundary through handoff_in / handoff_out, so this looks like it could be implemented as an optional transport rather than changing the layer-split design.
Suggested behavior
At startup:
- Check
cudaDeviceCanAccessPeer(). - If both stage devices support peer access, enable it with
cudaDeviceEnablePeerAccess(). - Allocate per-stage device hand-off buffers.
- Transfer with
cudaMemcpyPeerAsync()or a peer memcpy node inside the CUDA graph. - Synchronize the two stage streams with CUDA events.
- Automatically fall back to the existing mapped pinned-host hand-off when peer access is unavailable.
An opt-in implementation such as --split-p2p would make A/B testing easy while preserving the current portable path.
I can test patches and provide before/after prefill, decode, hand-off latency, stability and VRAM measurements on this dual RTX 2080 Ti NVLink system if useful.
Activity
Niko1221 commentedon Oct 1, 2026
Thanks! A measured opt-in would be welcome:
--split-p2pwith acudaDeviceCanAccessPeercheck and the normal path as fallback. #294 measured the hand-off at ~1.1 ms per window on PCIe, so please include numbers from your two 2080 Tis. We have no NVIDIA multi-GPU rig.Niko1221 commentedon Oct 1, 2026
Thanks for the detailed request and the numbers. A P2P/NVLink hand-off is reasonable as an opt-in (
--split-p2p, with the mapped host path kept as the default), but the hand-off is one small transfer per verify window, so we expect only a small gain. We have no NVLink pair to measure it on. I've noted it in our tracking list, and a PR is welcome if it measures faster on your 2080 Ti pair (decode and prompt, with the default unchanged). Closing the issue to keep the tracker for bugs.yeungyim commentedon Oct 2, 2026
Measured it on an NVLink pair, since you mentioned you have none to test on. I implemented
--split-p2pto the spec in your comment (cudaDeviceCanAccessPeer+cudaDeviceEnablePeerAccess, the mapped host path kept as the default and as the fallback on any failure) and A/B'd it.Correctness: bit-exact. The hand-off carries the same floats either way, so I gated on output identity. A seeded cold prompt produced the same set of outputs under both modes (mapped: A, B, A across three restarts; p2p: B, A) — the variation is run-to-run and comes from the adaptive expert cache re-rounding, not from the transport. The receiving kernel needed no change: its source pointer is already
const volatile float4*, which is correct for peer memory too.Speed on 2x RTX 3090 (NVLink NV4): 2-3% slower.
--split-p2pWhy it doesn't pay here: the mapped path is already zero-copy — each stage's GPU reads and writes the hand-off buffer in host memory directly, so there is no intermediate copy for peer access to remove. It only swaps "read host memory" for "read the peer's device memory", and on this board (both cards behind the CPU's PCIe lanes, Gen3 x8 each) that is not faster; the NVLink bridge is not chosen for this access pattern.
So your "we expect only a small gain" was right, and the sign is slightly worse than that.
Rig: 2x RTX 3090 24 GB, driver 595.91.07,
nvidia-smi topo -mreportsNV4, Linux, CUDA 13.1, engine 0.1.35 built from source, Qwen3.8-Flash-Next IQ3_XXS, layer split auto (K=26: layers 0-25 / 26-47), int8 KV, 262144 context. Measured with cold prefill — a unique leading salt per request, because the conversation cache otherwise serves a repeated prompt in ~0.4 s and you measure the cache, not the engine — temperature 0, max_tokens 256, median of 3.The patch is ~49 lines: the hand-off buffer allocated with
cudaMallocon the sending stage instead ofcudaHostAllocMapped, plus the capability check and fallback. Happy to post it if you want it on the record despite the negative result.vaskes commentedon Oct 29, 2026
Closing the loop on the numbers asked for here: the 2x RTX 2080 Ti 22 GB with NVLink (
NV2), Qwen3.8-Flash-Next GSQ-RCO IQ3_XXS (IQ3_XXS),--layer-split auto(K=26),--kv int8 --kv-resident 32768, context 262144, CUDA0 on a root-complex x16 slot and CUDA1 under a downstream port at x4.I implemented the prompt path's hand-off over peer access on the current
main(engine 0.1.53) and measured it:The reason is measurable and, I think, general for this topology:
nvidia-smi dmonduring a 16K prompt shows CUDA0 (x16) pulling 6.8 GB/s from the host (67% of its 10.1 GB/s link) and CUDA1 (x4) pulling 3.2 GB/s (97% of its 3.3 GB/s link, saturated). The prompt is bound by the pair's host traffic - the streamed experts - and not by the inter-card hand-off, which is ~36 ms of a 14.4 s prompt (0.24%). The NVLink removes those 36 ms but cannot widen the path to the host; and moving the x4 card's streamed experts over the NVLink instead would need the same bytes over the x16 card's link (~3 GB/s spare against the ~4 GB/s the x4 card pulls), which is a wash too.Full write-up, the implementation notes and the environment are in the discussion of PR #1099 (the same proposal): #1099 (comment)
vaskes commentedon Oct 29, 2026
Correction to the numbers above (the same correction is in the PR #1099 thread): the chunk's rows carry the hyper-connection's four heads, so a row is
D = N * HC = 10240floats, and one 8192-row chunk is 335.5 MB, not 83.9 MB. The host path therefore moves ~143 ms of PCIe time per chunk (the producer's download over its x16 at 8.1 GB/s: 41.7 ms, the consumer's upload over its x4 at 3.3 GB/s: 101.7 ms) against 7.16 ms (46.9 GB/s) over the NVLink - ~20x, not ~5x. The end-to-end numbers and the reason for them do not change (the copies are overlapped with the compute; the pair is bound by the host traffic).#1099 (comment)