Architecture & Measurement-Design Analysis

Demystifying NVSHMEM: A System-Level Analysis on Symmetric Memory and Device-Initiated Operations in GPU Communication

Source: Ma, Y.; Shen, S.; Chen, T.; Langer, A.; Kraus, J.; Glick, B.; Belusar, C.; Hammond, J.; Hoefler, T. arXiv:2606.05951v1 [cs.DC], 4 Jun 2026. Affiliations: ETH Zürich (Ma, Shen, Chen, Hoefler) + NVIDIA Corporation (Langer, Kraus, Glick, Belusar, Hammond). Analyzed version: NVSHMEM 3.3.9 (public source release). Reader: Direct PDF read via Claude Read tool (gemini-reader quota exhausted: RESOURCE_EXHAUSTED on gemini-2.0-flash; codex-reader unavailable on account). 12-page PDF read in a single pages=1-12 request. Analyst: Vishwakarma Date: 2026-06-17

Table of Contents

1. System Architecture (the "specimen": NVSHMEM as a layered runtime)
2. System-Under-Test Architecture (the H200 testbed + the dual-interface API surface)
3. Design-Space Diagram (axes swept in the microbenchmarks, axes held fixed)
4. Algorithm / Control-Flow Diagrams (symmetric-heap setup, fast vs slow RMA path, collective dispatch, DeepEP HT pipeline)
5. Quantitative Results — Empirical Findings by Regime
6. Configuration-Regime Trade-off Tables
7. Bottlenecks & Insights Surfaced by the Measurements
8. Limitations of the Methodology
9. Note on NCCL Tuning
10. Analogy

1. System Architecture (the "specimen")

NVSHMEM is not a benchmark harness — it is the system under analysis. The paper is a source-level dissection of NVIDIA's OpenSHMEM-based PGAS library, so the "architecture" here is the library's own layered runtime, and the "harness" (Section 2) is a thin microbenchmark wrapper bolted on top to expose its behavior. The defining design commitment is device-initiated, one-sided communication over symmetric memory: a CUDA thread, warp, or thread block issues put/get/atomic operations directly from inside a kernel, against a symmetric heap that every processing element (PE) allocates with identical type, size, and layout. The initiator names a (source, destination, target-PE) triple; the target performs no matching receive. This is the inverse control structure of NCCL, where the host enqueues a collective onto a CUDA stream and the CPU stays in the control path.

+-------------------------------------------------------------------+
|                    NVSHMEM Runtime (v3.3.9)                       |
|                                                                   |
|  +-------------------------------------------------------------+  |
|  |  DEVICE-CALLABLE API SURFACE  (issued from inside kernels)  |  |
|  |    nvshmem_<TYPE>_put / get   nvshmem_<TYPE>_atomic_<op>    |  |
|  |    nvshmem_fence / quiet      nvshmem_wait / test          |  |
|  |    *_warp  /  *_block  threadgroup-scoped variants          |  |
|  +----------------------------+--------------------------------+  |
|                               |                                   |
|  +----------------------------v--------------------------------+  |
|  |  HOST-CALLABLE API SURFACE  (setup + on-stream operations)  |  |
|  |    nvshmem_init / finalize    nvshmem_malloc / align / free|  |
|  |    nvshmemx_put_on_stream     nvshmemx_quiet_on_stream     |  |
|  |    team create / destroy      collective launchers         |  |
|  +----------------------------+--------------------------------+  |
|                               |                                   |
|  +----------------------------v--------------------------------+  |
|  |  COLLECTIVE LAYER  (built ON TOP of RMA + signaling)        |  |
|  |    Broadcast  AlltoAll  FCollect(=AllGather)               |  |
|  |    Reduce(=AllReduce)  ReduceScatter  Barrier  Sync        |  |
|  |    rule-based decision tree picks algo per call            |  |
|  +----------------------------+--------------------------------+  |
|                               |                                   |
|  +----------------------------v--------------------------------+  |
|  |  ONE-SIDED COMMUNICATION CORE                               |  |
|  |   FAST PATH: direct SM load/store to mapped peer heap       |  |
|  |   SLOW PATH: IBGDA (GPU-posted RDMA) OR host-proxy thread   |  |
|  +----------------------------+--------------------------------+  |
|                               |                                   |
|  +----------------------------v--------------------------------+  |
|  |  SYMMETRIC-MEMORY MANAGER                                   |  |
|  |   VA reservation (cuMemAddressReserve) eager                |  |
|  |   PA commit on demand (cuMemCreate/Map/SetAccess)          |  |
|  |   first-fit host allocator over 3 std::map structures      |  |
|  |   peer_heap_base_p2p_[] (fast) / peer_heap_base_remote_[]  |  |
|  +----------------------------+--------------------------------+  |
|                               |                                   |
|  +----------------------------v--------------------------------+  |
|  |  TRANSPORT LAYER (pluggable)                                |  |
|  |   P2P: NVLink / NVSwitch / MNNVL  |  IBGDA  |  IBRC/IBDEVX |  |
|  |   UCX | libfabric (cxi on HPE Slingshot)                   |  |
|  +-------------------------------------------------------------+  |
+-------------------------------------------------------------------+
▲ Fig 1: NVSHMEM as a layered runtime. The two API surfaces (device,
  host) sit above a collective layer that is itself built on the RMA
  core. The architectural pivot is the RMA core's fork into a fast
  path (SM-issued loads/stores) and a slow path (IBGDA or host proxy).

The single most consequential design choice is the fast-path / slow-path fork in the RMA core. Whether a put becomes a plain GPU store or a posted RDMA descriptor is decided by one predicate: is the target PE's heap P2P-mappable into the local GPU's virtual address space? If yes (same node, NVLink/NVSwitch-reachable, or MNNVL), the symmetric offset is dereferenced directly by a Streaming Multiprocessor — no NIC, no proxy, no host. If no (inter-node, or non-P2P intra-node), the same computed remote address is handed to a transport that either lets the GPU post the RDMA work itself (IBGDA) or stages it through a host proxy thread. This fork is the source of nearly every regime crossover in the measurement section: the fast path is latency-bound by SM-issue overhead, the slow path by NIC and synchronization cost.

A second structural decision is the dual host/device interface, which the authors call NVSHMEM's "key strength" and simultaneously the reason the library is "substantially more complex than communication libraries built around a single control path." The split is not clean: setup and memory management remain host-managed, ordering and synchronization are mostly device-side, and only RMA / atomics / collectives are genuinely available on both. The naming convention encodes the boundary — nvshmem_ is standard OpenSHMEM, nvshmemx_ is NVSHMEM-specific extension (stream-ordered and threadgroup-scoped forms), and the internal nvshmemi_ prefix marks routines that are not part of the public API.

2. System-Under-Test Architecture (the testbed + the API surface)

Because the paper is the dissection, the "system under test" splits into two complementary specimens: the hardware testbed the microbenchmarks run on (Section VII), and the API surface the source analysis maps (Table I). Both are reproduced here.

2.1 Hardware testbed (CoreWeave H200 cluster)

+--------- CoreWeave cluster: NVLink-4 intra-node + IB inter-node ---+
|                                                                    |
|   Node 0                                  Node 1                   |
|  +--------------------------------+      +----------------------+  |
|  |  8x NVIDIA H200 SXM5           |      |  8x H200 SXM5        |  |
|  |  144 GB HBM3e per GPU         |      |  144 GB HBM3e/GPU    |  |
|  |                                |      |                      |  |
|  |  NVLink-4 fabric:             |      |  NVLink-4 fabric     |  |
|  |   900 GB/s bidir per GPU      |      |   (intra-node P2P)   |  |
|  |   NVLS / NVLink-SHARP multicast|      |                      |  |
|  |   (in-network reduction)      |      |                      |  |
|  +---------------+----------------+      +----------+-----------+  |
|                  |                                  |              |
|             8x ConnectX-7 IB NICs            8x ConnectX-7 IB     |
|             (50 GB/s per NIC,                (8 NICs/node)        |
|              ~400 GB/s aggregate/node)                            |
|                  |                                  |              |
|                  +==================================+              |
|                       InfiniBand inter-node fabric                |
|                                                                    |
|   Reference upper bounds (red dashed lines in Figs 4-5):          |
|     - per-GPU NVLink:        450 GB/s (uni) / 900 GB/s (bidir)    |
|     - inter-node aggregate:  400 GB/s (8 NICs x 50 GB/s)         |
|                                                                    |
|   Software: CUDA 13.0.88, NVSHMEM 3.3.9, NCCL via nccl-tests      |
+--------------------------------------------------------------------+
▲ Fig 2: SUT hardware. NVLink-4 makes intra-node ~10x richer than
  inter-node IB, so the fast-path/slow-path fork in Fig 1 maps almost
  exactly onto the intra-node/inter-node boundary on this testbed.

This testbed is the opposite of the "flat interconnect" regime: NVLink-4 at 900 GB/s bidirectional per GPU versus an aggregate ~400 GB/s of IB per node means intra-node P2P is roughly an order of magnitude richer than inter-node RDMA. The fast-path/slow-path architectural fork therefore lines up almost perfectly with the intra-node/inter-node boundary, which is exactly why the microbenchmarks split every figure into an "intra-node" panel and an "inter-node" panel.

2.2 API surface (Table I, reproduced as a layered stack)

+-------------------------------------------------------------------+
|  API GROUP            AVAILABILITY      ROLE                       |
|-------------------------------------------------------------------|
|  Setup / exit         Mixed            init, finalize, bootstrap,  |
|                       (mostly host)    global_exit (also on dev)   |
|-------------------------------------------------------------------|
|  Memory management    Host-side        malloc / align / free over  |
|                                        the symmetric heap; buffer  |
|                                        registration                |
|-------------------------------------------------------------------|
|  Team management      Mixed            define PE subgroups; PE-id   |
|                       (mostly host)    translation; create/destroy |
|-------------------------------------------------------------------|
|  One-sided RMA        BOTH             put / get / scalar p,g /     |
|                                        strided iput,iget; the core |
|-------------------------------------------------------------------|
|  Atomics              BOTH             fetch, add, compare_swap,    |
|                                        swap, bitwise on symmetric   |
|-------------------------------------------------------------------|
|  Memory ordering      BOTH             fence (order), quiet         |
|                                        (order + completion)         |
|-------------------------------------------------------------------|
|  Synchronization      Mixed            wait / test on signal vars;  |
|                       (mostly device)  host via on-stream variants  |
|-------------------------------------------------------------------|
|  Collectives          BOTH             Barrier, Sync, Broadcast,    |
|                                        AlltoAll, FCollect, Reduce,  |
|                                        ReduceScatter                |
+-------------------------------------------------------------------+
▲ Fig 3: The eight API groups (Table I). "BOTH" rows are where the
  device-initiated model lives; "Host-side" rows (memory, most of
  setup/teams) reveal that control of the heap itself never leaves
  the CPU even though data movement does.

The mapping is the architecture's honest self-portrait: the data plane (RMA, atomics, ordering, collectives) is genuinely device-callable, but the control plane (heap allocation, team membership, bootstrap) stays host-resident. NVSHMEM's PGAS abstraction therefore gives the GPU a flat remote-memory view for moving data, but the shape of that memory — what is allocated where, which PEs are in which team — is still decided by host code. The paper contrasts this with NCCL's device API, which makes the scale-up vs scale-out distinction explicit and adds hierarchical composition through communicators; NVSHMEM hides that distinction behind one symmetric heap shared by all PEs in an instance.

3. Design-Space Diagram (axes swept, axes held fixed)

The microbenchmarks in Section VII are deliberately not a comprehensive study — the authors state performance evaluation "is not the primary focus." They sweep a narrow but well-chosen design space to expose a few characteristics. The dimensions actually varied, and those pinned, are:

                   DESIGN SPACE (swept vs held-fixed)
  +---------------------------------------------------------------+
  |                                                               |
  |  Axis 1: OPERATION KIND (4 levels)                            |
  |    [bulk put]  [bulk get]  [scalar p]  [scalar g]            |
  |    (+ AllReduce as the one collective benchmarked)           |
  |                                                               |
  |  Axis 2: MESSAGE SIZE (full sweep, the x-axis of every fig)  |
  |    2^? ... 2^? bytes  -> small-msg latency vs large-msg BW   |
  |    crossover region of interest: ~16 KiB ... 64 KiB          |
  |                                                               |
  |  Axis 3: LOCALITY / PATH (2 levels = the fast/slow fork)     |
  |    [intra-node H200 P2P (NVLink fast path)]                  |
  |    [inter-node H200 (IBGDA slow path)]                       |
  |                                                               |
  |  Axis 4: TRANSPORT CONFIG for inter-node (tuned IBGDA)       |
  |    NVSHMEM_IBGDA_NUM_RC_PER_PE = 64                          |
  |    64 CTAs, 1024 threads/CTA  (max RDMA concurrency)         |
  |                                                               |
  |  Axis 5 (AllReduce only): EXECUTION MODEL (4 levels)         |
  |    NVSHMEM [On-stream]  (host multi-CTA + NVLS)             |
  |    NVSHMEM [Device Block] (single-CTA device path)          |
  |    NCCL [Ring]                                               |
  |    NCCL [NVLS]                                               |
  |                                                               |
  |  Held FIXED (no sweep):                                       |
  |    - GPU:   H200 SXM5 only (no A100 / GH200 / consumer)      |
  |    - Scale: 2 nodes / up to 16 GPUs (small; not 100s)        |
  |    - NVSHMEM version: 3.3.9 (one release)                    |
  |    - PE-per-GPU: 1 (multi-PE-per-GPU exists but not swept)   |
  |    - Datatype for AllReduce: float (sum)                     |
  |    - Topology: single CoreWeave cluster                      |
  |    - Trials: 8 per point, averaged (std dev too small to see)|
  |                                                               |
  +---------------------------------------------------------------+
▲ Fig 4: Design space of the microbenchmarks. The product that matters
  is (operation x message-size x locality): bulk-vs-scalar crossed
  with small-vs-large crossed with intra-vs-inter. AllReduce adds a
  4-way execution-model comparison against NCCL.

Two scoping decisions define the measurement reach. First, only AllReduce is benchmarked among collectives, because it is the most performance- critical and the one where both NVSHMEM and NCCL ship optimized implementations — the other six collectives are described analytically via Table II but never timed. Second, the scale is small (2 nodes, <=16 GPUs); this is a behavior-characterization study, not a scaling study, so the findings speak to per-operation efficiency rather than to large-scale strong/weak scaling.

4. Algorithm / Control-Flow Diagrams

Four procedures carry the paper's mechanistic content: symmetric-heap setup, the fast-path RMA, the slow-path RMA, and the collective dispatch decision. A fifth (DeepEP's HT pipeline) is the real-application case study.

4.1 Symmetric-heap setup (Fig. 1 of the paper, the 5-marker flow)

  nvshmem_init() on each PE
       │
       ▼
① [Reserve VA range]  cuMemAddressReserve(p2p_npes × heap_size)
       │   segment 0 = local heap; segments 1..k = peer mappings
       ▼
② [Create host mspace allocator]  — NO physical pages committed yet
       │
       ▼
③ [Exchange heap_base_ across PEs]  (post transport init)
       │
       ▼
   ── application calls nvshmem_malloc(), heap must grow ──
       │
       ▼
④ allocate_physical_memory_to_heap():
       ⒈ cuMemCreate()      → physical handle
       ⒉ cuMemMap() + cuMemSetAccess()  → map subrange for local
                                          + eligible P2P peers
       ⒊ register region with P2P transport
       ⒋ insert region into mspace allocator
       ⒌ register region with NETWORK transport (non-P2P peers)
       │
       ▼
⑤ [Populate peer base tables + barrier]
       peer_heap_base_p2p_[pe]    ← directly mappable peers (FAST)
       peer_heap_base_remote_[pe] ← network peers + handles (SLOW)
▲ Fig 5: Symmetric-heap lifecycle. VA is reserved EAGERLY (step ①),
  physical memory is committed LAZILY on first allocation overflow
  (step ④). The two peer-base tables produced at ⑤ are exactly the
  fork predicate the RMA path will read.

The VA-eager / PA-lazy split (built on CUDA's low-level Virtual Memory Management API) is what lets every PE reserve a heap large enough to cover all P2P-reachable peers' heaps at fixed offsets, without paying for the physical memory up front. That fixed-offset layout is the whole reason a remote address can be computed from a local one by simple base substitution (Equation 1 below).

4.2 Remote-address computation (Equation 1) and the fast/slow fork

  Given a local symmetric pointer  dest_local
  and the local heap base          heap_base_:

      offset      = dest_local − heap_base_
      dest_remote = peer_heap_base_*[remote_pe] + offset

           ┌─────────────────────────────────────────┐
           │   is remote_pe P2P-mappable locally?     │
           └───────────────┬──────────────┬───────────┘
                    YES (fast)         NO (slow)
                       │                  │
        peer_heap_base_p2p_[pe]   peer_heap_base_remote_[pe]
                       │                  │
                       ▼                  ▼
              ┌─────────────────┐  ┌──────────────────────────┐
              │ SM issues a      │  │ encode RDMA work request  │
              │ direct store/    │  │   ↓                       │
              │ load to the      │  │ IBGDA available?          │
              │ mapped address   │  │   YES → GPU posts RDMA    │
              │ (no NIC, no host)│  │         (nvshmemi_ibgda_  │
              └─────────────────┘  │          rma_*) to NIC     │
                                   │   NO  → write descriptor   │
                                   │         to pinned buffer;  │
                                   │         host PROXY thread  │
                                   │         consumes & posts   │
                                   └──────────────────────────┘
▲ Fig 6: The one predicate that governs every put/get. The offset is
  preserved identically on both paths; only the base table and the
  dereference mechanism differ. The slow path forks AGAIN on IBGDA
  availability (GPU-posted RDMA vs host-proxy).

4.3 Collective algorithm dispatch (Table II as a decision tree)

NVSHMEM selects collective algorithms at runtime via a per-collective rule-based decision tree — explicitly not an analytical cost model. Capability checks, datatype/scope constraints, scratch-space availability, and fixed message-size thresholds gate which algorithm is allowed and preferred; unsupported cases fall back to more general implementations.

  AllReduce(msg M, PEs N) called
       │
       ▼
① NVLS hardware (NVLink-SHARP) available + datatype OK?
       ├── YES → NVLS one-shot  (O(MN) vol, O(1) sync)  ── in-network
       │         or NVLS two-shot for larger M
       └── NO ─┐
               ▼
② P2P-connected team + scratch available?
       ├── YES → k-ary recursive exchange (O(MN·k·log_k N), O(log_k N))
       │         or hierarchical fcollect (inherits FCollect cost)
       └── NO ─┐
               ▼
③ fall back → (segmented) linear AllReduce
               direct LD/ST: O(1) sync ; segmented: O(SN) sync
               where S = ⌈M / B_seg⌉ scratch-limited segments
▲ Fig 7: AllReduce dispatch. The tree prefers in-network NVLS first,
  then P2P recursive-exchange, then a general linear fallback. LL /
  LL128 protocols (Fig 8) overlay this for SMALL messages on the
  algorithms whose Table-II "LL/LL128" column says they support it.

The LL and LL128 protocols (borrowed in design from NCCL) couple data movement with arrival notification to avoid a separate post-transfer barrier on small messages. LL packs two 4-byte data elements plus two flags into a single 16-byte atomic write; the receiver polls the flags. LL128 groups 120 bytes of data with an 8-byte flag into a 128-byte unit for better bandwidth utilization — but is NVLink-only, because it relies on 128-byte atomic stores not guaranteed on PCIe. Both protocols require extra pSync storage because the receive buffer must hold payload plus flags.

4.4 DeepEP high-throughput dispatch (Fig. 6 of the paper)

DeepEP (DeepSeek's expert-parallel MoE library) is the real-application case study. It does not use NVSHMEM collectives — it uses NVSHMEM only at the critical points of its own multi-stage transport pipeline (metadata exchange, chunked RDMA puts, atomic credit updates), building everything else on top.

  HT DISPATCH (two-node, warp-specialized, 16 warps/SM)
       │
       ▼
① notify_dispatch: exchange per-rank/per-expert token counts as
   METADATA (not payload); compute prefix matrices, channel shares
       │
       ▼
② RDMA-side SM (7 sender warps + 1 sender-coordinator + 8 NVL recv):
     senders place tokens into per-peer RDMA RING BUFFERS on the
     symmetric heap
       │
       ▼
③ sender-coordinator warp batches writes into larger RDMA transfers
     via  nvshmemi_ibgda_put_nbi_warp
     and updates the remote tail with
          nvshmemi_ibgda_amo_nonfetch_add   (atomic credit)
       │  (cross-node tokens go only to GPUs with the SAME slot index)
       ▼
④ paired SM (8 RDMA→NVLink forwarder warps + 8 NVL-receiver warps):
     forwarders poll arrivals, decode metadata, copy tokens into
     intra-node NVLink ring buffers
       │
       ▼
⑤ NVLink-receiver warps place tokens into the output tensor
       │
       ▼
   combine kernel = the same structure in reverse (NVLink → RDMA → home)
▲ Fig 8: DeepEP HT path. NVSHMEM is the cross-node RDMA SUBSTRATE
  (steps ②③ via IBGDA put + atomic), while DeepEP layers its own
  warp-specialized two-stage (RDMA-then-NVLink) pipeline on top.

The low-latency (LL) DeepEP path for inference is structurally simpler: it drops the intra-node NVLink-forwarding stage, uses a single global NVSHMEM world team (instead of eight strided teams), partitions SMs by local expert rather than by token range, and reduces critical-path NVSHMEM activity to a single IBGDA put for payload plus one atomic update for the count.

5. Quantitative Results — Empirical Findings by Regime

5.1 One-sided RMA bandwidth (Fig. 4) — intra vs inter, bulk vs scalar

The single sharpest result is the scalar-g collapse: device-initiated scalar get stays below 9 GB/s intra-node and at 1.28 GB/s inter-node, because every g requires the issuing thread to wait for the returned value before completing — serializing each thread to one outstanding operation and killing pipelining. Scalar put (p) does far better (172 GB/s intra, 15.6 GB/s inter) precisely because remote stores can be buffered and overlapped without waiting for a return value. The bulk paths land much closer to the ceiling: intra-node bulk put reaches ~70% of NVLink, and with tuned IBGDA the inter-node bulk paths hit ~96% of a single 50 GB/s rail.

5.2 RMA latency (Fig. 4 insets)

The insets explain why small messages are not bandwidth-efficient: at the sub-microsecond-to-microsecond scale, fixed per-operation overheads (address translation, synchronization, work partitioning, protocol setup) dominate. The takeaway the authors draw is explicit — NVSHMEM is most effective for bulk or aggregated write-style RMA, and scalar operations should be treated as latency/control primitives and batched when bandwidth matters.

5.3 AllReduce: NVSHMEM vs NCCL (Fig. 5)

The intra-node story is a near-tie: NVSHMEM's host-side multi-CTA on-stream NVLS path (264 GB/s) closely tracks NCCL's NVLS (276 GB/s) and is latency-competitive on small messages. The inter-node story is a rout: NVSHMEM AllReduce falls below 0.20 GB/s across nodes while NCCL reaches 180–252 GB/s, and NVSHMEM's inter-node AllReduce latency grows into the milliseconds by 64 KiB while NCCL stays in the tens of microseconds. The cause the authors identify is that NVSHMEM's optimized collective effort has been concentrated on NVLS-based intra-node / MNNVL algorithms; its inter-node AllReduce is comparatively unoptimized. (Several inter-node NVSHMEM points are missing from Fig. 5 entirely due to timeouts.)

5.4 The multi-CTA gap (the root cause behind 5.3)

NVSHMEM's device-side collective is a single-participating-threadgroup collective by design — there are deliberately no public _grid variants, because a grid-scoped operation would require synchronization across CTAs inside a running kernel, which is more expensive than ending the kernel and launching a stream-ordered communication kernel. Multi-CTA execution is exposed only through host-side _on_stream wrappers, only for FCollect / AllReduce / ReduceScatter, and is gated on NVLS availability. This is the single-CTA bottleneck visible as the ~30 GB/s "Device Block" floor in Fig. 5.

6. Configuration-Regime Trade-off Tables

6.1 Communication path: fast (P2P) vs slow (IBGDA / proxy)

Winner is regime-dependent: within a node (or MNNVL domain) the fast path is strictly preferable; across nodes IBGDA is preferred over the host proxy because it keeps the CPU out of the control path and reaches ~96% of a single IB rail on bulk transfers.

6.2 Operation kind: bulk vs scalar, put vs get

Winner: bulk put for write-dominated movement; scalar g should be avoided in bandwidth-critical loops and reserved for genuine control reads. The asymmetry between p (172 GB/s) and g (<9 GB/s) intra-node is the clearest "use put, not get, where you can" signal in the paper.

6.3 AllReduce execution model

Winner is split: for intra-node AllReduce NVSHMEM's on-stream NVLS path is competitive with NCCL; for inter-node AllReduce NCCL (Ring or NVLS Tree) is decisively better and NVSHMEM should not be used as the AllReduce engine across nodes. NVSHMEM's value is as a fine-grained one-sided substrate, not as a drop-in collective replacement.

6.4 Protocol selection for small messages

Winner is message-size- and fabric-dependent: Simple for large messages, LL for small messages on any fabric, LL128 for small messages on NVLink where its larger transfer unit recovers bandwidth without violating the 128-byte atomic-store requirement.

7. Bottlenecks & Insights Surfaced by the Measurements

7.1 The scalar-get serialization wall

The headline bottleneck. Scalar g collapses to <9 GB/s intra-node and 1.28 GB/s inter-node because each operation forces the issuing thread to wait for its return value before retiring, capping the number of outstanding operations at one per thread and eliminating pipelining. Scalar p does not suffer this — buffered stores overlap freely. The architectural lesson: one-sided write primitives pipeline; one-sided read primitives that return a value into the issuing thread do not. This is a property of the completion semantics, not of the transport.

7.2 The single-CTA collective ceiling

NVSHMEM's deliberate choice to make device collectives single-threadgroup — to avoid expensive in-kernel cross-CTA synchronization — produces the ~30 GB/s "Device Block" AllReduce floor. Multi-CTA parallelism exists only via host-side on-stream wrappers and only when NVLS hardware is present. The authors flag "limited multi-CTA support remains a weakness of NVSHMEM's built-in collective implementation." This is the clearest gap between NVSHMEM and NCCL on the device side.

7.3 Inter-node collective immaturity

NVSHMEM's optimization investment has gone into NVLS/MNNVL intra-node algorithms, leaving inter-node AllReduce orders of magnitude behind NCCL (<0.20 GB/s vs 180–252 GB/s; millisecond vs microsecond latency). The fast/slow fork in the architecture is mirrored in the maturity of the two paths: the fast (P2P) path is highly tuned, the slow (network-collective) path is not.

7.4 The small-message fixed-cost floor

Across every operation, sub-64-KiB transfers are dominated by fixed per-operation overhead: address translation (Equation 1 + base-table lookup), synchronization, work partitioning, and protocol setup. This is exactly the regime LL/LL128 protocols exist to attack, and exactly why DeepEP batches many tokens into larger RDMA transfers via a sender-coordinator warp rather than issuing one RDMA per token.

7.5 The convergence-with-NCCL signal (DeepEP)

NVIDIA's published DeepEP comparison shows NCCL GIN matches NVSHMEM "typically within about 1–2%" while offering the same device-initiated style inside NCCL's runtime. The insight: the programming model NVSHMEM pioneered (device-initiated, symmetric-memory, one-sided) is being absorbed into NCCL's device API, so NVSHMEM's durable advantage is narrowing to applications that need fine-grained one-sided RMA and direct device-level control that NCCL's collective-centric model still does not expose.

7.6 The fence-vs-quiet ordering distinction

fence orders previously-issued operations to the same destination PE but guarantees nothing about completion; quiet orders and waits for completion and visibility at the destination. barrier carries stronger ordering than sync because it enforces completion (via quiet or __threadfence_system) before entering the dissemination rounds. These three completion tiers are the synchronization-cost knob behind small-message performance — choosing fence where quiet is not required avoids a completion wait.

8. Limitations of the Methodology

The most consequential limitation is the narrow collective coverage crossed with small scale: the paper convincingly characterizes per-operation RMA behavior and the intra-node AllReduce tie with NCCL, but cannot speak to how NVSHMEM's collectives behave at production scale or how the six untimed collectives perform. The inter-node AllReduce collapse is established, but its shape at scale is left open by the timeouts.

9. Note on NCCL Tuning

NVSHMEM's LL / LL128 / Simple protocols and its runtime rule-based algorithm decision tree are the same shape of selection problem that any NCCL configuration tuner faces: pick the protocol that couples notification with transfer for small messages (LL), the wider-unit variant only on fabrics that support the required atomic width (LL128 on NVLink), and the bulk protocol for large messages (Simple). The paper's empirical crossover at ~16–64 KiB between latency-bound and bandwidth-bound behavior, and its sharp intra-node/inter-node fork, are precisely the regime boundaries a protocol/ algorithm selector must learn rather than hard-code. The scalar-p-vs-g asymmetry also reinforces a general tuning prior: write-style, buffered primitives admit pipelining and respond well to higher concurrency, while return-value reads do not — a distinction worth respecting when choosing how aggressively to parallelize a given collective phase.

10. Analogy

NVSHMEM is a shared warehouse with pre-assigned, identical shelf layouts in every branch. Every branch (PE) builds the same rack plan (the symmetric heap), so a worker who knows where item X sits on his own shelf (dest_local − heap_base_) knows exactly where it sits on any other branch's shelf (peer_heap_base + offset) — no phone call, no order form, no clerk confirming receipt. That is one-sided communication. The crucial twist is how the worker reaches the remote shelf. If the other branch is in the same building connected by an internal conveyor (NVLink P2P), the worker just walks over and grabs the item himself — fast, no manager involved (the fast path, SM direct load/store). If the branch is across town (inter-node), he either dispatches a courier the warehouse robot can summon directly (IBGDA, GPU-posted RDMA) or, lacking that, fills out a slip and hands it to a dispatch clerk who calls the courier (the host proxy). The warehouse's brutal lesson, measured in the data, is that fetching an item and waiting to carry it back yourself (scalar get) jams the aisle — you can only hold one item at a time — whereas dropping off items (scalar put) lets you keep moving and stack many in flight. And the warehouse is superb at moving goods within a building but still clumsy at coordinating a synchronized inventory count across all branches at once (inter-node AllReduce), which is why for that one task the established courier company (NCCL) still wins. NVSHMEM's contribution is making the self-service shelf model work directly from the loading dock (the GPU), so specialized operations like DeepEP can build their own multi-stage logistics pipeline on top of it.