docs(cuda-oxide): state what sync_threads orders at compile time

[alejandro-soto-franco]

The shared-memory chapter states the hardware rule for sync_threads() and says nothing about what the compiler may move across it. Users cannot tell whether raw-pointer or ptx_asm! shared accesses stay on their side of the barrier, so speakrs wraps every shared access in a side-effecting ptx_asm!.

• Adds a table of access forms to the synchronisation section: SharedArray indexing, as_raw_mut_ptr, DynamicSharedArray, ptx_asm! with no options, with clobber("memory") and with options(register_only), cp_async_*, and a reference that outlives the barrier.
• Records how each backend emits the barrier: llvm.nvvm.barrier.cta.sync.aligned.all on the LLVM NVPTX route, and asm sideeffect "bar.sync 0;" with ~{memory} on the libNVVM route.
• Adds the cp_async_commit_group and cp_async_wait_group requirement before the barrier.
• Links the table from the ptx_asm! row of the feature list.

Closes #1463

Test: a throwaway probe example (not in this diff) with 25 kernels ran 21 of them on an RTX 5060 Laptop (sm_120) through both backends, and 20 passed on each route. The failure is the cp.async kernel with no cp_async_wait_group(0), which read the stale value in 128 of 128 threads. The same kernel with the wait read 0 stale values. PTX from llc and from libNVVM shows st.shared, bar.sync, ld.shared, bar.sync, st.shared, bar.sync, ld.shared for every access form, in straight-line code and in loops with a runtime trip count. The four compile-only kernels show that an unused register_only ld.shared leaves no PTX instruction and that a default one stays. sphinx-build -W on the book, check-book-api-names.sh, check-book-import-order.sh and check-source-references.sh pass.

[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]

The table is right. The paragraph about references is not, and I have a kernel that shows it:

• I rebuilt every row as a kernel and ran it through llc and libNVVM on an RTX 5090. Indexing, as_raw_mut_ptr, DynamicSharedArray, default ptx_asm! and clobber("memory") keep store, barrier, load order. An unused register_only load leaves no PTX. cp.async without cp_async_wait_group(0) read the stale value in all 128 threads.
• "The compiler reloads through such a reference today" is false. rustc merges reads through a &T with no interior mutability before our backend runs. This helper is one ld.shared on both routes, and all 128 threads saw the pre-barrier value for the second read:

#[device]
fn read_twice(r: &u32, p: *mut u32, tid: usize) -> (u32, u32) {
    let a = *r;
    thread::sync_threads();
    unsafe { p.add(tid).write(a + 1) }
    thread::sync_threads();
    (a, *r) // same value as a
}

• &UnsafeCell<u32> and the BlockAtomic* types do reload. That is the sound way to share one slot through a reference, so the paragraph should not call it UB. The tip at the top of the page already says SharedArray wraps UnsafeCell for this reason.
• Two smaller things: ptxas drops a dead default snippet from the SASS even when the PTX keeps it, and a register_only snippet can move across memory operations, not only disappear when unused.

To land this: rewrite the "Keep references out of the barrier" paragraph with the rule above and drop the lowering sentence; add the as_raw_mut_ptr example I asked for on #1463 (write your own slot, sync_threads(), read a neighbour's); list as_raw_mut_ptr in the API block; say register_only may move; scope "keeps it" to the PTX.

[nihalpasham]

Reviewed d5bdd5dd: changes requested. The landing list is in #1493 (comment).

[alejandro-soto-franco]

Addressed in ea51cc6:

• Rewrote the references paragraph with your rule and dropped the lowering sentence.
• Added the as_raw_mut_ptr example: write your own slot, call sync_threads(), read a neighbour's. All 128 threads read the right value on both routes on sm_120.
• Listed as_raw_mut_ptr in the API block.
• The register_only row now says it may move across memory operations.
• The default ptx_asm! row now says the PTX keeps the snippet and ptxas may drop it from the SASS.

Your read_twice repro gives the same result here. A plain &u32 is stale in all 128 threads on both routes. &UnsafeCell<u32> and BlockAtomicU32 read the new value.