DispatchInterNodeRecv derives the chunk sub-slice offset from blockId instead of bid (latent; masked when rdma_block_num % 8 == 0)

Start in src/ops/dispatch_combine/internode_v1.cpp at DispatchInterNodeRecv and compare its loop with CombineInterNodeTyped around line 928. Check the existing dispatch/combine tests and the rdma_block_num values in tests/python/ops/test_dispatch_combine_internode_v1.py; done means the receive loop uses the per-iteration block index consistently, with static verification if multi-node hardware is unavailable.

Summary

In DispatchInterNodeRecv, the per-chunk sub-slice offset is derived from blockId (the loop-invariant starting block) instead of bid (the loop variable). The combine path performs the identical computation with bid.

This is latent on every configuration in the repo — it only diverges when rdma_block_num % 8 != 0, and every value used in tests, examples and docs is a multiple of 8. Filing it because the fix is one line and the failure mode, if it ever fires, is silently wrong output rather than a crash.

The inconsistency

src/ops/dispatch_combine/internode_v1.cpp:378-381 — the dispatch receive loop strides bid across blocks and derives the chunk index k and node index i from it:

  for (int bid = blockId; bid < numRecvBlock * maxChunkNum * (nNodes - 1);
       bid += args.rdmaBlockNum) {
    int k = bid / (numRecvBlock * (nNodes - 1));
    int i = (bid / numRecvBlock) % (nNodes - 1);

but line 406, inside that same loop body, switches to blockId:

    for (int j = startTokenIdx + (blockId % numRecvBlock) * warpNum + warpId; j < endTokenIdx;
         j += numRecvBlock * warpNum) {

The combine twin, CombineInterNodeTyped at :928, uses bid for the same expression:

            for (int j = startTokenIdx + (bid % numRecvBlock) * warpNum + warpId; j < endTokenIdx;
                 j += numRecvBlock * warpNum) {

Its surrounding loop (:896) and index derivations (:906-907) are structurally identical to :378-381, so the two paths appear intended to partition work the same way. Within the dispatch loop body, blockId at :406 is the only index not derived from bid.

When it diverges

numRecvBlock is constexpr int numRecvBlock = 8 (:368). The outer loop advances bid += args.rdmaBlockNum, so on iteration n:

bid % 8  ==  (blockId + n * rdmaBlockNum) % 8

That equals blockId % 8 for all n exactly when rdmaBlockNum % 8 == 0. Otherwise the sub-slice offset stops tracking the chunk being processed while k and i continue to, so within a chunk some token sub-slices are visited more than once and others not at all.

Why it is currently latent

Every rdma_block_num I can find in the tree is a multiple of 8:

Location	Value
tests/python/ops/test_dispatch_combine_internode_v1.py	64, 128
tests/python/ops/test_dispatch_combine_jax.py	16
tests/python/ops/test_dispatch_combine_routing_handle.py	64
examples/ops/dispatch_combine/test_dispatch_combine_internode.py	64
docs/MORI-EP-GUIDE.md	64, 32

So this cannot be reached by the current test suite, and I am not claiming it explains any open issue.

The reason I think it is still worth fixing rather than leaving: rdma_block_num is caller-supplied, documented as a tunable in docs/MORI-EP-GUIDE.md:129, and the autotuner selects it programmatically (examples/ops/dispatch_combine/test_dispatch_combine_internode.py:135, "rdma_block_num": best_disp_config[2]). If the search space ever includes a non-multiple of 8, the result is wrong tokens rather than a failure.

Fix

--- a/src/ops/dispatch_combine/internode_v1.cpp
+++ b/src/ops/dispatch_combine/internode_v1.cpp
@@ -403,7 +403,7 @@ inline __device__ void DispatchInterNodeRecv(EpDispatchCombineArgs<T>& args) {
     int endTokenIdx = startTokenIdx + thisChunkTokenNum;
 
-    for (int j = startTokenIdx + (blockId % numRecvBlock) * warpNum + warpId; j < endTokenIdx;
+    for (int j = startTokenIdx + (bid % numRecvBlock) * warpNum + warpId; j < endTokenIdx;
          j += numRecvBlock * warpNum) {

This is a no-op for every configuration currently exercised, since bid % 8 == blockId % 8 whenever rdma_block_num % 8 == 0.

If the intent really is to key the sub-slice off the physical block rather than the logical chunk, a comment at :406 saying so would help — the asymmetry with :928 reads as a typo otherwise.

Verification

• Read against master @ 35e2eff; all line numbers and quotes verified there.
• Not executed. I do not have AMD multi-node hardware, so I have not run the dispatch/combine path either before or after the change. Analysis is static.

Happy to send the one-line PR if you would like it, though given it is unreachable from the current test matrix you may prefer to fold it into other work.