Skip to content

Emit LLVM noalias / readonly for Rust references with proven semantics #1088

Open
Open
Emit LLVM noalias / readonly for Rust references with proven semantics#1088
Labels
DependsBlocked on another PR/issue or an upstream dependencyIR-loweringLowering between dialects (dialect-mir to LLVM dialect)codegenDevice code-generation pipeline (Rust MIR to IR to PTX)llvm-exportTextual LLVM/NVVM IR exporter (llvm-export crate)perfPerformance of generated code or of the compiler itselfquestionFurther information is requestedsoundness-auditCorrectness case under active maintainer audit; not cleared for merge yet

Description

@uurl
uurl
opened on Aug 20, 2026

Motivation

Part of #413.

After #1062 preserves Rust pointer/reference kind through dialect MIR, the compiler can distinguish:

  • &T → SharedRef
  • &mut T → UniqueRef
  • *const T → RawConst
  • *mut T → RawMut
  • compiler-generated / erased pointers → Erased

This makes it possible to attach LLVM parameter attributes when the corresponding Rust semantics are actually proven, instead of inferring them from the lowered pointer type.

This matters for optimization. In particular, alias information can enable substantially better load/store optimization and vectorization for reference-based kernels.

Proposed semantics

Follow rustc's conservative reference rules:

Rust parameter Condition LLVM attributes
&T, &[T] pointee is Freeze noalias readonly
&mut T, &mut [T] pointee is Unpin noalias
*const T — none
*mut T — none
erased / compiler-generated pointer — none

In particular:

  • &UnsafeCell<T> must not receive readonly
  • shared references should not receive aliasing attributes unless the Freeze requirement is satisfied
  • mutable references should not receive noalias unless the Unpin requirement is satisfied
  • raw pointers must not receive reference-derived attributes
  • DisjointSlice should remain uninterpreted for this change

Implementation outline

  1. During MIR import, inspect the original rustc parameter type while the semantic information is still available.
  2. Record only proven per-source-argument facts on MirFuncOp.
  3. Preserve those facts through MIR lowering.
  4. Remap source arguments to their physical LLVM parameters after ABI flattening.
    • For slices, attributes apply only to the pointer component, not the length field.
  5. Make llvm-export mechanically serialize the already-proven facts.
    • The exporter must not infer aliasing from pointer type, mutability, address space, or DisjointSlice.

Acceptance criteria

  • &T / &[T] with Freeze emit noalias readonly
  • &UnsafeCell<T> emits neither attribute
  • &mut T / &mut [T] with Unpin emit noalias
  • &mut !Unpin emits no reference-derived attribute
  • raw pointers emit no reference-derived attributes
  • slice metadata parameters are never annotated
  • invalid attribute combinations are rejected
  • legacy typed-pointer LLVM export continues to work
  • existing CUDA examples continue to build
  • end-to-end generated LLVM IR contains the expected attributes

Out of scope

  • dereferenceable
  • recursive alias analysis
  • deriving attributes from DisjointSlice
  • deriving attributes from raw pointers
  • alias inference based on LLVM-level pointer shapes

Activity

  1. nihalpasham
    added
    codegenDevice code-generation pipeline (Rust MIR to IR to PTX)
    DependsBlocked on another PR/issue or an upstream dependency
    IR-loweringLowering between dialects (dialect-mir to LLVM dialect)
    llvm-exportTextual LLVM/NVVM IR exporter (llvm-export crate)
    perfPerformance of generated code or of the compiler itself
    questionFurther information is requested
    soundness-auditCorrectness case under active maintainer audit; not cleared for merge yet
    on Aug 24, 2026

  2. nihalpasham commented on Aug 24, 2026

    @nihalpasham
    nihalpasham
    on Aug 24, 2026
    Collaborator

    Keeping this open as a design question, but implementation is blocked for now.

    An incorrect optimizer promise can silently miscompile valid code, so the full proof is required under the current acceptance criteria.

  3. nihalpasham commented on Sep 3, 2026

    @nihalpasham
    nihalpasham
    on Sep 3, 2026
    Collaborator

    Building on #1186: the pointer kinds are in, minted from rustc_public only, but a kind is a classification, not an aliasing proof (the gap #1089 was closed for).

    One rule for the design: every fact comes through rustc_public. No TyCtxt reads, no Debug-string parsing.

    What rustc_public can prove today, and what it cannot:

    attribute              needs                                             rustc_public today
nonnull, align N       reference kind (#1186) + layout                   yes: RigidTy::Ref, Ty::layout()
noalias, readonly      Freeze for &T, Unpin (+UnsafeUnpin) for &mut T    no auto-trait query
dereferenceable(N)     same facts, via rustc's nofree                    no; Instance::fn_abi() has the bits, but only as Opaque strings
    • So the order flips:
      • the nonnull / align half (Emit LLVM nonnull / align / dereferenceable for proven Rust references #1098) can be built now from kind + layout. noalias / readonly / dereferenceable
      • wait for rustc_public to expose typed ArgAttributes in PassMode (today Opaque).
      • that is a small upstream change worth driving, I'll take it on; once it lands, we can copy rustc's bits, never re-derive Freeze/Unpin.
    • Flow once a fact exists:
      • typed attribute on mir.func per source argument;
      • mir-lower puts bits on the slice data pointer only (len bare; by-value structs and DisjointSlice nothing);
      • llvm-export prints them on kernel entry params, NVVM path only if nvvmVerifyProgram accepts.
    • Evidence and controls:

    @uurl - Do you want to take the nonnull/align half now? An approach note here first would be great.

  4. nihalpasham commented on Sep 3, 2026

    @nihalpasham
    nihalpasham
    on Sep 3, 2026
    Collaborator

    Filed the upstream ask for typed ArgAttributes in rustc_public: rust-lang/rustc_public#132.

  5. uurl commented on Sep 25, 2026

    @uurl
    uurl
    on Sep 25, 2026
    ContributorAuthor

    @nihalpasham — I saw rust-lang/rust#163242 is now up for review.

    Once that lands and cuda-oxide picks up the corresponding rustc version, are you planning to take the downstream #1088 / remaining #1098 integration, or would you like me to take it?

    I'd be happy to build on the reference-attribute plumbing from #1244 and keep the importer as a straight copy of rustc's typed ArgAttributes decisions, without re-deriving Freeze / Unpin semantics.

  6. praveenperera commented on Oct 10, 2026

    @praveenperera
    praveenperera
    on Oct 10, 2026

    A data point for this, from a speakrs repro on main (bafffe45), sm_80, ptxas 13.0.88:

    A kernel that copies 16-byte pixels from an input: &[f32] parameter by casting to *const U32x4 emits ld.global.v4.b32 (LDG.E.128), but never ld.global.nc. The same happens with *const f32 and with vector::as_vectors::<F32x4>(input). #[readonly] is a function-level marker for extern device functions, so there's currently no way to get the non-coherent path for a read-only input.

    In speakrs's FP16 convolution kernels (https://github.com/avencera/speakrs, crates/speakrs-cuda-kernels/src/wideconv/h16.rs), inputs that no kernel writes during the launch, such as weights and staged activations, are read with ptx_asm! using ld.global.nc.f32, ld.global.nc.v2.u32 and ld.global.nc.v4.u32. If &[T] with a Freeze pointee got noalias readonly as proposed here, NVPTX could select ld.global.nc itself, and we could drop that asm. An explicit .nc load intrinsic for raw pointers would cover the cases where the input isn't a reference.

  7. uurl commented on Oct 11, 2026

    @uurl
    uurl
    on Oct 11, 2026
    ContributorAuthor

    @nihalpasham — now that rust-lang/rust#163242 has landed, and given the new speakrs data point above, are you planning to take the downstream #1088 / remaining #1098 integration once cuda-rust picks up a nightly containing ArgAttributes::regular(), or would you like me to take it?

    I’d be happy to build on #1244 and keep the importer as a straight copy of rustc’s typed ArgAttributes decisions (no_alias, read_only, no_free, pointee_size), without re-deriving Freeze / Unpin.

  8. nihalpasham
    added a commit that references this issue on Oct 11, 2026

    cuda-oxide: update to nightly-2026-10-11 (#1482)

    91a0066

  9. nihalpasham commented on Oct 11, 2026

    @nihalpasham
    nihalpasham
    on Oct 11, 2026
    Collaborator

    The remaining explicit readonly work has an upstream blocker: LLVM #230960. With noalias readonly, NVPTX drops volatility and two required reads can become one machine load. This reproduces on LLVM 22.1.8, 23.1.0, and 23.1.3.

    #1483 adds rustc-proven noalias while deferring explicit kernel readonly. Keeping this issue open for that remaining work; the report includes the reduced IR and the working control.

  10. nihalpasham
    added a commit that references this issue on Oct 11, 2026

    cuda-oxide: preserve rustc reference attributes on kernels (#1483)

    d141118
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

No one assigned

    Labels

    DependsBlocked on another PR/issue or an upstream dependencyIR-loweringLowering between dialects (dialect-mir to LLVM dialect)codegenDevice code-generation pipeline (Rust MIR to IR to PTX)llvm-exportTextual LLVM/NVVM IR exporter (llvm-export crate)perfPerformance of generated code or of the compiler itselfquestionFurther information is requestedsoundness-auditCorrectness case under active maintainer audit; not cleared for merge yet

    Type

    No type

    Projects

    No projects

      Milestone

      No milestone

      Relationships

      None yet

      Development

      No branches or pull requests

      Issue actions