feat(m2n): quantized tensor support (scale forwarding + on-the-fly FP8) (#7) - #11
Open
kwen2501 wants to merge 434 commits into
Open
feat(m2n): quantized tensor support (scale forwarding + on-the-fly FP8) (#7)#11kwen2501 wants to merge 434 commits into
kwen2501 wants to merge 434 commits into
Conversation
Replace the strict sm < 90 error in both the Makefile and CMakeLists.txt with a per-CUDA-version selection trimmed to Hopper and newer, mirroring the patterns in makefiles/common.mk and the root CMakeLists.txt: - Makefile: guard the override with $(origin NVCC_GENCODE)=file so a user-supplied env/command-line value passes through unchanged while common.mk's ?= default is replaced with the trimmed sm_90+ list. - CMakeLists.txt: guard with NOT DEFINED CMAKE_CUDA_ARCHITECTURES so -D overrides pass through, and apply CMAKE_CUDA_ARCHITECTURES on the target the same way src/CMakeLists.txt does. When the user supplies an override that contains an arch below sm_90, let nvcc surface the error rather than pre-validating. Signed-off-by: Xiakun Lu <xiakunl@nvidia.com>
If user provides NCCL windows for the output token tensor and NVLink-only configuration is detected, write tokens directly into the output locations over NVLink, skipping the receive-stage copy. Only token headers are written into the staging buffer for token source information tracking. Signed-off-by: Artem Y. Polyakov <artemp@nvidia.com>
Take advantage of the z-copy path in LL/Dispatch Signed-off-by: Artem Y. Polyakov <artemp@nvidia.com>
## Summary Expanding testing matrix for NCCL EP CI, add testing for different layouts and batch size ## Implementation **ep_test** | Invocation | Algorithm | Layout | Tokens | Hidden | |-------------|-----------|----------------|--------|--------| | -a ll | LL | EM (hardcoded) | 128 | 7168 | | -a ht -L fl | HT | FLAT | 4096 | 7168 | | -a ht -L em | HT | EM | 4096 | 7168 | **ep_bench** Low-latency: | Layout | 128 | 256 | 1K | 4K | 8K | |--------------|-----|-----|----|----|----| | expert-major | ✓ | ✓ | ✓ | ✓ | X | | rank-major | ✓ | ✓ | ✓ | ✓ | X | High-throughput | Layout | 128 | 256 | 1K | 4K | 8K | |--------------|-----|-----|----|----|----| | flat | ✓ | ✓ | ✓ | ✓ | X | | expert-major | ✓ | ✓ | ✓ | ✓ | X | 8K batch size was disabled for now due to the failure on EOS: https://gitlab-master.nvidia.com/nccl/nccl/-/jobs/328613355 ## Testing CI ## Release notes (Optional.) This serves as a base when compiling the release notes. If you feel this should be mentioned then you can add a tentative text here. Otherwise indicate it should not be mentioned. ## Tracking https://nvbugspro.nvidia.com/bug/6169264 add coverages Signed-off-by: Chris Cao <chricao@nvidia.com> See merge request nccl/nccl!2789 Signed-off-by: Chris Cao (SW-GPU) <chricao@nvidia.com>
Stage libnccl_ep.so plus the headers its JIT runtime needs (nccl_ep.h,
ep_enums.h, common.hpp, device/*.cuh) into nccl/ep/{lib,include}/ so the
nccl4py wheel ships them as package data. A new build_assets/build_native.sh
builds libnccl.so (link-time dep, not shipped) and libnccl_ep.so under a
sterile env, then copies the artifacts into the package tree. For production
wheels, build_wheels.sh runs the script inside a manylinux_2_28 container
selected by uname -m, so the .so has glibc-2.28 symbol versions.
auditwheel repair excludes libnccl.so.2 (from nvidia-nccl-cu13) and
libcuda.so.1 (driver, not in container). nccl/ep/__init__.py points
NCCL_EP_JIT_SOURCE_DIR / NCCL_EP_JIT_BUILD_INCLUDE_DIR at the bundled
headers before any binding load, so HT-mode JIT finds them post-install.
contrib/nccl_ep no longer bakes a local rpath into libnccl_ep.so: the
Makefile drops -Xlinker -rpath -Xlinker \$(NCCL_LIBDIR), and CMakeLists
sets BUILD_WITH_INSTALL_RPATH=TRUE with INSTALL_RPATH="". Matches
libnccl.so's no-RUNPATH convention; the cu13 wheel's pre-dlopen of
libnccl.so.2 with RTLD_GLOBAL is what satisfies the NEEDED at runtime.
Signed-off-by: Xiakun Lu <xiakunl@nvidia.com>
Fix mis-allocation of the rdma_inter_node_group_flags. This aray is used in both Dispatch and Combine in HT mode. The indexing is `(node_id * max_num_of_chunks_per_rank + chunk_idx)` However, at buffer allocation the size of this array is calculated as flags_sz = align_size(rdma_team_size * sizeof(uint64_t), GIN_ALIGNMENT); Thanks to GIN_ALIGNMENT, the actual size is no less than 4KB At max batch size of 8192, chunk=128 → max_chunks=64 The 4KB allocation It works up to 8 nodes: 7 (nodes) × 64 (max chunk) × 8 = 3584 B However, for 9 nodes it will be overflown: 9 × 64 × 8 = 4608 B Signed-off-by: Artem Y. Polyakov <artemp@nvidia.com>
Signed-off-by: Bharath Ramesh <bhramesh@nvidia.com>
This commit fixes a problem when number of nodes exceed the number of G2S stages (NUM_OF_STAGES_G2S). The RDMA portion of multi-node path in inter_node_G2S_warp_group_device_function was not accounting for exceeding the NUM_OF_STAGES_G2S. Up to 4 nodes this would not happen, but starting from 5 nodes and topK of 5 it is possible to exceed the number of stages and thus data should be loaded in batches of NUM_OF_STAGES_G2S. Thank you to Yi Liu (@yliu.16) from ByteDance to reporting and investigating this fix. Signed-off-by: Artem Y. Polyakov <artemp@nvidia.com>
Currently, the Combine send buffer size is calculated based on the worst-case expert-major requirements. For 16 rank, 512 experts, 4K batch size and 16KB tokens this results in 512 x 4K x 16 KB = 35 GB buffer. Rank-majer, however, only depends on number of ranks thus 16 x 4K x 16KB = 1GB Signed-off-by: Artem Y. Polyakov <artemp@nvidia.com>
As we are getting a lot of concerns about memory consumption, extending ep_bench to report the consumption Signed-off-by: Artem Y. Polyakov <artemp@nvidia.com>
Currently, the layout size is calculated by separately picking max of the send and recv buffers for dispatch/combine. This is not optimal, as dispatch has larger receive and combine has larger send buffer. Optimize the calculation by selecting the max of the send+recv buffers sizes. Signed-off-by: Artem Y. Polyakov <artemp@nvidia.com>
1. Honor User-provided rdma buffer size and allocate it unconditionally Fail in the handle init if the buffer space is insufficient 2. Optimize the AUTO scenario. Allocating inside the EP Group creation is unoptimal as topK is not known and we have to use MAX=32. This significantly impacts the combine receive buffer size. Delay actual allocation to the first Handle init. Signed-off-by: Artem Y. Polyakov <artemp@nvidia.com>
Update documentation to reflect the consequences of lazy staging buffer allocation for LL mode. For `rdma_buffer_size == NCCL_EP_AUTO`, Handle init operation may become collective and reallocate staging buffer (rendering undefined behavior of any in-flight communication). Signed-off-by: Artem Y. Polyakov <artemp@nvidia.com>
Decouple the LL/active mask synchronization buffer from the staging buffer. The staging buffer is lazily allocated and may change it's shape. Active mask management code currently assumed that layout is allocated for the worst-case which would cause buffer overflow for other cases. Signed-off-by: Artem Y. Polyakov <artemp@nvidia.com>
…copy Use epGetTensorData to resolve dispatch output pointers in HT validation paths. tensor->data is null for window-backed tensors, causing validation to fail with a null deref. Signed-off-by: Subhadeep Bhattacharya <subhadeepb@nvidia.com>
Templatize the LL dispatch/combine kernels and host wrappers on TopkIdxT (int32_t or int64_t). The C-side API gates in ncclEpUpdateHandle now accept either dtype for LL; HT keeps strict ncclInt64. Callers select the kernel specialisation through the cached handle->topk_idx.datatype at dispatch / combine time. ep_bench gains a --topk-idx-int32 flag for validating the new specialisation. Caller is responsible for ensuring expert ids fit in the chosen integer width. Eventual landing target: separate MR. Signed-off-by: Ludwig Schneider <lschneider@nvidia.com>
Signed-off-by: Ludwig Schneider <lschneider@nvidia.com>
Signed-off-by: Ludwig Schneider <lschneider@nvidia.com>
Signed-off-by: Ludwig Schneider <lschneider@nvidia.com>
Signed-off-by: Phuong Nguyen <phuonguyen@nvidia.com>
Signed-off-by: Phuong Nguyen <phuonguyen@nvidia.com>
Catches the bug introduced by 7732e3d ("HT handle memory optimization"), which dropped the per-iteration memset of global_routing_map. With the memset gone, the convert kernel writes only rows 0..num_tokens-1, but ncclAllGather ships the full max_tokens rows — stale tail rows from a prior iteration get sent to peers and interpreted as live routing. The test drives ncclEpUpdateHandle twice on the same handle: iter 1: num_tokens=max_tokens=4, routes every expert across 4 ranks (max-dirties the routing map) iter 2: num_tokens=1, single token routed to expert 0 (convert writes row 0; rows 1..3 retain iter-1 bits) After iter-2's update, the test asserts per-rank ncclEpHandle_test_getNumRecvTokens against the with-fix expected {rank0: 4, rank1: 0, rank2: 0, rank3: 0}. Without the fix the counts are deterministically {6, 4, 2, 4} - every rank fails its EXPECT_EQ. The detector is independent of any CUDA OOB / page fault, so it works at the small unit-test fixture sizes (max_tokens=4, kHidden=16, kMaxRecvSlots=16) where the original illegal-address signature was absent. The test fails on this commit, demonstrating the bug. The follow-up commit restores the cudaMemsetAsync fix and the test passes. Signed-off-by: Hsin-Chun Yin <hsinchuny@nvidia.com>
Alternative fix to 0632add70 ("nccl-ep: zero local routing-map slot in
HT update_handle"): instead of a separate cudaMemsetAsync on the local
routing send slot, extend convert_topk_to_routing_map_kernel to launch
over max_tokens rows and zero the tail rows [num_tokens, max_tokens)
from the convert kernel itself. The downstream ncclAllGather over the
full max_tokens stride then ships clean tail bytes to peers.
Trade-offs vs the memset fix:
* Saves one kernel launch (~5-10 us) per HT dispatch.
* Costs nothing in the uniform num=max case -- the grid is already
sized to max_tokens and tail threads do zero work.
* In non-uniform / num<max cases, the extra work is (max_tokens -
num_tokens) rows of num_experts_packed bytes at HBM speed (sub-us
for the vLLM Qwen3-MoE config: max_tokens=8192, packed=16).
* Restores the kernel's pre-existing "the caller does not pre-zero"
invariant -- ownership of the routing map's zero-state stays with
the kernel that writes it.
Mechanical changes:
* hybridep_adapter.cu: add max_tokens parameter to kernel and host
wrapper, change early-return bound to max_tokens, add an early
return for tail threads after row zeroing, and resize the launch
grid to ceil(max_tokens / block_size).
* hybridep_adapter.cuh: thread max_tokens through the declaration.
* nccl_ep.cc: pass max_tokens at the ncclEpUpdateHandle call site
(max_tokens is already in scope as group->config.max_dispatch_tokens_per_rank).
Validated by the same test_ht_stale_routing_map regression test added
in 99bace8cd -- the test asserts the iter-2 per-rank recv counts and
will reproduce the stale-tail bug if the tail zeroing is removed.
Signed-off-by: Hsin-Chun Yin <hsinchuny@nvidia.com>
Asymmetric per-rank token counts (e.g. [32, 0, 16, 0] across an
ep_size=4 group) aborted any rank whose input was empty. Three
over-strict checks combined to make a legitimate zero-token
configuration unrepresentable:
1. tensorHasBinding rejected descriptors with data == nullptr, but
any zero-extent CUDA tensor has no element to address. The
binding check is the right contract for non-empty buffers; for
empty ones it should be waived.
2. ncclEpDispatch and ncclEpCombine had a lazy "if num_tokens == 0,
re-derive from inputs->tokens" branch that conflated "not yet
set via ncclEpUpdateHandle" with a legitimate zero, then fell
back to tensor_required(inputs->tokens) and tripped (1).
3. resolveTensorWindowBinding called ncclWinGetUserPtr on a NULL
win_hdl for empty tensors with no external window.
Fixes:
- tensorIsEmpty() (any sizes[i] == 0); tensorHasBinding() returns
true for empty tensors.
- bool num_tokens_set on ncclEpHandle, set in ncclEpUpdateHandle;
the two lazy branches now key off !num_tokens_set.
- resolveTensorWindowBinding short-circuits on empty tensors.
The dispatch / combine kernels are already empty-input safe: sender
loops in device/low_latency.cu degrade to no-ops and grid sizing is
independent of numTokens.
Signed-off-by: Ludwig Schneider <lschneider@nvidia.com>
The non-uniform sampler skipped 0 (sampled in [1, max_tokens] and
hardcoded the last rank to 1), so the asymmetric zero-token path
in libnccl_ep.so was never exercised at the C level. Three
bench-side issues prevented it from running cleanly:
1. generateRandomTopkIndicesLL computed a uniform distribution over
[0, num_tokens - 1]. At num_tokens = 0 the unsigned subtraction
underflowed to UINT_MAX and the masking write indexed far past
the empty topk_idx_host buffer.
2. epGetTensorData returned ncclInvalidUsage when the descriptor's
data pointer was NULL (cudaMalloc(0) yields nullptr) and the
tensor was not in the optional alloc-state map. Empty tensors
legitimately have nothing to address.
3. The "Non-uniform tokens in [1, %u]" startup banner advertised a
range the sampler no longer uses.
Fixes:
- Guard the LL masking loop on num_tokens > 0.
- epGetTensorData hands back data=nullptr/success when any
sizes[i] == 0; consumers must pass count=0 to CUDA APIs (the
bench already does).
- Sample non-uniform counts in [0, max_tokens] and pin the last
rank to 0 so the regression case is always exercised.
- Update the banner to "[0, %u] (last rank forced to 0)".
Signed-off-by: Ludwig Schneider <lschneider@nvidia.com>
As of v0.1, NCCL EP does not support quantization. Remove remainings of it from ep_bench to avoid confusion. The support is planned to be added and expected in future releases Signed-off-by: Artem Y. Polyakov <artemp@nvidia.com>
Signed-off-by: Phuong Nguyen <phuonguyen@nvidia.com>
The zero-copy mode enum (dc26057) introduced an opt-in for skipping library-owned staging, but on the LL dispatch path the new mode was only consumed as an implicit branch on recv_x->win_hdl. When the caller requested ON but did not satisfy the opportunistic conditions (nvlink-only topology, rank-major layout, !use_fp8, windowed recv_x), the library silently fell back to a state where the kernel either staged through buffers it expected to be elided, or addressed windows that were not registered. The result was undefined writes rather than a clean diagnostic. Add explicit gating on LL dispatch: zero_copy=ON errors with ncclInvalidArgument if any of nvlink_only, rank-major layout, !use_fp8_outer, or recv_x->win_hdl set is violated. AUTO/OFF stay opportunistic on HT (master parity) and on LL dispatch (matches the dc26057 opportunistic branch); no warnings emitted. LL combine is out of scope. The combine kernel reads inputs->tokens directly via pointer and uses windows only to translate peer recv-buffer pointers in ncclGetP2pPtr; the kernel-level zeroCopy flag is dispatch-shaped (sender-side P2P writes) and plumbing config.zero_copy through it on combine silently picks the wrong send path. Revert the LL combine plumbing to a hardcoded false (matches origin/master) and treat config.zero_copy as a dispatch-only switch. Signed-off-by: Ludwig Schneider <lschneider@nvidia.com>
Align the README with the implementation before v0.1 release Signed-off-by: Artem Y. Polyakov <artemp@nvidia.com>
Avoid outdated info in README.md due to dynamic nature of the test. Signed-off-by: Artem Y. Polyakov <artemp@nvidia.com>
Reshard collectives and call-time resource allocation are not safe during CUDA graph capture, which could otherwise fail late or replay stale state. The public APIs now reject capture immediately with an actionable invalid-usage error. Signed-off-by: Ching-Hsiang Chu <chichu@nvidia.com>
Staging allocations could grow without a fixed bound and reservations could remain occupied after failures. Fixed-size buckets now bound per-communicator staging memory, while failed or abandoned reservations are released safely. Signed-off-by: Ching-Hsiang Chu <chichu@nvidia.com>
Reshard staging now supports a default direct path and an opt-in packed-window path selected at runtime. Unsupported selections are rejected clearly instead of being silently ignored. Signed-off-by: Ching-Hsiang Chu <chichu@nvidia.com>
Partial-group reshards now use communicators containing only participating ranks. Adaptive selection reduces rank-varying inputs so all ranks agree on whether to split and on the required resources. Signed-off-by: Ching-Hsiang Chu <chichu@nvidia.com>
M2N calls can now be batched in groups and submitted together at the outermost group end, while an abort discards an incomplete group. Group execution reports which recorded call failed. Signed-off-by: Ching-Hsiang Chu <chichu@nvidia.com>
Correctness-only tests could not detect staged-group fusion regressing to one-at-a-time execution. Testing instrumentation now counts fused submissions so grouped tests can verify fusion without affecting production builds. Signed-off-by: Ching-Hsiang Chu <chichu@nvidia.com>
CI verified only that the library and header existed, so missing public symbols could go unnoticed. The export check now derives public API declarations from the header and verifies that each is defined and exported by the shared library. Signed-off-by: Ching-Hsiang Chu <chichu@nvidia.com>
Signed-off-by: Ke Wen <kwen@nvidia.com>
Signed-off-by: Ching-Hsiang Chu <chichu@nvidia.com>
Signed-off-by: Ke Wen <kwen@nvidia.com>
The default copy algorithm was DIRECT here and PACKWINDOW upstream, so an unconfigured caller got a different path than intended. Both the inline default and the runtime config reset now select PACKWINDOW. Signed-off-by: Ching-Hsiang Chu <chichu@nvidia.com>
Chunk granularity, GIN context count and the cross-dim transpose threshold were fixed at their compile-time constants, so the environment variables documented for them had no effect. Each now reads a runtime value that falls back to the same constant. Co-authored-by: Kaushik Kandadi Suresh <kkandadisure@nvidia.com> Co-authored-by: Sreeram Potluri <spotluri@nvidia.com> Signed-off-by: Ching-Hsiang Chu <chichu@nvidia.com>
An early return between reserving the transpose buffer and recording its completion event left the reservation held, so later calls on that communicator failed with a spurious concurrency error. The reservation is now released by a scope guard on both the success and error paths. Co-authored-by: Kaushik Kandadi Suresh <kkandadisure@nvidia.com> Co-authored-by: Sreeram Potluri <spotluri@nvidia.com> Signed-off-by: Ching-Hsiang Chu <chichu@nvidia.com>
Co-authored-by: Kaushik Kandadi Suresh <kkandadisure@nvidia.com> Co-authored-by: Sreeram Potluri <spotluri@nvidia.com> Signed-off-by: Ching-Hsiang Chu <chichu@nvidia.com>
The shipped matrix was missing regression groups that cover cache churn across fresh streams and two pack-window paths. The cases and the MPI stream-churn driver they need are restored. Co-authored-by: Kaushik Kandadi Suresh <kkandadisure@nvidia.com> Co-authored-by: Sreeram Potluri <spotluri@nvidia.com> Signed-off-by: Ching-Hsiang Chu <chichu@nvidia.com>
Signed-off-by: Ke Wen <kwen@nvidia.com>
… first failure Signed-off-by: Ke Wen <kwen@nvidia.com>
Signed-off-by: Ke Wen <kwen@nvidia.com>
Signed-off-by: Ke Wen <kwen@nvidia.com>
Signed-off-by: Ke Wen <kwen@nvidia.com>
Signed-off-by: Ke Wen <kwen@nvidia.com>
Signed-off-by: Ke Wen <kwen@nvidia.com>
Signed-off-by: Ching-Hsiang Chu <chichu@nvidia.com>
Signed-off-by: Ching-Hsiang Chu <chichu@nvidia.com>
`NCCL_RESHARD_NUM_CTAS` can select an invalid GIN context when the CTA count is not divisible by the context count. The change affects CTA-to-context assignment for reshard kernels. No other behavior changes are intended. Signed-off-by: Ching-Hsiang Chu <chichu@nvidia.com>
Signed-off-by: Ching-Hsiang Chu <chichu@nvidia.com>
Signed-off-by: Ke Wen <kwen@nvidia.com>
Signed-off-by: Ke Wen <kwen@nvidia.com>
Signed-off-by: Ching-Hsiang Chu <chichu@nvidia.com>
Port scale forwarding and on-the-fly FP8 quantization from the nccl-rl prototype onto nccl-extensions.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
14 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
nccl-rl(kwen/m2n-scale-forwarding) ontonccl-extensions.ncclReshardScaled/ncclReshardScaledWithWindowfor coupled (payload, scales) reshard with block-alignment validation.ncclReshardQuantizedfor on-the-fly FP8 (E4M3 / MXFP8) wire compression with keep-quantized or dequantize modes.ScalePlane,QuantSpec) and private CI unit tests for scale-plane validation + quantize API/numerics.Test plan
make -C nccl_m2n libwith an existing NCCL 2.30.x buildunit_tests --gtest_filter=ScalePlane*→ 20/20 passed (host)unit_tests --gtest_filter=Quantize*→ 10 API tests passed; 10 GPU kernel tests skipped (no device on login host)sync_nccl_ext_issue7.sh)