[GH-ISSUE #647] [hybrid_dispatch] Question on release scope of local-rank Rail bypass path #381

New issue

Closed

opened 2026-10-26 09:59:47 +00:00 by bsakaguchi · 2 comments

bsakaguchi commented

2026-10-26 09:59:47 +00:00

Owner

Originally created by @huakainomore on GitHub (May 26, 2026).
Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/647

Summary

hybrid_dispatch.cuh calls gin.red_add_rel<ncclTeamTagRail> to update the scaleout_channel_signaled_tail slot. On the local-rank bypass path this dispatches to ptx::red_add_rel_sys in handle.cuh, emitting red.release.sys.global.add.u64.
After tracing through get_sym_ptr<ncclTeamTagRail>, we are wondering whether .gpu might also be sufficient here, or whether there is a specific reason .sys is required that we are missing?

Question about the existing comment

the commentNOTES: the "release" scope will be sys for the local rank (we may involve NVLink so not gpu) mentions NVLink, but the only NVLink-bearing path goes through ncclTeamTagLsa/ncclTeamTagWorld, not ncclTeamTagRail. Is there a code path we missed where Rail can land on an NVLink peer?

Proposed change

Add a red_add_rel_gpu overload (or a scope tag) to NCCLGin so the call site can opt into .gpu release:

  // handle.cuh — new overload alongside existing red_add_rel
  template <typename team_t, typename dtype_t>
  __device__ __forceinline__
  void red_add_rel_gpu(dtype_t* sym_ptr, const dtype_t& value,
                       const int& dst_rank_idx,
                       const int& extra_options = 0) const {
      const auto dst_ptr = get_sym_ptr<team_t>(sym_ptr, dst_rank_idx);
      if (dst_ptr != nullptr) {
          ptx::red_add_rel_gpu(dst_ptr, value);   // <-- .gpu instead of
  .sys
      } else {
          // RDMA path unchanged: atomic provides system-scope ordering
          gin.signal(...);
      }
  }

The same change applies to hybrid_combine.cuh by the same reasoning.

Questions

Is there a non-local consumer of scaleout_channel_signaled_tail_ptr we might have missed ?
Is the NVLink mention in the comment a leftover from an earlier design, or does Rail ever land on an NVLink peer in some configuration?

update

~30% improvement when ep16 with 12 SMs
5% improvement when ep16 with 16 SMs

Originally created by @huakainomore on GitHub (May 26, 2026). Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/647 ## Summary `hybrid_dispatch.cuh` calls `gin.red_add_rel<ncclTeamTagRail>` to update the `scaleout_channel_signaled_tail` slot. On the local-rank bypass path this dispatches to `ptx::red_add_rel_sys` in `handle.cuh`, emitting `red.release.sys.global.add.u64`. After tracing through `get_sym_ptr<ncclTeamTagRail>`, we are wondering whether `.gpu` might also be sufficient here, or whether there is a specific reason `.sys` is required that we are missing? ## Question about the existing comment the comment`NOTES: the "release" scope will be `sys` for the local rank (we may involve NVLink so not gpu)` mentions NVLink, but the only NVLink-bearing path goes through ncclTeamTagLsa/ncclTeamTagWorld, not ncclTeamTagRail. Is there a code path we missed where Rail can land on an NVLink peer? ## Proposed change Add a red_add_rel_gpu overload (or a scope tag) to NCCLGin so the call site can opt into .gpu release: ``` // handle.cuh — new overload alongside existing red_add_rel template <typename team_t, typename dtype_t> __device__ __forceinline__ void red_add_rel_gpu(dtype_t* sym_ptr, const dtype_t& value, const int& dst_rank_idx, const int& extra_options = 0) const { const auto dst_ptr = get_sym_ptr<team_t>(sym_ptr, dst_rank_idx); if (dst_ptr != nullptr) { ptx::red_add_rel_gpu(dst_ptr, value); // <-- .gpu instead of .sys } else { // RDMA path unchanged: atomic provides system-scope ordering gin.signal(...); } } ``` The same change applies to hybrid_combine.cuh by the same reasoning. ## Questions 1. Is there a non-local consumer of scaleout_channel_signaled_tail_ptr we might have missed ? 2. Is the NVLink mention in the comment a leftover from an earlier design, or does Rail ever land on an NVLink peer in some configuration? ### update 1. ~30% improvement when ep16 with 12 SMs 2. 5% improvement when ep16 with 16 SMs

bsakaguchi closed this issue

2026-10-26 09:59:47 +00:00

bsakaguchi commented

2026-10-26 09:59:48 +00:00

Author

Owner

@gangxie112 commented on GitHub (May 26, 2026):

@sphish

@gangxie112 commented on GitHub (May 26, 2026): @sphish

bsakaguchi commented

2026-10-26 09:59:48 +00:00

Author

Owner

@polarstormx commented on GitHub (Sep 18, 2026):

This has been addressed in https://github.com/deepseek-ai/DeepEP/pull/674 (commit github.com/deepseek-ai/DeepEP@099d5f2bad). Closing as resolved.

@polarstormx commented on GitHub (Sep 18, 2026): This has been addressed in https://github.com/deepseek-ai/DeepEP/pull/674 (commit https://github.com/deepseek-ai/DeepEP/commit/099d5f2bad488b9c534ea785062b12f2e91d1d41). Closing as resolved.

bsakaguchi referenced this issue

2026-10-27 21:03:50 +00:00

[PR #381] [MERGED] Fix duplicate calculation of num_tokens_per_rank #544