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

まず、SH_RECHAIN_N と SH_CHAIN_NA 用に出力された PTX の抜粋を調べ、次に 02-shared-generic/ の候補 reproducer をカーネル全体の証拠と比較します。2 つの汎用 st.u32 store を再現できるか確認し、その原因となっている rustc_codegen_nvvm のパスを特定します。shared アドレスへの store が st.shared に lowering され、reproducer またはカーネル全体が無効な書き込みなしで memcheck を通過すれば完了です。

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.