feat(cuda-oxide): carry inline intent through the device pipeline

[ssoj13]

What this adds

Rust lets a function say how it wants to be inlined: #[inline] as a hint, #[inline(always)] as a demand, #[inline(never)] as a refusal. The device pipeline passed only the demand through to LLVM. The hint was dropped, and the refusal stopped Rust's own inliner but not LLVM's, so a helper that wanted to stay a separate GPU function never did. The three requests now reach LLVM as the matching function attributes, the mapping rustc's own backend uses, and #[inline(never)] means one real PTX function.

#[inline(never)] fn never_add(a: f32, b: f32) -> f32     #[inline] fn hint_mul(..)     #[inline(always)] fn always_sub(..)

before:  define float @never_add(float, float) #0        // opt inlines it anyway; the PTX entry has 0 calls
         define float @hint_mul(float, float) #0         // the hint is lost
after:   define float @never_add(float, float) noinline #0      -> call.uni never_add in the PTX
         define float @hint_mul(float, float) inlinehint #0
         define float @always_sub(float, float) alwaysinline #0 -> no function left after opt
         generic #[device] helper scale_by::<4>: no keyword, the factor still folds to an immediate

#[inline(never)]
pub fn never_add(a: f32, b: f32) -> f32 { a + b }   // a real call boundary in the PTX

#[cuda_module]
mod kernels {
    #[inline(never)]
    #[device]
    pub fn never_dev(a: f32, b: f32) -> f32 { a * 2.0 + b }   // also a real call
}

Contract

Source	mir.func attribute	LLVM attribute	In the PTX
#[inline]	inline_intent = hint	inlinehint	opt decides, as on the host
#[inline(always)]	inline_intent = always	alwaysinline	no function left
#[inline(never)]	inline_intent = never	noinline	a .func, reached with call.uni unless opt can forward a returned argument and drop the call, as the cluster example documents
no attribute	none	none	opt decides

• A kernel entry never carries an intent; the dialect verifier rejects one, and it rejects a value of the wrong attribute type under the key.
• The #[inline(never)] that #[device] adds to a helper with its own type or const parameters is not carried. It keeps each monomorphization visible to the host-side collector and nothing else, so const-generic arguments keep folding into their callers. A user's own inline attribute on such a helper still does not reach device code; that is [BUG]: #[device] does not obey #[inline] on a generic function #1471 and stays open.

What we fixed

• #[inline] and #[inline(never)] were dropped on the way to LLVM; only #[inline(always)] survived, through a string attribute on mir.func that the exporter matched by name.
• The intent is now a typed InlineIntentAttr on mir.func, set by the backend from codegen_fn_attrs, lowered to the LLVM function attribute in one place in mir-lower, and verified: a request on a kernel entry or a wrongly typed value under the key is an error, not a silent "none".
• A noinline helper that returns a shared-memory pointer is now a real function, and libNVVM declares that 32-bit return slot as .param .b64 while the body stores 32 bits and every call site reads a 32-bit retval0, so ptxas rejected addressof_sharedarray on the --emit-nvvm-ir route. The legalizer now gives every non-kernel definition that returns a bare non-generic pointer, inline attribute or not, a generic return and restores the address space with an addrspacecast on both sides of each call. Parameters, aggregates, kernels and declarations are untouched; a function whose address is taken is reported. The libNVVM side is tracked internally (NVBug 6946922); the pass goes once it ships.
• The drop_glue example (fix(mir-importer): preserve non-terminating drops #1361, landed after this branch was cut) relied on #[inline(never)] being dropped for its store-then-self-loop PTX check; test(drop_glue): keep destructor MIR with -Zinline-mir=no, not #[inline(never)] #1549 moved it to package-scoped -Zinline-mir=no ahead of this PR, so the full smoketest is green with every example.
• kernel_reference_validity kept its helper route inlined with #[inline(never)] and a comment saying LLVM would inline it anyway; it now says what it means with -Zinline-mir=no on its own package, and its two volatile loads stay in the entry.

Implementation

• rustc-codegen-cuda/src/device_codegen.rs: maps rustc_attr_ir::InlineAttr (Hint, Always/Force, Never) to InlineIntentAttr before entering the stable_mir context; is_generic_device_collector_boundary recognizes the macro's own #[inline(never)] on a prefixed device symbol with its own generics and drops it. The test keys on the device symbol name plus own generics, which holds while the macro writes its attribute first; once fix(cuda-macros): obey #[inline] on generic device functions #1472 lets a user attribute replace the macro's, it must key on the attribute's span instead.
• dialect-mir: InlineIntentAttr (mir.inline_intent, a pliron_attr enum) and MirFuncOp::{set_inline_intent, inline_intent} under the inline_intent key; the verifier rejects the attribute on a kernel and a foreign type under the key.
• mir-importer: FunctionInput.inline replaces is_inline_always; translator/body.rs sets the attribute on the new mir.func.
• mir-lower/src/lowering.rs: lower_inline_intent_attr is the only place that names inlinehint, alwaysinline and noinline; it writes them into llvm_func_attrs, which llvm-export already renders. The old by-name alwaysinline special case in the exporter is gone.
• nvvm-transforms/src/legalize.rs: legalize_pointer_returns runs for both NVVM dialects before the legacy pass; retype_pointer_return casts to generic before each ret, cast_call_result_back rebuilds each direct call and casts the result back. reserved-oxide-symbols gains the device-symbol test the backend uses.
• Examples: inline_intent (new; one helper of each kind, a generic #[device] helper, host and device compute the same values; verify-code-shape.sh checks the keyword on each pre-opt define, two call.uni in the entry, no always_sub function, and the folded immediate; MIR inlining is off for its package so every helper reaches LLVM), kernel_reference_validity (-Zinline-mir=no), cluster (shape script comment: opt still removes the call because it forwards the returned argument). scripts/smoketest.sh lists inline_intent among the examples rebuilt unoptimized for their shape check.
• Book: the #[device] helper row in the supported-features appendix states the mapping and the exception.

Verification

Checked at d0bcc80f8 merged with main 5b606fddd, #1509 and #1510 (nightly-2026-10-11, two RTX 5090, CUDA 13.4):

• ✅ cargo test --workspace --all-targets in cuda-oxide: 5596 passed, 0 failed; the backend crate's own 33 tests pass. Clippy with -D warnings on both workspaces, rustdoc, cargo oxide fmt --check, the lock check, the 18 guards and cargo deny pass. All 55 CI checks passed.
• ✅ cargo oxide run inline_intent: SUCCESS: all 256 elements match the host reference, and verify-code-shape.sh passes: noinline, inlinehint and alwaysinline on the right pre-opt defines, no keyword on the const-generic #[device] helper, two call.uni in the entry, no always_sub function left, scale_by::<4> folded to 0f40800000.
• ✅ kernel_reference_validity passes with its shape check (two volatile loads stay in the entry under -Zinline-mir=no); const_generic and cluster pass on sm_120.
• ✅ libNVVM route: cargo oxide run addressof_sharedarray --emit-nvvm-ir --arch sm_120 passes with the legalized generic return; main's backend inlines the helper and passes too. memcheck on the libNVVM-route and direct-route binaries: 0 errors.
• ✅ PTX unchanged where no inline attribute is involved: gemm_sol, tcgen05_matmul, wgmma, mma_mxf8f6f4 and vecadd built with main's backend and with this tree's are identical after hash normalization.
• ✅ Full GPU smoketest on the stacked tree: 237 of 238. The one failure was drop_glue's shape check from fix(mir-importer): preserve non-terminating drops #1361, which assumed #[inline(never)] is ignored; with test(drop_glue): keep destructor MIR with -Zinline-mir=no, not #[inline(never)] #1549 applied the lane and its neighbors (cluster, inline_intent, kernel_reference_validity) pass. The ten LTOIR-lane examples ran through libNVVM with --emit-nvvm-ir in that smoketest.

Caveats

• On the libNVVM route a device helper that returns a shared-memory pointer now returns a generic pointer with casts on both sides; the SASS is unchanged where opt inlines the helper, and the cast pair is what libNVVM's own CUDA front end emits for such returns. The pass is a workaround for the libNVVM declaration defect and can go when the fix ships.
• #[inline(never)] on a closure is unreachable on stable Rust, so the closure case named in the original description is not covered.
• The generic #[device] exception means a user's inline attribute on a helper with its own generics still does not reach device code ([BUG]: #[device] does not obey #[inline] on a generic function #1471).
• Not reproduced on sm_86 (the author's board) or with CUDA 13.3; the author's own runs cover those.
• The mapping covers the #[inline] family. rustc's own backend also emits noinline for #[optimize(none)] and inlinehint for shims and drop glue that require inlining; neither is touched here.
• cuda-device's own #[inline(never)] helpers (cluster_sync, the __cluster_config marker) now reach the pre-opt IR as noinline definitions; opt removes them as unreferenced, and the compared examples emit identical PTX on the direct route.
• The committed inline_intent/Cargo.lock predates Tile libraries: call Rust device functions from Tile kernels #1545's lock re-alignment, so the first build adds four entries (itoa, oxide-tile-export, serde_json, zmij); the follow-up that adds the book figure re-aligns it.

Related: #188, #189, #1471. Port of ansidium#11.

[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]

Thanks for this, the fix is right and it is going in. One snag: your fork is a fork of ansidium/cuda-rust-windows, not of NVIDIA/cuda-rust, so "allow edits from maintainers" gives push rights to ansidium, not to us, and I could not push my follow-ups onto your branch.

I reviewed and validated this on the 5090 and added the follow-ups on top of your commit, on branch landing/pr-1511 in https://github.com/nihalpasham/cuda-rust. Could you fork NVIDIA/cuda-rust directly, push that branch to your fork, and open the PR from there? Then it lands with you as the author. I will close this one once the new PR is up.

[ssoj13]

Thanks @nihalpasham! Since GitHub only permits one fork per repository network (ssoj13/cuda-rust-windows is already in this fork network via ansidium/cuda-rust-windows), I fetched your exact landing/pr-1511 branch (c9e4241366ab05d5028ba3362ce86c3f6fdcc3c5) and pushed it directly onto this PR's head branch (nvidia/feat/inline-intent), as well as ssoj13/cuda-rust-windows:landing/pr-1511. Ready to merge!

[nihalpasham]

/ok to test d0bcc80

[nihalpasham]

Thanks, this is the right direction and it matches what rustc's own backend does: #[inline] becomes inlinehint, #[inline(always)] becomes alwaysinline, #[inline(never)] becomes noinline. Also thanks for the pliron 0.20 .clone() fix in your merge. What I added before landing:

• The intent rides on mir.func as a typed InlineIntentAttr instead of a string, so the verifier can reject a request on a kernel entry and nothing downstream has to parse text.
• A noinline helper that returns a shared-memory pointer is now a real function, and libNVVM mis-declares that 32-bit return as 64 bits, so ptxas rejected addressof_sharedarray on the --emit-nvvm-ir route. The libNVVM legalizer (both NVVM dialects) now returns non-generic-address-space pointers, shared among them, from device helpers as generic pointers and casts back at the call site. The libNVVM side is tracked internally (NVBug 6946922); the pass goes once it ships.
• kernel_reference_validity keeps its two volatile loads with -Zinline-mir=no instead of #[inline(never)], which now means what it says.
• A new inline_intent example runs one helper of each kind and checks the attribute in the pre-opt IR and the result in the PTX: call.uni for the two noinline helpers, no function left for alwaysinline, and the const-generic factor still folded to an immediate.
• The generic #[device] exception stays as you had it; [BUG]: #[device] does not obey #[inline] on a generic function #1471 remains open for the user attribute that macro shadows (fix(cuda-macros): obey #[inline] on generic device functions #1472 is the fix in flight).
• One example that landed after your branch, drop_glue from fix(mir-importer): preserve non-terminating drops #1361, leaned on #[inline(never)] being ignored for its PTX check. test(drop_glue): keep destructor MIR with -Zinline-mir=no, not #[inline(never)] #1549 moved it to -Zinline-mir=no first, the same thing I did for kernel_reference_validity.

At d0bcc80f8 merged with today's main plus #1509 and #1510: the whole cuda-oxide workspace suite, clippy, rustdoc, fmt, the guards and all CI checks pass, and the full GPU smoketest on the 5090 is green with #1549 in place. The tensor-core examples emit identical PTX on the direct route, and the ten LTOIR-lane examples run through libNVVM.