[YeonwooSung]

What this adds

ManagedBarrier::init_by and inval_by elect the thread that runs mbarrier.init and mbarrier.inval by its block-linear id, threadIdx.x + blockDim.x * (threadIdx.y + blockDim.y * threadIdx.z), instead of by threadIdx.x alone. init() and inval() still select id 0, which is thread (0, 0, 0) in every block layout, and 1D blocks keep the same thread as before.

block (32, 4, 1), every thread calls init_by(128, 127)
before: threadIdx.x == 127 matches nobody; 128 threads arrive on an uninitialized word; error 719
after:  block-linear id 127 is thread (31, 3, 0); one mbarrier.init; all 128 threads complete

block (32, 4, 1), every thread calls init()
before: four threads (x = 0, y = 0..3) each run mbarrier.init on the same word
after:  one thread, (0, 0, 0)

// launched with block_dim = (32, 4, 1); every thread calls all of these
let bar = ManagedBarrier::<Uninit, GeneralBarrier>::from_static(&raw mut BAR);
let bar = unsafe { bar.init_by(128, 127) }; // block-linear id 127 = thread (31, 3, 0)
let t = bar.arrive();
bar.wait(t);
let _ = unsafe { bar.inval_by(127) };

What we fixed

• The election compared threadIdx.x with the selector. In a 2D or 3D block that admits one thread per row for an in-range x and no thread at all for a selector at or past blockDim.x; both are undefined per the PTX ISA, and the second faults with error 719.
• The election now goes through thread::block_linear_thread_id_here(), the helper fix(cuda-device): select tcgen05 alloc warp by linear thread id #1351 added and the numbering this_thread_block().thread_rank() uses, so warp election and thread election agree on how a block is numbered.
• The thread-requirement docs described a thread-0-only protocol: from_static(), init() and inval() are all-threads calls, and the from_static example called init() inside if tid == 0, a block-wide sync inside a divergent branch.

Contract

• init_by(count, id) and inval_by(id): all threads of the block call them; the thread whose block-linear id equals id runs the instruction. id must be the same value in every thread and below blockDim.x * blockDim.y * blockDim.z; an out-of-range id selects no thread, so the barrier is never initialized or invalidated.
• init() and inval() select id 0, thread (0, 0, 0).

Implementation

• barrier.rs: init_by and inval_by compare crate::thread::block_linear_thread_id_here() with the selector; these two sites are the only elections in the file, and init/inval delegate to them. The private copy of the formula and its tests are gone.
• thread.rs tests: every in-range id selects one thread across 1D, 2D, 3D and non-power-of-two layouts, and an id equal to the thread count selects nobody.
• future_apis: test_managed_barrier_2d launches (32, 4, 1); every thread calls init_by(128, 127), arrive, wait, inval_by(127); thread 0 fills the barrier word with all-ones first so a missing init faults instead of depending on what shared memory held; each thread writes 42 when its row-major index equals its thread_rank(), and the host returns an error on any other value.
• The generated mbarrier_init and mbarrier_inval intrinsics and the catalog are untouched; the defect was in the wrapper's election.

Verification

Checked at fd2bb3902 (nightly-2026-10-11, RTX 5090, sm_120a):

• ✅ cargo test -p cuda-device --all-targets passes with the new property test; cargo clippy --workspace --all-targets -- -D warnings is clean (the submitted head failed on clippy::identity_op in the removed tests); rustdoc with warnings as errors, fmt, the 18 guards and the lock check pass.
• ✅ Fails before, passes after: the same future_apis source built against main's cuda-device fails Test 6 with DriverError(719); this branch passes Tests 1 to 6 with memcheck, synccheck and racecheck clean.
• ✅ Scratch kernels, this branch against main with the same backend: (32,4,1), (8,4,2), (24,5,1), (1,64,1), (1,1,64), (3,5,7), (33,3,1) and two 1024-thread layouts fail on main with error 719 and pass here, through the llc route, the libNVVM route and with optimization off; 1D controls pass on both.
• ✅ PTX: one predicated mbarrier.init and one mbarrier.inval per kernel, guarded by the block-linear id; main guards on tid.x alone.

Caveats

• A 2D caller that passed an x coordinate as the selector now gets the block-linear id; no in-tree caller is 2D, and the old behaviour was already undefined.
• The four-initializer case (init() in a 2D block on main) runs silently on sm_120 and the sanitizers stay quiet; the PTX ISA rule is the evidence for it, not a runtime failure.
• The older future_apis kernels still call inval() inside if is_thread0 after a sync_threads; that predates this change and is a follow-up.

[copy-pr-bot]

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

[nihalpasham]

/ok to test fd2bb39

[nihalpasham]

Thanks, this is the right fix: mbarrier.init and mbarrier.inval must run on one thread, and in a 2D or 3D block that thread is the block-linear id, the numbering thread_rank() uses. Four things on top:

• Routed the election through thread::block_linear_thread_id_here(), which fix(cuda-device): select tcgen05 alloc warp by linear thread id #1351 landed after this branch was written, and removed the private predicate with its tests. The 4 * 1 literal in those tests also failed the clippy gate CI runs with -D warnings.
• Fixed the thread-requirement docs in the same file: from_static(), init() and inval() are all-threads calls, linear thread 0 acts, and the selector must be the same value in every thread and below the thread count. The old from_static example called init() inside if tid == 0, a block-wide sync inside a divergent branch.
• Added a (32, 4, 1) regression to future_apis: every thread calls init_by(128, 127) and inval_by(127), and the host returns an error on any wrong value. On main no thread has threadIdx.x == 127, nothing runs mbarrier.init, and the launch fails with error 719; here all 128 threads complete.
• One test next to the shared helper, and a one-page figure in the book.

At fd2bb3902, workspace clippy, the cuda-device tests, the guards and all CI checks passed. On the 5090 the 2D and 3D kernels fail on main with error 719 for nine layouts and pass here through both routes, with memcheck, synccheck and racecheck clean. 1D kernels keep the same thread.