Architecture & Compiler-Design Analysis
GC3: An Optimizing Compiler for GPU Collective Communication
Source: Cowan, M.; Maleki, S.; Musuvathi, M.;
Saarikivi, O.; Xiong, Y. GC3: An Optimizing Compiler for GPU
Collective Communication. arXiv preprint arXiv:2201.11840v3
[cs.DC], 19 Jul 2022 (revision 3). 13 pages, 8 figures.
URL: https://arxiv.org/abs/2201.11840
Code: https://github.com/microsoft/msccl-tools
and https://github.com/microsoft/msccl
Authors: Microsoft Research (Redmond) + Microsoft
Research Asia. Reader: Direct PDF read (gemini-reader
RESOURCE_EXHAUSTED on free-tier daily quota; codex-reader rejected
gpt-5.1-codex-mini for ChatGPT free-tier accounts; pages
1-13 read directly via the Read tool with pages parameter).
Analyst: Vishwakarma Date:
2026-05-04
Table of Contents
- Lineage Note — GC3 is the Preprint of MSCCLang
- Compiler/System Architecture (the "instrument" — the GC3 pipeline)
- Target-Hardware Architecture (Type-A 8xA100/HDR-IB and Type-B 16xV100/DGX-2)
- Design-Space Diagram (collectives x topologies x message sizes x knobs swept)
- The GC3 DSL — Chunk-Oriented Programming Model
- The Three IRs — Chunk DAG, Instruction DAG, GC3-IR
- Algorithm / Control Flow Diagrams (Tracing, Lowering, Fusion, Scheduling, Runtime Interpreter)
- Quantitative Results — Empirical Findings by Regime
- Configuration-Regime Trade-off Tables
- Bottlenecks & Insights Surfaced by the Measurements
- Limitations of the Methodology
- What to Borrow for DynamICCL — Compile-Time Graph Fixing vs Run-Time Knob Selection
- Analogy
1. Lineage Note — GC3 is the Preprint of MSCCLang
GC3 (arXiv:2201.11840, January 2022 first version, July 2022 v3) is the research-preprint identity of the system that was published one year later as MSCCLang at ASPLOS '23 (paper 0033 in this corpus). The authors, the architecture (DSL -> Chunk DAG -> Instruction DAG -> GC3-IR -> interpreter), and the core algorithms (Ring AllReduce, All-Pairs AllReduce, Hierarchical AllReduce, Two-Step AllToAll, AllToNext) are identical. The terminology rename is one-to-one:
+---------- GC3 (arXiv 2022) <--> MSCCLang (ASPLOS 2023) ----------+
| |
| GC3 DSL <==> MSCCLang DSL |
| GC3 program <==> MSCCLang program |
| GC3-IR <==> MSCCL-IR |
| GC3 runtime <==> MSCCL runtime |
| github.com/microsoft/msccl == same repo (compiler renamed) |
| github.com/microsoft/msccl-tools == same repo (DSL frontend) |
+------------------------------------------------------------------+
^ Fig 0: Identity map between the GC3 preprint and the MSCCLang
ASPLOS paper. The GC3 preprint includes one named collective —
AllToNext — that the ASPLOS version retains under the same name.This block-diagram analysis treats GC3 as a paper in its own right, using the GC3 paper's own terminology and figure references throughout. Where the framing benefits from explicit lineage to SCCL [4] and TACCL, we use citation [4] (Cai et al.) for SCCL and refer to TACCL by name. The "What to Borrow for DynamICCL" section below specializes the compile-time-vs-runtime composition argument to GC3's specific measurements and language affordances.
2. Compiler/System Architecture (the "instrument" — the GC3 pipeline)
GC3 is a three-stage compiler with an interpreter runtime, glued together by a Python-embedded DSL that is traced (not parsed) into a graph IR. Where SCCL [4] synthesizes schedules from a topology and collective spec via SMT, and TACCL guides synthesis with a human-supplied sketch, GC3 takes a fundamentally different stance: the algorithm is hand-written by the user in a high-level Python DSL, and the compiler's job is purely to lower it correctly and efficiently onto GPU thread blocks. Synthesis is out of scope; programmability is in scope. This trade is what lets the same compiler emit code for both intra-node Ring AllReduce on 8xA100 and inter-node Two-Step AllToAll on 256xA100 without re-running an SMT solver per topology.
+---------------------------------------------------------------------+
| GC3 Pipeline |
| |
| +-----------------------------+ +----------------------------+ |
| | GC3 DSL (Python eDSL) | | Scheduling Directives | |
| | - chunk(rank,buffer,index, | | - parallelize(N) | |
| | count) -> c | | - channel ch=k | |
| | - c1.copy(rank,buf,idx) | | - aggregation count=N | |
| | - c1.reduce(c2) | | - protocol (LL/LL128/Simple)| |
| | - signature, postcondition | | (Sec 5.1) | |
| | - hierarchical AllReduce in | | | |
| | ~30 LOC (Fig 3a) | | | |
| +--------------+---------------+ +-------------+--------------+ |
| | | |
| v v |
| +---------------------------------------------------------------+ |
| | STAGE 1: TRACING (Sec 4.1) | |
| | Sequentially execute the Python program; record every | |
| | source/copy/reduce as a Chunk DAG node. | |
| | Edges = chunk-data dependencies (true) + buffer-index reuse | |
| | (false dependencies). | |
| | Output: Chunk DAG (global parallelism view, exact semantics) | |
| +-------------------------+-------------------------------------+ |
| | |
| v |
| +---------------------------------------------------------------+ |
| | STAGE 2: LOWERING (Sec 4.2 -- "Instruction Generation") | |
| | Each chunk-DAG node -> 1..2 instruction nodes: | |
| | local copy -> copy(srcBuf,srcInd,dstBuf,dstInd) | |
| | local reduce -> reduce(srcBuf,srcInd,dstBuf,dstInd) | |
| | remote copy -> [send] + [recv] | |
| | remote reduce-> [send] + [recvReduceCopy (rrc)] | |
| | Adds COMMUNICATION EDGES (cross-rank receive-after-send) | |
| | preserves PROCESSING EDGES (intra-rank order from Chunk DAG).| |
| | Output: Instruction DAG | |
| +-------------------------+-------------------------------------+ |
| | |
| v |
| +---------------------------------------------------------------+ |
| | STAGE 2b: INSTRUCTION FUSION (Sec 4.2, peephole) | |
| | rcs : recv + send-on-same-chunk -> receiveCopySend | |
| | rrcs: rrc + send -> receiveReduceCopySend | |
| | rrs : if rrc result not used locally -> receiveReduceSend | |
| | -> eliminates intermediate global-memory writes; fused | |
| | instructions transfer values through GPU registers. | |
| +-------------------------+-------------------------------------+ |
| | |
| v |
| +---------------------------------------------------------------+ |
| | STAGE 3: SCHEDULING (Sec 5) | |
| | 1. Channel Assignment (lowest legal channel by default) | |
| | 2. Thread Block Creation (one TB per (sendPeer,recvPeer,ch))| |
| | 3. Priority Sort (depth + reverse-depth heuristic) | |
| | 4. TB Assignment (greedy ready-first into TBs) | |
| | 5. Cross-TB Sync (processing edges that cross TBs | |
| | -> semaphore dependencies) | |
| | Output: GC3-IR (XML-like tree, per-GPU per-TB instr lists) | |
| +-------------------------+-------------------------------------+ |
| | |
| v |
| +---------------------------------------------------------------+ |
| | GC3 RUNTIME (Sec 6, Fig 5) | |
| | Single CUDA cooperative-kernel launch, one thread block per | |
| | GC3-IR thread block, interpreter loop: | |
| | for tile in chunkTiling : | |
| | for instr in instrList : | |
| | if instr.hasDep : wait(semaphore, depBid, depStep) | |
| | switch (instr.opCode) { send / recv / rrc / rrs ... } | |
| | if instr.hasDep : threadfence + sync_threads + set sem | |
| | - Uses NCCL FIFO slots (s = 1..8 of size 512 KB..5 MB) | |
| | - Reuses NCCL transport (NVLink, PCIe, SHM, IB GPUDirect RDMA| |
| | over HDR IB, TCP) | |
| | - Picks Simple / LL128 / LL protocol per buffer-size range | |
| | - Falls back to NCCL's built-in algorithms for sizes outside | |
| | user-configured ranges (Sec 6: "the user can hyper-optimize| |
| | GC3 programs to a specific use case"). | |
| +---------------------------------------------------------------+ |
+---------------------------------------------------------------------+
^ Fig 1: GC3 pipeline. Three stages (trace -> lower -> schedule) emit
GC3-IR which the runtime interpreter executes as a single cooperative
CUDA kernel. The runtime extends NCCL 2.8.4-1 — GC3 is a co-tenant
of NCCL, not a replacement. The compiler is **API-compatible with
NCCL** so existing ML workloads can switch over with no code change
(Sec 1).Two architectural choices in this pipeline are load-bearing for the
rest of the analysis. First, GC3 is a traced eDSL, not a parsed
DSL. The user writes Python that is executed, and the
trace is captured as the Chunk DAG. This means programs benefit from
Python's control flow (loops, list comprehensions,
with parallelize(N): context managers) without the compiler
having to understand any of it — the compiler only sees the resulting
graph. The same architectural choice that powers PyTorch's eager-mode
tracing is reused here. Second, the runtime is an interpreter,
not a code generator. GC3-IR is parsed once at init and the
per-thread-block instruction list is walked by a CUDA switch statement
(Fig 5 of the paper). This is the opposite of SCCL [4], which emits a
custom CUDA kernel per algorithm. The interpreter overhead is amortized
across many tiles (the inner pipelining loop) and the flexibility wins
are large: the same kernel runs every algorithm.
Lower-layer +-------------------------+ Upper-layer
inputs | GC3 Compiler+Runtime | output
+-------------+-----------+
|
v
+-------------------------+
| NCCL transport API | "GC3 is built on top
| - GPUDirect P2P / NVLink| of NCCL-2.8.4-1"
| - GPUDirect RDMA / IB | (Sec 7: Experimental
| - SHM / PCIe / TCP | Setup)
| - Simple / LL128 / LL |
| - FIFO slots (1<=s<=8) |
+-------------+-----------+
|
v
+-------------------------+
| Type-A: 8xA100 + 12 NVL3|
| + 6 NVSwitch + |
| + 2 HDR IB NICs |
| Type-B: 16xV100 (DGX-2) +|
| + 6 NVL2 + 6 NVS|
| + 1 HDR IB NIC |
+-------------------------+
^ Fig 2: GC3's stack position. The compiler+runtime pair sits *above*
NCCL's transport layer and *replaces* NCCL's algorithm layer.
Critically, GC3 inherits NCCL's protocol selector (LL/LL128/Simple)
and FIFO machinery — meaning the same NCCL knob space DynamICCL
operates on is *available* to GC3 programs (and is exposed through
DSL scheduling directives).The verbatim relationship to NCCL is established in Sec 1:
"Lastly, the runtime is API-compatible with NCCL allowing existing ML workloads to easily switch over to GC3, inherit NCCL's support of diverse set GPUs and interconnections, and safely fall over to NCCL kernels for yet unimplemented algorithms in GC3 to enable safe operation in production."
This is the strongest statement in the paper for DynamICCL's purposes. GC3 explicitly preserves NCCL's full transport layer (NVLink, PCIe, SHM, IB, TCP), inherits its protocol selector (LL/LL128/Simple), and falls back to NCCL's algorithms for cases where no GC3 program is provided. DynamICCL — which lives below the algorithm layer, inside NCCL's tuner-plugin slot — therefore composes with GC3 by default: when GC3 dispatches into NCCL fallback, DynamICCL tunes the NCCL algorithm; when GC3 runs its own program, DynamICCL can override the runtime-configurable knobs (protocol, tile size, parallelize factor) that GC3's DSL exposes as defaults.
3. Target-Hardware Architecture
GC3 is evaluated on two distinct multi-GPU systems (Sec 7, "Experimental Setup"). The two single-node systems differ on every consequential dimension: GPU generation, NVLink topology, NVSwitch count, and IB NIC count. This makes the evaluation's regime axis explicit: A100/NVSwitch is the modern "flat" intra-node interconnect; V100/NVLink-fabric is the older DGX-2-style interconnect with two boards bridged by NVSwitch crossbars.
3.1 Type-A — 8xA100 per node + 2 HDR IB NICs per pair
+---- Type-A node: 8 NVIDIA A100 + 6 NVSwitch + 2 HDR IB NICs --------+
| |
| A100-0 A100-1 A100-2 A100-3 A100-4 A100-5 A100-6 A100-7 |
| | | | | | | | | |
| +-------+-------+-------+---+---+-------+-------+-------+ |
| | |
| 12 third-generation NVLinks per GPU |
| | |
| +---+---+---+---+---+---+ v +---+---+---+---+---+---+ |
| | NVS-0 | NVS-1 | NVS-2 | NVS-3 | NVS-4 | NVS-5 | |
| +---+---+---+---+---+---+---+---+---+---+---+---+ |
| | |
| Total bisection: 600 GB/s bi-directional intra-node |
| | |
| +-----------------+ +-----------------+ |
| | PCIe Switch (4G)| | PCIe Switch (4G)| ... |
| +--------+--------+ +--------+--------+ |
| | | |
| +--------+-------+ +--------+-------+ |
| | 2 x HDR IB NIC | | 2 x HDR IB NIC | per GPU pair |
| | 25 GB/s each | | 25 GB/s each | |
| +-----------------+ +-----------------+ |
| |
+---------------------------------------------------------------------+
|
Cross-node: HDR InfiniBand (25 GB/s per NIC)
|
+---------------- Type-A cluster ------------------+
| 16 nodes x 8 GPU = 128 GPU at 1-node tests |
| Multi-node tests: 2-node 16xA100, 3-node 24xA100,|
| 16-node 256xA100 |
+--------------------------------------------------+
^ Fig 3: Type-A hardware. 8 A100 GPUs per node, fully connected by 12
third-generation NVLinks each via 6 NVSwitches yielding 600 GB/s
bi-directional bandwidth per GPU. Each PAIR of GPUs shares a single
PCIe Switch connected to 2 HDR IB NICs (so each GPU has effective
access to 25 GB/s of cross-node bandwidth via shared NIC pair). The
largest cluster experiment uses 16 such nodes for 256 A100 GPUs.3.2 Type-B — 16xV100 DGX-2 with two-board NVLink fabric
+---- Type-B node: 16xV100 in two boards of 8 (DGX-2 style) ---------+
| |
| Board 0: V100-0..V100-7 Board 1: V100-8..V100-15 |
| |
| +----------------------+ +----------------------+ |
| | V100-0 ... V100-7 | | V100-8 ... V100-15 | |
| | | 6 NVLink2 each | | | 6 NVLink2 each | |
| | v | | v | |
| | NVS-A0..NVS-A5 (6) | | NVS-B0..NVS-B5 (6) | |
| +-----------+----------+ +----------+-----------+ |
| | | |
| +-----+ +---------------+ |
| v v |
| Each NVSwitch on board 0 has 8 NVLinks to its |
| counterpart NVSwitch on board 1 (the DGX-2 "fabric") |
| | |
| +--------------+---------------+ |
| | PCIe Switch (per GPU pair) | |
| | 1 HDR IB NIC, 25 GB/s | (1 per pair, shared) |
| +------------------------------+ |
| |
+---------------------------------------------------------------------+
|
Cross-node: HDR InfiniBand (25 GB/s per IB NIC)
|
+---------------- Type-B cluster ------------------+
| 1-node 16xV100, 2-node 32xV100, 4-node 64xV100 |
+--------------------------------------------------+
^ Fig 4: Type-B hardware. 16 V100 GPUs split across two boards of 8.
Each GPU has 6 second-generation NVLinks to 6 NVSwitches on its own
board. Cross-board comm goes through 8 NVLinks per NVSwitch to the
counterpart NVSwitch on the other board — i.e., the V100 fabric is
*NOT* fully flat (intra-board faster than cross-board). This is the
hierarchical-interconnect regime that benefits most from
"Hierarchical AllReduce" in Sec 7.2.3.3 Software stack
| Layer | Component |
|---|---|
| Application | ML workloads: production language model (8xA100), |
| Mixture-of-Experts (256xA100) | |
| Compiler+Runtime | GC3 (built on NCCL-2.8.4-1) [Sec 7] |
| Collective lib | NCCL 2.8.4-1 |
| Transports | NVLink (gen2/gen3), PCIe, SHM, GPUDirect RDMA over IB |
| GPU runtime | CUDA + cooperative kernel launch |
| Hardware | A100+NVSwitch+HDR-IB / V100+NVSwitch+HDR-IB |
The two-cluster setup is designed to exercise fundamentally different intra-node topology shapes. A100 (Type-A) is a single fully-connected NVSwitch fabric; V100 (Type-B) is a two-board hierarchical fabric. GC3 must therefore prove that the same DSL can express algorithms whose schedule is optimal on both topologies — which it does by exposing channel and parallelize directives that the user configures per-topology. This compile-time topology specialization is the design choice that DynamICCL must compose with: DynamICCL's runtime topology fingerprint signal must distinguish flat-A100 from hierarchical-V100, because the optimal NCCL knobs differ between the two even when the GC3 program is identical.
4. Design-Space Diagram (collectives x topologies x message sizes x knobs)
The independent variables form a five-dimensional sweep. Every panel of Fig 7 in the paper fixes the (collective, scale, hardware) tuple and varies (algorithm, parallelize factor r, protocol, buffer size); every named-collective row in Sec 7 fixes a collective and reads off the regimes where each algorithm dominates.
GC3 DESIGN SPACE (5 axes + held-fixed)
+---------------------------------------------------------------+
| |
| Axis 1: COLLECTIVE (4 levels) |
| [AllReduce] |
| [AllToAll] |
| [AllToNext] <- novel custom collective |
| (AllGather) <- used as helper inside Hierarchical AR |
| |
| Axis 2: ALGORITHM (per collective) |
| AllReduce : {Ring, All-Pairs, Hierarchical} |
| AllToAll : {Two-Step, naive baseline} |
| AllToNext : {GC3 r-parallel, hand CUDA baseline} |
| |
| Axis 3: HARDWARE x SCALE (8 cells from Fig 7) |
| 1 node x 8 A100 (Fig 7a) |
| 1 node x 16 V100 (Fig 7b) |
| 2 node x 16 A100 (Fig 7c) |
| 2 node x 32 V100 (Fig 7d) |
| 16 node x 256 A100 (Fig 7e -- AllToAll) |
| 4 node x 64 V100 (Fig 7f -- AllToAll) |
| 3 node x 24 A100 (Fig 7g -- AllToNext) |
| 4 node x 64 V100 (Fig 7h -- AllToNext) |
| |
| Axis 4: BUFFER SIZE (16-22 levels, log-spaced) |
| AR small : 1 KB, 2 KB, 4 KB ... 32 MB |
| AR large : 256 KB ... 4 GB (multi-node) |
| A2A : 256 KB ... 4 GB |
| AllToNext: 4 KB ... 256 MB |
| |
| Axis 5: KNOBS (within-program, per Fig 7 legend) |
| parallelize r in {1, 2, 4, 8, 16, 24} |
| channel ch in {1, 4, 8} |
| protocol in {LL, LL128, Simple} |
| |
| Held FIXED (no sweep): |
| - NCCL version: 2.8.4-1 (fall-back paths) |
| - Cluster topology: 2 node types only (Type-A, Type-B) |
| - GPU model: A100 (Type-A) or V100 (Type-B) only |
| - Transport: NVLink + IB; no Ethernet/RoCE measured |
| - Synchronization: BSP (no SSP/ASP) |
| - Compression: NONE (lossless only) |
| - Datatype: not explicitly varied |
| - Optimizer: workload-defined (e.g. Adam for MoE) |
| - Tile size: discussed but not explicitly swept in plots |
| |
+---------------------------------------------------------------+
^ Fig 5: 5-axis design space — 4 x ~5 x 8 x ~20 x ~6 = ~19,200 logical
cells (though the paper plots a curated subset of ~150 -- the most
informative regime curves). NCCL's internal algorithm/protocol/nCh/
numThreads are NOT swept independently inside fallback; they are
surfaced into GC3 as DSL directives the user manually tunes.Two absences define the paper's measurement scope. First, GC3 fixes NCCL 2.8.4-1 as the fallback baseline — this is the only NCCL version compared against, so any speedup numbers in this paper compare to that specific NCCL release. Second, the buffer-size dispatch ranges that select among GC3 program variants are user-configured at compile time, not learned at runtime. The user empirically measures the crossover points in Fig 7 and bakes them into a static lookup table in the runtime. This is exactly the manual analog of what an RL agent like DynamICCL would learn dynamically. Sec 6 makes this explicit:
"All GC3-IR generated by our compiler is guaranteed to be correct, but some programs might only be performant for a range of buffer sizes. Therefore, the runtime dynamically selects the right algorithm to invoke based on user configurable size ranges and falls back to NCCL's built-in algorithms otherwise. This allows a user to hyper-optimize GC3 programs to a specific use case." (Sec 6)
For DynamICCL, this means the very mechanism GC3 exposes for buffer-size-conditional algorithm selection is the natural insertion point for an RL policy: replace the "user-configurable size ranges" with a learned policy that conditions on a richer state vector (msg size + topology fingerprint + recent timing window).
5. The GC3 DSL — Chunk-Oriented Programming Model
The DSL is described in Sec 3. Its key design choice is to manipulate chunk references rather than chunks directly, with the compiler enforcing that only the latest reference for any (rank, buffer, index) location can be operated on. This makes programs data-race free by construction (Sec 3.3), which lets the compiler do aggressive scheduling without correctness checks.
+------------------------------------------------------------------+
| GC3 DSL Surface |
| |
| Buffers (3 per rank): |
| +------------+ +------------+ +------------+ |
| | input | | output | | scratch | |
| | (init data)| | (result --| | (temp, | |
| | | | empty at | | uninit) | |
| | | | start) | | | |
| +------------+ +------------+ +------------+ |
| |
| Chunk forms (Sec 3.1): |
| - Input chunks : initial elements of input buffers |
| - Reduction chunks : combinations of input chunks via |
| point-wise operator (e.g. addition) |
| - Uninitialized chunks: type-only, hold no data |
| |
| Operations (Table 1): |
| chunk(rank,buf,idx,count=C) -> c (reference, count=1 def) |
| c1.copy(rank2,buf,idx2) -> c2 (returns copy reference) |
| c1.reduce(c2) -> c3 (in-place into c1's idx) |
| |
| Collective signature & postcondition (Sec 3.2): |
| signature = (#input chunks, #output chunks, in_place?) |
| postcondition = for each output chunk, either: |
| - an input chunk (passthrough), OR |
| - a reduction chunk (combination) |
| -> compiler validates that the prospective algorithm |
| actually implements the collective claim (correctness gate)|
| |
| Scratch buffer size: not user-supplied; GC3 deduces from |
| highest indices accessed in the program. |
+------------------------------------------------------------------+
^ Fig 6: GC3 DSL surface from Sec 3. Three primitives (chunk, copy,
reduce) plus pre/post-condition contracts give a small, total
language for chunk routing — small enough to be data-race-free
by construction, expressive enough to encode every algorithm in
Sec 7.The hierarchical AllReduce in Fig 3a (paper) takes ~30 LOC including its helper functions (Fig 3b). It has four phases — intra-node ReduceScatter, inter-node ReduceScatter, inter-node AllGather, intra- node AllGather — each implemented by a Ring algorithm. This same program is compiled and run on both Type-A (8xA100) and Type-B (16xV100) clusters; only the topology-specific scheduling directives (channel, parallelize) differ.
6. The Three IRs — Chunk DAG, Instruction DAG, GC3-IR
GC3's three IRs at decreasing levels of abstraction (Fig 4 of the paper) capture the lowering pipeline cleanly:
+-----------------------------------------------------------------+
| IR 1: CHUNK DAG (output of Sec 4.1 Tracing) |
| +-----------------------------------------------------------+ |
| | Nodes : copy / reduce operations (one per DSL call) | |
| | Edges : | |
| | - true (data) deps from chunk movement | |
| | - false deps from buffer-index reuse | |
| | Source nodes = all input chunks | |
| | Captures: global parallelism view of program semantics | |
| | Audience: humans reasoning about the algorithm | |
| +-----------------------------------------------------------+ |
| | |
| v |
| IR 2: INSTRUCTION DAG (output of Sec 4.2 Lowering + Fusion) |
| +-----------------------------------------------------------+ |
| | Nodes : low-level instructions | |
| | {send, recv, copy, reduce, rrc, rcs, rrcs, rrs} | |
| | Edges : | |
| | - communication edges (send <-> matching recv) | |
| | - processing edges (preserved Chunk DAG order) | |
| | Captures: how chunks become wire-level operations, | |
| | with cross-rank communication explicit | |
| | Audience: compiler optimization passes | |
| +-----------------------------------------------------------+ |
| | |
| v |
| IR 3: GC3-IR (output of Sec 5 Scheduling, Fig 4 right panel) |
| +-----------------------------------------------------------+ |
| | Tree structure: | |
| | Program | |
| | | | |
| | +--> GPU 0 | |
| | | +--> Threadblock 0 | |
| | | | +--> Connections: send peer, recv peer, | |
| | | | channel | |
| | | | +--> Instructions [list, sequential] | |
| | | +--> Threadblock 1 | |
| | | +--> Threadblock 2 | |
| | +--> GPU 3 | |
| | | +--> Threadblock 0, 1, ... | |
| | +--> GPU 5 | |
| | Constraints (Sec 5): | |
| | - At most 1 send + 1 recv connection per TB | |
| | - At most 1 sending TB + 1 receiving TB per connection| |
| | - Total TB count <= available SMs (cooperative kernel) | |
| | Audience: GC3 runtime interpreter | |
| +-----------------------------------------------------------+ |
+-----------------------------------------------------------------+
^ Fig 7: The three IRs. Chunk DAG is the algorithm spec; Instruction
DAG is the wire-level expansion; GC3-IR is the executable thread-
block schedule. Each lowering step adds *strictly more constraint*
while preserving the semantic invariants.The three-IR structure mirrors a classic compiler decomposition: high-level intent (Chunk DAG) -> mid-level operations (Instruction DAG) -> low-level scheduling (GC3-IR). Crucially, scheduling decisions (channel, TB assignment, sync points) are deferred to the lowest IR; this means the user's high-level program is portable across topologies, and only the Sec 5 scheduling directives need to be re-tuned per hardware. For DynamICCL, this is a clean separation between the algorithm and the parameters: the algorithm (Chunk DAG and Instruction DAG) is fixed at compile time; the parameters (channel assignment, parallelize factor, tile size, protocol) appear at the GC3-IR layer and below — exactly where DynamICCL's action space lives.
7. Algorithm / Control Flow Diagrams
7.1 Tracing — Chunk DAG generation (Sec 4.1)
START (a GC3 program written in Python)
|
v
(1) GC3 imports DSL primitives, intercepts copy/reduce calls
|
v
(2) Sequentially execute the Python program body
|
v
(3) For each chunk(rank,buf,idx,count) call:
create a Chunk DAG source node OR retrieve existing
|
v
(4) For each c.copy(...) or c.reduce(...) call:
- validate that c is the LATEST reference (data-race check)
- create a Chunk DAG node
- add true-data edges from operand source nodes
- add false-buffer-index-reuse edges to overwriters
- return a fresh reference to the new chunk
|
v
(5) Drop unused references; finalize Chunk DAG
|
v
END --> Chunk DAG ready for lowering
^ Fig 8: Tracing control flow. The Python program is executed sequen-
tially; the compiler intercepts every DSL call and grows the Chunk
DAG one node at a time. Crucially, control flow (loops, if-then-
else) is handled by Python at trace time -- the compiler never sees
it. This is the same trick PyTorch eager-mode uses.7.2 Lowering — Instruction Generation + Fusion (Sec 4.2)
INPUT: Chunk DAG (from tracing)
|
v
(1) For each chunk-DAG node, classify by locality:
- local copy -> 1 instr: copy
- local reduce -> 1 instr: reduce
- remote copy -> 2 instr: send + recv
- remote reduce-> 2 instr: send + recvReduceCopy (rrc)
|
v
(2) Add COMMUNICATION EDGES between matching send/recv pairs
(cross-rank dependencies in the new Instruction DAG)
|
v
(3) Preserve PROCESSING EDGES from Chunk DAG
(intra-rank order from the DSL trace)
|
v
(4) PEEPHOLE FUSION pass (Fig 4 highlighted in green):
scan adjacent instr pairs:
recv + send-of-same-chunk -> rcs (recvCopySend)
rrc + send-of-result -> rrcs (recvReduceCopySend)
rrc with no local consumer + send -> rrs (recvReduceSend)
each fusion REMOVES intermediate global-mem write/read
values transferred via GPU registers instead
|
v
(5) If multiple sends depend on a single recv, fuse the longest-
path send into rcs first (paper Sec 4.2 last paragraph)
|
v
END --> Instruction DAG (with fused instructions where applicable)
^ Fig 9: Lowering + fusion control flow. Step 4 is the core
performance pass: each fused instruction saves one global-memory
round-trip per chunk, which is the dominant cost on saturated
NVLink/IB links.7.3 Scheduling — TB Assignment + Sync Insertion (Sec 5.2)
INPUT: Instruction DAG + scheduling directives (channel, parallelize)
|
v
(1) CHANNEL ASSIGNMENT
for each communication edge:
if user-specified channel directive: use it
elif edge is from a parallelize fragment: use its set channels
elif edge is in a fused chain: use lowest channel of chain
else: use lowest channel that doesn't violate the
"1 send peer + 1 recv peer per TB" invariant
|
v
(2) THREAD BLOCK CREATION
scan all instructions per GPU
for each unique (sendPeer, recvPeer, channel) tuple:
create one thread block
|
v
(3) PRIORITY SORT
compute dependency depth d(v) = longest path from source
compute reverse dep depth r(v) = longest path to sink
sort instructions by (d ascending, r descending) into a
GLOBAL TOPOLOGICAL ORDER
|
v
(4) TB ASSIGNMENT (greedy)
for each instruction in priority order:
find the set of TBs whose (sP,rP,ch) match the instr
if multiple candidates: pick the TB whose latest assigned
instruction is earliest in the topological order
assign instruction to that TB
|
v
(5) SYNCHRONIZATION INSERTION
for each processing edge that crosses TB boundaries:
insert a semaphore-based wait/set pair in GC3-IR
(intra-TB processing edges are implicitly sequential)
|
v
END --> GC3-IR (deadlock-free by construction; Sec 5.2 last para)
^ Fig 10: Scheduling control flow. The greedy TB assignment is
deadlock-free because the topological order respects all comm and
processing edges; "thread block sequential execution" cannot
introduce cycles into the IR.The deadlock-freedom argument (Sec 5.2 last paragraph) is the correctness keystone for the whole compiler:
"All Instruction DAGs are guaranteed to have a global topological order because it was generated by sequentially tracing the GC3 program. By assigning instructions to thread blocks in a topological order that respects communication and processing edges, all implicit dependencies introduced by thread block sequential execution cannot produce cycles so that the GC3-IR does not have deadlocks." (Sec 5.2)
This is the same argument SCCL [4] makes with SMT, but achieved via a much cheaper greedy heuristic that exploits the trace order.
7.4 Runtime Interpreter (Sec 6.2, Fig 5)
INPUT: GC3-IR file + buffer pointers
|
v
(1) PARSE GC3-IR; load instruction lists per TB into GPU memory
|
v
(2) Initialize NCCL P2P connections for every (sendPeer, recvPeer,
channel) tuple required by the program (Sec 6.1)
Allocate NCCL FIFO slots: s slots of size b
512 KB <= b <= 5 MB ; 1 <= s <= 8
|
v
(3) Launch single CUDA cooperative kernel
cudaLaunchCooperativeKernel(GC3_interpreter, ...)
|
v
(4) Inside each thread block (bid = blockIdx.x):
OUTER LOOP -- chunk tiling:
for (t=0; t<chunkSize; t += tileSize) {
for each instr in instrList[bid] {
[a] if instr.hasDep:
wait(semaphore[instr.depBid[tid]],
instr.depStep[tid])
[b] switch (instr.opCode):
case SEND: send(srcPtr+srcOff, count*tileSize)
case RECV: recv(...)
case RRC : recvReduceCopy(...)
case RCS : recvCopySend(...)
case RRS : recvReduceSend(...)
case RRCS: recvReduceCopySend(...)
case COPY: locallyCopy(...)
case REDUCE: locallyReduce(...)
[c] if instr.hasDep:
threadfence + sync_threads + set sem to step s
}
}
|
v
(5) Kernel exits; output buffers contain collective result
|
v
END
^ Fig 11: Runtime interpreter (Fig 5 of the paper). The outer chunk-
tiling loop is what enables PIPELINING: tile k of phase i overlaps
in time with tile k-1 of phase i+1 (Fig 6 of paper). Tile size is
bounded above by the FIFO slot size b.The runtime is the smallest piece of GC3 by line count but the most
performance-critical: every fused instruction's bandwidth savings
multiplies by the number of tiles, and every cross-TB semaphore wait
serializes pipelining. The choice of b (NCCL FIFO buffer
size) and tileSize together determine pipelining depth —
and both are runtime-configurable knobs that DynamICCL
can tune.
8. Quantitative Results — Empirical Findings by Regime
GC3 evaluates four collectives across eight (cluster, scale) cells in Fig 7. Results are reported as speedup over a baseline that varies by regime (NCCL, hand-optimized CUDA, or both).
8.1 AllReduce — Single Node (Fig 7a-b)
1-node 8xA100, AllReduce (Fig 7a):
buffer 1KB - 8KB : All Pairs r=2 LL ~= NCCL (no clear winner)
buffer 16KB-256KB : GC3 Ring ch=8 LL up to 1.9x faster than NCCL
buffer 512KB-1MB : GC3 All Pairs r=4 LL128 best at ~1.5x
buffer 2MB - 8MB : crossover region; GC3 Ring ch=4 LL near tie
buffer >32MB : NCCL Ring matches GC3 Ring (saturated link)
1-node 16xV100, AllReduce (Fig 7b):
buffer 16KB-128KB : GC3 All Pairs peak ~3.0x faster than NCCL
buffer 256KB-1MB : GC3 Ring ch=8 LL128 strong (~2.0x at 256KB)
buffer >2MB : NCCL closer to GC3The peak speedup of 1.9x at 32 KB - 256 KB on A100 is the paper's headline AllReduce number. The mechanism is explained in Sec 7.1.1:
"Our ring implementation distributes a single logical ring across multiple channels by varying the channel of copy and reduce operations. We tune the number of channels per ring, parallelization, and protocol for the system. We compare our Ring implementations against NCCL's Ring implementation in Figure 7a. While examining NCCL's codebase, we found and experimentally validated that NCCL's Ring schedule is roughly equivalent to scheduling a logical ring onto one channel, parallelizing the entire program 24 times, and varying the protocol based on the buffer size."
So GC3's win comes from putting one logical ring across
ch=k channels with a smaller r
factor, instead of NCCL's "one logical ring on one channel with
r=24". For sizes >32 MB, NCCL recovers parity because chunk
parallelization can't help once links are saturated.
"The GC3 ring implementation outperforms NCCL by up to 1.9x when the buffer size is between 32KB and 3MB. ... For buffer sizes greater than 32MB, more parallelization is required, and the best GC3 configurations matched NCCL's performance by scheduling a logical ring onto one channel and parallelizing the program 24 times." (Sec 7.1.1)
8.2 AllReduce — All-Pairs (small-buffer specialist)
"For buffer sizes from 1KB to 1MB, All Pairs is up to 1.8x faster than NCCL, depending on the number of instances used to optimize the program." (Sec 7.1.2)
The All-Pairs AllReduce is a 2-step algorithm (vs Ring's 2R-2 steps, where R = ranks): every rank receives a chunk from every other rank, sums, and broadcasts to all. It's bandwidth-equivalent to Ring but latency-optimal — and GC3's contribution is enabling it via the DSL with only ~2 fewer steps, which dominates at small sizes where per-step latency is the cost.
8.3 AllReduce — Hierarchical (multi-node specialist; Sec 7.2, Fig 7c-d)
2-node 16xA100, AllReduce (Fig 7c):
buffer 1KB - 256KB : GC3 LL r=1 best, up to 1.4x over NCCL
(NCCL Hierarchical close behind)
buffer 256KB - 1GB : GC3 Simple r=4 ~ ties NCCL Hierarchical
buffer >1GB : GC3 11% faster than NCCL Hierarchical
2-node 32xV100, AllReduce (Fig 7d):
buffer 256KB - 4MB : GC3 Simple r=4 peak ~1.6x over NCCL
(NCCL Hierarchical also strong here)
buffer >32MB : NCCL recovers parityThe headline statement (Sec 7.2):
"For small sizes we are up to 1.4x faster than NCCL. For large buffers, greater than 1GB, our implementation is up to 11% faster than NCCL."
The Hierarchical algorithm is specifically what the V100 DGX-2's two-board fabric and the multi-node IB topology demand: separate intra-node and inter-node phases avoid mixing the two interconnect types in a single ring, which would clog the slower link. A naive Ring across 32 GPUs spanning 2 nodes is bottlenecked by the IB hop; Hierarchical confines IB hops to the smaller inter-node ReduceScatter
- AllGather phases.
8.4 AllToAll — Two-Step (Sec 7.3, Fig 7e-f)
The Two-Step AllToAll algorithm aggregates cross-node sends into a single coalesced IB transfer per pair:
GC3 Two-Step AllToAll (Fig 8 of paper, ~16 LOC):
Phase 1 (intra-node copies): for each (n, g) -> (m, g):
c.copy(m, 'sc', g) for all m != n [Line 11]
c.copy(n, 'in', m=n) [Line 9]
Phase 2 (cross-node coalesced send + intra-node distribute):
c = chunk(m,g,'sc',n*G,sz=G) [Line 14]
c.copy(n, g, 'out', m*G) [Line 16]
16-node 256xA100, AllToAll (Fig 7e):
buffer 256KB - 4MB : GC3 Two-Step LL128 peak >2x over CUDA TS
but with 2-64MB high-variance region
(cloud IB cross-tenant noise)
buffer >512MB : GC3 1.3x faster than hand-optimized
GC3 1.2x faster than NCCL
Hand-optimized SLOWER than NCCL at >512MB
4-node 64xV100, AllToAll (Fig 7f):
buffer 4MB - 32MB : GC3 ~1.2x over hand-optimized
buffer >256MB : GC3 ~ ties hand-optimizedThe headline statement (Sec 7.3):
"At large sizes the GC3 implementation is up to 1.3x faster than the hand-optimized implementation. ... Note, at smaller sizes between 2MB-64MB there are large fluctuations in speedup caused by congestion in the IB network which is shared with other cloud tenants; however the general trends show that GC3's optimizations improve performance."
The mechanism (Sec 7.3 last paragraph):
"The hand-optimized version is implemented using point-to-point primitives exposed by NCCL, but lacks scheduling decisions made by the compiler that decides communication across multiple parallel thread blocks. The GC3 seamless handles aggregating chunks in the scratch buffer (Line 12), while the handwritten implementation requires a separate kernel that copies and contiguously arranges chunks in a scratch buffer for the aggregated IB send resulting in extra synchronization overhead. Furthermore, the GC3 implementation is much succinct and requires only 15 lines of code while the hand optimized kernel requires roughly 70 lines of code."
15 LOC vs 70 LOC, with a 1.3x performance win — this is the productivity-and-performance dual claim of the paper.
8.5 AllToNext — Custom Collective (Sec 7.4, Fig 7g-h)
AllToNext is GC3's novel custom collective: each GPU i
sends a buffer to GPU i+1 (last GPU sends nothing). This
pattern arises in pipeline-parallel inference where a
model is sharded across GPUs.
3-node 24xA100, AllToNext (Fig 7g):
buffer 4KB - 64KB : GC3 r=4 best (small init cost)
buffer 1MB - 16MB : GC3 r=8 surpasses CUDA baseline
buffer 32MB - 256MB: GC3 r=16 up to 14.5x over hand CUDA
4-node 64xV100, AllToNext (Fig 7h):
similar trends; up to ~5x at 32MB - 256MBThe headline statement (Sec 7.4):
"AllToNext begins to show improvement over the baseline, and is ultimately up to 14.5x for a large buffers. The best performing selection of r depends buffer sizes. For small buffer sizes, less parallelization provide better performance, as the benefit from parallelizing communication doesn't offset the cost of initializing extra resources."
The mechanism: a single GPU-to-GPU send is bottlenecked by the
single PCIe-shared IB NIC between the sender's PCIe
switch and the receiver's PCIe switch. By splitting the buffer across
all G GPUs in the sender's node, GC3 routes pieces over
G different IB NICs to corresponding GPUs on the receiver
node, then a final intra-node copy assembles the result. This is
fan-out across IB NICs that the naive 1-NIC baseline leaves
unused — and explains the 14.5x figure (the cluster has up to
~16 IB NICs available depending on topology specifics).
8.6 End-to-End Production Workloads (Sec 7.5)
"GC3 is currently used in inferencing a public facing language model on cloud service provider X on 8xA100 GPUs with 1.22-1.29x total GPU time speed up depending on the batch size used. GC3 is also used for training a large Mixture-of-Experts model for speech, language, and vision on 256xA100 GPU providing 1.10-1.89x speed up depending on the model architecture."
These are two real production workloads — a public LM (inference) and a giant MoE (training) — with concrete time-to-solution speedups. The 1.22-1.29x for LM inference comes from accelerating AllReduce inside model-parallel decoding; the 1.10-1.89x for MoE training comes from accelerating AllToAll for expert routing.
8.7 Summary of headline speedups
| Collective | Cluster | Buffer regime | GC3 speedup vs baseline |
|---|---|---|---|
| AllReduce Ring | 1n 8xA100 | 32 KB - 3 MB | up to 1.9x vs NCCL |
| AllReduce Ring | 1n 8xA100 | > 32 MB | matches NCCL (saturated) |
| AllReduce All-P | 1n 8xA100 | 1 KB - 1 MB | up to 1.8x vs NCCL |
| AllReduce Hier | 2n 16xA100 | small (< 1 MB) | up to 1.4x vs NCCL |
| AllReduce Hier | 2n 16xA100 | > 1 GB | 11% faster vs NCCL |
| AllToAll 2-Step | 16n 256xA100 | > 512 MB | 1.3x vs hand CUDA |
| AllToAll 2-Step | 16n 256xA100 | > 512 MB | 1.2x vs NCCL |
| AllToNext | 3n 24xA100 | 32 MB - 256 MB | up to 14.5x vs hand CUDA |
| LM inference E2E | 8xA100 | (mixed) | 1.22-1.29x |
| MoE training E2E | 256xA100 | (mixed) | 1.10-1.89x |
9. Configuration-Regime Trade-off Tables
9.1 Algorithm choice (within AllReduce)
| Dimension | Ring | All-Pairs | Hierarchical | Winner (DynamICCL) |
|---|---|---|---|---|
| Latency-optimal | 2R-2 steps | 2 steps | 4 phases x ring | All-Pairs |
| Bandwidth-optimal | yes | yes | yes (intra+inter split) | tie |
| Small msg (< 16 KB) | weak | best | weak | All-Pairs |
| Medium msg (32 KB - 3 MB) | best with ch>1 | strong | strong on multi-node | Ring (1n) / Hier (2n+) |
| Large msg (> 32 MB) | matches NCCL | weak | best on multi-node | Hier |
| Multi-node necessity | OK | weak (R^2 sends) | designed for it | Hier |
| Code complexity | low | very low | medium | -- |
For DynamICCL, prefer regime-aware algorithm
dispatch. The crossover between All-Pairs (small), Ring (medium
1-node), and Hierarchical (large multi-node) is exactly the kind of
state-conditional optimal action that justifies an RL policy. Static
defaults miss the crossover; an agent that conditions on
(msg_size_bin, is_multi_node, scale) can pick the right
side of each flip.
9.2
Parallelization factor r (Chunk Parallelization)
| Buffer size | Best r | Reason | Winner (DynamICCL) |
|---|---|---|---|
| 1 KB - 16 KB | r = 1 - 2 | startup cost dominates; r competes for BW | low r |
| 32 KB - 1 MB | r = 2 - 4 | sweet spot — fills NVLink without contention | medium r |
| 1 MB - 32 MB | r = 4 - 8 | larger transfers benefit from saturating BW | high r |
| > 32 MB | r = 16-24 | matches NCCL's "ring on 1 channel x r=24" | very high r |
| AllToNext small | r = 1 - 4 | extra IB-NIC init overhead not worth it | low r |
| AllToNext large | r = 8 - 16 | up to 14.5x by fan-out across IB NICs | high r |
For DynamICCL, prefer a non-monotonic-in-r action
axis. The optimal r is small at both extremes
(init dominates at small msg; contention dominates at very large msg)
and peaks in the middle — a Pareto-frontier traversal that an RL agent
can learn from reward signal alone.
9.3 Protocol choice (LL / LL128 / Simple)
| Buffer size | Best protocol | Reason | Winner (DynamICCL) |
|---|---|---|---|
| 1 KB - 16 KB | LL | 64-byte chunks, lowest latency | LL |
| 32 KB - 256 KB | LL or LL128 | LL128 begins to amortize init | LL/LL128 (regime) |
| 256 KB - 1 MB | LL128 | balance of latency + bandwidth | LL128 |
| 1 MB - 8 MB | LL128 / Simple | crossover region | LL128/Simple (regime) |
| > 8 MB | Simple | bandwidth-optimal; large chunks | Simple |
For DynamICCL, prefer a 3-level protocol action axis with a buffer-size-conditional policy. The Fig 7 plots show all three protocols crossing in the 32 KB - 1 MB range — exactly where the exploration budget should be densest.
9.4 Channel count (number of NCCL channels per logical ring)
| Buffer size | Best ch | Reason | Winner (DynamICCL) |
|---|---|---|---|
| 1 KB - 16 KB | ch = 1 - 4 | small msg doesn't fill 1 channel | low ch |
| 32 KB - 256 KB | ch = 4 - 8 | distribute across NVSwitches | medium-high ch |
| 256 KB - 8 MB | ch = 8 | fully saturate NVLink fabric | high ch (8) |
| > 32 MB | ch = 1 | single channel + r=24 wins (NCCL) | low ch + high r |
For DynamICCL, prefer channel-and-r as a coupled
action. The paper shows that (ch=1, r=24) (NCCL's
default) wins at very large buffers and (ch=8, r=4) wins at
medium buffers — they trade off but cannot both be high (each uses GPU
TB resources). Agent-2's joint action over (ch, r) is a constrained 2D
selection.
9.5 Hardware sensitivity (Type-A flat vs Type-B hierarchical)
| Algorithm | Type-A (8xA100, flat NVSwitch) | Type-B (16xV100, 2-board) | Winner (DynamICCL) |
|---|---|---|---|
| Ring AllReduce | strong (single contiguous link) | strong intra-board, weaker x-board | Ring on Type-A |
| All-Pairs | strong (every pair via NVSwitch) | weaker (board-crossings) | All-Pairs on Type-A |
| Hierarchical AR | matches Ring (no advantage) | specifically designed for | Hier on Type-B |
| Two-Step AllToAll | native fit (1 IB NIC pair) | native fit (1 IB NIC pair) | tie |
| AllToNext | fan-out across 2 NICs/pair | fan-out across 1 NIC/pair | A100 has more NICs |
For DynamICCL, prefer a topology fingerprint feature that distinguishes flat-NVSwitch from hierarchical-2-board. The optimal algorithm differs on these two even when scale and buffer size are identical.
10. Bottlenecks & Insights Surfaced by the Measurements
10.1 NCCL's "1 channel x r=24" implicit policy is exposed
The paper's deepest empirical insight (Sec 7.1.1) is that NCCL's internal Ring schedule is equivalent to "single logical ring on one channel, with the entire program parallelized 24 times, with protocol chosen per buffer size." This is essentially reverse-engineering NCCL's hidden policy. For DynamICCL, this is gold:
NCCL's implicit policy (uncovered by GC3 paper):
channel : 1 (single channel per logical ring)
parallelize: 24 (the whole ring repeated 24 times)
protocol : f(buffer size) -- LL/LL128/Simple
GC3's empirical finding:
this policy WINS at > 32 MB
this policy LOSES (by up to 1.9x) at 32 KB - 3 MB
-- where (channel=8, r=4) is preferredFor DynamICCL, the implication is concrete: Agent-2's policy must differ from NCCL's hidden default in the 32 KB - 3 MB regime, and the agent's reward signal will drive it toward (ch=8, r=4) territory without needing to see the NCCL code. The paper essentially gives the agent a "ceiling" to beat in this regime.
10.2 Single TB cannot saturate A100 NVLink
"Our experience has shown that a single thread block in an NVIDIA A100 GPU is not capable of saturating the bandwidth of an outgoing NVLink. The user should carefully choose the parallelization factor as increasing it beyond a certain point will reduce performance due to competition for bandwidth." (Sec 5.1)
The architectural fact that one TB cannot saturate one outgoing
NVLink explains why r > 1 is needed even on a
fully-connected NVSwitch. But it also explains why the curve is
non-monotonic: too many TBs compete for a finite number of NCCL FIFO
slots and PCIe channels. For DynamICCL, encode
link_utilization < threshold as a state feature that
triggers an r increase, and treat link contention
(detected via slowdown) as a state feature that triggers an
r decrease.
10.3 Cross-tenant cloud-IB noise at 2-64 MB
"Note, at smaller sizes between 2MB-64MB there are large fluctuations in speedup caused by congestion in the IB network which is shared with other cloud tenants." (Sec 7.3)
This is the empirical fingerprint of multi-tenant cloud IB: variance
in the 2-64 MB regime is high because messages this size traverse the
fabric long enough to encounter congestion from neighbors but not long
enough to amortize it. For DynamICCL, a
is_shared_cloud_ib flag (detectable via collective time
variance) should trigger reward smoothing (rolling-mean over W=10
calls) so the policy doesn't update on noise.
10.4 Single-IB-NIC underutilization is a 14.5x leverage point
The AllToNext result (Sec 7.4, Fig 7g) shows that the naive
implementation — every GPU sends its full buffer to its successor
through one IB NIC — leaves up to 14.5x of bandwidth on the
table. The fix is fan-out across IB NICs. For
DynamICCL, recognize "single-link underutilization" via aggregate IB BW
telemetry as a state feature, and prefer fan-out actions (high
r, splitting one collective across multiple
channels) to recover the orphaned bandwidth.
10.5 Hand-optimized CUDA kernels are NOT the ceiling
"The GC3 implementation is up to 1.3x faster than the hand-optimized implementation. ... Furthermore, the GC3 implementation is much succinct and requires only 15 lines of code while the hand optimized kernel requires roughly 70 lines of code." (Sec 7.3)
GC3 beats a hand-tuned CUDA AllToAll by 1.3x. The gap is from cross-kernel optimizations the compiler does and the human author forgets — fusion across kernel boundaries, scratch-buffer aggregation without an extra copy kernel, pipelined tiling. This is a crucial framing for DynamICCL's evaluation: treat NCCL_default as the RL reward baseline (the "do nothing" arm) and treat hand-optimized CUDA as a sanity-check ceiling — but not as a hard target, because GC3 has already shown the ceiling can be beaten.
10.6 Custom collectives are first-class — AllToNext as proof
GC3 lets users define arbitrary new collectives with
the same DSL that expresses standard ones. The AllToNext collective is
novel — it's not in the MPI standard, has no NCCL implementation — and
yet GC3 supports it natively, optimizes it to 14.5x over a hand CUDA
baseline, and delivers it in production. For DynamICCL, the
implication is that the action space must include a "custom collective"
arm, with a learned policy that recognizes pipeline-parallel
patterns (GPU i -> GPU i+1) and dispatches
to AllToNext rather than a generic AllReduce or AllToAll.
10.7 The conclusion the paper hints at but does not act on
"Recent works [44, 15, 20, 29] have shown the advantage of overlapping computation and communication when optimizing distributed ML workloads. While our focus here is on specifying communication collectives, extending GC3 to further specify the scheduling of computation is an interesting future work." (Sec 8)
GC3 explicitly notes that adding compute-comm overlap into
the DSL is future work. DynamICCL operates one layer below the
algorithm graph, so the comm-comm overlap signal (i.e., whether
the caller itself is overlapping compute with this collective) is an
exogenous state feature DynamICCL should read — not an action
it should take. Encode caller_is_overlapping_compute = bool
in state.
11. Limitations of the Methodology
| Limitation | Implication for DynamICCL |
|---|---|
| No automatic algorithm synthesis | DynamICCL handles within-algorithm tuning; orthogonal |
| User must hand-write GC3 programs | Algorithm choice is exogenous; agent picks among given progs |
| Tile size requires manual configuration | DynamICCL can subsume this knob automatically |
| Buffer-size dispatch ranges set by user | DynamICCL's policy learns the dispatch ranges from rewards |
| Peephole fusion limited to listed patterns | Compiler may miss higher-order fusions |
| Cooperative kernel constraint (TB <= SMs) | Caps maximum parallelization factor — also caps DynamICCL r |
| Evaluation models: only LM + MoE production | Limited model-architecture coverage (no CNN, GNN, RecSys) |
| Hardware: A100 + V100 only | No H100, MI300, TPU regimes — DynamICCL must extrapolate |
| IB-only inter-node (no RoCE / Ethernet) | Transport coverage thinner than HPC variety |
| Cloud-IB cross-tenant noise unaddressed in 2-64MB | Reward-noise model needed before DynamICCL trains on cloud |
| SCCL direct-copy protocol not implemented | Author-acknowledged gap; protocol coverage incomplete |
| 50-iteration measurement window | Reasonable for end-to-end but thin for tail-latency analysis |
| 20-iteration warmup (Sec 7) | Same protocol DynamICCL should reuse |
| No formal cost-model for compiler choices | Heuristic scheduling — could miss optimal TB assignment |
| No adaptive runtime — all decisions compile-time | EXACTLY the gap DynamICCL fills — runtime adaptation |
| Anonymous cloud provider "X" | Reproduction blocked; cluster details partially redacted |
| Single-collective focus per cell | No multi-collective interaction (which DynamICCL must model) |
| Datatypes not varied | fp32 vs fp16 vs bf16 effect on optimal config not measured |
The most consequential limitation for DynamICCL is the same as for MSCCLang: GC3 makes all scheduling decisions at compile time, with the only runtime adaptivity being size-based dispatch among pre-compiled programs. Within a chosen program, the values of (channel, tile size, parallelize factor, protocol) are baked in by the user. DynamICCL fills exactly this gap — it adapts these values at runtime based on observed state. The composition is therefore clean: GC3 fixes the algorithmic graph; DynamICCL tunes the parameters inside that graph.
12. What to Borrow for DynamICCL — Compile-Time Graph Fixing vs Run-Time Knob Selection
GC3 and DynamICCL operate at adjacent layers of the collective- communication stack. GC3 is a programmable algorithm layer with compile-time scheduling; DynamICCL is a configurable parameter layer with runtime selection. The two compose, and the borrowable patterns from GC3 fall into five clusters: layer composition, state-vector features, exploration-budget allocation, the two crossovers Agent-2 must discover, and reward-shaping inputs.
12.1 GC3 is composable with DynamICCL — explicit composition
+--- Layer composition (GC3 + DynamICCL) ----------------------------+
| |
| User code: PyTorch DDP / Megatron / DeepSpeed |
| | |
| v |
| GC3 program (compile time) |
| - chunk routing |
| - algorithm shape (Ring / All-Pairs / Hier / Two-Step / ...) |
| - thread block layout (1 send peer + 1 recv peer per TB) |
| - default ch, default tile size, default r, default protocol |
| | |
| v |
| GC3-IR (deployed binary) |
| | |
| v |
| +-----------------------------------------------+ |
| | DynamICCL Tuner Plugin (runtime, NCCL slot) | |
| | observes: msg_size, model_intensity I, | |
| | local_batch_size, topology fp, | |
| | recent collective LSTM window | |
| | chooses among the run-time-tunable knobs: | |
| | - which GC3 program to dispatch | |
| | (subsumes user-set buffer-size ranges) | |
| | - protocol override (LL / LL128 / Simple) | |
| | - tile size override | |
| | - r override (within compiled bounds) | |
| | - nChannels override | |
| +-----------------------------------------------+ |
| | |
| v |
| GC3 runtime interpreter (executes one program, one CUDA kernel) |
| | |
| v |
| NCCL transport (NVLink / IB GPUDirect RDMA / SHM / PCIe / TCP) |
+--------------------------------------------------------------------+
^ Fig 12: Layer composition — explicit. GC3 occupies the algorithm
layer; DynamICCL occupies the parameter layer immediately above
NCCL transport. The composition is clean because every GC3
scheduling directive (channel, parallelize, protocol) is exposed
with a runtime override path. **GC3's contribution is composable
with DynamICCL — not in tension.**The GC3 paper itself states the composition possibility (Sec 6):
"All GC3-IR generated by our compiler is guaranteed to be correct, but some programs might only be performant for a range of buffer sizes. Therefore, the runtime dynamically selects the right algorithm to invoke based on user configurable size ranges and falls back to NCCL's built-in algorithms otherwise. This allows a user to hyper-optimize GC3 programs to a specific use case." (Sec 6)
The "user-configurable size ranges" is precisely the manual hand-tuned table that DynamICCL replaces with a learned policy. Where the GC3 paper says "the user must measure and configure," DynamICCL says "the policy will learn online from reward signal." This is not in tension with GC3 — it is GC3's explicitly-anticipated automation path.
12.2 Composability vs tension — the explicit assessment
Composability dimensions (GC3 and DynamICCL multiply):
| Dimension | Why composable |
|---|---|
| Algorithm shape (Chunk DAG) | Fixed by user at compile time; DynamICCL doesn't touch it |
| Thread block layout | Fixed by Sec 5.2 scheduling; deadlock-freedom invariant |
| Cross-TB synchronization | Fixed; DynamICCL doesn't move sync points |
| Buffer-size dispatch table | DynamICCL replaces static table with learned policy |
| Protocol selection | GC3 exposes as DSL directive; DynamICCL overrides at RT |
| Parallelize factor r | GC3 exposes as DSL directive; DynamICCL overrides at RT |
| Channel assignment | GC3 exposes as DSL directive; DynamICCL overrides at RT |
| Tile size | GC3 deduces from buffer/FIFO; DynamICCL can override |
| Custom-collective dispatch | GC3 supports custom collectives; DynamICCL learns when |
Tension dimensions (potential conflicts):
| Dimension | Why potential tension |
|---|---|
| Static deadlock-freedom proof | DynamICCL must NOT move sync edges -- would break GC3 IR |
| User-baked "tuned for use case" | DynamICCL might override user intent if not gated |
| Compile-time-fixed algorithm | DynamICCL cannot pick an algorithm GC3 didn't compile |
| Cooperative-kernel TB count | DynamICCL's r override must respect TB <= SMs invariant |
The four tension dimensions all resolve to invariant constraints DynamICCL must respect: it cannot move sync points, cannot exceed TB count, cannot pick algorithms outside the compiled set. As long as DynamICCL's action space is constrained to the runtime-configurable subset (protocol, tile, r within bounds, channel within bounds, and program dispatch among compiled programs), the systems are strictly composable. The conclusion is: GC3 and DynamICCL multiply; they do not conflict.
12.3 State-vector features the GC3 paper validates as predictive
Add to Agent-2 state vector s_t (informed by GC3):
+-----------------------------------------------------------------+
| msg_size_bin : enum (already there; log-binned) |
| algo_program_id : enum (which GC3 program is dispatched|
| -- the structural prior) |
| topology_fingerprint : enum |
| ({flat-NVSwitch (Type-A 8xA100), |
| hier-2-board (Type-B 16xV100), |
| multi-node-IB, |
| multi-node-IB-shared-cloud}) |
| num_ib_nics_per_pair : int (1 on V100, 2 on A100 -- sets |
| bandwidth ceiling for AllToNext) |
| is_intra_node_only : bool (1-node vs multi-node) |
| is_shared_cloud_ib : bool (variance-detected; triggers |
| reward smoothing) |
| mean_link_utilization : float (recent-window aggregate; |
| flags single-TB undersaturation)|
| caller_is_overlapping_compute : bool (Sec 8 future work signal:|
| overlap is exogenous) |
| is_pipeline_parallel : bool (caller is using GPU-i to GPU- |
| i+1 pattern -- AllToNext fit) |
| is_compute_in_reduce : bool (true for Ring/Hier; false for |
| All-Pairs -- changes r-sensit.)|
| collective_kind : enum ({AR, A2A, AN, AG, RS}) |
+-----------------------------------------------------------------+
^ Fig 13: Borrowed state features. The first three are the GC3
paper's central exogenous topology features. The next three
capture observed runtime conditions (link util, IB sharing, IB
NIC count). The last three capture caller intent that the agent
cannot control but should condition on.12.4 Empirical findings that constrain the policy's prior
PRIOR: Agent-2 should be CONSERVATIVE (low-exploration) in:
+--------------------------------------------------------------+
| Regime Reason |
|--------------------------------------------------------------|
| > 32 MB AllReduce on 1-node 8xA100 NCCL Ring matches GC3 |
| (saturated single-channel + r=24) (Sec 7.1.1 final para)|
| |
| Large AllReduce > 1 GB on 2-node Hier only +11% over |
| (Hierarchical at scale) NCCL (Sec 7.2) |
| |
| Small AllToNext (< 256 KB) r=1 is best (init |
| cost dominates) |
| |
| AllToAll < 256 KB (16-node) Two-Step roughly ties |
| baselines |
+--------------------------------------------------------------+
PRIOR: Agent-2 should be AGGRESSIVE (high-exploration) in:
+--------------------------------------------------------------+
| Regime Reason |
|--------------------------------------------------------------|
| 16 KB - 256 KB AllReduce 1.5-1.9x gap to NCCL |
| (1-node 8xA100, ch=8 r=4 LL) from channel choice |
| (Fig 7a) |
| |
| 1 KB - 1 MB AllReduce All-Pairs vs Ring |
| (1-node 16xV100) crossover at small |
| sizes (Sec 7.1.2, |
| Fig 7b: 3x peak) |
| |
| Small msg AllReduce on 2-node Hier 1.4x over NCCL |
| (16xA100; < 1 MB) (Sec 7.2) |
| |
| AllToAll 256+ MB on 256xA100 1.3x over hand CUDA, |
| 1.2x over NCCL |
| |
| AllToNext > 16 MB up to 14.5x over CUDA |
| (3-node 24xA100) baseline (Sec 7.4) |
| |
| 2-64 MB on shared cloud IB High variance, possib.|
| for large gain when |
| congestion is low |
+--------------------------------------------------------------+
^ Fig 14: Where to allocate exploration budget. Conservative regions
are where GC3 ties or marginally beats NCCL/CUDA -- limited room
for further gain. Aggressive regions are where GC3's wins are
1.4x-14.5x -- meaning there is *real* leverage that DynamICCL can
extract through parameter selection on top of GC3's algorithm
graph (or, equivalently, on top of NCCL's algorithms when the GC3
program is not compiled for that buffer-size range).12.5 The two crossovers Agent-2 must discover
Crossover A — Algorithm choice flips with buffer
size. All-Pairs wins small (1 KB - 1 MB), Ring wins medium (32
KB - 3 MB on 1-node), Hierarchical wins large multi-node (> 1 GB).
The GC3 paper's Sec 7.1 + 7.2 explicitly partition the buffer-size axis
into ranges where different algorithms dominate. DynamICCL
Agent-2 must learn this partition from reward signal alone, conditional
on the same (msg_size_bin, scale, topology) axes.
The GC3 paper provides ground-truth break-points the trained agent's
policy should match.
Crossover B — Parallelization factor r flips
with buffer size in a non-monotonic way. r = 4
wins at 32 KB; r = 24 wins at 32 MB; intermediate sizes
have intermediate optima (Sec 7.1.1). The crossover is
non-monotonic in r at fixed buffer size, which is
the empirical fingerprint of a competition effect (more r
helps until contention dominates). Agent-2 must learn a 5-level
discrete-action axis over r and capture its
buffer-size-conditional optimum — a Pareto-frontier
traversal.
12.6 Reward-shaping inputs from the paper
| Paper finding | DynamICCL reward-shaping rule |
|---|---|
| 1.9x peak vs NCCL at 32-256 KB (Fig 7a) | Define
excess_speedup = wall(NCCL_default) / wall(chosen) and
reward log(excess_speedup) -- encourages large absolute
wins, dampens marginal noise |
| Cloud-IB variance at 2-64 MB (Sec 7.3) | Apply rolling-mean reward over W=10 calls when state.is_shared_cloud_ib = true |
| End-to-end 1.10-1.89x range (Sec 7.5) | Track end-to-end model-step-time as a slow-loop reward in addition to per-collective time -- captures cross-collective interactions |
| Hand-optimized beaten by 1.3x (Sec 7.3) | Use NCCL_default as standard baseline; treat hand-optimized as a "ceiling check" rather than reward target |
| AllToNext 14.5x via fan-out (Sec 7.4) | Recognize "single-link underutilization" as state and amplify reward when agent successfully fans out |
12.7 Exploration budget allocation
The GC3 paper measures 4 collectives x 8 (cluster, scale) cells x ~5 algorithm/parallelize variants x ~22 buffer sizes x ~3 protocols = ~10,560 logical cells, but plots a curated subset of ~150 in Fig 7. Each cell uses 20 warmup + 50 measurement iterations (Sec 7 "Experimental Setup"). This is the right ballpark for an RL agent: a few hundred iterations per (collective, hardware) combination is enough to identify the regime structure.
For DynamICCL, allocate ~50 reward samples per
(msg_size_bin x topology x algo) cell during the
high-exploration phase, then narrow the sampling budget to
high-leverage regimes (Fig 14 lower box) for the exploitation phase. The
two production end-to-end workloads (LM inference on 8xA100, MoE
training on 256xA100) should serve as the gold-standard validation set:
if Agent-2's learned policy matches or beats GC3's hand-tuned 1.22-1.29x
and 1.10-1.89x end-to-end speedups, the agent is good enough for
production.
12.8 The composition principle stated cleanly
GC3 (compile time) DynamICCL (run time)
------------------------------ -----------------------------
- Pick algorithm graph - Pick which graph to dispatch
- Default ch, tile, r, proto - Override ch, tile, r, proto
- Insert TB sync points - (no override -- invariant)
- Static dispatch ranges - Dynamic dispatch by state
- Compile-time correctness gate - Run-time adaptation
- Frozen at deploy - Re-trained per-cluster
Together: the graph structure is fixed by the human author who
proves it correct via GC3's pre/post-condition contracts; the
parametric values inside the graph are tuned by RL on observed
hardware. Neither subsumes the other; they multiply.This is the cleanest framing of how GC3 and DynamICCL co-exist. The paper's authors explicitly make space for runtime adaptation in the form of "user configurable size ranges" — which is the manual, static version of what DynamICCL automates dynamically.
12.9 Methodological patterns to reuse
| Pattern (GC3) | DynamICCL adoption |
|---|---|
| 20 warmup + 50 measurement iterations | Same protocol per (algo, ch, r, tile) cell |
| Speedup against NCCL default as primary metric | Same; NCCL default = the "do nothing" RL baseline |
| Speedup against hand-optimized as ceiling check | Optional gating signal: agent must beat NCCL default but not need to beat hand-optimized to be useful |
| Production end-to-end workloads (LM, MoE) | Validate trained policy on real LLM serving + training |
| Open-source compiler + tools | DynamICCL tuner plugin must be open-source (NCCL plugin) |
| Three IRs at decreasing abstraction | DynamICCL's policy network can mirror this: high-level (collective type) -> mid-level (algo + protocol) -> low-level (ch, r, tile) |
| Cooperative-kernel single-launch model | DynamICCL works inside NCCL's same single-launch model -- no kernel-launch knob in action space |
| Buffer-size-conditional algorithm dispatch | Replace static lookup table with state-conditioned policy |
| 30 LOC programs (Fig 3a hierarchical AR) | DynamICCL's action API must be small enough to express in similar LOC for runtime overrides |
12.10 Compiler-runtime knob taxonomy — full mapping
The cleanest contribution GC3 makes to DynamICCL's design is forcing an explicit taxonomy of which knobs are compile-time-fixed vs run-time- tunable. The table below resolves every NCCL/GC3 parameter to one of three categories:
| Knob | Category | Why |
|---|---|---|
| Algorithm shape (Ring/AP/Hier/Two-Step) | Compile-time-structural | Defined by chunk routing in DSL |
| Chunk-DAG topology | Compile-time-structural | Inherent to user program |
| Thread block layout (1sP+1rP) | Compile-time-structural | Determined by Sec 5 scheduling |
| Cross-TB sync points | Compile-time-structural | Computed from processing edges; deadlock-free |
| Channel assignment | Compile-time-default | DSL ch=k directive; runtime can override |
Parallelize factor r |
Compile-time-default | DSL parallelize(N) directive; runtime override |
| Protocol (LL/LL128/Simple) | Compile-time-default | DSL directive; runtime can override per buffer |
| Tile size | Run-time-configurable | Set by runtime based on buffer size + FIFO slot |
| nChannels (NCCL) | Run-time-tunable | DynamICCL's primary action axis |
| numThreads (NCCL) | Run-time-tunable | DynamICCL's primary action axis |
| chunkSize (NCCL) | Run-time-tunable | Maps to GC3's tile size |
| Algorithm dispatch | Run-time-tunable | DynamICCL picks among pre-compiled programs |
| Buffer-size dispatch table | Run-time-tunable | Replaces user's static lookup with learned policy |
For DynamICCL, the action space is exactly the union of the last six rows of the table — nChannels, numThreads, chunkSize, protocol, "which GC3 program to dispatch," and the buffer-size threshold for that dispatch. All other knobs are either invariants of the user's program (compile-time-structural) or sensible defaults the user provided (compile-time-default that DynamICCL is permitted to override).
13. Analogy
GC3 is to collective communication what CUDA shader languages are to graphics. Before HLSL/GLSL/CUDA, every graphics programmer wrote hand-tuned assembly per GPU per pipeline stage; the result was fast but unportable and intractable to maintain. Shader languages introduced a high-level domain-specific abstraction (vertex shader, fragment shader, compute shader) with strong correctness guarantees (no out-of-bounds writes, deterministic execution), let a compiler emit optimal kernels, and reduced the typical shader from 1000s of hand-written assembly lines to ~30 lines of HLSL. Productivity won; optimality matched or beat the hand-written baseline because the compiler captured cross-stage optimizations a human programmer forgets.
In GC3, the same shift happens in the collective space. Before GC3, every custom collective was hand-written CUDA against the NCCL P2P primitives; correct programs were rare, and cross-kernel optimizations were absent. GC3 introduces a chunk-oriented DSL (chunks, ranks, channels, threadblocks) with strong correctness guarantees (data-race- free by construction, postcondition-verified), lets a compiler emit GC3-IR, and reduces the typical hierarchical AllReduce from hundreds of CUDA lines to ~30 lines of Python (Fig 3a). The 1.9x AllReduce speedup over NCCL, the 1.3x AllToAll speedup over hand-optimized, and the 14.5x AllToNext speedup over hand CUDA are the analogs of "shader compilers beat hand-written assembly" — and the rationale is the same: the compiler sees the whole program at once and applies cross-step fusion (rrcs, rrs, rcs) and pipelined tiling that no human keeps straight.
DynamICCL fits into this analogy as the GPU driver's runtime shader optimizer — the layer that takes a compiled shader and, based on observed device state (clock speed, register pressure, recent kernel timings), picks the right scheduling parameters for this invocation on this GPU. DynamICCL doesn't synthesize the shader; it tunes the dispatch. The GC3 program is the shader, the user-supplied scheduling directives (channel, parallelize, protocol) are the shader's compile- time hints, and DynamICCL's RL policy is the runtime hint adapter that overrides those hints based on what the cluster is actually doing right now. The "user configurable size ranges" the GC3 runtime exposes (Sec 6) is exactly the manual hand-tuned table that the RL policy replaces.
The two together are the same architectural pattern that makes modern graphics pipelines work: a high-level correctness-preserving compiler upstream + a state-aware runtime tuner downstream. GC3 upstream gives you correctness, portability, and cross-kernel optimization. DynamICCL downstream gives you state-conditional parameter selection that no static deploy can match. Neither is sufficient on its own. Together, they cover both the "I want to write a correct collective in 30 lines of Python" problem and the "I want this collective to run optimally on whatever cluster you deploy it to" problem.
The key composability claim — to be stated unambiguously — is this: GC3's contribution is composable with DynamICCL, not in tension with it. GC3 fixes the algorithmic graph at compile time via a DSL and a correctness-preserving compiler; DynamICCL tunes the parameters inside that graph at run time via an RL policy that conditions on observed cluster state. The four NCCL knobs DynamICCL controls (algorithm selection at the program-dispatch granularity, protocol, nChannels, chunkSize/tile, parallelize factor) are all surfaced as either runtime APIs in GC3 or as user-configured defaults that GC3's runtime explicitly permits override of. The deadlock-free invariant GC3 maintains (Sec 5.2) does constrain DynamICCL: the agent cannot move sync points or pick algorithms outside the compiled set. But within that envelope, the multiplication is clean — and the GC3 paper's own "user configurable size ranges" mechanism is the gateway through which DynamICCL inserts its learned policy.
Summary of Borrowed Patterns
| Pattern from Cowan et al. (GC3 / arXiv 2201.11840v3, 2022) | DynamICCL application |
|---|---|
| Three-IR pipeline (Chunk DAG -> Instr DAG -> GC3-IR) | Mirror in policy network: collective-type -> algo+protocol -> nCh/r/tile |
| Chunk-oriented DSL with pre/post contracts | Constrain Agent-2's action space to satisfy NCCL correctness invariants |
| Peephole fusion (rcs / rrcs / rrs) | Recognize fused-instruction states as low-overhead in the LSTM encoding |
| Cooperative-kernel single-launch interpreter | DynamICCL inherits this -- no kernel-launch knob needed |
| Pipelined tile execution (Sec 6.2, Fig 6) | Tile size = NCCL chunkSize; agent learns size-conditional optimum |
| 1 TB cannot saturate A100 NVLink (Sec 5.1) | link_utilization < 0.9 as state feature triggering
nChannels increase |
| Algorithm vs implementation strategy gap (Sec 7.2 NCCL coll path) | Confirms DynamICCL operates above the kernel-chaining trap (single launch) |
| Buffer-size-conditional algorithm choice (Sec 7.1) | Action axis: "which GC3 program to dispatch" -- discrete 3-5 levels |
| Parallelize factor r non-monotonic with buffer size (Sec 7.1.1) | Pareto-frontier traversal action axis r in {1,2,4,8,16,24} |
| LL / LL128 / Simple protocol size-conditional (Fig 7a-d) | Protocol is a 3-level discrete action; learn size-conditional optimum |
| Cross-tenant cloud IB noise at 2-64 MB (Sec 7.3) | Reward-smoothing flag: rolling mean over W=10 when shared_cloud_ib=true |
| AllToNext 14.5x via all-IB-NIC fan-out (Sec 7.4) | Recognize "single-link underutilization"; prefer fan-out actions |
| End-to-end LM 1.22-1.29x, MoE 1.10-1.89x (Sec 7.5) | End-to-end model-step-time as slow-loop reward |
| Compile-time vs runtime knob taxonomy | Action space = run-time-configurable knobs only (last 6 rows of taxonomy) |
| 20 warmup + 50 measurement (Sec 7) | Reuse exact protocol for DynamICCL's per-cell sweep |
| Open-source release (msccl + msccl-tools) | DynamICCL tuner plugin must be open-source |
| Hand-optimized as ceiling check, not reward target | NCCL_default is the reward baseline; hand-optimized is a sanity gate |
| Layer composition: graph fixed + parameters tuned | DynamICCL's mission statement: tune within GC3's structural prior |
| NCCL's hidden policy = (1ch, r=24, proto=f(size)) (Sec 7.1.1) | Concrete "default" the agent must beat at 32 KB - 3 MB; ground truth ceiling |
| AllToNext custom collective | Agent's program-dispatch action must include custom-collective option |
| Type-A flat NVSwitch vs Type-B hier 2-board | Topology fingerprint feature with at least 4 levels (Type-A, Type-B, multi-node IB, shared cloud IB) |
| Sec 8 future work on compute-comm overlap | Encode caller_is_overlapping_compute as exogenous state
feature |