NVIDIA SHARP: In-Network All-Reduce, NVLS and CollNet

The all-reduce cost problem

A data-parallel training step ends with an all-reduce: every rank holds a gradient buffer of size S and must end up with the element-wise sum of all n buffers. In NCCL's ring algorithm, GPU streaming multiprocessors (SMs) do that arithmetic, and the data crosses the fabric twice, once for reduce-scatter and once for all-gather. The nccl-tests performance notes put the ideal ring time at t = (S/B) * (2*(n-1)/n) for link bandwidth B, and define bus bandwidth as busbw = algbw * (2*(n-1)/n) so that results can be compared against hardware peak ^[1]. The factor 2(n-1)/n is the problem: each rank sends, and receives, almost twice its buffer.

Table 1. Ring all-reduce traffic per rank. Derived from the nccl-tests formula; S = 1 GB is an illustrative buffer size.

The ring also has 2(n-1) dependent steps, so per-hop latency dominates small messages and any slow rank stalls the rest, and its kernels occupy SMs that could run overlapping compute. NVIDIA's NCCL 2.27 post states that "traditional ring-based implementations can consume 16 or more SMs" ^[2]. SHARP addresses all three costs (bytes, steps and SMs) by moving the summation out of the GPUs and into the switches.

What SHARP does

The Scalable Hierarchical Aggregation and Reduction Protocol "improves the performance of MPI and Machine Learning collective operation, by offloading collective operations from CPUs and GPUs to the network and eliminating the need to send data multiple times between endpoints" ^[3]. The logic is "incorporated in the switch ASIC" and performs reductions "directly within the network fabric". According to NVIDIA, moving the collective into the switch fabric "reduces the amount of transferred data by half and minimizes jitter" ^[4].

Each endpoint sends its contribution once up an aggregation tree of switches; each switch sums its children's streams and forwards one partial result, and the root's result is replicated back down to every endpoint. Data volume shrinks "as aggregation nodes are reached" ^[3]. Each endpoint therefore sends and receives about S instead of about 2S, which is where the "half" comes from.

Two domains: NVLink SHARP and InfiniBand SHARP

SHARP exists in two places, which NCCL treats as different algorithms.

NVLink SHARP (NVLS) runs inside NVSwitch. Third-generation NVSwitch added "hardware acceleration for collective operations with multicast and NVIDIA SHARP in-network reductions". Embedded ALUs offer "up to 400 GFLOPS of FP32 throughput", and a SHARP controller "can manage up to 128 SHARP groups in parallel" ^[5]. NCCL documents NVLS as available "in third-generation NVSwitch systems (NVLink4) with Hopper and later GPU architectures" ^[6].

InfiniBand SHARP runs in Quantum-family InfiniBand switches. It reduces across nodes, and NCCL reaches it through the CollNet path and an external network plugin.

NCCL 2.18.1 combined the two domains. It added the inter-node algorithms "NVLink SHARP + IB SHARP (NVLS)" and "NVLink SHARP + Tree (NVLSTREE)" ^[7]. In the hierarchical case, NVSwitch reduces within the node, IB switches reduce across nodes, and the result is multicast back down both levels.

Worked example: NVLS inside one eight-GPU H100 node

NVIDIA's NVSwitch post counts the operations. A ring-style NCCL all-reduce costs "2N reads and 2N writes at each GPU interface, or 4N total operations", while with SHARP it costs "2N+2 operations, or approximately halving the number of read/write operations" ^[5]. In units of S/N per GPU, that is (N+1)/N × S in each direction for NVLS, compared with 2(N-1)/N × S for the ring.

Assumptions (illustrative): N = 8 H100 GPUs, S = 1 GB, and 450 GB/s per direction per GPU. The 450 GB/s is derived as half of the 900 GB/s NVLink 4 bandwidth per GPU ^[8] and matches NVIDIA's "450 GB/s of bandwidth for reduction operations" for an eight-H100 system ^[5].

1. Ring traffic per GPU = 2 × 7/8 × 1 GB = 1.75 GB. Time = 1.75 / 450 = 3.89 ms.
2. NVLS traffic per GPU = 9/8 × 1 GB = 1.125 GB. Time = 1.125 / 450 = 2.50 ms.
3. Ratio = 3.89 / 2.50 ≈ 1.56× (derived, wire-time only).
4. nccl-tests would report algbw = 1 GB / 2.50 ms = 400 GB/s for NVLS. It would then report busbw = 400 × 1.75 = 700 GB/s.

Because nccl-tests applies the ring factor whatever algorithm ran, an NVLS or CollNet run can show busbw above the per-direction link rate; that indicates in-network reduction, not a measurement error. All four figures are upper bounds that ignore protocol overhead and synchronization.

Worked example: IB SHARP across 128 nodes

Assumptions (illustrative): 128 nodes × 8 GPUs = 1,024 ranks, one 400 Gb/s NIC per GPU (8 rails per node), treated as 50 GB/s per NIC with encoding and protocol overhead ignored, and S = 1 GB.

1. Ring. Every ring channel enters and leaves each node once, and each link in a channel carries 2(n-1)/n of that channel's slice. Summed over channels, a node's network egress is 2 × 1023/1024 × 1 GB ≈ 1.998 GB, or ≈ 0.250 GB per NIC. Time ≈ 0.250 / 50 = 5.0 ms.
2. Hierarchical IB SHARP. An intra-node reduce-scatter over NVLink leaves each GPU holding a reduced 1/8 slice. Each NIC sends 0.125 GB up its rail's aggregation tree and receives 0.125 GB back. Time ≈ 0.125 / 50 = 2.5 ms, plus the intra-node phases on the faster NVLink domain.
3. The inter-node byte reduction is ≈ 2×, which matches NVIDIA's "by half" statement. The step count also changes: from 2(n-1) = 2,046 dependent ring steps to a round trip through a few switch tiers.

NCCL normally picks its own Tree algorithm for small messages, so SHARP's real latency advantage is smaller than the ring comparison suggests; see the guide to NCCL algorithms and protocols.

SHARP generations

Table 2. SHARP generations. Sources: NVIDIA SHARP blog, SHARP User Manual v3.0, Quantum-X800 datasheet.

SHARPv4 needs a status check because NVIDIA's wording changed. The October 2024 SHARP blog says SHARPv4 "will be released with the NVIDIA Quantum-X800 XDR InfiniBand switch platforms" ^[4]. The March 2024 launch release claimed "a 9x increase of 14.4Tflops of In-Network Computing" with SHARPv4 and said the platform would be available "next year" ^[9]. Later documents use the present tense:

• The NVIDIA datasheet (AUG25, hosted by Dell) says the fourth generation "adds support for FP8 precision and new collective operations such as ReduceScatter and ScatterGather" ^[10].
• The product page says the switch "includes hardware-based In-Network Computing with SHARP v4" ^[11].
• NVIDIA's validated-configuration matrix for XDR clusters lists Q3400-RA and Q3200-RA switches alongside SHARP software versions 3.10.3 through 3.16.3 ^[12].
• The SHARP software's supported-device list includes Quantum-3, with the footnote "Support limited to NVIDIA Quantum-3 based XDR Q34xx-RA InfiniBand Switch" ^[13].

As of 2026-10-28, Quantum-X800 is documented as a shipping, validated platform with SHARP software support. This guide found no NVIDIA document stating which SHARPv4 features (FP8 in-switch reduction, the new collectives) a given SHARP and NCCL release exposes. The newest public SHARP change log, 3.15 (June 2026), adds a beta parallel All-Gather/Reduce-Scatter over InfiniBand multicast and does not mention FP8 ^[14]. Confirm against the release notes of the exact build.

NCCL integration

NCCL reaches the two domains through different controls. The defaults below come from the NCCL 2.32.3 environment-variable reference (NCCL v2.32.3-1, released 2026-09-17 and still the latest release on 2026-10-28; NCCL environment variables).

Table 3. NCCL controls relevant to SHARP. Source: NCCL environment variables.

Three details about NCCL_NVLS_ENABLE matter in practice. In 2.27 and 2.28, the default value 2 tried to fall back silently when NVLS allocation failed. That "could result in a hang if the allocation failed on just some of the ranks", so 2.29 reverted to failing, as value 1 does. The same reference states that NVLS is not compatible with multiple ranks of one communicator sharing a GPU. In that case, value 1 fails initialization and the default 2 silently disables NVLS ^[6].

IB SHARP requires the NCCL-SHARP plugin. NVIDIA's SHARP manual says the plugin ships with HPC-X and loads automatically in HPC-X module environments. It can also be built from the nccl-rdma-sharp-plugins repository ^[15][16]. Streaming aggregation is supported on one communicator at a time, so NCCL_SHARP_DISABLE (per process group) and NCCL_SHARP_GROUP_SIZE_THRESH choose where it is used, and NCCL_IBEXT_DISABLE falls back to native transports. SHARP_COLL_LOCK_ON_COMM_INIT=1 and SHARP_COLL_NUM_COLL_GROUP_RESOURCE_ALLOC_THRESHOLD=0 reserve resources at initialization and are on by default from plugin 2.1.x.

Buffer registration determines how much SM work disappears. NVLS user-buffer registration has been available since 2.19. It requires ncclMemAlloc() or a qualified VMM allocator and identical buffer offsets across ranks. IB SHARP registration has been available since 2.21 and "can reduce the number of NCCL SM usage down to 1". It is currently limited to communicators with one rank per node, and network registration requires NCCL_PXN_DISABLE=1 because PXN is on by default ^[17]. NCCL 2.27 extended SHARP to AllGather and ReduceScatter "for both NVLink and IB fabrics". It noted that these collectives "are now preferred over AllReduce" for large LLM training ^[2].

What the published gains are

Table 4. Vendor-published SHARP results, with their context.

These are vendor numbers on vendor-chosen configurations, and two figures often quoted alongside them are not SHARP results. The NCCL 2.27 "up to 9x reduction in latency for small message sizes" and "up to 2.5x higher performance" on NVL8 come from the symmetric-memory low-latency kernels, not from in-network reduction ^[2]. Application-level gains such as the 17% BERT figure depend on how much of the step is exposed communication.

Precision caveats

In-network reduction moves the arithmetic to an ALU the framework does not control. For the NCCL 2.27 symmetric kernels, NVIDIA states that "reductions are computed using FP32 accumulators (or FP16 for FP8 on NVLink Switch (NVLS) Systems)" ^[2]. An FP8 all-reduce on NVLS therefore accumulates at lower precision. On the IB side, SHARP 3.14 added a sharp_am parameter, rounding_method, defaulting to 0 (round to nearest even), with 4 selecting LFSR stochastic rounding, for all SHARP jobs ^[18]. That setting is fabric-wide and owned by the fabric administrator, not the job. The NCCL 2.18.1 notes list a hang for bfloat16 reductions with IB SHARP "on systems with less than 1 NIC per GPU" ^[7]. A prudent operator compares loss curves, or bitwise outputs where determinism matters, with SHARP on and off before relying on it for low-precision collectives.

Requirements and deployment checklist

Table 5. Deployment checklist. Sources: SHARP General Information, Using SHARP with NCCL, Changes and New Features, SHARP blog, NCCL env, Buffer registration.

Rail grouping interacts with fabric design, so plan the SHARP tree layout together with the rail and oversubscription sizing rather than added afterwards.

When SHARP is not worth it

IB SHARP is an InfiniBand feature. The SHARP supported-device list names only Quantum-family switches ^[13], and NVIDIA markets Quantum InfiniBand as "the world's only fully offloadable, In-Network Computing platform" ^[19]. A RoCE or other Ethernet scale-out fabric therefore runs GPU-based algorithms between nodes, though NVLS can still help inside the NVLink domain. Whether that difference justifies InfiniBand is one input among many in the InfiniBand vs Ethernet for AI decision.

SHARP also adds little in several other situations:

• Single-node or two-node jobs. CollNet is not enabled below NCCL_COLLNET_NODE_THRESHOLD, which defaults to 2.
• Compute-bound steps. Where communication is already hidden behind backward compute, little is left to save.
• Multiple ranks per GPU. NVLS is disabled in this configuration.
• Multi-tenant clusters on SHARPv2 hardware. Only one job at a time gets aggregation.
• PXN dependence. Software stacks that rely on PXN cannot also use network buffer registration.

There, better compute and communication overlap and NCCL's own hierarchical algorithms are the levers.

Validation and troubleshooting

Run A/B comparisons on the same nodes and the same NCCL build:

1. Run all_reduce_perf from nccl-tests with NCCL_COLLNET_ENABLE=0 and then =1, and with NCCL_NVLS_ENABLE=0 and then the default. Sweep message sizes from kilobytes to gigabytes.
2. Set NCCL_DEBUG=INFO and NCCL_DEBUG_SUBSYS=INIT,NVLS,TUNING,REG to confirm which algorithm ran and whether registration succeeded. Add SHARP_COLL_LOG_LEVEL=3 for the plugin.
3. Read busbw with the arithmetic above in mind. A busbw above the per-direction link rate suggests an in-network path. A busbw identical to the baseline suggests that SHARP was never selected.
4. Measure step time and exposed communication in the real training job. Freed SMs only pay off if the framework uses them.

An NVLS initialization failure usually points to resource allocation, and partial-rank failures on 2.27 and 2.28 could hang instead of erroring; those cases call for the NCCL hang triage procedure rather than further tuning.

GPU Smith can review SHARP and NVLS enablement as part of a broader fabric assessment.

Sources

1. Advancing Performance with NVIDIA SHARP In-Network Computing - NVIDIA Technical Blog (Oct 2024), generations v1-v4, halving data and jitter, 17% BERT, Azure latency, MVAPICH2 results (accessed 2026-10-28)
2. Enabling Fast Inference and Resilient Training with NCCL 2.27 - NVIDIA Technical Blog (Jul 2025), SHARP for AllGather/ReduceScatter, SM figures, FP32/FP16 accumulators, symmetric-memory 9x/2.5x context (accessed 2026-10-28)
3. NVIDIA SHARP User Manual v3.0 - NVIDIA, SHARP definition, aggregation nodes, multiple aggregation trees in v3 (accessed 2026-10-28)
4. Using NVIDIA SHARP with NVIDIA NCCL - NVIDIA SHARP docs, plugin requirements, NCCL and SHARP variables, rail grouping (accessed 2026-10-28)
5. SHARP: General Information - NVIDIA SHARP docs, supported switches, HCAs, packages, Q34xx-RA footnote (accessed 2026-10-28)
6. SHARP: Changes and New Features - NVIDIA SHARP docs 3.14, rounding_method, stochastic rounding, DOCA-Host packaging change (accessed 2026-10-28); 3.15 edition (last updated 2026-06-03), parallel AG/RS beta: https://networking-docs.nvidia.com/sharpum/3150/changes-and-new-features
7. NCCL Environment Variables - NVIDIA NCCL docs, NVLS/CollNet/ALGO/DEBUG_SUBSYS defaults and semantics (accessed 2026-10-28)
8. NCCL User Buffer Registration - NVIDIA NCCL docs, NVLS and IB SHARP registration, PXN incompatibility, SM usage down to 1 (accessed 2026-10-28)
9. NCCL Release 2.18.1 - NVIDIA, NVLS/NVLSTree inter-node algorithms, bfloat16 IB SHARP known issue (accessed 2026-10-28)
10. nccl-tests PERFORMANCE.md - NVIDIA on GitHub, algbw/busbw definitions and 2(n-1)/n factor (accessed 2026-10-28)
11. Upgrading Multi-GPU Interconnectivity with the Third-Generation NVIDIA NVSwitch - NVIDIA Technical Blog (Aug 2022), NVSwitch SHARP ALUs, 128 groups, 2N+2 vs 4N operations, 450 GB/s reduction bandwidth (accessed 2026-10-28)
12. NVIDIA Hopper Architecture In-Depth - NVIDIA Technical Blog (Mar 2022), NVSwitch in-network reductions, up to 2x small-block claim (accessed 2026-10-28)
13. NVIDIA NVLink and NVLink Switch - NVIDIA, NVLink bandwidth per GPU by generation (accessed 2026-10-28)
14. NVIDIA Announces New Switches Optimized for Trillion-Parameter GPU Computing and AI Infrastructure - NVIDIA Newsroom (Mar 2024), SHARPv4 14.4 TFLOPS / 9x claim, availability "next year" (accessed 2026-10-28)
15. NVIDIA Quantum-X800 InfiniBand Platform - NVIDIA, SHARP v4 in-network computing, 144 x 800 Gb/s ports (accessed 2026-10-28)
16. NVIDIA Quantum-X800 InfiniBand Switches Datasheet - NVIDIA datasheet (AUG25) hosted by Dell, 4th-gen SHARP, FP8 and new collectives (accessed 2026-10-28)
17. NVIDIA Quantum-X800 (XDR) Clusters - NVIDIA validated configuration, Q3400-RA/Q3200-RA with SHARP 3.10.3-3.16.3 (accessed 2026-10-28)
18. NVIDIA InfiniBand Networking - NVIDIA, "fully offloadable, In-Network Computing platform" (accessed 2026-10-28)
19. nccl-rdma-sharp-plugins - NVIDIA/Mellanox on GitHub, NCCL SHARP plugin source and build requirements (accessed 2026-10-28)