`dynamic_shared_mem` overlaps a kernel's static shared memory (the extern symbol is declared with a size)

该错误位于 crates/cuda_std/src/shared.rs 中的 extern 声明 DYN_SHARED，其带大小的类型 u128 导致 PTX 输出 [16] 而不是 []。请阅读该声明和报告中的 PTX 片段以查看差异，然后决定如何使该符号不带大小。当输出的 PTX 中 _Zcuda_std_dyn_shared 声明没有大小，并且复现结果报告动态区域位于静态数组之后（偏移 0x800）时，即完成。运行需要 NVIDIA GPU 和 CUDA 工具链。

Summary

cuda_std::shared::dynamic_shared_mem returns a pointer into memory that also holds the kernel's static
#[address_space(shared)] arrays. A kernel that uses both kinds of shared memory corrupts its static arrays when it
writes the dynamic region. The extern symbol behind dynamic_shared_mem is declared as a sized u128, so the PTX
declares it with a size ([16]) where nvcc declares dynamic shared memory unsized ([]). PTX treats a sized
.extern .shared symbol as an ordinary 16-byte shared variable and lays it out among the kernel's static shared
variables, so the "dynamic" region starts inside the static data instead of after it.

Declaring the symbol as a zero-length array (UnsafeCell<[u128; 0]>) makes the PTX declaration unsized and fixes it
(verified below).

Environment

• Rust-CUDA 3a82c35a7196731fb60eebdac2a40d2b1bf0ba95, the head of main at the time of the run;
  crates/cuda_std/src/shared.rs there is unchanged in this respect.
• Rust nightly-2026-04-02 (the repository's rust-toolchain.toml).
• rustc_codegen_nvvm with the llvm21 feature: LLVM 21.1.8, from the project's prebuilt linux-x86_64 archive.
• CUDA toolkit 13.3 (nvcc V13.3.73) for libNVVM; NvvmArch::Compute120.
• NVIDIA RTX PRO 6000 Blackwell Server Edition (sm_120), driver 595.91.07 (CUDA 13.2). The PTX's .version was
  lowered from 9.3 to 9.2 to load on this driver; nothing else changed.
• Ubuntu 24.04, x86_64.

Reproduction

One block of 256 threads: fill a static shared array, write a different pattern to 1 KiB of dynamic shared memory,
then copy the static array out. Under CUDA's semantics the static array is untouched.

Kernel (cuda_std at the commit above):

#![feature(asm_experimental_arch)]
use core::mem::MaybeUninit;
use cuda_std::{address_space, kernel, thread};

#[inline(always)]
fn barrier() {
    unsafe { core::arch::asm!("barrier.sync 0;") };
}

#[kernel]
#[allow(improper_ctypes_definitions)]
pub unsafe fn repro(out: *mut u32, dyn_words: u32) {
    #[address_space(shared)]
    static mut S: [MaybeUninit<u32>; 256] = [MaybeUninit::uninit(); 256];
    let s = core::ptr::addr_of_mut!(S) as *mut u32;
    let d = cuda_std::shared::dynamic_shared_mem::<u32>();
    let t = thread::thread_idx_x() as usize;
    unsafe { *s.add(t) = 0xA000_0000 | t as u32 };
    barrier();
    let mut i = t;
    while i < dyn_words as usize {
        unsafe { *d.add(i) = 0xD000_0000 | i as u32 };
        i += 256;
    }
    barrier();
    unsafe { *out.add(t) = *s.add(t) };
}

(barrier.sync is written as inline PTX only to keep this report to one issue.)

Host (cust, PTX built by cuda_builder at NvvmArch::Compute120):

let module = Module::from_ptx(PTX, &[])?;
let stream = Stream::new(StreamFlags::NON_BLOCKING, None)?;
let out = DeviceBuffer::from_slice(&[0u32; 256])?;
unsafe { launch!(module.repro<<<1, 256, 1024, stream>>>(out.as_device_ptr(), 256u32))? };
stream.synchronize()?;
let mut s = [0u32; 256];
out.copy_to(&mut s[..])?;
let bad: Vec<usize> = (0..256).filter(|&i| s[i] != 0xA000_0000 | i as u32).collect();
println!("static array intact: {} ({} of 256 words overwritten)", bad.is_empty(), bad.len());

The reproduction also prints the shared-window offsets of both regions (cvta.to.shared). I can share the full
crate, with its lockfiles and the CUDA C++ control below, if useful.

Expected and actual

Expected (CUDA C++, the same kernel with __shared__ uint32_t S[256] and extern __shared__ uint32_t dyn[], built
by nvcc 12.8 and by nvcc 13.3):

static array at shared offset 0x400, dynamic region at 0x800
static array intact: true (0 of 256 words overwritten)

Actual (Rust-CUDA at 3a82c35):

static array at shared offset 0x410, dynamic region at 0x400
static array intact: false (252 of 256 words overwritten)
  S[0] = 0xd0000004, expected 0xa0000000
  S[1] = 0xd0000005, expected 0xa0000001
  S[2] = 0xd0000006, expected 0xa0000002
  S[3] = 0xd0000007, expected 0xa0000003

The dynamic region begins at the 16-byte symbol, and the static array is placed right after that symbol, so dynamic
word 4 lands on S[0].

PTX

Rust-CUDA (3a82c35):

.extern .shared .align 16 .b8 _Zcuda_std_dyn_shared[16];
	.shared .align 4 .b8 _RNvNvCs53lP8RTauJO_13repro_kernels5repro1S[1024];

nvcc (the CUDA C++ control):

.extern .shared .align 16 .b8 dyn[];
	.shared .align 4 .b8 _ZZ5reproPjS_jE1S[1024];

Cause

crates/cuda_std/src/shared.rs:

unsafe extern "C" {
    #[nvvm_internal::addrspace(3)]
    #[allow(improper_ctypes)]
    #[link_name = "_Zcuda_std_dyn_shared"]
    static DYN_SHARED: ::core::cell::UnsafeCell<u128>;
}

The symbol's type is a sized u128, so it is emitted as .b8 _Zcuda_std_dyn_shared[16]. Only an unsized
.extern .shared array denotes the dynamic shared memory placed after the kernel's static shared memory; a sized one
is laid out like any other 16-byte shared variable.

Suggested fix (verified)

Declare the symbol as a zero-length array, which keeps the 16-byte alignment and emits an unsized declaration:

-        static DYN_SHARED: ::core::cell::UnsafeCell<u128>;
+        static DYN_SHARED: ::core::cell::UnsafeCell<[u128; 0]>;

With only this change, the PTX reads .extern .shared .align 16 .b8 _Zcuda_std_dyn_shared[]; and the reproduction
prints:

static array at shared offset 0x400, dynamic region at 0x800
static array intact: true (0 of 256 words overwritten)

which matches nvcc. Another option would be for the codegen to emit [] for this one symbol whatever its type; the
change above needs no codegen work.

Workaround

Until it is fixed, take each kernel's shared memory from one kind only: either all from #[address_space(shared)]
statics, or all from dynamic_shared_mem, carving any scratch space out of the dynamic region. A kernel that uses only
the dynamic region is unaffected.