Hacktoberfest 2026 : les issues que les mainteneurs ont marquées pour octobre, ouvertes et accessibles aux débutants. Parcourir les issues Hacktoberfest

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

Ouverte Adaptée aux débutants
  • Quelque chose ne va pas ?
#416 3 commentaires 0 réactions 0 personnes assignées Voir sur GitHub

@shubhamkapoor01 y travaille déjà.

Depuis le 25/10/2026.

  • #417 par @shubhamkapoor01 — ouverte

Évaluation

Difficulté
2/5
Temps estimé
1-3 heures
Accessibilité débutants
72/100
Type d'issue
Bug
Clarté
Clairement spécifiée
Activité
Active
Stack technique
rust
Domaine
compilers

Piste de recherche

Le bogue se situe dans la déclaration extern DYN_SHARED de crates/cuda_std/src/shared.rs, dont le type dimensionné u128 fait que le PTX émet [16] au lieu de []. Lis cette déclaration et les extraits PTX du rapport pour voir la différence, puis décide comment rendre le symbole sans taille. C'est terminé lorsque le PTX émis déclare _Zcuda_std_dyn_shared sans taille et que la reproduction signale la région dynamique après le tableau statique (décalage 0x800). L'exécution nécessite un GPU NVIDIA et la chaîne d'outils CUDA.

Rédigé par le modèle d'indexation à partir du texte de l'issue.

Description

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.

Langage dominant
Rust
Étoiles
5.4k
Forks
251
Merge moyen
30 min
PR mergées (30 j)
1

Préparer son environnement

Ouvrir dans Codespaces

Lance le conteneur de développement du projet dans votre navigateur, avec votre propre compte GitHub.

Par où commencer

  1. Lisez l'issue en entier, puis le guide de contribution du projet.
  2. Signalez en commentaire que vous la prenez — cela évite que deux personnes fassent le même travail.
  3. Forkez le dépôt et travaillez sur une branche.
  4. Ouvrez une pull request qui référence le numéro de l'issue.

Autres issues de Rust-GPU/rust-cuda

Toutes les issues de Rust-GPU/rust-cuda

Issues similaires

Plus d'issues Rust

Recevez les nouvelles issues par e-mail

Un résumé court des issues GitHub adaptées aux débutants.