diff --git a/cuda-oxide/cuda-oxide-book/advanced/shared-memory-and-synchronization.md b/cuda-oxide/cuda-oxide-book/advanced/shared-memory-and-synchronization.md index c068c8af4f..aa2afd6b47 100644 --- a/cuda-oxide/cuda-oxide-book/advanced/shared-memory-and-synchronization.md +++ b/cuda-oxide/cuda-oxide-book/advanced/shared-memory-and-synchronization.md @@ -140,6 +140,7 @@ impl SharedArray { 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::::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` 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 = 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` 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 diff --git a/cuda-oxide/cuda-oxide-book/appendix/supported-features.md b/cuda-oxide/cuda-oxide-book/appendix/supported-features.md index c038225a85..5c89c5a15c 100644 --- a/cuda-oxide/cuda-oxide-book/appendix/supported-features.md +++ b/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 `. | ---