Skip to content

fix: resolve the GIN context count on the single-node path (kNumQPs == 0 corrupts combine) #4

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
KeitaW wants to merge 1 commit into
base: main
Choose a base branch
from
+23 −3
Open

fix: resolve the GIN context count on the single-node path (kNumQPs == 0 corrupts combine) #4

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 2 additions & 3 deletions csrc/kernels/backend/nccl.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -126,16 +126,15 @@ NCCLSymmetricMemoryContext::NCCLSymmetricMemoryContext(const int64_t& nccl_comm,
if (scaleout_active) {
gin_config = elastic::gin_alloc::make_gin_resources(resolve_gin_context_cnt());
} else {
gin_config.gin_context_cnt = this->num_allocated_qps;
gin_config.gin_context_cnt = resolve_gin_context_cnt();
gin_config.gin_indexed_signals_cnt = 0;
}

EP_HOST_ASSERT(gin_config.gin_indexed_signals_cnt >= (num_rdma_ranks - 1) and
"GIN indexed-signal budget cannot give each peer rail team a dedicated "
"signal; reduce num_allocated_qps to raise the per-context signal count");

if (scaleout_active)
this->num_allocated_qps = gin_config.gin_context_cnt;
this->num_allocated_qps = gin_config.gin_context_cnt;

if (get_env<int>("EP_BUFFER_DEBUG"))
printf("GIN layout: gin_context_cnt=%d, gin_indexed_signals_cnt=%d, num_qp=%d\n",
Expand Down
5 changes: 5 additions & 0 deletions deep_ep/buffers/elastic.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -377,6 +377,11 @@ def __init__(self,
self.explicitly_destroy)

self.num_allocated_qps = self.runtime.get_num_allocated_qps()
# The unordered hybrid path intentionally passes 0 down and relies on C++ to
# resolve it (see the automatic QP count above). A 0 read back here means that
# delegation did not happen; it would surface far downstream as `kNumQPs == 0`.
assert self.num_allocated_qps >= 1, \
f'runtime returned num_allocated_qps={self.num_allocated_qps}: the QP count was never resolved'

# Logical rank indices
self.num_scaleout_ranks, self.num_scaleup_ranks = self.get_logical_domain_size()
Expand Down
6 changes: 6 additions & 0 deletions deep_ep/include/deep_ep/common/comm.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -77,6 +77,12 @@ __device__ __forceinline__ void timeout_while(const func_t& func, const int64_t&
template <int kNumSMs, int kNumQPs, int kNumChannelsPerSM, bool kWithNotifyWarps = false>
__device__ __forceinline__ std::pair<int, ncclGinResourceSharingMode> get_qp_mode(
const int& sm_idx, const int& channel_in_sm_idx, const bool& is_notify_warp = false) {
static_assert(kNumQPs >= 1,
"kNumQPs must be >= 1. With 0 the kNumQPs == 1 fast path is skipped, "
"kNumSMs <= kNumAvailableQPs is false, and balanced_partition() is called "
"with q == 0, i.e. n / 0. Compiled for the host this specialization exits on "
"SIGFPE; in a GPU run it produced corrupted output and no reported fault.");

constexpr auto kSharingCTA = NCCL_GIN_RESOURCE_SHARING_CTA;
constexpr auto kSharingGrid = kNumSMs == 1 ? NCCL_GIN_RESOURCE_SHARING_CTA : NCCL_GIN_RESOURCE_SHARING_GPU;

Expand Down
10 changes: 10 additions & 0 deletions deep_ep/include/deep_ep/common/qp_mapping.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -62,6 +62,11 @@ QP_MAPPING_HD constexpr QPSlot balanced_partition(int idx, int n, int q) {
template <int kNumSMs, int kNumQPs, int kNumChannelsPerSM, bool kWithNotifyWarps>
QP_MAPPING_HD constexpr int channel_to_qp(int sm_idx, int channel_in_sm_idx,
bool is_notify_warp = false) {
static_assert(kNumQPs >= 1,
"kNumQPs must be >= 1. With 0 the kNumQPs == 1 fast path is skipped, "
"kNumSMs <= kNumAvailableQPs is false, and balanced_partition() is called "
"with q == 0, i.e. n / 0. Compiled for the host this specialization exits on "
"SIGFPE; in a GPU run it produced corrupted output and no reported fault.");
// Only one QP
if constexpr (kNumQPs == 1)
return 0;
Expand Down Expand Up @@ -95,6 +100,11 @@ QP_MAPPING_HD constexpr int channel_to_qp(int sm_idx, int channel_in_sm_idx,
// `signal_id < ceil(channels / qps)` -- within the tuner's signal budget.
template <int kNumSMs, int kNumQPs, int kNumChannelsPerSM, bool kWithNotifyWarps>
QP_MAPPING_HD constexpr int channel_to_signal_id(int sm_idx, int channel_in_sm_idx) {
static_assert(kNumQPs >= 1,
"kNumQPs must be >= 1. With 0 the kNumQPs == 1 fast path is skipped, "
"kNumSMs <= kNumAvailableQPs is false, and balanced_partition() is called "
"with q == 0, i.e. n / 0. Compiled for the host this specialization exits on "
"SIGFPE; in a GPU run it produced corrupted output and no reported fault.");
if constexpr (kNumQPs == 1)
return sm_idx * kNumChannelsPerSM + channel_in_sm_idx;

Expand Down