CUDA illegal memory access in internode_ll::dispatch+0x1050 on B300/RoCE (Blackwell SM100, 16 ranks, hidden=7168)

Summary

Running DeepEP V1 on B300/RoCE produces a CUDA error: an illegal memory access was encountered inside the deep_ep::internode_ll::dispatch kernel. compute-sanitizer pinpoints the failing instruction at offset +0x1050 of the dispatch kernel — the access is 47 GB out of bounds from a small 8-byte allocation, consistent with a pointer-arithmetic miscomputation. NVSHMEM IBGDA's data plane is functional on the same cluster (verified independently via NVSHMEM's reduction_on_stream perftest — full latency table below).

Environment

Item | Value
-- | --
Hardware | NVIDIA B300 SXM6, 8× GPUs/node, x86_64
GPU arch | SM 10.0 (Blackwell B300)
Network | RoCE (Ethernet link layer) on Mellanox ConnectX-7 mlx5_0, RoCEv2 with PFC priority 3 / DSCP 24
OS / Container | Ubuntu 24.04, CUDA 13.0.88, PyTorch 2.9.1+cu130, gcc 13.3.0 (SGLang deepseek-v4-b300 family container)
DeepEP | 1.2.1
NVSHMEM | recent NVSHMEM build (host + IBGDA transport) — both NVSHMEM init and IBGDA transport selection succeed; the data plane is verified independently below
Test | tests/test_low_latency.py --num-processes 8 (2 nodes × 8 GPUs = 16 ranks total)
Hidden size | 7168 (default per tests/utils.py)

Symptom

NVSHMEM IBGDA transport initializes successfully on all 16 ranks (Successfully initialized: IBGDA × 16), buffer construction succeeds, and test_main enters the dispatch call. The test then crashes at ~47 sec wall time — NCCL ProcessGroupNCCL watchdog catches the IMA on multiple ranks:

[rank3]:[E505 22:27:27] [PG ID 1 PG GUID 1 Rank 3] Process group watchdog

  thread terminated with exception: CUDA error: an illegal memory access

  was encountered

[rank4]:  same

[rank5]:  same

[rank6]:  same

[rank7]:  same

[rank11..15]: same

10 of 16 ranks reported the error directly; the rest were SIGTERM'd by torch.multiprocessing before logging their own. No bandwidth measurement is reached — EOFError: Ran out of input from torch/multiprocessing/spawn.py:217 is the parent observing children that died before serializing tracebacks.

compute-sanitizer trace pinpoints the failing instruction

========= COMPUTE-SANITIZER

========= Invalid __global__ read of size 8 bytes

=========     at deep_ep::internode_ll::dispatch<(bool)0, (bool)0, (int)7168>(...)+0x1050

=========     by thread (224,0,0) in block (8,0,0)

=========     Access to 0x71ca0ddfa8d8 is out of bounds

=========     and is 46,747,262,417 bytes after the nearest allocation at 0x71bf2b850b00 of size 8 bytes

=========     Saved host backtrace up to driver entry point at kernel launch time

=========         Host Frame: cudaLaunchKernelExC [...]

=========         Host Frame: deep_ep::internode_ll::dispatch(...) [0x7a689] in deep_ep_cpp.cpython-312-x86_64-linux-gnu.so

=========         Host Frame: deep_ep::Buffer::low_latency_dispatch(...) in deep_ep.cpp:1184

=========         Host Frame: low_latency_dispatch in buffer.py:585

=========         Host Frame: test_main in test_low_latency.py:58

=========         Host Frame: test_loop in test_low_latency.py:179

The "47 GB out of bounds from a small 8-byte allocation" pattern strongly suggests pointer-arithmetic gone wrong. The kernel's offset computation — per csrc/kernels/internode_ll.cu source around the dispatch kernel:

dst_ptr = (uint64_t)rdma_recv_x

        + dst_expert_local_idx * num_ranks * num_max_dispatch_tokens_per_rank * num_bytes_per_msg

        + rank * ...

        + slot_idx * num_bytes_per_msg

appears to be computing into unmapped memory on B300 (SM100 / 16 ranks / hidden=7168). Likely candidates:

• dst_expert_local_idx, num_max_dispatch_tokens_per_rank, or num_bytes_per_msg having a value that produces an out-of-stride offset

• Dereferencing a small per-PE counter pointer (the 8-byte allocation) as a buffer base

• Symmetric heap base pointer being computed from a stale or wrongly-typed source

Independent NVSHMEM IBGDA validation on the same cluster

To rule out NVSHMEM as the source, I ran NVSHMEM's own reduction_on_stream perftest from the NVSHMEM tree (perftest/host/coll/reduction_on_stream.cpp) on the same B300/RoCE cluster (16 PEs, 2 nodes, same RoCE env). Full latency table printed end-to-end:

size_B   elems   type op   avg_us       min_us     max_us     avg_GB/s   max_GB/s

128      32      int  sum  36.451840    32.800     39.968     0.004      0.007

256      64      int  sum  36.391040    32.832     39.968     0.007      0.013

...

1048576  262144  int  sum  74.136961    71.168     81.952    14.144     26.520

Cross-node int-sum reduction works for all sizes 128 B → 1 MB. Peak 14 GB/s. Wall time 43 sec. Exit 0 on all 16 ranks. NVSHMEM IBGDA's symmetric heap allocation, on-stream collective, and IBGDA-RDMA writes are functional on this cluster — the bug is therefore inside DeepEP's internode_ll::dispatch kernel, not in NVSHMEM.

Asks for the DeepEP team

I am not asking the team to reproduce the issue from scratch (the B300/RoCE software stack and getting NVSHMEM IBGDA up cleanly on Blackwell B300 has its own moving pieces — happy to coordinate on that separately if useful). What would help most:

• A look at internode_ll::dispatch for any 16-rank / hidden=7168 / num_max_dispatch_tokens_per_rank assumption that could miscompute on B300 / SM100

• Whether the kernel was tested on B300/RoCE specifically (not B300/IB)

• If V1 isn't going to support B300, an explicit mention in docs/legacy.md would help future users avoid this pitfall

• If the fix is non-trivial, V2 (NCCL Gin backend, no NVSHMEM) is the migration path — but most current SGLang inference deployments are still on V1, so a V1 fix or workaround would have immediate impact

Related upstream issues

• #550 (Blackwell SM100 support tracker)

• #608, #590 (HybridEP on B300 — different code path; their bugs not directly related to ours but useful context for B300 support state)

• DeepEP V1 is officially Hopper-only per docs/legacy.md. We acknowledge B300 / SM100 is unofficial in V1; this report aims to flag a specific kernel-level issue that's reachable for users who pin V1 against a Blackwell cluster.

Bundle / evidence

I can share full compute-sanitizer logs (csan-rank0.txt, ~291 KB), Slurm run logs, and the NVSHMEM-side validation evidence (reduction_on_stream PASSED on the same cluster) on request.

Thanks for DeepEP — would love to see this resolved so V1 + Blackwell B300 + RoCE works end-to-end while the V2 migration is in progress.