Purlin: Separating Orchestration from the Datapath of Collectives

Separating the concerns

• Purlin targets the challenge of adapting collectives to changing GPU mechanisms and application regimes, which existing implementations often make difficult by coupling semantics, orchestration, and datapath.Applications span throughput-oriented prefill, latency-sensitive decode, and device-resident building blocks that can fuse with computation.Source: §1 · pp. 1–2

Purlin’s properties

• Purlin aims to derive orchestration for new collectives and let applications customize datapaths or codesign policies without rebuilding orchestration.Existing frameworks still require developers to translate new collective semantics into orchestration, while coupled orchestration and movement can make datapath customization require revisiting orchestration.Source: §1 · pp. 2–3

Where we are today

• Existing systems trade portability against specialization: portable libraries are inconvenient to evolve, while generation-specific implementations can abandon support for still-deployed GPUs.The introduction cites delayed adoption of hardware mechanisms and specialized implementations that omit older Ampere GPUs as examples.Source: §1 · p. 2

Purlin’s three-layer design

• Purlin describes collectives through input and output layouts and a consume operation, then uses SNAC for shared coordination and Atom for hardware-specific copying or reduction.SNAC schedules synchronization, manages buffer lifetimes, and initiates transfers through Atom; the layers separate the communication description from orchestration and movement.Source: Figure 1 · p. 2, §1 · p. 3

Why this separation is hard

• Purlin supports multiple GPU generations with distinct Atom implementations and uses layout-derived coordination, protocol specializations, and compile-time selection to address collective-specific execution needs.The framework provides a broadly compatible vectorized-load/store Atom alongside Ampere, Hopper, and Blackwell specializations; its shared SNAC protocol handles differing layouts and latency- or throughput-bound transfers.Source: §1 · p. 3

Results

• Across seven collectives, Purlin achieves latency speedups of up to 5.14× and bandwidth improvements of up to 4.50× over existing baselines.In SGLang, offline LLM serving interactivity improves by 1.13× on average and up to 1.37×; online interactivity improves by 1.26× on average and up to 2.85×, with the largest gain under overload; diffusion generation speeds up by up to 1.13×.Source: §1 · p. 3

2 Background and Motivation

• Table 2 summarizes representative communication patterns in distributed inference, with (V) denoting a variable-length variant.Source: Table 2 · p. 3

2.1 Stable Semantics, Evolving Application Demands

• Collective contracts specify input and output conditions independently of communication steps and hardware, while AllGather, ReduceScatter, AllReduce, and AllToAll have distinct data-distribution semantics.Variable-length variants allow contributions or partitions to differ in size.Source: §2.1 · p. 3
• Prefill’s large payloads emphasize throughput, while decode’s small transfers emphasize latency; fine-grained fusion also calls for device-side primitives that applications can schedule and compose.A high-bandwidth strategy can incur setup and synchronization costs that dominate small messages.Source: §2.1 · p. 4
• Communication accounts for 30–32% of end-to-end latency at concurrencies one and four; replacing SGLang’s backend with Purlin improves request latency by 1.23–1.25×.Purlin also improves TPOT by the same gain and TTFT by 1.10×.Source: §2.1 · p. 4

2.2 Hardware Evolution and Communication Cost

• GPU datapaths have evolved from vectorized loads and stores to asynchronous transfers, TMA, and NVLink SHARP, alongside bandwidth growth from 300 GB/s on Ampere to 900 GB/s on Blackwell.The section argues that datapaths should evolve separately from collective semantics and orchestration, without sacrificing performance, while supporting application regimes and fine-grained primitives.Source: §2.2 · pp. 4–5
• Across A100, H200, and B200, increasing interconnect bandwidth raises the bandwidth-delay product from roughly 512 KiB to 2 MiB.Measured NVLink round-trip latencies are approximately 2.2 µs, 2.1 µs, and 2.8 µs, respectively.Source: §2.2 · p. 4
• Atom::copy reaches comparable or higher bandwidth than tuned NCCL Device API with fewer SMs; on B200, both approach 760 GB/s with a fourfold difference in SM allocation.NVSHMEM remains below that bandwidth plateau across the measured SM range.Source: §2.2 · p. 5

3 Purlin Design

• Purlin separates collective layout specifications, SNAC coordination, and Atom data movement, allowing orchestration to be reused as datapaths and applications are customized.Applications can also invoke Atom directly and tune its pipeline.Source: §3 · p. 5

3.1 Layouts

• Layouts describe rank-relative input and output regions; Purlin uses packed, scattered, and transposed layouts, and composes compatible transformations into one GPU kernel.For example, AllReduce is ReduceScatter followed by AllGather, composed as a scattered-to-scattered transformation.Source: §3.1 · pp. 5–6

3.2 Stage, Notify, and Consume (SNAC)

• SNAC derives rank dependencies from input and output layouts, then stages, notifies, and consumes chunks using the collective’s copy or reduction operation.For composed AllReduce, staging, reduction, and gathering CTAs operate on chunks so staging can overlap with reduction and gathering.Source: §3.2 · pp. 6–8
• Up to 1.98× lower latency with push and up to 7% higher bandwidth with pull guide SNAC’s choice of push for latency-bound execution and pull for throughput.The latency path combines payload and notification in an atomic 16-byte packet, doubling network data volume, so codesign policies choose the crossover to the throughput path.Source: Figure 7 · p. 8, §3.2 · p. 8

3.3 The Atom

• Atom provides copy and reduction operations through a unified pipeline that uses CTA-local shared memory to overlap asynchronous data fetching with processing.The processing phase is configurable; on Ampere and above, Atom uses cp.async for asynchronous global-to-shared data movement.Source: §3.3 · pp. 8–9
• SCORE makes Atom reductions deterministic for a fixed source ordering and hardware architecture by assigning each output element to one thread and accumulating sources in order.It pipelines source fetches and uses multiple independent accumulators to improve throughput while preserving ordered accumulation.Source: §3.3 · p. 10
• SCORE reductions are deterministic by construction, and SNAC’s latency path computes each output once in source order.The latency path reduces packets while maintaining both stated invariants: one thread computes each output element, in source order.Source: §3.3 · p. 10
• NVLS offers hardware-accelerated reduction on Hopper and newer GPUs, but its unspecified order lacks SCORE’s determinism guarantee, so Purlin lets users choose either path.The paper compares the implementations in §B.Source: §3.3 · p. 10
• Pipeline tuning aims to bring in-flight capacity Qcap near the bandwidth–delay product while minimizing CTA count C.Increasing per-CTA stage size or stage count raises shared-memory use; increasing C uses more SMs. Policies are specialized by collective and hardware generation, and direct Atom callers can tune configurations.Source: §3.3 · p. 10

4 Purlin Implementation

• Purlin comprises 6,200 lines of CUDA C++ and 704 lines of Python, depending on libcu++ and CUB for prefix sums.Its Python API JIT-compiles an architecture-specific module and selects corresponding Atom implementations and codesign policies.Source: §4 · pp. 10–11

4.1 Staging Buffers

• Each GPU allocates two 2 MiB staging buffers for latency-bound execution and two 256 MiB buffers for throughput-bound execution.Users can configure buffer capacity to trade memory use for performance; larger throughput footprints cycle through slots.Source: §4.1 · p. 11

4.2 Codesign Policies and Tuning

• Codesign policies configure SNAC parameters, including latency–throughput crossover, chunk sizes, CTA allocations, and Atom pipeline depths.The authors tune these offline for each collective and GPU generation, while applications can define their own policies.Source: §4.2 · p. 11

4.3 Epoch Bookkeeping

• A monotonic 64-bit epoch selects between staging buffers by parity and also supplies SNAC’s notification value.Epoch increments use the invocation’s chunk count rounded to an odd number, or the maximum chunk count across ranks for variable-length collectives.Source: §4.3 · p. 11

5.1 Collective Microbenchmarks

• 1.65–4.36× geometric-mean latency improvements over NCCL span twelve collective–platform combinations, with a 5.14× maximum over NCCLX for 2 KiB B200 AllReduce.Against MSCCL++, AllGather improves 1.23–1.56×, while AllReduce ranges from 0.98× to 1.06× in geometric-mean speedup.Source: §5.1 · p. 13
• 1.14–1.53× geometric-mean bandwidth improvements over NCCL vary by GPU generation, while H200 NCCL leads Purlin on all four collectives at 1 GiB.For B200 AllReduce over 1 MiB–1 GiB, Purlin reaches 432.9 GB/s versus MSCCL++ at 371.5 GB/s; the B200 peer-buffer comparison also reports Purlin AllGather at 718.7 GB/s versus MSCCL++ at 592.4 GB/s.Source: §5.1 · p. 13

Can we increase bandwidth without allocating more

• 6.1× more pull bandwidth on B200 results from moving from one to eight pipeline stages with the same sixteen CTAs.Push bandwidth rises from 239 GB/s at one stage to 710 GB/s at four and 713 GB/s at eight, whereas pull rises from 477 to 761 GB/s between four and eight stages.Source: Can we increase bandwidth without… · pp. 14–15

5.3 Application Performance

• 1.13× geometric-mean gains across offline serving comparisons and 1.26× across online arrival rates accompany image-latency gains of approximately 1.09×–1.13× over stock SGLang in the reported resolution range.Offline gains reach 1.37×; online gains reach 2.85×, with the largest gain under overload.Source: §5.3 · p. 15
• At 1.5 requests/s, online serving end-to-end latency falls from 121.3 to 41.5 seconds versus stock SGLang, but achieved throughput plateaus at approximately 1.2–1.4 requests/s.The authors hypothesize that reduced queuing delay for backlogged requests contributes to the larger overload gain.Source: §5.3 · p. 15

5.4 Output Quality

• For Qwen-Image, outputs were pixel-identical across all 100 Qwen-Image-Bench prompts at 1024^2 resolution with matched seeds; pass@16 and majority@16 scores were identical.Source: §5.4 · p. 16

6 Limitations and Future Work

• Evaluation covers only eight-GPU NVLink domains; larger domains may need different orchestration, while multi-node RDMA deployments would need a scale-out datapath.Both directions are left to future work.Source: §6 · p. 16
• Although SNAC expressions are shown for Broadcast, Reduce, Scatter, and Gather, the prototype implements only non-rooted collectives.Implementing and tuning the rooted collectives remains future work.Source: §6 · p. 16
• The Atom supports device-initiated mechanisms such as cp.async and TMA, but not host-initiated mechanisms such as the copy engine.The authors identify device-compatible interfaces for host-side mechanisms as a possible future direction.Source: §6 · p. 16

7 Related Work

• SCCL, TACCL, and OptCCL synthesize collective algorithms, while MSCCLang compiles chunk operations into execution schedules; Purlin instead develops a shared orchestration protocol for compile-time descriptions.The synthesis systems cover topology-guided search, communication sketches, or concurrent-collective optimization.Source: §7 · p. 16
• NVSHMEM, NCCL’s device API, and MSCCL++ expose programmable communication, while ParallelKittens and fused-kernel work target communication-computation overlap.Purlin uses the same Atom interface for direct kernel use and collectives derived through SNAC.Source: §7 · p. 16
• UCCL-Tran and UCCL-EP separate control and datapaths for RDMA settings, whereas Purlin applies decoupling within a scale-up domain.The separation is intended to let Purlin evolve its datapath without rewriting orchestration or sacrificing compatibility.Source: §7 · p. 16
• CUTLASS uses composable hardware-aware components and CuTe layouts for computation; Purlin layouts describe distributed input and output arrangements for collectives.The authors compare deriving collectives with SNAC and Atom to deriving high-performance computation with CUTLASS.Source: §7 · p. 16

8 Conclusion

• Purlin separates collective semantics, orchestration, and the hardware datapath, supporting independent evolution while preserving regime specialization and hardware-aware tuning.Source: §8 · p. 16

A.1 Rooted Collectives in SNAC

• Rooted collectives use a compile-time Root parameter and runtime root rank to specify which ranks produce or consume data, beyond the layouts’ data arrangement.Source: §A.1 · p. 21
• Broadcast can compose Scatter followed by AllGather, while Reduce can compose ReduceScatter followed by Gather; SNAC fuses and orchestrates each composition end-to-end.Source: §A.1 · p. 21
• For rooted Scatter, only the root stages and notifies; for Gather, all ranks stage and notify but only the root consumes, and tuning is still required.Source: §A.1 · p. 21

• Up to 1.47× higher algorithm bandwidth is achieved by Purlin with NVLS compared to SCORE, which guarantees deterministic reduction.Source: Figure 13 · p. 21, B SCORE vs. NVLink SHARP · p. 22
• At n = 8, SCORE moves 1.56× more bytes than NVLS: 1.75S versus 1.125S per rank.Source: B SCORE vs. NVLink SHARP · p. 22
• On H200, the corresponding peak algorithm bandwidths are 257 GB/s for SCORE and 400 GB/s for NVLS, using 450 GB/s nominal unidirectional NVLink bandwidth.Source: Table 7 · p. 21, B SCORE vs. NVLink SHARP · p. 22
• NVLS delivers higher bandwidth but lacks SCORE’s deterministic-reduction guarantee, so Purlin exposes both techniques and defers the choice to the application.Source: B SCORE vs. NVLink SHARP · p. 22

• The appendix covers configurations for inference experiments and collective microbenchmarks.Source: C Additional Evaluation Details · p. 22
• Additional microbenchmark results examine peer-accessible buffers.Source: C Additional Evaluation Details · p. 22
• The final appendix subsection presents results for variable-length collectives.Source: C Additional Evaluation Details · p. 22

C.1 Software and Application Configuration

• Application integrations share a specified SGLang revision and use PyTorch 2.13.0, CUDA 12.9, FlashInfer 0.6.18, and Transformers 5.12.1.Source: §C.1 · p. 22
• Offline serving uses fixed-length random requests in three input/output token configurations, while online serving uses the first 1,000 Mooncake conversation requests.Source: §C.1 · p. 22
• Image generation uses Ulysses sequence parallelism across eight GPUs for 50 denoising steps; AIME26 enables thinking and allows up to 131,072 output tokens.Source: §C.1 · p. 22

C.2 Collective Buffer Configuration

• For collective buffers, NCCL uses symmetric peer-accessible allocation settings when applicable, while B200 ordinary-buffer MSCCL++ skips a zero-copy algorithm under CUDA Graphs.Source: §C.2 · p. 22

C.3 Additional Peer-Accessible Buffer Results

• Figure 14 covers AllGather and AllReduce on eight A100 GPUs with peer-accessible buffers, showing latency and algorithm bandwidth.Latency is in the left column and algorithm bandwidth in the right.Source: Figure 14 · p. 22
• Figure 15 covers AllGather and AllReduce on eight H200 GPUs with peer-accessible buffers, reporting latency and bandwidth.Lower latency and higher bandwidth are better.Source: Figure 15 · p. 22
• Figure 16 covers variable-length collectives on eight B200, H200, and A100 GPUs using ordinary application buffers and a Zipf partition with s = 0.125.Message size is total T as defined in Eq. (4).Source: Figure 16 · p. 23

C.4 Variable-Length Collectives

• Partition sizes are passed at runtime; NCCL baselines emulate AllGatherV with broadcasts, ReduceScatterV with reductions, and AllToAllV with sends, receives, and a local copy.NCCLX uses TorchComms variable-length APIs.Source: §C.4 · p. 23
• The benchmark assigns rank r a Zipf-split share of total T across eight ranks with exponent s = 0.125; for T ≥ 8 KiB, the largest partition is 1.125× the mean.Splits preserve the total; 1 and 2 KiB totals use uniform partitions.Source: §C.4 · p. 23
• AllGatherV uses each rank’s assigned share as its contribution, ReduceScatterV produces that share as its output, and AllToAllV rotates partitions so each rank sends and receives T bytes.The same partitions are used across systems, and T is plotted for all three operations.Source: §C.4 · p. 23
• Latency improves over both baselines at every point in the 1 KiB–512 KiB sweep, with geometric-mean speedups of 1.91–2.73× over NCCL and 2.06–3.03× over NCCLX.These ranges cover nine collective–platform combinations.Source: §C.4 · p. 23
• Across platforms, bandwidth gains over NCCL range from 1.23–1.65× for AllGatherV, 1.09–1.29× for AllToAllV, and 1.57–2.00× for ReduceScatterV.Bandwidth speedups depend more strongly on the collective and message size.Source: §C.4 · p. 23
• At 2 MiB on A100, ReduceScatterV reaches 87.0 GB/s versus NCCL’s 20.4 GB/s and NCCLX’s 19.3 GB/s, gains of 4.26× and 4.50×.This is the largest individual gain reported.Source: §C.4 · p. 23
• At 1 GiB, NCCL leads ReduceScatterV on all three platforms; on H200, Purlin achieves 318.5 GB/s against NCCL’s 391.4 GB/s.The passage identifies this as the largest gap.Source: §C.4 · p. 23
• The results establish performance under mild skew; stronger imbalance requires tuning and a separate evaluation left for future work.Source: §C.4 · p. 23