From bb19e68e4f692b1dff2e6f22b9d8b0deae28862e Mon Sep 17 00:00:00 2001
From: Bizuayeu <30091157+Bizuayeu@users.noreply.github.com>
Date: Fri, 16 Oct 2026 02:43:42 +0900
Subject: [PATCH] glm: a split start exchanges one value with every peer before
 the weights load, refusing a pair that cannot reach each other

---
 zig/src/families/glm/cuda_engine.zig |  1 +
 zig/src/families/glm/cuda_ranks.zig  | 18 ++++++++++++++++++
 2 files changed, 19 insertions(+)

diff --git a/zig/src/families/glm/cuda_engine.zig b/zig/src/families/glm/cuda_engine.zig
index ee38e00b3e..f8cd1fbf2f 100644
--- a/zig/src/families/glm/cuda_engine.zig
+++ b/zig/src/families/glm/cuda_engine.zig
@@ -118,6 +118,7 @@ pub const Engine = struct {
             checkpoint,
         );
         try ranks.agree(d, lib, comm, stream, fp, try a.alloc(ranks.Fingerprint, world));
+        if (o.settings.split) try ranks.probeSplit(d, lib, comm, stream, rank, world);
         const share = try shard.Share.of(&c, rank, world);
         const spans = try gpa.alloc(shard.Span, world);
         errdefer gpa.free(spans);
diff --git a/zig/src/families/glm/cuda_ranks.zig b/zig/src/families/glm/cuda_ranks.zig
index ab5b01b770..d7682c4896 100644
--- a/zig/src/families/glm/cuda_ranks.zig
+++ b/zig/src/families/glm/cuda_ranks.zig
@@ -203,6 +203,24 @@ pub fn agree(
     }
 }
 
+/// Before loading, split's peers: one value each way with each, in rank order (no pair waits on a later one).
+pub fn probeSplit(d: *const cuda.Driver, lib: *const cuda.nccl.Library, comm: cuda.nccl.Comm, stream: cuda.Stream, rank: usize, world: usize) !void {
+    var dev = try cuda.DeviceBuffer.alloc(d, 8);
+    defer dev.free();
+    for (0..world) |peer| if (peer != rank) {
+        try lib.check(lib.api.ncclGroupStart(), "ncclGroupStart");
+        try lib.check(lib.api.ncclSend(dev.ptr, 1, .f32, @intCast(peer), comm, stream.handle), "ncclSend");
+        try lib.check(lib.api.ncclRecv(dev.ptr + 4, 1, .f32, @intCast(peer), comm, stream.handle), "ncclRecv");
+        lib.check(lib.api.ncclGroupEnd(), "ncclGroupEnd") catch return unreachablePeer(rank, peer);
+        stream.synchronize() catch return unreachablePeer(rank, peer);
+    };
+}
+
+fn unreachablePeer(rank: usize, peer: usize) error{SplitUnreachable} {
+    std.debug.print("REFUSE TF_GLM_PREFILL_REDUCE=split: rank {d} cannot exchange with rank {d}; split sends every rank's rows to every other, and a fabric where some ranks reach only their neighbours (a ring of four without a switch) needs TF_GLM_PREFILL_REDUCE=gather\n", .{ rank, peer });
+    return error.SplitUnreachable;
+}
+
 fn explain(f: Fingerprint.Field) []const u8 {
     return switch (f) {
         .layout => "a different build of this check",