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

cuda-oxide/cuda-oxide-book/advanced/shared-memory-and-synchronization.md

@@ -140,6 +140,7 @@ impl<T, const N: usize, const ALIGN: usize> SharedArray<T, N, ALIGN> {
     pub const fn len() -> usize;
     pub fn as_ptr(&self) -> *const T;
     pub fn as_mut_ptr(&mut self) -> *mut T;
+    pub unsafe fn as_raw_mut_ptr(shared: *mut Self) -> *mut T;
 }
 
 // Indexing (unsafe access via static mut)
@@ -272,6 +273,105 @@ that not all threads take. Every thread in the block must reach the same
 control flow, restructure so the barrier is outside the branch.
 :::
 
+(compiler-ordering-around-the-barrier)=
+
+### Compiler ordering around the barrier
+
+The hardware guarantee above covers visibility between threads. A second
+question is whether the compiler keeps your own accesses on the side of the
+barrier where you wrote them. Both backends emit the barrier as an operation
+that the compiler treats as reading and writing all memory:
+
+- The LLVM NVPTX route calls `llvm.nvvm.barrier.cta.sync.aligned.all`. The
+  call is `convergent` and has no `memory(...)` attribute.
+- The libNVVM route emits `bar.sync 0;` as `asm sideeffect` with a `~{memory}`
+  clobber.
+
+| Access form                                              | Kept on its side of `sync_threads()` | Notes |
+| :------------------------------------------------------- | :----------------------------------- | :---- |
+| `SharedArray` indexing, `S[i]`                           | Yes | A plain load or store in address space 3. |
+| `SharedArray::as_raw_mut_ptr(&raw mut S)`                | Yes | The same instructions as indexing. Prefer it to `as_mut_ptr()` and `as_ptr()`. Those methods borrow `S`, and the `static_mut_refs` lint rejects a borrow of a `static mut` under edition 2024. |
+| `DynamicSharedArray::<T>::get()` and `offset()`          | Yes | A plain load or store through the returned pointer. |
+| `ptx_asm!` with no options                               | Yes | The call is `asm sideeffect` with no memory attribute. The PTX keeps it in order against the barrier, plain loads and stores, and other side-effecting asm. ptxas may drop a snippet with an unused result from the SASS. |
+| `ptx_asm!` with `clobber("memory")`                      | Yes | Adds `~{memory}` to the constraint string. The order matches the row above in the current lowering. Without the clobber that order comes from the `sideeffect` flag alone. That flag does not tell the compiler that the snippet touches memory. Add it to every snippet that loads or stores through an address. |
+| `ptx_asm!` with `options(register_only)`                 | No | The call has no `sideeffect` flag. The compiler may move it across memory operations and the barrier, and deletes it when its result is unused. Use it for snippets that touch registers only. |
+| `cp_async_*` copies, `cp_async_commit_group()`, `cp_async_wait_group(n)` | Yes, but see below | Issue, commit and wait stay in order. The barrier does not wait for the copy to land. |
+| A `&T` or `&[T]` to slots another thread writes, live across the barrier | No | rustc merges reads through it before either backend runs. See below. |
+| A `&UnsafeCell<T>` to such a slot | Yes | A plain load that is re-read after each barrier. Two adjacent reads with nothing between them may still become one load. |
+| A reference to a `BlockAtomic*` type | Yes | Each `load()` is side-effecting asm with a memory clobber. The compiler keeps every one. |
+
+The indexing, `as_raw_mut_ptr`, `DynamicSharedArray` and both side-effecting
+`ptx_asm!` rows were checked on both backends with a producer and consumer
+kernel: store, barrier, load, barrier, store, barrier, load. The PTX keeps that
+order for each of them. It keeps the same order when the loads sit in a loop
+whose trip count is a kernel argument. A run on an RTX 5060 Laptop (sm_120)
+read the expected neighbour values. A `register_only` `ld.shared` whose result
+is unused leaves no instruction in the PTX. The same snippet with no options
+stays. The `cp.async` and reference rows have their own tests below.
+
+The table covers compiler ordering only. The rule under "When to sync" still
+decides where the barriers belong.
+
+#### Raw pointers across the barrier
+
+Each thread derives its pointer from the raw address of the allocation and
+creates no reference to the whole array. A thread writes its own slot
+before the barrier and reads a neighbour's slot after it:
+
+```rust
+static mut TILE: SharedArray<u32, 128> = SharedArray::UNINIT;
+
+let tid = thread::threadIdx_x() as usize;
+let tile = unsafe { SharedArray::as_raw_mut_ptr(&raw mut TILE) };
+unsafe { tile.add(tid).write(tid as u32 + 1) }; // this thread's slot
+thread::sync_threads();
+let next = unsafe { tile.add((tid + 1) % 128).read() }; // a neighbour's slot
+```
+
+A kernel built from this snippet read the neighbour's value in all 128 threads
+on both backends.
+
+#### `cp.async` needs a wait before the barrier
+
+`bar.sync` does not wait for a pending `cp.async`. Each thread waits for its
+own copies and then every thread meets at the barrier:
+
+```rust
+unsafe {
+    cp_async_ca_4(dst, src);
+    cp_async_commit_group();
+    cp_async_wait_group(0); // this thread's copies have landed
+}
+thread::sync_threads(); // other threads may now read `dst`
+```
+
+Reading the destination before the copy completes is undefined behaviour. A
+test kernel that drops `cp_async_wait_group(0)` and reads a neighbour's slot
+after the barrier can return the stale value: in a plain run on sm_120 all 128
+threads did, and under compute-sanitizer every copy had landed first. The same
+kernel with the wait returned the copied value in all 128 threads.
+
+#### Share a slot through a reference only with `UnsafeCell`
+
+When `T` has no interior mutability, a `&T` into shared memory tells rustc
+that the target does not change while the reference is live. rustc merges two
+reads through it into one load before either backend runs. In a test kernel
+each thread read its neighbour's slot through a `&u32`, passed two barriers
+while the neighbour wrote a new value, and read again. The second read
+returned the old value in all 128 threads on both backends. A `&[T]` built
+over the tile with `from_raw_parts` is the same case: two indexed reads across
+the barrier compile to one load. Keeping such a reference live across a barrier
+while another thread writes the target is undefined behaviour. So is a `&mut T`
+to such a slot: the other thread's write breaks its exclusive access, even
+though no pass merges reads through it today.
+
+A `&UnsafeCell<T>` is re-read after each barrier, and every `BlockAtomic*`
+load is an instruction the compiler keeps. The same kernel with either form
+read the new value in all 128 threads.
+Share a slot through one of these, or read through a raw pointer after each
+barrier. The tip under "Declaration rules" gives the same reason for
+`SharedArray` wrapping `UnsafeCell`.
+
 ---
 
 ## Shared memory vs. other approaches

cuda-oxide/cuda-oxide-book/appendix/supported-features.md

@@ -192,7 +192,7 @@ Anonymous promoted allocations remain unsupported.
 
 | Feature | Status | Description |
 |:--------|:-------|:------------|
-| `ptx_asm!` Macro | **Partial** | CUDA inline PTX with `%0` operands, `in`, `out`, and `inout`; up to 16 output operands across `out` with `=`-prefixed constraints and `inout` with `+`-prefixed constraints; up to 16 explicit inputs; CUDA register constraints `h`, `r`, `l`, `q`, `f`, and `d`; immediate integer constraint `n`; compile-time string constraint `C`; `clobber("memory")`; and `options(register_only)` for pure register snippets. With multiple output operands, the macro writes tuple results back in declaration order. By default, snippets are treated as side-effecting and stay inside their current control flow. Use `options(register_only, may_diverge)` only for pure snippets that are safe to move across divergent control flow; **never** use it for `.sync` instructions or collectives. More than 16 output operands are not implemented yet. |
+| `ptx_asm!` Macro | **Partial** | CUDA inline PTX with `%0` operands, `in`, `out`, and `inout`; up to 16 output operands across `out` with `=`-prefixed constraints and `inout` with `+`-prefixed constraints; up to 16 explicit inputs; CUDA register constraints `h`, `r`, `l`, `q`, `f`, and `d`; immediate integer constraint `n`; compile-time string constraint `C`; `clobber("memory")`; and `options(register_only)` for pure register snippets. With multiple output operands, the macro writes tuple results back in declaration order. By default, snippets are treated as side-effecting and stay inside their current control flow. Use `options(register_only, may_diverge)` only for pure snippets that are safe to move across divergent control flow; **never** use it for `.sync` instructions or collectives. More than 16 output operands are not implemented yet. For how each option orders against `sync_threads()`, see {ref}`Compiler ordering around the barrier <compiler-ordering-around-the-barrier>`. |
 
 ---