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

Lỗi nằm ở khai báo extern DYN_SHARED trong crates/cuda_std/src/shared.rs, nơi kiểu có kích thước u128 khiến PTX xuất ra [16] thay vì []. Hãy đọc khai báo đó và các đoạn PTX trong báo cáo để thấy sự khác biệt, sau đó quyết định cách làm cho symbol không có kích thước. Hoàn thành khi PTX được xuất ra khai báo _Zcuda_std_dyn_shared không có kích thước và bản tái hiện báo cáo vùng động nằm sau mảng tĩnh (offset 0x800). Để chạy cần có GPU NVIDIA và bộ công cụ 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.