From 3b0093a3534ee67f1216dfc8f79d0e57e907bc81 Mon Sep 17 00:00:00 2001
From: Bizuayeu <30091157+Bizuayeu@users.noreply.github.com>
Date: Fri, 16 Oct 2026 07:16:15 +0900
Subject: [PATCH 1/2] glm: TF_GLM_PREFILL_REDUCE=ring passes split's exchanges
 along the ring, for a fabric where a rank reaches only its neighbours

---
 zig/src/families/glm/cuda.zig         |   2 +
 zig/src/families/glm/cuda_engine.zig  |   4 +-
 zig/src/families/glm/cuda_forward.zig |  38 +++---
 zig/src/families/glm/cuda_ranks.zig   |  24 +++-
 zig/src/families/glm/cuda_relay.zig   | 167 ++++++++++++++++++++++++++
 5 files changed, 211 insertions(+), 24 deletions(-)
 create mode 100644 zig/src/families/glm/cuda_relay.zig

diff --git a/zig/src/families/glm/cuda.zig b/zig/src/families/glm/cuda.zig
index af9254de61..5df3099331 100644
--- a/zig/src/families/glm/cuda.zig
+++ b/zig/src/families/glm/cuda.zig
@@ -13,6 +13,7 @@ pub const mlp = @import("cuda_mlp.zig");
 pub const shard = @import("cuda_shard.zig");
 pub const weights = @import("cuda_weights.zig");
 pub const forward = @import("cuda_forward.zig");
+pub const relay = @import("cuda_relay.zig");
 pub const config = @import("config.zig");
 pub const quant = @import("cuda_quant.zig");
 pub const q4 = @import("cuda_q4.zig");
@@ -37,6 +38,7 @@ test {
     _ = shard;
     _ = weights;
     _ = forward;
+    _ = relay;
     _ = quant;
     _ = q4;
     _ = ranks;
diff --git a/zig/src/families/glm/cuda_engine.zig b/zig/src/families/glm/cuda_engine.zig
index ee38e00b3e..f29ba18303 100644
--- a/zig/src/families/glm/cuda_engine.zig
+++ b/zig/src/families/glm/cuda_engine.zig
@@ -187,7 +187,7 @@ pub const Engine = struct {
 
     fn view(e: *Engine, b: *fw.Buffers, out: *[1]fw.Rank) fw.Forward {
         out.* = .{.{ .w = &e.w, .st = &e.cur.st, .b = b }};
-        return .{ .d = e.d, .m = &e.m, .stream = e.stream, .ranks = out, .exchange = e.exchange(), .pieces = if (b == &e.bp) &e.pieces else null, .split = e.settings.split, .kda_wide = e.settings.kda_wide };
+        return .{ .d = e.d, .m = &e.m, .stream = e.stream, .ranks = out, .exchange = e.exchange(), .pieces = if (b == &e.bp) &e.pieces else null, .split = e.settings.split, .ring = e.settings.ring, .kda_wide = e.settings.kda_wide };
     }
 
     /// Binds slot `slot`'s sequence for the steps that follow, making it (a zeroed state) on first use.
@@ -417,7 +417,7 @@ pub const Engine = struct {
             rows += p.ids.len;
         }
         var one: [1]fw.Rank = .{.{ .w = &e.w, .st = segs[0].st, .b = b, .segs = segs[0..parts.len] }};
-        const f: fw.Forward = .{ .d = e.d, .m = &e.m, .stream = e.stream, .ranks = &one, .exchange = e.exchange(), .pieces = null, .split = e.settings.split, .kda_wide = e.settings.kda_wide };
+        const f: fw.Forward = .{ .d = e.d, .m = &e.m, .stream = e.stream, .ranks = &one, .exchange = e.exchange(), .pieces = null, .split = e.settings.split, .ring = e.settings.ring, .kda_wide = e.settings.kda_wide };
         try f.run(@ptrCast(ids[0..rows]));
         for (parts, segs[0..parts.len]) |p, g| {
             try e.bind(p.slot);
diff --git a/zig/src/families/glm/cuda_forward.zig b/zig/src/families/glm/cuda_forward.zig
index 1e57e55a08..007a2c8e71 100644
--- a/zig/src/families/glm/cuda_forward.zig
+++ b/zig/src/families/glm/cuda_forward.zig
@@ -18,6 +18,7 @@ const quant = @import("cuda_quant.zig");
 const q4 = @import("cuda_q4.zig");
 const state = @import("cuda_state.zig");
 const shared = @import("cuda_shared.zig");
+const relay = @import("cuda_relay.zig");
 
 pub const prefill_rows = state.prefill_rows;
 pub const window_rows = state.window_rows;
@@ -158,6 +159,7 @@ pub const Forward = struct {
     kda_wide: bool = true, // TF_GLM_KDA_DECODE_WIDE (the Python family's default): decode windows on KDA's three-kernel chain
     pieces: ?*const Pieces = null, // a prompt chunk's exchanges in row pieces on a second stream (null: whole)
     split: bool = true, // TF_GLM_PREFILL_REDUCE=split (the default with send/recv): each rank glues its own rows
+    ring: bool = false, // TF_GLM_PREFILL_REDUCE=ring: split, its exchanges passed along the ring (cuda_relay)
 
     pub fn set(f: *const Forward) glue.Set {
         return .{ .glue = &f.m.glue };
@@ -275,7 +277,7 @@ pub const Forward = struct {
     }
 
     /// reduce.shares: [lo, hi) of the rows of a piece of `n` each rank owns: about equal, the first n % world one more.
-    fn share(n: usize, world: usize, rank: usize) [2]usize {
+    pub fn share(n: usize, world: usize, rank: usize) [2]usize {
         const base = n / world;
         const extra = n % world;
         const lo = rank * base + @min(rank, extra);
@@ -324,13 +326,15 @@ pub const Forward = struct {
                     const me = r.w.share.rank;
                     const mine = share(n, world, me);
                     const own = mine[1] - mine[0];
-                    try x.lib.check(x.lib.api.ncclGroupStart(), "ncclGroupStart");
-                    for (0..world) |peer| if (peer != me) {
-                        const theirs = share(n, world, peer);
-                        try x.lib.check(x.lib.api.ncclSend(r.b.part + (lo + theirs[0]) * D * 4, (theirs[1] - theirs[0]) * D, .f32, @intCast(peer), x.comm, p.side.handle), "ncclSend");
-                        try x.lib.check(x.lib.api.ncclRecv(room(r, lo) + peer * own * D * 4, own * D, .f32, @intCast(peer), x.comm, p.side.handle), "ncclRecv");
-                    };
-                    try x.lib.check(x.lib.api.ncclGroupEnd(), "ncclGroupEnd");
+                    if (f.ring) try relay.partials(x.lib, x.comm, p.side, .{ .world = world, .rank = me, .n = n }, room(r, lo), r.b.part, lo, D) else {
+                        try x.lib.check(x.lib.api.ncclGroupStart(), "ncclGroupStart");
+                        for (0..world) |peer| if (peer != me) {
+                            const theirs = share(n, world, peer);
+                            try x.lib.check(x.lib.api.ncclSend(r.b.part + (lo + theirs[0]) * D * 4, (theirs[1] - theirs[0]) * D, .f32, @intCast(peer), x.comm, p.side.handle), "ncclSend");
+                            try x.lib.check(x.lib.api.ncclRecv(room(r, lo) + peer * own * D * 4, own * D, .f32, @intCast(peer), x.comm, p.side.handle), "ncclRecv");
+                        };
+                        try x.lib.check(x.lib.api.ncclGroupEnd(), "ncclGroupEnd");
+                    }
                     try f.copy(p.side, room(r, lo) + me * own * D * 4, r.b.part + (lo + mine[0]) * D * 4, own * D * 4);
                 },
             }
@@ -401,14 +405,16 @@ pub const Forward = struct {
                     const r = f.ranks[0];
                     const me = r.w.share.rank;
                     const rows_of = if (last) r.b.hidden else r.b.normed;
-                    const mine = share(n, world, me);
-                    try x.lib.check(x.lib.api.ncclGroupStart(), "ncclGroupStart");
-                    for (0..world) |peer| if (peer != me) {
-                        const theirs = share(n, world, peer);
-                        try x.lib.check(x.lib.api.ncclSend(rows_of + (lo + mine[0]) * D * 2, (mine[1] - mine[0]) * D, .bf16, @intCast(peer), x.comm, p.side.handle), "ncclSend");
-                        try x.lib.check(x.lib.api.ncclRecv(rows_of + (lo + theirs[0]) * D * 2, (theirs[1] - theirs[0]) * D, .bf16, @intCast(peer), x.comm, p.side.handle), "ncclRecv");
-                    };
-                    try x.lib.check(x.lib.api.ncclGroupEnd(), "ncclGroupEnd");
+                    if (f.ring) try relay.glued(x.lib, x.comm, p.side, .{ .world = world, .rank = me, .n = n }, rows_of, lo, D) else {
+                        const mine = share(n, world, me);
+                        try x.lib.check(x.lib.api.ncclGroupStart(), "ncclGroupStart");
+                        for (0..world) |peer| if (peer != me) {
+                            const theirs = share(n, world, peer);
+                            try x.lib.check(x.lib.api.ncclSend(rows_of + (lo + mine[0]) * D * 2, (mine[1] - mine[0]) * D, .bf16, @intCast(peer), x.comm, p.side.handle), "ncclSend");
+                            try x.lib.check(x.lib.api.ncclRecv(rows_of + (lo + theirs[0]) * D * 2, (theirs[1] - theirs[0]) * D, .bf16, @intCast(peer), x.comm, p.side.handle), "ncclRecv");
+                        };
+                        try x.lib.check(x.lib.api.ncclGroupEnd(), "ncclGroupEnd");
+                    }
                 },
             }
             try p.shared[q].record(p.side);
diff --git a/zig/src/families/glm/cuda_ranks.zig b/zig/src/families/glm/cuda_ranks.zig
index ab5b01b770..71866d5f0c 100644
--- a/zig/src/families/glm/cuda_ranks.zig
+++ b/zig/src/families/glm/cuda_ranks.zig
@@ -7,7 +7,8 @@ const Sha256 = std.crypto.hash.sha2.Sha256;
 
 /// The forward's knobs a rank reads from its environment (the Python family's names and defaults).
 pub const Settings = struct {
-    split: bool = true, // TF_GLM_PREFILL_REDUCE: split (each rank glues its own rows) or gather
+    split: bool = true, // TF_GLM_PREFILL_REDUCE: split (each rank glues its own rows), ring or gather
+    ring: bool = false, // TF_GLM_PREFILL_REDUCE=ring: split, its exchanges passed along the ring (cuda_relay)
     kda_wide: bool = true, // TF_GLM_KDA_DECODE_WIDE: 1 (decode windows on KDA's three-kernel chain) or 0
     pieces: u32 = 3, // TF_GLM_OVERLAP_PIECES: a prompt chunk's exchanges in this many row pieces (1: whole).
     heat: ?heat.Bands = null, // TF_HEAT_HIGH and TF_HEAT_LOW (upstream's names): a prompt chunk waits on every rank's hottest zone
@@ -17,8 +18,11 @@ pub const Settings = struct {
     pub fn read(get: *const fn ([:0]const u8) ?[]const u8, problem: *[]const u8) !Settings {
         var s: Settings = .{};
         if (get("TF_GLM_PREFILL_REDUCE")) |v| {
-            if (std.mem.eql(u8, v, "split")) s.split = true else if (std.mem.eql(u8, v, "gather")) s.split = false else {
-                problem.* = "TF_GLM_PREFILL_REDUCE: split or gather";
+            if (std.mem.eql(u8, v, "split")) s.split = true else if (std.mem.eql(u8, v, "ring")) {
+                s.split = true;
+                s.ring = true;
+            } else if (std.mem.eql(u8, v, "gather")) s.split = false else {
+                problem.* = "TF_GLM_PREFILL_REDUCE: split, ring or gather";
                 return error.BadSetting;
             }
         }
@@ -63,7 +67,7 @@ pub const Fingerprint = extern struct {
     layout: u32 = 5, // this struct's version
     world: u32,
     rank: u32,
-    split: u32,
+    split: u32, // TF_GLM_PREFILL_REDUCE: 0 gather, 1 split, 2 ring
     kda_wide: u32,
     pieces: u32,
     layers: u32, // TF_GLM_LAYERS (0: all)
@@ -105,7 +109,7 @@ pub const Fingerprint = extern struct {
         var f: Fingerprint = .{
             .world = @intCast(world),
             .rank = @intCast(rank),
-            .split = @intFromBool(s.split),
+            .split = if (!s.split) 0 else if (s.ring) 2 else 1,
             .kda_wide = @intFromBool(s.kda_wide),
             .pieces = s.pieces,
             .layers = s.layers,
@@ -230,6 +234,11 @@ fn fake(name: [:0]const u8) ?[]const u8 {
     return null;
 }
 
+fn ringed(name: [:0]const u8) ?[]const u8 {
+    if (std.mem.eql(u8, name, "TF_GLM_PREFILL_REDUCE")) return "ring";
+    return null;
+}
+
 fn bad(name: [:0]const u8) ?[]const u8 {
     if (std.mem.eql(u8, name, "TF_GLM_PREFILL_REDUCE")) return "sum";
     return null;
@@ -240,7 +249,10 @@ test "settings from the environment: the defaults, the choices, and a refusal na
     const s = try Settings.read(&fake, &problem);
     try std.testing.expect(!s.split and s.kda_wide and s.pieces == 2 and s.heat.?.high_c == 92 and s.heat.?.low_c == 88);
     try std.testing.expectError(error.BadSetting, Settings.read(&bad, &problem));
-    try std.testing.expectEqualStrings("TF_GLM_PREFILL_REDUCE: split or gather", problem);
+    try std.testing.expectEqualStrings("TF_GLM_PREFILL_REDUCE: split, ring or gather", problem);
+    const r = try Settings.read(&ringed, &problem);
+    try std.testing.expect(r.split and r.ring);
+    try std.testing.expectEqual(@as(u32, 2), Fingerprint.of(2, 0, r, 8192, true, @splat(0), @splat(0), "/c").split);
 }
 
 test "fingerprints: agreeing ranks pass; the first differing field and rank are named" {
diff --git a/zig/src/families/glm/cuda_relay.zig b/zig/src/families/glm/cuda_relay.zig
new file mode 100644
index 0000000000..9f12339114
--- /dev/null
+++ b/zig/src/families/glm/cuda_relay.zig
@@ -0,0 +1,167 @@
+//! TF_GLM_PREFILL_REDUCE=ring: split's exchanges passed along the ring, for fabrics where ranks reach only neighbours.
+
+const std = @import("std");
+const cuda = @import("cuda");
+const fw = @import("cuda_forward.zig");
+
+/// One piece's relay: each owner gets every rank's block in split's slot, so it sums split's bytes in split's order.
+pub const Relay = struct {
+    world: usize,
+    rank: usize,
+    n: usize, // the piece's rows
+
+    pub fn next(x: Relay) usize {
+        return (x.rank + 1) % x.world;
+    }
+    pub fn prev(x: Relay) usize {
+        return (x.rank + x.world - 1) % x.world;
+    }
+    /// Whose blocks this rank sends at step s (its own at step 1).
+    pub fn sentFrom(x: Relay, s: usize) usize {
+        return (x.rank + x.world + 1 - s) % x.world;
+    }
+    /// Whose blocks this rank takes at step s.
+    pub fn takenFrom(x: Relay, s: usize) usize {
+        return (x.rank + x.world - s) % x.world;
+    }
+    /// The i-th owner after this rank: step s sends owners 0 .. world - s, and keeps 0 .. world - s - 1 to send on.
+    pub fn owner(x: Relay, i: usize) usize {
+        return (x.rank + 1 + i) % x.world;
+    }
+    /// Owner p's rows of the piece (split's shares).
+    pub fn rows(x: Relay, p: usize) [2]usize {
+        return fw.Forward.share(x.n, x.world, p);
+    }
+    /// Rows into a transit where owner(i)'s block begins.
+    pub fn at(x: Relay, i: usize) usize {
+        var o: usize = 0;
+        for (0..i) |k| {
+            const r = x.rows(x.owner(k));
+            o += r[1] - r[0];
+        }
+        return o;
+    }
+    /// Rows into a piece's room where transit t (by step parity) begins: past every slot split fills.
+    pub fn transit(x: Relay, t: usize) usize {
+        const most = (x.n + x.world - 1) / x.world;
+        return x.world * most + t * (x.world - 2) * most;
+    }
+};
+
+/// A piece's partials (`part`: this rank's fp32 rows) into `room`'s slots on `side`; NCCL keeps a pair's posting order.
+pub fn partials(lib: *const cuda.nccl.Library, comm: cuda.nccl.Comm, side: cuda.Stream, x: Relay, room: u64, part: u64, lo: usize, dims: usize) !void {
+    const mine = x.rows(x.rank);
+    const own = mine[1] - mine[0];
+    for (1..x.world) |s| {
+        const from = x.transit((s - 1) % 2); // what step s - 1 kept (step 1 sends from this rank's partial)
+        const into = x.transit(s % 2);
+        try lib.check(lib.api.ncclGroupStart(), "ncclGroupStart");
+        for (0..x.world - s) |i| {
+            const p = x.rows(x.owner(i));
+            const src = if (s == 1) part + (lo + p[0]) * dims * 4 else room + (from + x.at(i)) * dims * 4;
+            try lib.check(lib.api.ncclSend(src, (p[1] - p[0]) * dims, .f32, @intCast(x.next()), comm, side.handle), "ncclSend");
+        }
+        try lib.check(lib.api.ncclRecv(room + x.takenFrom(s) * own * dims * 4, own * dims, .f32, @intCast(x.prev()), comm, side.handle), "ncclRecv");
+        for (0..x.world - s - 1) |i| {
+            const p = x.rows(x.owner(i));
+            try lib.check(lib.api.ncclRecv(room + (into + x.at(i)) * dims * 4, (p[1] - p[0]) * dims, .f32, @intCast(x.prev()), comm, side.handle), "ncclRecv");
+        }
+        try lib.check(lib.api.ncclGroupEnd(), "ncclGroupEnd");
+    }
+}
+
+/// A piece's glued rows (`rows_of`, bf16) to every rank, each step sending on what the last one brought in.
+pub fn glued(lib: *const cuda.nccl.Library, comm: cuda.nccl.Comm, side: cuda.Stream, x: Relay, rows_of: u64, lo: usize, dims: usize) !void {
+    for (1..x.world) |s| {
+        const out = x.rows(x.sentFrom(s));
+        const got = x.rows(x.takenFrom(s));
+        try lib.check(lib.api.ncclGroupStart(), "ncclGroupStart");
+        try lib.check(lib.api.ncclSend(rows_of + (lo + out[0]) * dims * 2, (out[1] - out[0]) * dims, .bf16, @intCast(x.next()), comm, side.handle), "ncclSend");
+        try lib.check(lib.api.ncclRecv(rows_of + (lo + got[0]) * dims * 2, (got[1] - got[0]) * dims, .bf16, @intCast(x.prev()), comm, side.handle), "ncclRecv");
+        try lib.check(lib.api.ncclGroupEnd(), "ncclGroupEnd");
+    }
+}
+
+test "Relay: each owner's slot q holds rank q's block of its rows, as split's direct sends leave it; the transits fit" {
+    const gpa = std.testing.allocator;
+    for (2..7) |w| for ([_]usize{ 0, 1, w - 1, 2 * w + 1, 683 }) |extra| {
+        const n = w + extra; // a split piece has a row for every rank
+        // rank q's partial of piece row j is q << 16 | j (one value a row); every rank's room is the piece's [w * n]
+        const part = try gpa.alloc(u32, w * n);
+        defer gpa.free(part);
+        const rooms = try gpa.alloc(u32, w * w * n);
+        defer gpa.free(rooms);
+        @memset(rooms, std.math.maxInt(u32));
+        for (0..w) |q| for (0..n) |j| {
+            part[q * n + j] = @intCast(q << 16 | j);
+        };
+        var wire: std.ArrayList(u32) = .empty; // what each rank sent its next one this step, rank after rank
+        defer wire.deinit(gpa);
+        var cut: [8]usize = undefined;
+        for (1..w) |s| {
+            wire.clearRetainingCapacity();
+            for (0..w) |r| { // every rank's sends first: a step's receives write nothing its sends read
+                const x: Relay = .{ .world = w, .rank = r, .n = n };
+                const room = rooms[r * w * n ..][0 .. w * n];
+                cut[r] = wire.items.len;
+                for (0..w - s) |i| {
+                    const p = x.rows(x.owner(i));
+                    const len = p[1] - p[0];
+                    try wire.appendSlice(gpa, if (s == 1) part[r * n + p[0] ..][0..len] else room[x.transit((s - 1) % 2) + x.at(i) ..][0..len]);
+                }
+            }
+            for (0..w) |r| {
+                const x: Relay = .{ .world = w, .rank = r, .n = n };
+                const room = rooms[r * w * n ..][0 .. w * n];
+                var got = wire.items[cut[x.prev()]..if (x.prev() + 1 < w) cut[x.prev() + 1] else wire.items.len];
+                const mine = x.rows(r);
+                const own = mine[1] - mine[0];
+                @memcpy(room[x.takenFrom(s) * own ..][0..own], got[0..own]);
+                got = got[own..];
+                for (0..w - s - 1) |i| {
+                    const p = x.rows(x.owner(i));
+                    const len = p[1] - p[0];
+                    @memcpy(room[x.transit(s % 2) + x.at(i) ..][0..len], got[0..len]);
+                    got = got[len..];
+                }
+                try std.testing.expectEqual(@as(usize, 0), got.len);
+            }
+        }
+        for (0..w) |p| {
+            const x: Relay = .{ .world = w, .rank = p, .n = n };
+            const mine = x.rows(p);
+            const own = mine[1] - mine[0];
+            const room = rooms[p * w * n ..][0 .. w * n];
+            @memcpy(room[p * own ..][0..own], part[p * n + mine[0] ..][0..own]); // the own block, as split copies it
+            for (0..w) |q| for (0..own) |k| try std.testing.expectEqual(@as(u32, @intCast(q << 16 | (mine[0] + k))), room[q * own + k]);
+            try std.testing.expect(x.transit(1) + (w - 2) * ((n + w - 1) / w) <= w * n);
+        }
+    };
+}
+
+test "Relay: a piece's glued rows reach every rank, each step sending what the last one brought in" {
+    for (2..7) |w| for ([_]usize{ 0, 1, w - 1, 2 * w + 1 }) |extra| {
+        const n = w + extra;
+        var rows: [6][64]u32 = undefined; // rank r's copy of the piece's rows: its own glued, the rest unknown
+        for (0..w) |r| {
+            const mine = (Relay{ .world = w, .rank = r, .n = n }).rows(r);
+            for (0..n) |j| rows[r][j] = if (j >= mine[0] and j < mine[1]) @intCast(j) else std.math.maxInt(u32);
+        }
+        for (1..w) |s| {
+            var sent: [6][64]u32 = undefined; // every rank's send of step s, before any receive lands
+            for (0..w) |r| {
+                const x: Relay = .{ .world = w, .rank = r, .n = n };
+                const out = x.rows(x.sentFrom(s));
+                @memcpy(sent[r][out[0]..out[1]], rows[r][out[0]..out[1]]);
+            }
+            for (0..w) |r| {
+                const x: Relay = .{ .world = w, .rank = r, .n = n };
+                const y: Relay = .{ .world = w, .rank = x.prev(), .n = n };
+                const got = x.rows(x.takenFrom(s));
+                try std.testing.expectEqual(y.rows(y.sentFrom(s)), got); // rank - 1 sends the block this rank takes
+                @memcpy(rows[r][got[0]..got[1]], sent[x.prev()][got[0]..got[1]]);
+            }
+        }
+        for (0..w) |r| for (0..n) |j| try std.testing.expectEqual(@as(u32, @intCast(j)), rows[r][j]);
+    };
+}

From 98a5dc6578a8a776eb5ffb33a180ab4779af8211 Mon Sep 17 00:00:00 2001
From: Bizuayeu <30091157+Bizuayeu@users.noreply.github.com>
Date: Fri, 16 Oct 2026 16:52:04 +0900
Subject: [PATCH 2/2] nccl: Library.group ends a group whose post failed, so
 later calls stay out of it; the ring relay posts through it, and a fake NCCL
 fails each posting in turn

---
 zig/src/cuda/nccl.zig               |  11 +++
 zig/src/families/glm/cuda_relay.zig | 141 +++++++++++++++++++++++-----
 2 files changed, 127 insertions(+), 25 deletions(-)

diff --git a/zig/src/cuda/nccl.zig b/zig/src/cuda/nccl.zig
index b0d527ce75..018634916b 100644
--- a/zig/src/cuda/nccl.zig
+++ b/zig/src/cuda/nccl.zig
@@ -59,11 +59,22 @@ pub const Library = struct {
 
     pub fn check(self: *const Library, r: Result, what: []const u8) Error!void {
         if (r == 0) return;
+        if (@import("builtin").is_test) return error.NcclFailed; // a test's injected failure: the runner fails error logs
         const s = self.api.ncclGetErrorString(r);
         std.log.err("{s}: {s} ({d})", .{ what, if (s) |p| std.mem.span(p) else "?", r });
         return error.NcclFailed;
     }
 
+    /// `post(ctx)`'s sends and receives as one group; a failed post still ends it, so later calls stay out of it.
+    pub fn group(self: *const Library, ctx: anytype, comptime post: fn (@TypeOf(ctx)) anyerror!void) anyerror!void {
+        try self.check(self.api.ncclGroupStart(), "ncclGroupStart");
+        post(ctx) catch |err| {
+            _ = self.api.ncclGroupEnd();
+            return err;
+        };
+        try self.check(self.api.ncclGroupEnd(), "ncclGroupEnd");
+    }
+
     /// NCCL's version as 10000 * major + 100 * minor + patch.
     pub fn version(self: *const Library) Error!c_int {
         var v: c_int = 0;
diff --git a/zig/src/families/glm/cuda_relay.zig b/zig/src/families/glm/cuda_relay.zig
index 9f12339114..95e8f4a2c5 100644
--- a/zig/src/families/glm/cuda_relay.zig
+++ b/zig/src/families/glm/cuda_relay.zig
@@ -50,37 +50,56 @@ pub const Relay = struct {
 
 /// A piece's partials (`part`: this rank's fp32 rows) into `room`'s slots on `side`; NCCL keeps a pair's posting order.
 pub fn partials(lib: *const cuda.nccl.Library, comm: cuda.nccl.Comm, side: cuda.Stream, x: Relay, room: u64, part: u64, lo: usize, dims: usize) !void {
-    const mine = x.rows(x.rank);
-    const own = mine[1] - mine[0];
-    for (1..x.world) |s| {
-        const from = x.transit((s - 1) % 2); // what step s - 1 kept (step 1 sends from this rank's partial)
-        const into = x.transit(s % 2);
-        try lib.check(lib.api.ncclGroupStart(), "ncclGroupStart");
-        for (0..x.world - s) |i| {
+    for (1..x.world) |s| try lib.group(Step{ .lib = lib, .comm = comm, .side = side, .x = x, .s = s, .base = room, .part = part, .lo = lo, .dims = dims }, Step.partials);
+}
+
+/// A piece's glued rows (`rows_of`, bf16) to every rank, each step sending on what the last one brought in.
+pub fn glued(lib: *const cuda.nccl.Library, comm: cuda.nccl.Comm, side: cuda.Stream, x: Relay, rows_of: u64, lo: usize, dims: usize) !void {
+    for (1..x.world) |s| try lib.group(Step{ .lib = lib, .comm = comm, .side = side, .x = x, .s = s, .base = rows_of, .lo = lo, .dims = dims }, Step.glued);
+}
+
+/// Step `s` of an exchange: its postings, which `Library.group` puts in one group.
+const Step = struct {
+    lib: *const cuda.nccl.Library,
+    comm: cuda.nccl.Comm,
+    side: cuda.Stream,
+    x: Relay,
+    s: usize,
+    base: u64, // the piece's room (partials) or the glued rows
+    part: u64 = 0, // this rank's partial (partials)
+    lo: usize,
+    dims: usize,
+
+    fn send(t: Step, src: u64, count: usize, ty: cuda.nccl.DataType) !void {
+        try t.lib.check(t.lib.api.ncclSend(src, count, ty, @intCast(t.x.next()), t.comm, t.side.handle), "ncclSend");
+    }
+    fn recv(t: Step, dst: u64, count: usize, ty: cuda.nccl.DataType) !void {
+        try t.lib.check(t.lib.api.ncclRecv(dst, count, ty, @intCast(t.x.prev()), t.comm, t.side.handle), "ncclRecv");
+    }
+    fn partials(t: Step) anyerror!void {
+        const x = t.x;
+        const mine = x.rows(x.rank);
+        const own = mine[1] - mine[0];
+        const from = x.transit((t.s - 1) % 2); // what step s - 1 kept (step 1 sends from this rank's partial)
+        const into = x.transit(t.s % 2);
+        for (0..x.world - t.s) |i| {
             const p = x.rows(x.owner(i));
-            const src = if (s == 1) part + (lo + p[0]) * dims * 4 else room + (from + x.at(i)) * dims * 4;
-            try lib.check(lib.api.ncclSend(src, (p[1] - p[0]) * dims, .f32, @intCast(x.next()), comm, side.handle), "ncclSend");
+            const src = if (t.s == 1) t.part + (t.lo + p[0]) * t.dims * 4 else t.base + (from + x.at(i)) * t.dims * 4;
+            try t.send(src, (p[1] - p[0]) * t.dims, .f32);
         }
-        try lib.check(lib.api.ncclRecv(room + x.takenFrom(s) * own * dims * 4, own * dims, .f32, @intCast(x.prev()), comm, side.handle), "ncclRecv");
-        for (0..x.world - s - 1) |i| {
+        try t.recv(t.base + x.takenFrom(t.s) * own * t.dims * 4, own * t.dims, .f32);
+        for (0..x.world - t.s - 1) |i| {
             const p = x.rows(x.owner(i));
-            try lib.check(lib.api.ncclRecv(room + (into + x.at(i)) * dims * 4, (p[1] - p[0]) * dims, .f32, @intCast(x.prev()), comm, side.handle), "ncclRecv");
+            try t.recv(t.base + (into + x.at(i)) * t.dims * 4, (p[1] - p[0]) * t.dims, .f32);
         }
-        try lib.check(lib.api.ncclGroupEnd(), "ncclGroupEnd");
     }
-}
-
-/// A piece's glued rows (`rows_of`, bf16) to every rank, each step sending on what the last one brought in.
-pub fn glued(lib: *const cuda.nccl.Library, comm: cuda.nccl.Comm, side: cuda.Stream, x: Relay, rows_of: u64, lo: usize, dims: usize) !void {
-    for (1..x.world) |s| {
-        const out = x.rows(x.sentFrom(s));
-        const got = x.rows(x.takenFrom(s));
-        try lib.check(lib.api.ncclGroupStart(), "ncclGroupStart");
-        try lib.check(lib.api.ncclSend(rows_of + (lo + out[0]) * dims * 2, (out[1] - out[0]) * dims, .bf16, @intCast(x.next()), comm, side.handle), "ncclSend");
-        try lib.check(lib.api.ncclRecv(rows_of + (lo + got[0]) * dims * 2, (got[1] - got[0]) * dims, .bf16, @intCast(x.prev()), comm, side.handle), "ncclRecv");
-        try lib.check(lib.api.ncclGroupEnd(), "ncclGroupEnd");
+    fn glued(t: Step) anyerror!void {
+        const out = t.x.rows(t.x.sentFrom(t.s));
+        const got = t.x.rows(t.x.takenFrom(t.s));
+        try t.send(t.base + (t.lo + out[0]) * t.dims * 2, (out[1] - out[0]) * t.dims, .bf16);
+        try t.recv(t.base + (t.lo + got[0]) * t.dims * 2, (got[1] - got[0]) * t.dims, .bf16);
     }
-}
+};
 
 test "Relay: each owner's slot q holds rank q's block of its rows, as split's direct sends leave it; the transits fit" {
     const gpa = std.testing.allocator;
@@ -139,6 +158,78 @@ test "Relay: each owner's slot q holds rank q's block of its rows, as split's di
     };
 }
 
+/// NCCL for the relay's postings: each call in order, and the call numbered `fail` returns an error.
+const Fake = struct {
+    const Kind = enum { start, end, send, recv };
+    const Call = struct { kind: Kind, peer: c_int = -1 };
+    var calls: [64]Call = undefined;
+    var n: usize = 0;
+    var fail: ?usize = null;
+
+    fn note(c: Call) cuda.nccl.Result {
+        if (n < calls.len) calls[n] = c;
+        n += 1;
+        return if (fail == n - 1) 1 else 0;
+    }
+    fn start() callconv(.c) cuda.nccl.Result {
+        return note(.{ .kind = .start });
+    }
+    fn end() callconv(.c) cuda.nccl.Result {
+        return note(.{ .kind = .end });
+    }
+    fn send(_: cuda.abi.DevicePtr, _: usize, _: cuda.nccl.DataType, peer: c_int, _: cuda.nccl.Comm, _: cuda.abi.Stream) callconv(.c) cuda.nccl.Result {
+        return note(.{ .kind = .send, .peer = peer });
+    }
+    fn recv(_: cuda.abi.DevicePtr, _: usize, _: cuda.nccl.DataType, peer: c_int, _: cuda.nccl.Comm, _: cuda.abi.Stream) callconv(.c) cuda.nccl.Result {
+        return note(.{ .kind = .recv, .peer = peer });
+    }
+    fn library() cuda.nccl.Library {
+        var api: cuda.nccl.Api = undefined;
+        api.ncclGroupStart = &start;
+        api.ncclGroupEnd = &end;
+        api.ncclSend = &send;
+        api.ncclRecv = &recv;
+        return .{ .lib = undefined, .api = api };
+    }
+    /// One exchange's whole schedule on `lib`: a piece's partials, or with `glue` its glued rows.
+    fn run(lib: *const cuda.nccl.Library, x: Relay, glue: bool) !void {
+        const side: cuda.Stream = .{ .d = undefined, .handle = null };
+        if (glue) try glued(lib, null, side, x, 0, 0, 8) else try partials(lib, null, side, x, 0, 0, 0, 8);
+    }
+};
+
+test "Relay: a posting that fails, at each posting in turn, ends its group and posts nothing after it" {
+    const lib = Fake.library();
+    for (2..7) |w| for (0..w) |r| for ([_]bool{ false, true }) |glue| {
+        const x: Relay = .{ .world = w, .rank = r, .n = 2 * w + 1 };
+        Fake.n = 0;
+        Fake.fail = null;
+        try Fake.run(&lib, x, glue);
+        const total = Fake.n;
+        const clean = Fake.calls; // each failing run below writes over the calls
+        for (clean[0..total]) |c| switch (c.kind) {
+            .send => try std.testing.expectEqual(@as(c_int, @intCast(x.next())), c.peer),
+            .recv => try std.testing.expectEqual(@as(c_int, @intCast(x.prev())), c.peer),
+            else => {},
+        };
+        for (0..total) |k| {
+            const failed = clean[k].kind;
+            Fake.n = 0;
+            Fake.fail = k;
+            try std.testing.expectError(error.NcclFailed, Fake.run(&lib, x, glue));
+            var opened: usize = 0; // a failed GroupStart opens nothing
+            var ended: usize = 0;
+            for (Fake.calls[0..Fake.n], 0..) |c, i| switch (c.kind) {
+                .start => opened += @intFromBool(i != k),
+                .end => ended += 1,
+                else => {},
+            };
+            try std.testing.expectEqual(opened, ended);
+            try std.testing.expectEqual(if (failed == .send or failed == .recv) k + 2 else k + 1, Fake.n);
+        }
+    };
+}
+
 test "Relay: a piece's glued rows reach every rank, each step sending what the last one brought in" {
     for (2..7) |w| for ([_]usize{ 0, 1, w - 1, 2 * w + 1 }) |extra| {
         const n = w + extra;