Hacktoberfest 2026: los issues que los mantenedores marcaron para octubre, abiertos y aptos para principiantes. Explorar issues de Hacktoberfest

A `#[address_space(shared)]` static is stored with a generic `st.u32` instead of `st.shared`, corrupting memory — two identically-declared siblings lower differently

Abierto
  • ¿Algo no está bien?
#402 0 comentarios 0 reacciones 0 asignados Ver en GitHub

Nadie ha tomado este issue todavía.

Evaluación

Dificultad
5/5
Tiempo estimado
Más de una semana
Aptitud para principiantes
30/100
Tipo de issue
Error
Claridad
Necesita aclaración
Estado de actividad
Tranquilo
Stack tecnológico
rust
Área
backend, compilers

Línea de trabajo

Empieza examinando los fragmentos de PTX emitidos para SH_RECHAIN_N y SH_CHAIN_NA y, a continuación, compara el reproducer candidato en 02-shared-generic/ con toda la evidencia del kernel. Confirma si se pueden reproducir los dos stores genéricos st.u32 e identifica la ruta de rustc_codegen_nvvm responsable. La tarea estará completa cuando los stores a direcciones shared se reduzcan a st.shared y el reproducer o el kernel completo supere memcheck sin escrituras no válidas.

Escrito por el modelo de indexación a partir del texto del issue.

Descripción

In device code compiled by rustc_codegen_nvvm, one #[address_space(shared)] static is written through a generic-space st.u32 on its raw shared offset, while an identically-declared sibling in the same function is correctly written with st.shared.u32. Dereferencing a shared offset in the generic window is an out-of-bounds access: compute-sanitizer memcheck reports Invalid __global__ write of size 4 bytes ... Access to 0x26c is out of bounds. It compiles with no error or warning.

This is not the cvta.shared issue one might expect — the module contains zero cvta.shared and needs none, because correct shared accesses use ld.shared/st.shared on the raw symbol offset. The defect is purely that these stores got the wrong state space.

The contrast that isolates it

Two shared statics, declared identically:

#[address_space(shared)]
static mut SH_RECHAIN_N: [MaybeUninit<u32>; WARPS_PER_BLOCK] = [MaybeUninit::uninit(); WARPS_PER_BLOCK];
#[address_space(shared)]
static mut SH_CHAIN_NA:  [MaybeUninit<u32>; WARPS_PER_BLOCK] = [MaybeUninit::uninit(); WARPS_PER_BLOCK];

Both emit the same PTX declaration (.shared .align 4 .b8 …[16]), and their per-warp slot addresses are computed by the same instruction form, from the same offset register, on adjacent lines:

21565:  add.s64  %rd42, %rd566, %rd562;   // %rd566 = SH_CHAIN_NA
21567:  add.s64  %rd44, %rd567, %rd562;   // %rd567 = SH_RECHAIN_N

Yet they lower differently:

// SH_CHAIN_NA — CORRECT
23035:  st.shared.u32  [%rd42], %r397;
23042:  ld.shared.u32  %r193, [%rd42];
24928:  st.shared.u32  [%rd42], %r418;

// SH_RECHAIN_N — WRONG (generic state space)
23206:  st.u32  [%rd44], %r206;
23470:  st.u32  [%rd44], %r194;

Within that one kernel there are 106 ld/st.shared and exactly 2 generic st.u32 — the two above. A second kernel in the same module reproduces the pattern exactly (%rd43 ← SH_RECHAIN_N, 2 generic stores).

%rd44 is defined once at 21567 and never redefined before either store; both stores are inside the same .visible .entry (which spans lines 21469–25225).

Source correspondence

The offending store is unambiguous — the Rust is

if lane == 0 {
    SH_RECHAIN_N[w.wib].write(if go { n_a as u32 } else { 0 });
}

and the PTX select maps one-to-one onto the if/else:

23204:  cvt.u32.u64  %r205, %rd286;
23205:  selp.b32     %r206, %r205, 0, %p395;   // if go { n_a } else { 0 }
23206:  st.u32       [%rd44], %r206;           // <-- generic store to shared address

Runtime effect

========= Invalid __global__ write of size 4 bytes
=========     at spine_chain_big+0x7c90
=========     by thread (96,0,0) in block (0,0,0)
=========     Access to 0x26c is out of bounds

0x26c = 620. For the six module-scope shared symbols this kernel uses, laid out in declaration order with their alignments (SH_SC[512] → 0, SH_J[32] → 512, SH_TLO[32] → 544, SH_STOP[16] → 576, SH_CHAIN_NA[16] → 592, SH_RECHAIN_N[16] → 608), warp 3's slot is 608 + 3×4 = 620. The faulting thread (96,0,0) is lane 0 of warp 3 — the lane that executes the guarded store, with a warp index matching the slot. All 226 reported errors reduce to this single invalid access plus cascading CUDA_ERROR_ILLEGAL_ADDRESS on subsequent API calls from the poisoned context.

Removing the shared broadcast (deriving the value warp-uniformly from global memory instead) takes memcheck to 0 errors on both kernels with no other change.

What distinguishes the two symbols

The only difference I can observe at the source level: SH_CHAIN_NA is written and read only by lane 0, whereas SH_RECHAIN_N is written by lane 0 and then read by every lane after a __syncwarp — it is a genuine broadcast. Both writes sit in lane-0-guarded blocks and use the same MaybeUninit::write API. I am reporting that as an observation, not a root cause; I have not located the responsible code in the backend.

Because the offending store sits inside a data-dependent branch, only warps whose data takes that path fault — so the failure looks intermittent and warp-dependent even though the miscompilation is static. That misdirection cost me a while, and is worth knowing when triaging similar reports.

Environment

Rust-CUDA 7fa76f3d717038a92c90bf4a482b0b8dd3259344, cuda_builder 0.3.0, nightly-2025-08-04, CUDA 13.2 libNVVM, driver 595.71.05, Tesla T4 (sm_75), -arch=compute_75. Kernel crate no_std, -Z saturating_float_casts=false, --override-libm on (default).

Reproducer status — please read

I do not have a minimal reproducer. The evidence above is from a real kernel; a synthetic kernel mirroring the pattern (lane-0 store + post-__syncwarp broadcast read, with a control array alongside) did not reproduce it, which suggests the trigger depends on surrounding context — function size, inlining, or register pressure. The PTX and memcheck excerpts are self-contained and verifiable as they stand, and I can supply the full 1 MB PTX module, the full memcheck log, or the offending source on request. A candidate (non-triggering) synthetic kernel is at https://github.com/nh13/rust-cuda-nvvm-repros (02-shared-generic/).

Related: #401 (a separate rustc_codegen_nvvm codegen defect in the bit-counting intrinsics).

Prepared with the assistance of Claude (Anthropic). Every PTX line number and address above was read directly from the emitted module and cross-checked; the layout arithmetic is derived, and the source-level distinction between the two symbols is an observation rather than a located root cause.

Lenguaje dominante
Rust
Estrellas
5.4k
Forks
251
Merge medio
4 d 15 h
PR fusionados (30 d)
4

Preparar el entorno

Abrir en Codespaces

Inicia el contenedor de desarrollo del proyecto en tu navegador, con tu propia cuenta de GitHub.

Primeros pasos

  1. Lee el issue completo y luego la guía de contribución del proyecto.
  2. Comenta en el issue que vas a ocuparte — evita que dos personas hagan lo mismo.
  3. Haz un fork del repositorio y trabaja en una rama.
  4. Abre un pull request que haga referencia al número del issue.

Más de Rust-GPU/rust-cuda

Todos los issues de Rust-GPU/rust-cuda

Issues similares

Más issues de Rust

Recibe los nuevos issues en tu correo

Un resumen breve de issues de GitHub para principiantes.