# B12X f3686b5 TP3 minimal backport

## Finding

Adding an `N=3` CUDA specialization to `f3686b5` is necessary but not
sufficient. The first all-reduce can deadlock because the eager IPC slab is
published before its signal region is initialized.

The old initialization sequence is:

1. Allocate the local signal/eager slab without zero-filling it.
2. Publish the CUDA IPC handle and open peer handles.
3. Zero the local signal region after handle exchange returns.

Ranks can return from the handle exchange at different times. A faster rank
can post its first arrival into a slower peer's `peer_counter`; the slower
rank's late `cudaMemset` then erases that arrival. The barrier waits for an
exact generation value, so the erased write is never recreated and all ranks
spin indefinitely.

Commit `97b3d642` initializes the slab before publishing its IPC handle. It
also moves eager input staging into the CUDA kernel, but that staging change
is not required to fix this observed TP3 deadlock.

## Controlled tests

| Candidate | Result |
| --- | --- |
| `f3686b5` plus TP3 validation/dispatch | Hangs on first FP16, 8-element all-reduce |
| Same plus a leading block `__syncthreads()` | Same first-operation hang |
| Same plus pre-publication signal initialization | Eager, graph, and multistream tests pass |
| Minimal candidate, 256-iteration soak | Passes in 11.97 seconds |

The soak executed 7,936 all-reduce launches per rank: 3,072 eager launches
covering FP16, BF16, FP32 and four message sizes; 4,352 CUDA-graph launches;
and 512 launches split across two streams. That is 23,808 launches across the
three ranks.

## Patch scope

`b12x-f3686b5-tp3-init-race.patch` changes only:

- Python and C++ supported-world-size validation.
- The C++ `N=3` template dispatch.
- Signal initialization ordering so zero-fill completes before IPC handle
  publication.

It does not backport staged-input transport, push mode, fused RMSNorm,
stream-pool changes, or unrelated B12X kernels.