From d5bdd5dd75138aa54dad5afc255a68bbe2cc1955 Mon Sep 17 00:00:00 2001 From: Alejandro Soto Franco Date: Sun, 11 Oct 2026 16:10:36 +0300 Subject: [PATCH 1/3] docs(cuda-oxide): state what sync_threads orders at compile time The shared-memory chapter explains the hardware rule for sync_threads and says nothing about what the compiler may move across it. Users could not tell whether raw-pointer or ptx_asm! shared accesses are safe to leave visible to the optimiser, and wrapped every access in a side-effecting ptx_asm!. Add a table of access forms to the synchronisation section. It covers SharedArray indexing, raw pointers, DynamicSharedArray, ptx_asm! with no options, with clobber("memory") and with register_only, cp.async, and references held across the barrier. Add the cp.async commit and wait requirement before the barrier. Link the table from the ptx_asm! row of the feature list. Closes #1463 Signed-off-by: Alejandro Soto Franco --- .../shared-memory-and-synchronization.md | 62 +++++++++++++++++++ .../appendix/supported-features.md | 2 +- 2 files changed, 63 insertions(+), 1 deletion(-) 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..3d0fd5fb00 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 @@ -272,6 +272,68 @@ 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 + +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()`, which needs `&mut S` and is rejected by the `static_mut_refs` lint 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 compiler keeps it, and keeps its order against the barrier and other side-effecting asm. | +| `ptx_asm!` with `clobber("memory")` | Yes | Adds `~{memory}` to the constraint string. The order is the same as with no options today. The clobber declares a memory access that the operands do not name. Add it to every snippet that loads or stores through an address. | +| `ptx_asm!` with `options(register_only)` | Not promised | The call has no `sideeffect` flag, and the compiler 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 `&mut T` into shared memory that outlives the barrier | Not promised | Undefined behaviour in Rust when another thread writes after the barrier. | + +The first six 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 every form and keeps it 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, while the same snippet with no options stays. +The `register_only` loads in the test kernels kept their order against the +barrier. The macro does not promise that order. Treat it as unspecified. + +The table covers compiler ordering only. The rule under "When to sync" still +decides where the barriers belong. + +#### `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` +``` + +A test kernel that drops `cp_async_wait_group(0)` and reads a neighbour's slot +after the barrier returned the stale value in all 128 threads on sm_120. The +same kernel with the wait returned the copied value in all 128. + +#### Keep references out of the barrier + +A `&T` into shared memory promises that the target does not change while the +reference is live. A store by another thread after the barrier breaks that +promise. Reading through the same reference before and after the barrier is +undefined behaviour under Rust's aliasing rules. The compiler reloads through +such a reference today, and the current lowering is the only reason. Read +through a fresh raw pointer or index after each barrier. + --- ## 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 0cd0412d38..6518fd32a6 100644 --- a/cuda-oxide/cuda-oxide-book/appendix/supported-features.md +++ b/cuda-oxide/cuda-oxide-book/appendix/supported-features.md @@ -181,7 +181,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 [Shared Memory and Synchronization](../advanced/shared-memory-and-synchronization.md). | --- From ea51cc66c1cd31fea652da2f69dc8184cefeccc8 Mon Sep 17 00:00:00 2001 From: Alejandro Soto Franco Date: Thu, 15 Oct 2026 06:07:55 -0600 Subject: [PATCH 2/3] docs(cuda-oxide): share a slot by reference only through UnsafeCell rustc merges two reads through a plain &T before either backend runs. The paragraph claiming the compiler reloads was wrong. A test kernel read the stale value in all 128 threads on both routes. &UnsafeCell and BlockAtomicU32 reloaded. Adds the as_raw_mut_ptr example from #1463: write your own slot, call sync_threads(), read a neighbour's. It passes on both backends. Lists as_raw_mut_ptr in the API block. register_only may move across memory operations. The default ptx_asm! row names the PTX as what keeps the snippet. The verification paragraph names the rows it covers. --- .../shared-memory-and-synchronization.md | 73 +++++++++++++------ 1 file changed, 51 insertions(+), 22 deletions(-) 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 3d0fd5fb00..7930afc7ce 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) @@ -287,26 +288,46 @@ that the compiler treats as reading and writing all memory: | 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()`, which needs `&mut S` and is rejected by the `static_mut_refs` lint under edition 2024. | +| `SharedArray::as_raw_mut_ptr(&raw mut S)` | Yes | The same instructions as indexing. Prefer it to `as_mut_ptr()`. That method needs `&mut S`, and the `static_mut_refs` lint rejects it 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 compiler keeps it, and keeps its order against the barrier and other side-effecting asm. | -| `ptx_asm!` with `clobber("memory")` | Yes | Adds `~{memory}` to the constraint string. The order is the same as with no options today. The clobber declares a memory access that the operands do not name. Add it to every snippet that loads or stores through an address. | -| `ptx_asm!` with `options(register_only)` | Not promised | The call has no `sideeffect` flag, and the compiler deletes it when its result is unused. Use it for snippets that touch registers only. | +| `ptx_asm!` with no options | Yes | The call is `asm sideeffect` with no memory attribute. The PTX keeps it in order against the barrier 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 declares no memory access. 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 `&mut T` into shared memory that outlives the barrier | Not promised | Undefined behaviour in Rust when another thread writes after the barrier. | - -The first six 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 every form and keeps it 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, while the same snippet with no options stays. -The `register_only` loads in the test kernels kept their order against the -barrier. The macro does not promise that order. Treat it as unspecified. +| A `&T` to a slot another thread writes, live across the barrier | No | rustc merges reads through it before either backend runs. See below. | +| A `&UnsafeCell` or a reference to a `BlockAtomic*` type | Yes | Each read reloads the slot. | + +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 @@ -325,14 +346,22 @@ A test kernel that drops `cp_async_wait_group(0)` and reads a neighbour's slot after the barrier returned the stale value in all 128 threads on sm_120. The same kernel with the wait returned the copied value in all 128. -#### Keep references out of the barrier - -A `&T` into shared memory promises that the target does not change while the -reference is live. A store by another thread after the barrier breaks that -promise. Reading through the same reference before and after the barrier is -undefined behaviour under Rust's aliasing rules. The compiler reloads through -such a reference today, and the current lowering is the only reason. Read -through a fresh raw pointer or index after each barrier. +#### 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. Keeping such a +reference live across a barrier while another thread writes the target is +undefined behaviour. + +A `&UnsafeCell` or a reference to a `BlockAtomic*` type reloads at each +read. 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`. --- From b83ea3ebaf86e6e9358c4f0927631bf92b272e67 Mon Sep 17 00:00:00 2001 From: nihalpasham Date: Mon, 19 Oct 2026 15:38:37 +0530 Subject: [PATCH 3/3] docs(book): tighten the compiler ordering rows at sync_threads Reword the "Compiler ordering around the barrier" section so each row says what the lowering guarantees. The clobber row said the sideeffect flag "declares no memory access", which reads as its opposite. The default ptx_asm! row now names plain loads and stores among the operations it is ordered against. The reference rule covers a &[T] over the tile and a &mut T to a shared slot. The UnsafeCell and BlockAtomic rows are split: a &UnsafeCell is re-read after each barrier, while every BlockAtomic load is kept. The cp.async paragraph leads with the ISA rule and states the no-wait measurement as a race. The as_raw_mut_ptr row names as_ptr() beside as_mut_ptr() under the static_mut_refs lint. Add an explicit MyST target above the section so the ptx_asm! row of the feature list links to the table instead of the page. Signed-off-by: nihalpasham --- .../shared-memory-and-synchronization.md | 37 ++++++++++++------- .../appendix/supported-features.md | 2 +- 2 files changed, 24 insertions(+), 15 deletions(-) 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 7930afc7ce..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 @@ -273,6 +273,8 @@ 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 @@ -288,14 +290,15 @@ that the compiler treats as reading and writing all memory: | 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()`. That method needs `&mut S`, and the `static_mut_refs` lint rejects it under edition 2024. | +| `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 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 declares no memory access. Add it to every snippet that loads or stores through an address. | +| `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` to a slot another thread writes, live across the barrier | No | rustc merges reads through it before either backend runs. See below. | -| A `&UnsafeCell` or a reference to a `BlockAtomic*` type | Yes | Each read reloads the slot. | +| 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 @@ -342,9 +345,11 @@ unsafe { thread::sync_threads(); // other threads may now read `dst` ``` -A test kernel that drops `cp_async_wait_group(0)` and reads a neighbour's slot -after the barrier returned the stale value in all 128 threads on sm_120. The -same kernel with the wait returned the copied value in all 128. +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` @@ -353,12 +358,16 @@ 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. Keeping such a -reference live across a barrier while another thread writes the target is -undefined behaviour. - -A `&UnsafeCell` or a reference to a `BlockAtomic*` type reloads at each -read. The same kernel with either form read the new value in all 128 threads. +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`. diff --git a/cuda-oxide/cuda-oxide-book/appendix/supported-features.md b/cuda-oxide/cuda-oxide-book/appendix/supported-features.md index c5e103264f..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. For how each option orders against `sync_threads()`, see [Shared Memory and Synchronization](../advanced/shared-memory-and-synchronization.md). | +| `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 `. | ---