[YunusMutlu]

What this adds

Static and dynamic shared memory now carry a checked, typed alignment contract. The MIR verifier rejects an ALIGN that is not a non-zero power of two, with the source position, and lowering declares every static __shared_mem_N global and the per-function __dynamic_smem pool at the larger of the requested alignment and the element type's ABI alignment. Before, the declared alignment came only from the request, or from an exporter fallback that printed 8 for everything it did not know; ptxas then packed under-declared arrays behind smaller statics and the typed stores faulted.

static mut V: SharedArray<u64, 2, 1> = SharedArray::UNINIT;   // touched before a u8 pad
static mut S: SharedArray<V4, 4> = SharedArray::UNINIT;       // #[repr(C, align(16))], natural
let p: *mut u64 = DynamicSharedArray::<u64, 1>::get();
static mut BAD: SharedArray<u64, 4, 3> = SharedArray::UNINIT;

before: .shared .align 1 .b8 V[16]    st.shared.b64     -> CUDA error 716, misaligned address
        .shared .align 8 .b8 S[64]    st.shared.v4.b32  -> CUDA error 716, misaligned address
        .extern .shared .align 1 .b8 __dynamic_smem_...[]
        BAD: llc: error: alignment is not a power of two
after:  .shared .align 8 .b8 V[16]
        .shared .align 16 .b8 S[64]
        .extern .shared .align 8 .b8 __dynamic_smem_...[]
        BAD: src/main.rs:20:13: MirSharedAllocOp alignment must be a non-zero power of two, found 3

#[repr(C, align(16))]
#[derive(Copy, Clone)]
struct V4([f32; 4]);

#[kernel]
pub fn fill(mut out: DisjointSlice<f32>) {
    static mut VS: SharedArray<V4, 4> = SharedArray::UNINIT;   // declared .align 16 now
    static mut PAD: SharedArray<u64, 1> = SharedArray::UNINIT;
    // write VS and PAD, sync, read back through `out`
}

What we fixed

• MirSharedAllocOp and MirExternSharedOp verify that an explicit alignment is a non-zero power of two; the error names the use site instead of surfacing as an llc failure.
• Static storage alignment is max(ALIGN, abi_align(T)), computed in mir-lower by required_storage_align from the Rust ABI raise (mir_type_abi_align) and the LLVM natural layout (llvm_type_size_align), so repr(align(N)) aggregates and LLVM leaf elements such as TmaDescriptor are both covered.
• The dynamic pool alignment is the maximum over the function's typed DynamicSharedArray uses, then combined with the launch-contract alignment as before.
• SharedArray<bool, N> compiles: the exporter fallback used to emit align 0 for i1 arrays and llc rejected it.

Alignment contract

• SharedArray<T, N, ALIGN>: ALIGN is 0 (natural) or a non-zero power of two. Storage is never declared below align_of::<T>(); a larger ALIGN is honored.
• DynamicSharedArray<T, ALIGN>: ALIGN is a non-zero power of two. The kernel's pool is declared at the strongest typed requirement across its get, get_raw and offset uses and the launch contract. The alignment of an offset(bytes) result remains the caller's responsibility.
• Natural-alignment aggregates, nested arrays and f16 now declare their natural .align (4 for [f32; N], 2 for f16) where the exporter used to print 8.

Implementation

• dialect-mir ops/memory.rs: verifier arms on both ops, tests asserting the diagnostic text; the attribute table names extern_byte_offset and extern_alignment.
• mir-lower convert/types/layout.rs: required_storage_align(ctx, mir_ty, llvm_ty) is max of the two layout helpers and None only when both are unknown. convert_shared_alloc_dc and the dynamic pre-scan in lowering.rs use it and report an unknown element with the op's location. The power-of-two re-checks in lowering are gone; the verifiers own the rule.
• The dialect keeps no public alignment oracle: the helper the PR had moved into dialect-mir is back as the private cast.rs helper it was, because dialect-mir cannot see LLVM leaf types.
• Tests: verifier reject and accept for both ops; lowering i64 1 to 8 (static and dynamic), an align-32 tuple 1 to 32, the maximum over two pointees, and an llvm.array<16 x i64> element and pointee to 8.
• Example sharedmem: two kernels that fault on main plus a PTX declaration check.

Verification

Checked at bbea3a2c6 (nightly-2026-10-11, RTX 5090, sm_120a):

• ✅ cargo test -p dialect-mir (70) and cargo test -p mir-lower (327 + 147) pass; clippy and rustdoc with warnings as errors on dialect-mir, mir-lower and cuda-device, cargo oxide fmt --check, the 18 guards and the lock check are clean.
• ✅ Fails before, passes after: the same sharedmem source built with main's backend fails Test 3 with DriverError(716, "misaligned address") and declares .align 1 and .align 8; this branch passes all five tests with .align 8 and .align 16, memcheck and synccheck clean.
• ✅ SharedArray<TmaDescriptor, 1> and DynamicSharedArray::<TmaDescriptor, 128> compile at align 8, as on main; they did not compile on the submitted head.
• ✅ ALIGN = 3 (static and dynamic) and DynamicSharedArray::<T, 0> fail in the verifier with the source position; on main they reach llc as align 3 and align 0.
• ✅ 19 shared-memory examples run or compile with unchanged shared declarations, except shared_slice_unsize, whose [f32; 128] tile moves from .align 8 to the correct .align 4 and still passes.

Caveats

• An ALIGN below align_of::<T>() is raised silently, as nvcc's __align__ and Rust's repr(align) do; it is not an error.
• TmaDescriptor is repr(C, align(64)) but the importer translates it to a plain [16 x u64], so its shared storage is declared align 8, as on main. That translation gap is separate from this change.
• Natural-alignment aggregates lose the accidental align 8, and LLVM then stops forming ld/st.shared.v2 on them. No in-tree example is affected.

Closes #1344.

[copy-pr-bot]

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

[nihalpasham]

/ok to test bbea3a2

[nihalpasham]

Thanks, this is the right route: verifier arms for the request and a typed raise where the storage is declared. Four things on top:

• Computed the required alignment in mir-lower from its own two layout helpers, the Rust ABI raise plus the LLVM natural layout, instead of a helper in dialect-mir. The dialect cannot see LLVM leaf types, so SharedArray<TmaDescriptor, N> had stopped compiling on this branch; it compiles again at align 8, with llvm.array lowering tests.
• Removed the power-of-two re-checks from lowering; the two verifier arms own the rule and run before conversion.
• Added a GPU regression to sharedmem: a SharedArray<u64, 2, 1> and a natural-alignment repr(C, align(16)) element, each touched before a smaller static. On main both fault with misaligned address; here they pass and the PTX declares .align 8 and .align 16.
• Documented the contract on SharedArray, DynamicSharedArray and the book page, with a one-page figure.

At bbea3a2c6, both crate suites, the 18 guards, the shared-memory examples and compute-sanitizer on the new kernels passed on the 5090, and all CI checks passed. This addresses #1344. Aggregate, nested-array and f16 shared arrays now declare their natural alignment where the exporter used to print a blanket 8.