diff --git a/crates/cuda-device/src/lib.rs b/crates/cuda-device/src/lib.rs index a3171afbc..ea8b7528f 100644 --- a/crates/cuda-device/src/lib.rs +++ b/crates/cuda-device/src/lib.rs @@ -82,7 +82,7 @@ pub use cusimd::{CuSimd, Float2, Float4, TmemRegs4, TmemRegs32}; pub use disjoint::{__LaunchContractDisjointSlice, __LaunchContractDisjointSliceAbi}; pub use disjoint::{DisjointSlice, SpaceLayout}; pub use fence::*; -pub use shared::{DynamicSharedArray, SharedArray}; +pub use shared::{DynamicSharedArray, SharedArray, SharedPtr}; pub use tcgen05::{ TensorMemoryHandle, TmemAddress, TmemDeallocated, TmemF32x4, TmemF32x32, TmemGuard, TmemReady, TmemUninit, diff --git a/crates/cuda-device/src/shared.rs b/crates/cuda-device/src/shared.rs index 5467ee7d3..9d1102551 100644 --- a/crates/cuda-device/src/shared.rs +++ b/crates/cuda-device/src/shared.rs @@ -220,6 +220,23 @@ impl SharedArray { let _ = shared; unreachable!("SharedArray::as_raw_mut_ptr called outside CUDA kernel context") } + + /// Returns a [`SharedPtr`] to the first element of a shared array. + /// + /// Unlike [`Self::as_raw_mut_ptr`], which returns a generic pointer, the + /// result keeps the array's CTA-shared address space in its type, so it + /// stays a shared-memory pointer when it is passed to helper functions. + /// + /// ```rust,ignore + /// static mut TILE: SharedArray = SharedArray::UNINIT; + /// let tile: SharedPtr = SharedArray::shared_ptr(&raw mut TILE); + /// ``` + #[inline(always)] + pub fn shared_ptr(shared: *mut Self) -> SharedPtr { + SharedPtr { + ptr: shared as *mut SharedSlot, + } + } } impl Index for SharedArray { @@ -493,6 +510,24 @@ impl DynamicSharedArray { let _ = byte_offset; unreachable!("DynamicSharedArray::offset called outside CUDA kernel context") } + + /// Like [`Self::get`], but returns a [`SharedPtr`], which keeps the + /// CTA-shared address space in its type. + #[inline(always)] + pub fn shared_ptr() -> SharedPtr { + SharedPtr { + ptr: Self::get() as *mut SharedSlot, + } + } + + /// Like [`Self::offset`], but returns a [`SharedPtr`], which keeps the + /// CTA-shared address space in its type. + #[inline(always)] + pub fn shared_ptr_at(byte_offset: usize) -> SharedPtr { + SharedPtr { + ptr: Self::offset(byte_offset) as *mut SharedSlot, + } + } } /// Convert a generic-address pointer into its raw `.shared` window offset. @@ -546,4 +581,136 @@ pub unsafe fn cvta_generic_to_shared_u32(ptr: *const u8) -> u32 { unreachable!("cvta_generic_to_shared_u32 called outside CUDA kernel context") } +// ============================================================================ +// SharedPtr - a pointer into the issuing CTA's shared memory, typed as such +// ============================================================================ + +/// The pointee of [`SharedPtr`]'s field. A pointer to `SharedSlot` is a +/// pointer into CTA-shared memory (`addrspace(3)`) by its type, as a pointer +/// to [`SharedArray`] is; `repr(transparent)` gives it `T`'s size, so pointer +/// arithmetic strides by `T`. +#[doc(hidden)] +#[repr(transparent)] +pub struct SharedSlot(T); + +/// A raw pointer into the shared memory of the issuing CTA (`.shared::cta`). +/// +/// A `*mut T` in device code is a generic pointer, which may address global, +/// local or shared memory, including another CTA's shared memory in the +/// cluster. `SharedPtr` instead carries its address space in its type: the +/// compiler knows it is a CTA-shared pointer wherever it goes, including +/// through calls to helper functions, and emits the `.shared::cta` forms of +/// instructions that take it without converting a generic address. +/// +/// It is created from shared memory the CTA owns, with +/// [`SharedArray::shared_ptr`] or [`DynamicSharedArray::shared_ptr`], or from +/// a generic pointer with the `unsafe` [`SharedPtr::from_raw`]. Like a raw +/// pointer, it may dangle; dereferencing it goes through [`SharedPtr::as_ptr`] +/// and `unsafe` code. +/// +/// Pointers from [`crate::cluster::map_shared_rank`], which may name another +/// CTA's shared memory, are not `SharedPtr`s. +/// +/// # Example +/// +/// ```rust,ignore +/// static mut TILE: SharedArray = SharedArray::UNINIT; +/// +/// fn store(tile: SharedPtr, i: usize, value: f32) { +/// // `tile` is still a shared pointer here, although `store` is a +/// // separate function. +/// unsafe { tile.add(i).as_ptr().write(value) }; +/// } +/// +/// let tile = SharedArray::shared_ptr(&raw mut TILE); +/// store(tile, thread::threadIdx_x() as usize, 1.0); +/// ``` +#[repr(transparent)] +pub struct SharedPtr { + ptr: *mut SharedSlot, +} + +impl Clone for SharedPtr { + #[inline(always)] + fn clone(&self) -> Self { + *self + } +} + +impl Copy for SharedPtr {} + +impl SharedPtr { + /// Converts a generic pointer to a `SharedPtr`, as CUDA C++'s + /// `__cvta_generic_to_shared` or CuTe's `make_smem_ptr` do. + /// + /// # Safety + /// + /// `ptr` must point into the shared memory of the issuing CTA. A pointer + /// to global or local memory, or to another CTA's shared memory, gives an + /// unspecified address. + #[inline(always)] + pub unsafe fn from_raw(ptr: *mut T) -> Self { + Self { + ptr: ptr as *mut SharedSlot, + } + } + + /// The generic pointer to the same memory, for ordinary loads and stores + /// and for APIs that take raw pointers. + #[inline(always)] + pub fn as_ptr(self) -> *mut T { + self.ptr as *mut T + } + + /// The pointer's 32-bit `.shared::cta` address, as taken by the + /// `_shared_u32` forms of CTA-shared instructions. + #[inline(always)] + pub fn addr_u32(self) -> u32 { + // SAFETY: `self` points into the issuing CTA's shared memory. + unsafe { cvta_generic_to_shared_u32(self.as_ptr() as *const u8) } + } + + /// Reinterprets the pointee type, keeping the address space. + #[inline(always)] + pub fn cast(self) -> SharedPtr { + SharedPtr { + ptr: self.ptr as *mut SharedSlot, + } + } + + /// The pointer `count` elements of `T` further on, as `<*mut T>::add`. + /// + /// # Safety + /// + /// As for `<*mut T>::add`. + #[inline(always)] + pub unsafe fn add(self, count: usize) -> Self { + Self { + ptr: unsafe { self.ptr.add(count) }, + } + } + + /// The pointer `count` elements of `T` away, as `<*mut T>::offset`. + /// + /// # Safety + /// + /// As for `<*mut T>::offset`. + #[inline(always)] + pub unsafe fn offset(self, count: isize) -> Self { + Self { + ptr: unsafe { self.ptr.offset(count) }, + } + } + + /// The pointer `count` bytes further on, as `<*mut T>::byte_add`. + /// + /// # Safety + /// + /// As for `<*mut T>::byte_add`. + #[inline(always)] + pub unsafe fn byte_add(self, count: usize) -> Self { + unsafe { self.cast::().add(count).cast::() } + } +} + include!("generated/shared_sreg.rs"); diff --git a/crates/mir-importer/src/translator/facts.rs b/crates/mir-importer/src/translator/facts.rs index 2be0ceafe..e585fac95 100644 --- a/crates/mir-importer/src/translator/facts.rs +++ b/crates/mir-importer/src/translator/facts.rs @@ -389,6 +389,15 @@ mod pointer_origin { } } + /// ABI RULE: `cuda_device::SharedPtr` stores a `*mut SharedSlot`, so + /// the shared pointer it lowers to is always `RawMut`. + pub(crate) fn abi_shared_ptr_field() -> PointerOrigin { + PointerOrigin { + kind: MirPointerKind::RawMut, + mutable: true, + } + } + /// ABI RULE: `cuda_device::DisjointSlice<'_, T>` stores its data pointer /// as `*mut T`, so its data/element pointers are always `RawMut`. pub(crate) fn abi_disjoint_slice_data_ptr() -> PointerOrigin { diff --git a/crates/mir-importer/src/translator/rvalue/place_addr.rs b/crates/mir-importer/src/translator/rvalue/place_addr.rs index 0b09daa48..f48b6d74b 100644 --- a/crates/mir-importer/src/translator/rvalue/place_addr.rs +++ b/crates/mir-importer/src/translator/rvalue/place_addr.rs @@ -64,6 +64,7 @@ pub(crate) fn translate_place_address( prev_op: Option>, loc: Location, ) -> TranslationResult>)>> { + let place = &*super::place_read::strip_shared_ptr_field_projections(body, place); let Some(slot) = value_map.get_slot(place.local) else { return Ok(None); }; diff --git a/crates/mir-importer/src/translator/rvalue/place_read.rs b/crates/mir-importer/src/translator/rvalue/place_read.rs index d1e6a0f39..b0fc482f0 100644 --- a/crates/mir-importer/src/translator/rvalue/place_read.rs +++ b/crates/mir-importer/src/translator/rvalue/place_read.rs @@ -80,6 +80,7 @@ pub fn translate_place( prev_op: Option>, loc: Location, ) -> TranslationResult<(Value, Option>)> { + let place = &*strip_shared_ptr_field_projections(body, place); match classify_place_read_strategy(ctx, place, value_map)? { PlaceReadStrategy::Address => { if let Some((value, last_op)) = translate_place_load_from_address( @@ -109,6 +110,51 @@ pub fn translate_place( } } +/// Drop the projections of a place onto the field of a `cuda_device::SharedPtr`. +/// +/// `SharedPtr { ptr: *mut SharedSlot }` is a scalar-lowered newtype: its +/// value is its one field, a CTA-shared pointer (see `types::translate_type`). +/// Projecting that field is therefore the identity. Without this, the place +/// walkers would read the pointer value of a `SharedPtr` nested in another +/// place, such as `(region.base).ptr`, as the address of an aggregate. +pub(crate) fn strip_shared_ptr_field_projections<'a>( + body: &mir::Body, + place: &'a mir::Place, +) -> std::borrow::Cow<'a, mir::Place> { + use rustc_public::ty::{RigidTy, TyKind}; + + if !place + .projection + .iter() + .any(|elem| matches!(elem, mir::ProjectionElem::Field(0, _))) + { + return std::borrow::Cow::Borrowed(place); + } + let mut place_ty = body.locals()[place.local].ty; + let mut kept = Vec::with_capacity(place.projection.len()); + for elem in &place.projection { + let is_shared_ptr = matches!( + place_ty.kind(), + TyKind::RigidTy(RigidTy::Adt(adt_def, _)) if facts::is_cuda_device_adt(&adt_def, "SharedPtr") + ); + let Ok(next_ty) = elem.ty(place_ty) else { + return std::borrow::Cow::Borrowed(place); + }; + if !(is_shared_ptr && matches!(elem, mir::ProjectionElem::Field(0, _))) { + kept.push(elem.clone()); + } + place_ty = next_ty; + } + if kept.len() == place.projection.len() { + std::borrow::Cow::Borrowed(place) + } else { + std::borrow::Cow::Owned(mir::Place { + local: place.local, + projection: kept, + }) + } +} + // ============================================================================ // Place Read Strategy and Address Path // ============================================================================ diff --git a/crates/mir-importer/src/translator/types.rs b/crates/mir-importer/src/translator/types.rs index ee5c23b8e..f860dbe82 100644 --- a/crates/mir-importer/src/translator/types.rs +++ b/crates/mir-importer/src/translator/types.rs @@ -364,6 +364,14 @@ fn translate_pointer_like( let elem = shared_array_element_type(ctx, &substs, "SharedArray")?; Ok(facts::mint_shared_ptr_type(ctx, elem, origin).into()) } + rustc_public::ty::TyKind::RigidTy(rustc_public::ty::RigidTy::Adt(adt_def, substs)) + if is_cuda_device_adt(&adt_def, "SharedSlot") => + { + // `*mut SharedSlot` is the field of `SharedPtr`: a pointer + // to a `T` in CTA-shared memory, by its type. + let elem = shared_array_element_type(ctx, &substs, "SharedSlot")?; + Ok(facts::mint_shared_ptr_type(ctx, elem, origin).into()) + } rustc_public::ty::TyKind::RigidTy(rustc_public::ty::RigidTy::Adt(adt_def, _substs)) if is_cuda_device_adt(&adt_def, "Barrier") => { @@ -718,6 +726,14 @@ pub fn translate_type( // ThreadIndex is a newtype around usize - translate to usize // The type safety is enforced at the Rust level, not the IR level Ok(get_usize_type(ctx).into()) + } else if is_cuda_device_adt(&adt_def, "SharedPtr") { + // `SharedPtr { ptr: *mut SharedSlot }` is a scalar-lowered + // newtype: the value is its field, a CTA-shared (`addrspace(3)`) + // pointer to `T`. Params, returns, locals and fields of type + // `SharedPtr` are therefore shared pointers by type, across + // function boundaries. + let elem = shared_array_element_type(ctx, &substs, "SharedPtr")?; + Ok(facts::mint_shared_ptr_type(ctx, elem, facts::abi_shared_ptr_field()).into()) } else if is_cuda_device_adt(&adt_def, "RangeToken") { Ok(dialect_iket::types::IketRangeTokenType::get(ctx).into()) } else if is_cuda_device_adt(&adt_def, "SharedArray") { diff --git a/crates/rustc-codegen-cuda/examples/shared_memory_helper/README.md b/crates/rustc-codegen-cuda/examples/shared_memory_helper/README.md index 8947f5a80..b764d8a0f 100644 --- a/crates/rustc-codegen-cuda/examples/shared_memory_helper/README.md +++ b/crates/rustc-codegen-cuda/examples/shared_memory_helper/README.md @@ -13,6 +13,26 @@ unsafe fn helper(smem: *mut u8, offset: usize) -> u64 { } ``` +Every build also runs the same probe through `SharedPtr`, a CTA-shared +pointer typed as such: + +```rust +let smem = DynamicSharedArray::::shared_ptr(); +shared_ptr_probe(Region { base: smem, offset }); + +#[inline(never)] +unsafe fn shared_ptr_probe(region: Region) -> u64 { + let full = region.base.cast::().byte_add(1024).cast::().as_ptr(); + // ... +} +``` + +`SharedPtr` lowers to an `addrspace(3)` pointer by its type, so the +generic helper's parameter, nested in a struct, is one too, and the call +needs no address-space adaptation. `verify-code-shape.sh` checks that the +helper takes `ptr addrspace(3)` and that no PTX converts a generic address to a +shared one. + The launch supplies 2048 dynamic shared bytes. Four independent CTAs each use one thread. Nine launches cover three data offsets and three input values, including wrapping arithmetic. Every CTA checks both a volatile shared-memory @@ -79,6 +99,9 @@ RTX 5090, NVIDIA driver 580.173.02, and CUDA tools 13.3. | Local helper, all three inline policies | MIR verification rejects the call | | Dependency helper, all three inline policies | Same rejection | +With no features, the `SharedPtr` probe runs after the direct control and +reports `PASS shared-ptr: 9 launches, 36 CTA data/barrier checks`. + Both Compute Sanitizer runs of the direct control reported: ```text diff --git a/crates/rustc-codegen-cuda/examples/shared_memory_helper/src/main.rs b/crates/rustc-codegen-cuda/examples/shared_memory_helper/src/main.rs index 3abc8dc97..44c1a0fd3 100644 --- a/crates/rustc-codegen-cuda/examples/shared_memory_helper/src/main.rs +++ b/crates/rustc-codegen-cuda/examples/shared_memory_helper/src/main.rs @@ -9,7 +9,8 @@ use cuda_core::simt::LaunchConfig; use cuda_core::{CudaContext, DeviceBuffer}; -use cuda_device::{DynamicSharedArray, cuda_module, kernel, launch_bounds, thread}; +use cuda_device::{DynamicSharedArray, SharedPtr, barrier::Barrier}; +use cuda_device::{cuda_module, kernel, launch_bounds, thread}; #[cfg(all(feature = "local-helper", feature = "cross-crate-helper"))] compile_error!("Select at most one helper feature; no features selects the direct control"); @@ -27,6 +28,45 @@ unsafe fn local_probe(smem: *mut u8, offset: usize) -> u64 { shared_memory_helper_lib::shared_probe_body!(smem, offset) } +/// The shared memory a typed probe works on, as a struct field. +#[derive(Clone, Copy)] +struct Region { + base: SharedPtr, + offset: usize, +} + +/// The probe body over a typed shared pointer, in a generic helper that is not +/// inlined. Unlike the raw `*mut u8` helpers above, it compiles: `SharedPtr` +/// is a CTA-shared pointer by its type, so the helper's parameter is one too, +/// and the barrier is addressed without converting a generic pointer. +/// +/// # Safety +/// As for `shared_probe_body!`: 2048 live shared bytes aligned to 1024, one +/// thread per CTA, and an initialized u64 at `offset`, below byte 1024. +#[inline(never)] +unsafe fn shared_ptr_probe(region: Region) -> u64 { + unsafe { + let bytes = region.base.cast::(); + let full: *mut Barrier = bytes.byte_add(1024).cast::().as_ptr(); + let data = bytes.byte_add(region.offset).cast::().as_ptr(); + cuda_device::barrier::mbarrier_init(full, 1); + let pending = !cuda_device::barrier::mbarrier_try_wait_parity(full, 0); + let _token = cuda_device::barrier::mbarrier_arrive(full); + let mut complete = 0u64; + let mut attempts = 0u32; + while attempts < 1024 { + if cuda_device::barrier::mbarrier_try_wait_parity(full, 0) { + complete = if pending { 1 } else { 2 }; + break; + } + attempts += 1; + } + let value = core::ptr::read_volatile(data); + core::ptr::write_volatile(data, value.wrapping_mul(3) ^ 0x1234_5678); + complete + } +} + #[cuda_module] mod kernels { use super::*; @@ -52,6 +92,23 @@ mod kernels { *output.add(block * 2 + 1) = completed; } } + + /// The same probe through `SharedPtr`, a typed CTA-shared pointer, passed + /// to a non-inlined generic helper inside a struct. + #[kernel] + #[launch_bounds(1)] + pub fn shared_ptr_memory_probe(output: *mut u64, seed: u64, byte_offset: u64) { + let smem = DynamicSharedArray::::shared_ptr(); + let block = thread::blockIdx_x() as usize; + let offset = byte_offset as usize; + unsafe { + let data = smem.byte_add(offset).cast::().as_ptr(); + core::ptr::write_volatile(data, seed.wrapping_add(block as u64)); + let completed = shared_ptr_probe(Region { base: smem, offset }); + *output.add(block * 2) = core::ptr::read_volatile(data); + *output.add(block * 2 + 1) = completed; + } + } } fn main() -> Result<(), Box> { @@ -79,41 +136,44 @@ fn main() -> Result<(), Box> { block_dim: (1, 1, 1), shared_mem_bytes: 2048, }; - let mut cases = 0; - for seed in [0u64, 17, u64::MAX] { - for offset in [0u64, 64, 248] { - // SAFETY: four independent CTAs, one thread each; output holds two - // u64s per CTA. Data offsets are aligned and disjoint from the - // barrier. Every launch receives a fresh shared-memory allocation. - unsafe { - module.shared_memory_probe( - &stream, - config, - output.cu_deviceptr() as *mut u64, - seed, - offset, - )?; - } - let got = output.to_host_vec(&stream)?; - for block in 0..BLOCKS { - let expected = seed.wrapping_add(block as u64).wrapping_mul(3) ^ 0x1234_5678; - assert_eq!( - got[block * 2], - expected, - "{variant}: data, seed={seed}, offset={offset}, block={block}" - ); - assert_eq!( - got[block * 2 + 1], - 1, - "{variant}: barrier, seed={seed}, offset={offset}, block={block}" - ); + for typed in [false, true] { + let variant = if typed { "shared-ptr" } else { variant }; + let mut cases = 0; + for seed in [0u64, 17, u64::MAX] { + for offset in [0u64, 64, 248] { + let out = output.cu_deviceptr() as *mut u64; + // SAFETY: four independent CTAs, one thread each; output holds + // two u64s per CTA. Data offsets are aligned and disjoint from + // the barrier. Every launch receives a fresh shared-memory + // allocation. + unsafe { + if typed { + module.shared_ptr_memory_probe(&stream, config, out, seed, offset)?; + } else { + module.shared_memory_probe(&stream, config, out, seed, offset)?; + } + } + let got = output.to_host_vec(&stream)?; + for block in 0..BLOCKS { + let expected = seed.wrapping_add(block as u64).wrapping_mul(3) ^ 0x1234_5678; + assert_eq!( + got[block * 2], + expected, + "{variant}: data, seed={seed}, offset={offset}, block={block}" + ); + assert_eq!( + got[block * 2 + 1], + 1, + "{variant}: barrier, seed={seed}, offset={offset}, block={block}" + ); + } + cases += 1; } - cases += 1; } + println!( + "PASS {variant}: {cases} launches, {} CTA data/barrier checks", + cases * BLOCKS + ); } - println!( - "PASS {variant}: {cases} launches, {} CTA data/barrier checks", - cases * BLOCKS - ); Ok(()) } diff --git a/crates/rustc-codegen-cuda/examples/shared_memory_helper/verify-code-shape.sh b/crates/rustc-codegen-cuda/examples/shared_memory_helper/verify-code-shape.sh new file mode 100755 index 000000000..8ffe705a1 --- /dev/null +++ b/crates/rustc-codegen-cuda/examples/shared_memory_helper/verify-code-shape.sh @@ -0,0 +1,30 @@ +#!/usr/bin/env bash +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 +# +# `SharedPtr` keeps a CTA-shared pointer typed `addrspace(3)` across a call: +# the non-inlined `shared_ptr_probe` helper takes one as its parameter, and +# no PTX converts a generic address to a shared one (`cvta.to.shared`). +set -euo pipefail + +dir="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +ll="${dir}/shared_memory_helper.ll" +ptx="${dir}/shared_memory_helper.ptx" + +for file in "${ll}" "${ptx}"; do + [[ -f "${file}" ]] || { echo "FAIL: ${file} not found; build the example first" >&2; exit 1; } +done + +if ! grep -Eq '^define .*shared_ptr_probe.*\(ptr addrspace\(3\) ' "${ll}"; then + echo "FAIL: shared_ptr_probe does not take a ptr addrspace(3) parameter" >&2 + grep -E '^define .*shared_ptr_probe' "${ll}" >&2 || true + exit 1 +fi + +if grep -q 'cvta\.to\.shared' "${ptx}"; then + echo "FAIL: PTX converts a generic address to a shared one" >&2 + grep -n 'cvta\.to\.shared' "${ptx}" >&2 + exit 1 +fi + +echo "PASS: SharedPtr stays addrspace(3) across shared_ptr_probe, no cvta.to.shared" diff --git a/cuda-oxide-book/advanced/shared-memory-and-synchronization.md b/cuda-oxide-book/advanced/shared-memory-and-synchronization.md index c068c8af4..8ffa7d009 100644 --- a/cuda-oxide-book/advanced/shared-memory-and-synchronization.md +++ b/cuda-oxide-book/advanced/shared-memory-and-synchronization.md @@ -233,6 +233,48 @@ total does not exceed `shared_mem_bytes` — there are no runtime guards. (Hopper+), use `DynamicSharedArray` to get the required 128-byte alignment. The alignment is encoded in the PTX `.align` directive. +### Typed shared pointers: `SharedPtr` + +`DynamicSharedArray::get()` and `SharedArray::as_mut_ptr()` return `*mut T`, +which in device code is a *generic* pointer: it may address global, local or +shared memory, including another block's shared memory in a cluster. Once +such a pointer is passed to a helper function, the compiler no longer knows it +points into this block's shared memory. + +`SharedPtr` is a raw pointer into the issuing block's shared memory +(`.shared::cta`) that carries that fact in its type. It lowers to an +`addrspace(3)` pointer wherever it goes, including helper parameters and +return values, struct fields and generic code, so instructions that take it +use their shared-memory forms without converting a generic address. + +```rust +use cuda_device::{DynamicSharedArray, SharedArray, SharedPtr}; + +#[inline(never)] +unsafe fn fill(tile: SharedPtr, i: usize, value: f32) { + // Still a shared-memory pointer here, although `fill` is not inlined. + unsafe { tile.add(i).as_ptr().write(value) }; +} + +static mut TILE: SharedArray = SharedArray::UNINIT; +let tile: SharedPtr = SharedArray::shared_ptr(&raw mut TILE); +let pool: SharedPtr = DynamicSharedArray::::shared_ptr_at(1024); +unsafe { fill(tile, thread::threadIdx_x() as usize, 1.0) }; +``` + +| Operation | Result | +|:---------------------------------------|:----------------------------------------------------| +| `SharedArray::shared_ptr(&raw mut S)` | `SharedPtr` to the first element of `S` | +| `DynamicSharedArray::::shared_ptr()` / `shared_ptr_at(bytes)` | Like `get()` / `offset(bytes)` | +| `unsafe SharedPtr::from_raw(p)` | Asserts that generic `p` points into this block's shared memory | +| `add`, `offset`, `byte_add`, `cast` | Pointer arithmetic that keeps the type | +| `as_ptr()` | The generic `*mut T`, for loads, stores and raw-pointer APIs | +| `addr_u32()` | The 32-bit `.shared::cta` address, for `_shared_u32` forms | + +`SharedPtr` cannot be a kernel parameter: the host has no shared-memory +address to pass. Pointers from `cluster::map_shared_rank` may name another +block's shared memory and are not `SharedPtr`s. + --- ## Synchronization: sync_threads() diff --git a/cuda-oxide-book/appendix/api-quick-reference.md b/cuda-oxide-book/appendix/api-quick-reference.md index 6a986c919..64ab9ee8f 100644 --- a/cuda-oxide-book/appendix/api-quick-reference.md +++ b/cuda-oxide-book/appendix/api-quick-reference.md @@ -310,6 +310,7 @@ pub fn dynamic(data: &[f32]) { | `SharedArray` | Compile-time sized, block-scoped shared memory | | `SharedArray` | With 128-byte alignment (required for TMA destinations) | | `DynamicSharedArray` | Runtime-sized shared memory (set via `LaunchConfig`) | +| `SharedPtr` | Pointer into this block's shared memory, `addrspace(3)` by type | Both are `!Sync` — concurrent access requires explicit barriers. @@ -647,7 +648,7 @@ debug::prof_trigger::<7>(); // Nsight profiler trigger | `thread` | Thread/block IDs, `index_1d`, `sync_threads` | All | | `config` | Compile-time policies, shapes, tiles, atoms, layouts, memory spaces, and scopes | All | | `disjoint` | `DisjointSlice` — typed writes completed by a launch proof | All | -| `shared` | `SharedArray`, `DynamicSharedArray` | All | +| `shared` | `SharedArray`, `DynamicSharedArray`, `SharedPtr` | All | | `warp` | Shuffle, vote, match, lane/warp ID | All | | `atomic` | Scoped atomics (device/block/system) | sm_70+ | | `debug` | `clock64`, `trap`, `breakpoint`, `gpu_printf!` | All |