deepseek-ai/DeepEP
1
0
Fork 0
mirror of https://github.com/deepseek-ai/DeepEP.git synced 2026-11-05 00:47:32 +00:00
Code Issues 308 Projects Packages Wiki Activity

[PR #478] Add barrier before nvshmem_team_split_strided to ensure proper initialization #593

New issue
Open
opened 2026-10-27 21:04:21 +00:00 by bsakaguchi · 0 comments
bsakaguchi commented 2026-10-27 21:04:21 +00:00
Owner
Copy link

📋 Pull Request Information

Original PR: https://github.com/deepseek-ai/DeepEP/pull/478
Author: @pierDipi
Created: 11/3/2025
Status: 🔄 Open

Base: main ← Head: barrier-low-latency

📝 Commits (1)

  • ce054ae Add barrier before nvshmem_team_split_strided to ensure proper initialization

📊 Changes

1 file changed (+3 additions, -1 deletions)

View changed files

📝 csrc/kernels/runtime.cu (+3 -1)

📄 Description

I'm getting

/tmp/nvshmem_src/src/host/team/team.cu:nvshmem_team_split_strided:63: NVSHMEM API called before NVSHMEM initialization has completed

Move nvshmem_barrier_all() to execute before nvshmem_team_split_strided rather than after. This is required because team split is a collective operation that must be called by all PEs in the parent team, and all PEs must reach this call in a synchronized manner.

Without the barrier after nvshmemx_init_attr(), ranks may complete initialization at different times, leading to race conditions where some PEs attempt to split teams before others have finished NVSHMEM initialization. This can cause undefined behavior and incorrect team formation.

The barrier after team split was unnecessary per NVSHMEM documentation: teams are immediately usable after creation without intervening synchronization.

🔄 This issue represents a GitHub Pull Request. It cannot be merged through Gitea due to API limitations.

## 📋 Pull Request Information **Original PR:** https://github.com/deepseek-ai/DeepEP/pull/478 **Author:** [@pierDipi](https://github.com/pierDipi) **Created:** 11/3/2025 **Status:** 🔄 Open **Base:** `main` ← **Head:** `barrier-low-latency` --- ### 📝 Commits (1) - [`ce054ae`](https://github.com/deepseek-ai/DeepEP/commit/ce054aeaa21b0b52d67a749ab4b8073eb316e093) Add barrier before nvshmem_team_split_strided to ensure proper initialization ### 📊 Changes **1 file changed** (+3 additions, -1 deletions) <details> <summary>View changed files</summary> 📝 `csrc/kernels/runtime.cu` (+3 -1) </details> ### 📄 Description I'm getting ``` /tmp/nvshmem_src/src/host/team/team.cu:nvshmem_team_split_strided:63: NVSHMEM API called before NVSHMEM initialization has completed ``` Move `nvshmem_barrier_all()` to execute before `nvshmem_team_split_strided` rather than after. This is required because team split is a collective operation that must be called by all PEs in the parent team, and all PEs must reach this call in a synchronized manner. Without the barrier after nvshmemx_init_attr(), ranks may complete initialization at different times, leading to race conditions where some PEs attempt to split teams before others have finished NVSHMEM initialization. This can cause undefined behavior and incorrect team formation. The barrier after team split was unnecessary per NVSHMEM documentation: teams are immediately usable after creation without intervening synchronization. --- <sub>🔄 This issue represents a GitHub Pull Request. It cannot be merged through Gitea due to API limitations.</sub>
Sign in to join this conversation.
No Branch/Tag specified
Branches Tags
main
hybrid-ep
chore/remove-clang-format-yapf
v2.5
nvDev
try-fix-cu128
mori-ep
antgroup-opt
compat-old-torch
add-linter
eager-rdma
v1.2.1
Labels
Clear labels   
pull-request
Mirrored from GitHub Pull Request

No labels
pull-request
Milestone
Clear milestone
No items
No milestone
Projects
Clear projects
No items
No project
Assignees
Clear assignees
No assignees
1 participant
Notifications
Due date
The due date is invalid or out of range. Please use the format "yyyy-mm-dd".

No due date set.

Dependencies

No dependencies set.

Reference
deepseek-ai/DeepEP#593
No description provided.