feat(cuda-device): add SharedPtr, a CTA-shared pointer typed addrspace(3)

[lucifer1004]

Part of #1337; groundwork for #1238.

What this adds

SharedPtr<T>: a raw pointer into the issuing CTA's shared memory (.shared::cta) that carries its address space in its type.

A *mut T in device code is a generic pointer. Once it is passed to a helper, nothing tells the compiler that it points into this CTA's shared memory; helpers are compiled as separate functions whatever their inline attribute. SharedPtr<T> lowers to an addrspace(3) pointer wherever it goes: parameters and returns of non-inlined helpers, struct fields, generic functions and closures. Instructions that take it can use their .shared::cta forms without converting a generic address, which is what #1238 needs for TMA on SM12x (see #1337).

static mut TILE: SharedArray<f32, 256> = SharedArray::UNINIT;

#[inline(never)]
unsafe fn fill(tile: SharedPtr<f32>, i: usize, value: f32) {
    unsafe { tile.add(i).as_ptr().write(value) };   // `tile` is addrspace(3) here
}

let tile = SharedArray::shared_ptr(&raw mut TILE);
unsafe { fill(tile, thread::threadIdx_x() as usize, 1.0) };

API
SharedArray::shared_ptr(&raw mut S)	the array's shared address, with no conversion
DynamicSharedArray::<T>::shared_ptr() / shared_ptr_at(bytes)	as get() / offset(bytes)
unsafe SharedPtr::from_raw(p)	the one explicit locality assertion, like CUDA C++ __cvta_generic_to_shared or CuTe make_smem_ptr
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 the _shared_u32 forms (#1236)

Only a mutable form is provided: shared memory is written through the same pointers it is read through, and as_ptr() gives a *mut T that coerces to *const T where an API wants one.

How it works

• SharedPtr<T> is #[repr(transparent)] struct SharedPtr<T> { ptr: *mut SharedSlot<T> }, with a #[doc(hidden)] #[repr(transparent)] SharedSlot<T>(T) pointee.
• The importer types a pointer to SharedSlot<T> as an addrspace(3) pointer to T, as it already does for pointers to SharedArray and Barrier. It lowers SharedPtr<T> itself to that same pointer, as a scalar-lowered newtype like ThreadIndex. Construction therefore reuses the existing newtype aggregate path, and reading the field is the existing scalar extract_field.
• One importer change beyond the type mapping: a Field(0) projection on a SharedPtr place is the identity, so the place walkers drop it. Without that, a SharedPtr nested in another place, such as region.base.ptr after rustc inlines cast, was read as the address of an aggregate.
• The methods are ordinary #[inline(always)] Rust: casts between *mut SharedSlot<_> stay in addrspace(3), as_ptr is the 3 → 0 addrspacecast that &raw mut S as *mut T already emits, and from_raw is the 0 → 3 cast.
• SharedPtr cannot be a kernel parameter. The exporter already rejects addrspace(3) entry parameters, and it does so for SharedPtr too: kernel ... parameter 0 is a pointer into shared memory (address space 3) ....
• Pointers from map_shared_rank, which may name another CTA's shared memory, are not SharedPtrs.
• Nothing existing changes: SharedArray::as_ptr / as_mut_ptr and DynamicSharedArray::get still return generic pointers.

Tests

• shared_memory_helper runs its probe a second time through SharedPtr, with no features required:

  • the probe passes a non-inlined generic helper a Region<T> { base: SharedPtr<T>, offset };
  • it addresses the barrier and data slots with byte_add / cast;
  • it reports PASS shared-ptr: 9 launches, 36 CTA data/barrier checks.

  The raw *mut u8 helpers of the same example still fail MIR verification, as the README documents.

• The new verify-code-shape.sh checks that shared_ptr_probe takes ptr addrspace(3) in the emitted LLVM IR, and that no PTX converts a generic address to a shared one (cvta.to.shared).

• In a scratch kernel, not committed, the remaining API also computed correct values on the GPU:

  • SharedArray::shared_ptr, from_raw and a negative offset;
  • a SharedPtr passed to a closure;
  • addr_u32 differences.

• A SharedPtr kernel parameter fails to compile with the exporter diagnostic above.

Validation

All runs used an RTX PRO 6000 (sm_120) with the CUDA 13.3 toolkit.

• just fmt-check, clippy, test and test-cuda pass (3963 and 1505 tests).

• just doc-check passes.

• just check-guards passes every script except two, whose results could not be obtained here:

  • cargo deny --locked check stops at the advisory-database lock (NFS do not support locking). cargo deny check bans licenses sources passes for all three workspaces.
  • check-example-license-policy.sh hits the same lock.

  This PR adds no dependency and changes no lockfile.

• Full scripts/smoketest.sh: 233/234 pass, the same as main. The one failure, sparse_mma_fp8, also fails on main.

• shared_memory_helper passes on both backends: LLVM NVPTX, and libNVVM with --emit-nvvm-ir --arch=sm_120. Each backend runs both the direct and the SharedPtr probe. On libNVVM, the helper is still define i64 @…shared_ptr_probe…(ptr addrspace(3) %v0, i64 %v1).

Follow-ups tracked in #1337: #1336 moves cp_async_bulk_tensor_*_g2s_cta onto SharedPtr<u8> on top of this PR; SharedPtr forms for cp.async, ldmatrix/stmatrix and the tcgen05/wgmma descriptors; a cluster-shared counterpart from map_shared_rank.

🤖 Generated with Claude Code

[nihalpasham]

Landed through #1336 at 8bb24a4: your commit is the first one on that branch, unchanged and with your sign-off, plus the review follow-ups (SharedArray::shared_ptr made unsafe like as_raw_mut_ptr, the aggregate kernel-parameter refusal with its fixture, the shared_ptr_api regressions, the limits in the book). Closing this one as already fixed. Thanks for the type, it is what the .shared::cta form needed.