From f2bdfa3fa5aa7ebfb6132bd64e14203f32a3c594 Mon Sep 17 00:00:00 2001 From: Yunus Mutlu Date: Tue, 29 Sep 2026 22:44:37 +0300 Subject: [PATCH 1/5] fix(mir-importer): preserve emitted diverging device calls Signed-off-by: Yunus Mutlu --- .../crates/mir-importer/src/pipeline.rs | 13 +++-- .../mir-importer/src/translator/facts.rs | 21 +++++++ .../src/translator/terminator/helpers.rs | 57 ++++++++++++------- .../src/translator/terminator/mod.rs | 46 +++++++++------ .../examples/function_item_call/README.md | 10 ++++ .../examples/function_item_call/src/main.rs | 25 +++++++- .../function_item_call/verify-code-shape.sh | 50 ++++++++++++++++ .../rustc-codegen-cuda/src/collector.rs | 14 ++--- 8 files changed, 184 insertions(+), 52 deletions(-) create mode 100644 cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/verify-code-shape.sh diff --git a/cuda-oxide/crates/mir-importer/src/pipeline.rs b/cuda-oxide/crates/mir-importer/src/pipeline.rs index 5d88f4a234..d405375f40 100644 --- a/cuda-oxide/crates/mir-importer/src/pipeline.rs +++ b/cuda-oxide/crates/mir-importer/src/pipeline.rs @@ -494,11 +494,16 @@ pub fn run_pipeline( config: &PipelineConfig, known_defs: crate::translator::facts::KnownDefs, ) -> Result { - // Install the driver-resolved lang-item ids for this run. Set (not - // merged) every entry: the ids are only valid inside the caller's - // `rustc_internal::run` context, so a stale set from a previous run on - // this thread must never survive. + // Install the driver-resolved lang-item ids and emitted-function identities + // for this run. Replace (rather than merge) both thread-local sets so + // stale state from a previous run on this thread cannot survive. crate::translator::facts::set_known_defs(known_defs); + crate::translator::facts::set_emitted_functions( + functions + .iter() + .map(|func| func.instance.mangled_name().to_string()) + .collect(), + ); prepare_output_dir(&config.output_dir)?; diff --git a/cuda-oxide/crates/mir-importer/src/translator/facts.rs b/cuda-oxide/crates/mir-importer/src/translator/facts.rs index 69302b9241..6e1a262f06 100644 --- a/cuda-oxide/crates/mir-importer/src/translator/facts.rs +++ b/cuda-oxide/crates/mir-importer/src/translator/facts.rs @@ -85,6 +85,12 @@ thread_local! { index_mut_trait: None, }) }; + // Exact monomorphized Rust functions that this pipeline run will emit. + // + // Store rustc's linkage identity rather than the legalised module symbol: + // call classification must not depend on Legaliser ordering or spelling. + static EMITTED_FUNCTIONS: std::cell::RefCell> = + const { std::cell::RefCell::new(std::collections::BTreeSet::new()) }; } /// Installs the driver-resolved lang-item ids for this pipeline run, @@ -99,6 +105,21 @@ pub(crate) fn known_defs() -> KnownDefs { KNOWN_DEFS.with(|cell| cell.get()) } +/// Replace the emitted-function set for the current pipeline run. +/// +/// Like [`set_known_defs`], this replaces rather than extends previous +/// thread-local state so one rustc invocation cannot leak into the next. +pub(crate) fn set_emitted_functions(names: std::collections::BTreeSet) { + EMITTED_FUNCTIONS.with(|set| *set.borrow_mut() = names); +} + +/// Whether this exact monomorphized Rust instance will have a `mir.func` +/// definition in the module being translated. +pub(crate) fn is_emitted_function(instance: &mir::mono::Instance) -> bool { + let identity = instance.mangled_name().to_string(); + EMITTED_FUNCTIONS.with(|set| set.borrow().contains(&identity)) +} + /// Validate the source reference and the fully specialized read-only storage /// contract using rustc's type system. MIR locals have erased regions, so the /// outer lifetime is read from the declaration before instantiation instead. diff --git a/cuda-oxide/crates/mir-importer/src/translator/terminator/helpers.rs b/cuda-oxide/crates/mir-importer/src/translator/terminator/helpers.rs index a85b623374..74a09706e3 100644 --- a/cuda-oxide/crates/mir-importer/src/translator/terminator/helpers.rs +++ b/cuda-oxide/crates/mir-importer/src/translator/terminator/helpers.rs @@ -316,10 +316,10 @@ pub fn bundle_generated_u32_results_as_array( /// 1. Translate all MIR arguments to Pliron IR values /// 2. At a foreign ABI boundary, adapt pointer address spaces to the exact /// declared parameter types without changing pointer kind or mutability -/// 3. Prepare the call destination using the current projection state +/// 3. For returning calls, prepare the destination using the current projection state /// 4. Create a `mir.call` operation carrying the callee's name attribute -/// 5. Store the result through the prepared destination -/// 6. Emit a zero-operand goto to the call's success target +/// 5. For returning calls, store the result and goto the success target +/// 6. For diverging calls, emit `mir.unreachable` after the call /// /// Reference arguments (`&mut local`) are handed the local's alloca slot /// pointer directly, so callee writes through the reference are observed by @@ -383,16 +383,21 @@ pub fn emit_function_call( } } - let (prepared_destination, prepared_last_op) = prepare_destination_write( - ctx, - body, - destination, - value_map, - block_ptr, - last_op, - loc.clone(), - )?; - last_op = prepared_last_op; + let prepared_destination = if target.is_some() { + let (prepared, prepared_last_op) = prepare_destination_write( + ctx, + body, + destination, + value_map, + block_ptr, + last_op, + loc.clone(), + )?; + last_op = prepared_last_op; + Some(prepared) + } else { + None + }; use pliron::builtin::attributes::StringAttr; @@ -423,11 +428,20 @@ pub fn emit_function_call( call_op }; + if target.is_none() { + return Ok(super::emit_unreachable_after( + ctx, + block_ptr, + Some(call_op), + loc, + )); + } + let result_value = call_op.deref(ctx).get_result(0); let goto_prev = finish_destination_write( ctx, - prepared_destination, + prepared_destination.expect("returning call prepared its destination"), result_value, value_map, block_ptr, @@ -435,14 +449,13 @@ pub fn emit_function_call( loc.clone(), )?; - if let Some(target_idx) = target { - Ok(emit_goto(ctx, *target_idx, goto_prev, block_map, loc)) - } else { - input_err!( - loc.clone(), - TranslationErr::unsupported("Call terminator without target not supported".to_string(),) - ) - } + Ok(emit_goto( + ctx, + target.expect("returning call has a success target"), + goto_prev, + block_map, + loc, + )) } /// Adapt a foreign-call pointer argument only when the address space is the diff --git a/cuda-oxide/crates/mir-importer/src/translator/terminator/mod.rs b/cuda-oxide/crates/mir-importer/src/translator/terminator/mod.rs index 15321134d3..a7f76a5b10 100644 --- a/cuda-oxide/crates/mir-importer/src/translator/terminator/mod.rs +++ b/cuda-oxide/crates/mir-importer/src/translator/terminator/mod.rs @@ -1254,25 +1254,37 @@ fn translate_call( return Ok(result); } - // Handle diverging calls (calls that never return, like unwrap_failed, panic, etc.) - // These have no target block because the function never returns. + // `Call { target: None }` means there is no normal successor. It does not + // mean an emitted callee can be erased: a collected Rust `-> !` function + // may perform observable work before it diverges. Preserve that call and + // let `emit_function_call` terminate the caller with `mir.unreachable`. // - // The callee is dropped and replaced by an immediate trap. That applies - // to every `-> !` callee reaching this point, including a user - // `#[device]` fn that legitimately diverges (e.g. `loop {}`); full Rust - // fidelity would emit the call for resolvable callees the way - // `translate_function_item_call` does. Dropping the call is a - // pre-existing semantic: the trap only makes it safe, where the bare - // `unreachable` previously emitted here was UB that let `opt` delete - // the whole panic path. - // - // Panic entry points additionally never get here with any statements - // translated ahead of them: `block::translate_block` recognizes the same - // shape via [`is_dropped_panic_call`] and emits the trap directly. + // For every diverging callee that is not part of this pipeline run's + // emitted-function set, retain the pre-existing trap fallback. Keep this + // before the unsupported-intrinsic and libm checks below so non-emitted + // diverging calls still take the same path as before this fix. if target_usize.is_none() { - // This is a diverging call (returns !) - emit trap + unreachable - // Examples: unwrap_failed(), panic!(), abort() - return Ok(emit_trap_unreachable_after(ctx, block_ptr, prev_op, loc)); + let resolved_instance = match func { + mir::Operand::Constant(constant) + if *constant.const_.kind() == ConstantKind::ZeroSized => + { + match constant.const_.ty().kind() { + rustc_public::ty::TyKind::RigidTy(rustc_public::ty::RigidTy::FnDef( + fn_def, + substs, + )) => rustc_public::mir::mono::Instance::resolve(fn_def, &substs).ok(), + _ => None, + } + } + _ => None, + }; + + if !resolved_instance + .as_ref() + .is_some_and(crate::translator::facts::is_emitted_function) + { + return Ok(emit_trap_unreachable_after(ctx, block_ptr, prev_op, loc)); + } } // A call to a rustc intrinsic that no dispatch arm above recognized can diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/README.md b/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/README.md index 62efd401c6..26bd973de6 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/README.md +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/README.md @@ -7,3 +7,13 @@ Rust lowers `apply_once(plus_seven, x)` through a callable trait method such as function-item body, and MIR import must emit a direct call name that matches the collector/export naming policy. Otherwise the generated device IR references a callee symbol that was never emitted. + +The diverging-call regression also covers an ordinary direct `#[device]` +function returning `!`. Its body performs an observable store before looping +forever. A MIR `Call` with no normal successor must therefore keep the call and +end the caller block with `unreachable`; replacing the call itself with a trap +would erase the store. + +`verify-code-shape.sh` pins that contract in the pre-optimization LLVM IR, +before the middle-end can inline the diverging helper. The divergent branches +are not launched; the host run exercises only the returning path. diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/src/main.rs b/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/src/main.rs index 68b442f644..c878dbbfc9 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/src/main.rs +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/src/main.rs @@ -10,8 +10,9 @@ //! back to the concrete function body instead of emitting a dangling trait-shim //! callee symbol. The regression also covers `Fn`/`FnMut`, references, //! multiple and nested-tuple arguments, tuple returns, and type/const-generic -//! function items. Diverging function items and closures also verify that a -//! direct Rust call returning `!` ends its MIR block with `unreachable`. +//! function items. Diverging function items, closures, and an ordinary direct +//! `-> !` device helper also verify that a non-returning Rust call executes +//! before its MIR block ends with `unreachable`. //! `call_once_decoy` proves that an ordinary function whose name contains //! trait-like text is not mistaken for a callable-trait shim. @@ -57,6 +58,18 @@ fn spin_forever(_value: u32) -> ! { loop {} } +// Unlike `spin_forever`, this is reached by an ordinary direct MIR `Call`. +// The store makes replacing that call with an immediate trap observably wrong. +#[allow(clippy::empty_loop)] +#[inline(never)] +#[device] +unsafe fn write_then_diverge(ptr: *mut u32) -> ! { + unsafe { + ptr.write(0x1111_1111); + } + loop {} +} + #[device] fn apply_once u32>(f: F, x: u32) -> u32 { f(x) @@ -148,6 +161,14 @@ mod kernels { #[allow(clippy::empty_loop)] apply_never(|_| loop {}, flag); } + if flag == 3 { + // The importer must preserve this ordinary direct call. The callee + // stores first and then diverges, so replacing it with a trap loses + // an observable side effect. + unsafe { + write_then_diverge(out.as_mut_ptr()); + } + } if let Some(slot) = out.get_mut(idx) { *slot = 0xd1ce; } diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/verify-code-shape.sh b/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/verify-code-shape.sh new file mode 100644 index 0000000000..f231d6d102 --- /dev/null +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/verify-code-shape.sh @@ -0,0 +1,50 @@ +#!/usr/bin/env bash +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +set -euo pipefail + +root="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +llvm="${root}/function_item_call.ll" + +test -s "${llvm}" + +extract_llvm_function() { + local needle="$1" + awk -v needle="${needle}" ' + /^define / && index($0, needle) != 0 { in_function = 1 } + in_function { print } + in_function && /^}/ { exit } + ' "${llvm}" +} + +kernel="$(extract_llvm_function '@diverging_callable(')" +test -n "${kernel}" + +call_count="$( + grep -Ec 'call void @[^ (]*write_then_diverge\(' <<<"${kernel}" || true +)" +if [[ "${call_count}" -ne 1 ]]; then + echo "diverging_callable must contain exactly one direct call to write_then_diverge" >&2 + exit 1 +fi + +call_line="$( + grep -nE 'call void @[^ (]*write_then_diverge\(' <<<"${kernel}" | + head -n1 | + cut -d: -f1 +)" +next_line="$(sed -n "$((call_line + 1))p" <<<"${kernel}")" +if ! grep -Eq '^[[:space:]]*unreachable[[:space:]]*$' <<<"${next_line}"; then + echo "direct diverging call must be immediately followed by LLVM unreachable" >&2 + exit 1 +fi + +helper="$(extract_llvm_function 'write_then_diverge')" +test -n "${helper}" +if ! grep -Eq 'store i32 286331153, ptr ' <<<"${helper}"; then + echo "write_then_diverge lost its observable sentinel store" >&2 + exit 1 +fi + +echo "function_item_call code shape: PASS" diff --git a/cuda-oxide/crates/rustc-codegen-cuda/src/collector.rs b/cuda-oxide/crates/rustc-codegen-cuda/src/collector.rs index 0b2bb713ab..a83078e2cf 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/src/collector.rs +++ b/cuda-oxide/crates/rustc-codegen-cuda/src/collector.rs @@ -2464,10 +2464,9 @@ impl<'tcx> DeviceCollector<'tcx> { /// "Symbol ... not found" error. /// /// Functions that are declared diverging (return type `!`) are left - /// alone: their call sites have no target block, so the translator - /// already lowers them to LLVM `unreachable` (note: NOT a trap; - /// the optimizer may delete paths that provably reach it, a known - /// gap tracked separately). + /// alone. If their body is emitted, the importer preserves the call and + /// terminates its caller with `unreachable`; if collection deliberately + /// omits the body, the importer retains the existing device-trap fallback. fn check_unreachable_callee( &self, resolved: Instance<'tcx>, @@ -2547,8 +2546,9 @@ impl<'tcx> DeviceCollector<'tcx> { // No stub marker: leave non-local functions to the existing silent // skip (this is what keeps the real `cuda_device` intrinsic - // placeholders and `core`'s cold panic wrappers working), and leave - // declared-diverging functions to the translator's `unreachable` lowering. + // placeholders and `core`'s cold panic wrappers working). Declared- + // diverging functions that are omitted here retain the translator's + // device-trap fallback rather than producing an undefined call. if !def_id.is_local() || mir.return_ty().is_never() { return; } @@ -2577,7 +2577,7 @@ impl<'tcx> DeviceCollector<'tcx> { .with_help( "annotate device helpers with `#[device]`; if the panic is \ intentional, declare the function as diverging (`-> !`) so the \ - call lowers to LLVM `unreachable`", + device pipeline can treat the path as non-returning", ) .emit() } From d5906d226fbec09beb088324572fccb8267943fb Mon Sep 17 00:00:00 2001 From: nihalpasham Date: Mon, 19 Oct 2026 15:38:37 +0530 Subject: [PATCH 2/5] test(mir-importer): cover the emitted-function set and name the retained traps Split is_emitted_function into a name-based is_emitted_name that it calls, so the per-run set can be tested without a rustc session, and add that test: the set installed for one pipeline run replaces the previous run's set instead of merging into it. A leftover name would keep a diverging call to a function the current run never emits. Make the translate_call comment at the diverging-call arm accurate. The arm fires for any visibly uninhabited destination type, because that is when rustc clears the call target, and two cases keep the trap: a call through a function pointer resolves no instance, and a callee whose path matches the panic-entry predicate is trapped by block translation before translate_call runs. Signed-off-by: nihalpasham --- .../mir-importer/src/translator/facts.rs | 34 +++++++++++++++++-- .../src/translator/terminator/mod.rs | 26 +++++++++----- 2 files changed, 50 insertions(+), 10 deletions(-) diff --git a/cuda-oxide/crates/mir-importer/src/translator/facts.rs b/cuda-oxide/crates/mir-importer/src/translator/facts.rs index 6e1a262f06..db134cb78d 100644 --- a/cuda-oxide/crates/mir-importer/src/translator/facts.rs +++ b/cuda-oxide/crates/mir-importer/src/translator/facts.rs @@ -116,8 +116,15 @@ pub(crate) fn set_emitted_functions(names: std::collections::BTreeSet) { /// Whether this exact monomorphized Rust instance will have a `mir.func` /// definition in the module being translated. pub(crate) fn is_emitted_function(instance: &mir::mono::Instance) -> bool { - let identity = instance.mangled_name().to_string(); - EMITTED_FUNCTIONS.with(|set| set.borrow().contains(&identity)) + is_emitted_name(&instance.mangled_name()) +} + +/// Whether `mangled_name`, rustc's linkage identity of one monomorphized +/// function, is in the current run's emitted-function set. The name-based +/// form lets the set's replace-not-merge contract be tested without a rustc +/// session. +pub(crate) fn is_emitted_name(mangled_name: &str) -> bool { + EMITTED_FUNCTIONS.with(|set| set.borrow().contains(mangled_name)) } /// Validate the source reference and the fully specialized read-only storage @@ -1125,4 +1132,27 @@ pub fn reference_validity( assert!(fact.pointee_alignment.is_power_of_two()); } } + + /// `set_emitted_functions` replaces the previous run's set instead of + /// merging into it. A name left over from an earlier rustc invocation on + /// the same thread would keep a diverging call to a function this run + /// never emits, and the module would reference an undefined symbol. + #[test] + fn emitted_function_set_is_replaced_per_run() { + let first_run = std::collections::BTreeSet::from(["first_run_die".to_string()]); + set_emitted_functions(first_run); + assert!(is_emitted_name("first_run_die")); + assert!(!is_emitted_name("second_run_spin")); + + let second_run = std::collections::BTreeSet::from(["second_run_spin".to_string()]); + set_emitted_functions(second_run); + assert!(is_emitted_name("second_run_spin")); + assert!( + !is_emitted_name("first_run_die"), + "an entry from the previous run must not survive the replacement" + ); + + set_emitted_functions(std::collections::BTreeSet::new()); + assert!(!is_emitted_name("second_run_spin")); + } } diff --git a/cuda-oxide/crates/mir-importer/src/translator/terminator/mod.rs b/cuda-oxide/crates/mir-importer/src/translator/terminator/mod.rs index cfc5478c1f..d88f37d4b0 100644 --- a/cuda-oxide/crates/mir-importer/src/translator/terminator/mod.rs +++ b/cuda-oxide/crates/mir-importer/src/translator/terminator/mod.rs @@ -1254,15 +1254,25 @@ fn translate_call( return Ok(result); } - // `Call { target: None }` means there is no normal successor. It does not - // mean an emitted callee can be erased: a collected Rust `-> !` function - // may perform observable work before it diverges. Preserve that call and - // let `emit_function_call` terminate the caller with `mir.unreachable`. + // `Call { target: None }` means there is no normal successor. rustc + // clears the target when the destination type is visibly uninhabited: + // `!`, `core::convert::Infallible`, an empty enum. That does not mean the + // callee can be erased. An emitted Rust function with such a return type + // may do observable work before it diverges. When the resolved instance + // is in this pipeline run's emitted-function set, keep the call and let + // `emit_function_call` end the block with `mir.unreachable`, which is + // what rustc's own backend emits for a call without a destination. // - // For every diverging callee that is not part of this pipeline run's - // emitted-function set, retain the pre-existing trap fallback. Keep this - // before the unsupported-intrinsic and libm checks below so non-emitted - // diverging calls still take the same path as before this fix. + // Every other diverging call keeps the pre-existing trap fallback. That + // covers callees the collector skipped, such as core helpers without MIR + // and local `-> !` bodies that are nothing but a panic, and a call + // through a function pointer, whose `func` operand is not a constant and + // so resolves no instance. Panic entry points never reach this arm: for + // a callee whose path matches `is_panic_entry_path`, + // `block::translate_block` lowers the whole block to the trap through + // `is_dropped_panic_call` before any call is classified. Keep this arm + // before the unsupported-intrinsic and libm checks below so the trapped + // cases take the same path as before. if target_usize.is_none() { let resolved_instance = match func { mir::Operand::Constant(constant) From a619f2d5644841adf2727eae5202c01b6f1e36f3 Mon Sep 17 00:00:00 2001 From: nihalpasham Date: Mon, 19 Oct 2026 15:38:37 +0530 Subject: [PATCH 3/5] test(examples): check every diverging call site in function_item_call The shape check asserted one of the three call sites the fix changes in diverging_callable, read only the pre-optimization IR, had no control for the trap fallback, and rejected the `!dbg` form of `unreachable` that a --device-debug build emits. It now asserts the two apply_never calls and the write_then_diverge call, each followed by `unreachable`, accepts the `!dbg` form, and, when the direct route wrote PTX, reads the kernel entry: the sentinel store or a call to the helper must survive opt and llc, with one trap for the control. The control is panic_only, a `-> !` function whose body is only a panic: the collector does not emit it, so its call site must still lower to the device trap. Built with the previous backend the check exits 1; with the fix it passes on the direct, NVVM and device-debug builds. The module doc now says the diverging calls are emitted, not executed, the README describes what the check covers, and the script carries the executable bit like most of its siblings. Signed-off-by: nihalpasham --- .../examples/function_item_call/README.md | 13 ++- .../examples/function_item_call/src/main.rs | 30 +++++-- .../function_item_call/verify-code-shape.sh | 82 +++++++++++++++---- 3 files changed, 103 insertions(+), 22 deletions(-) mode change 100644 => 100755 cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/verify-code-shape.sh diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/README.md b/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/README.md index 26bd973de6..3a06d7c5af 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/README.md +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/README.md @@ -14,6 +14,13 @@ forever. A MIR `Call` with no normal successor must therefore keep the call and end the caller block with `unreachable`; replacing the call itself with a trap would erase the store. -`verify-code-shape.sh` pins that contract in the pre-optimization LLVM IR, -before the middle-end can inline the diverging helper. The divergent branches -are not launched; the host run exercises only the returning path. +`verify-code-shape.sh` checks that contract in the pre-optimization LLVM IR, +before the middle-end can inline the diverging helper. It covers the three +call sites in `diverging_callable` that used to lower to a trap: two through +`apply_never` (a function item and a closure) and the direct call to +`write_then_diverge`. Each kept call must be followed by `unreachable`. When +the build also wrote PTX, the kernel entry must keep the sentinel store or a +call to the helper. The control is `panic_only`, a `-> !` function whose body +is nothing but a panic: the collector does not emit it, so its call site must +still lower to the device trap, in the IR and in the PTX. The divergent +branches are not launched; the host run exercises only the returning path. diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/src/main.rs b/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/src/main.rs index c878dbbfc9..08e0784125 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/src/main.rs +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/src/main.rs @@ -11,8 +11,10 @@ //! callee symbol. The regression also covers `Fn`/`FnMut`, references, //! multiple and nested-tuple arguments, tuple returns, and type/const-generic //! function items. Diverging function items, closures, and an ordinary direct -//! `-> !` device helper also verify that a non-returning Rust call executes -//! before its MIR block ends with `unreachable`. +//! `-> !` device helper also verify that a non-returning Rust call is emitted +//! before its MIR block ends with `unreachable`. `panic_only` is the control: +//! its `-> !` body is nothing but a panic, so the collector never emits it and +//! its call site must still lower to the device trap. //! `call_once_decoy` proves that an ordinary function whose name contains //! trait-like text is not mistaken for a callable-trait shim. @@ -70,6 +72,18 @@ unsafe fn write_then_diverge(ptr: *mut u32) -> ! { loop {} } +// Control for the retained trap. A `-> !` body that is nothing but a panic is +// a placeholder to the collector, so this function is never emitted and its +// call site must keep the device trap. `#[inline(never)]` keeps the call in +// the kernel's MIR, so the importer classifies this callee rather than the +// panic entry point inside it. The name must not contain a `panicking` module +// segment: that would trap the whole block before the call is classified. +#[inline(never)] +#[device] +fn panic_only() -> ! { + panic!("not collected") +} + #[device] fn apply_once u32>(f: F, x: u32) -> u32 { f(x) @@ -146,9 +160,10 @@ mod kernels { } } - /// Compile both diverging callable paths without entering them on-device. - /// A Rust call returning `!` has no successor block, so the importer must - /// emit the direct call followed by `mir.unreachable`. + /// Compile the diverging call paths without entering them on-device. + /// A Rust call returning `!` has no successor block. When the callee's + /// body is emitted, the importer must keep the direct call followed by + /// `mir.unreachable`; when it is not, the call site keeps the device trap. #[kernel] pub fn diverging_callable(flag: u32, mut out: DisjointSlice) { let idx = thread::index_1d(); @@ -169,6 +184,11 @@ mod kernels { write_then_diverge(out.as_mut_ptr()); } } + if flag == 4 { + // Retained-trap control: the callee's body is only a panic, so the + // collector does not emit it and this call site must still trap. + panic_only(); + } if let Some(slot) = out.get_mut(idx) { *slot = 0xd1ce; } diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/verify-code-shape.sh b/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/verify-code-shape.sh old mode 100644 new mode 100755 index f231d6d102..a866ee8104 --- a/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/verify-code-shape.sh +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/verify-code-shape.sh @@ -6,6 +6,7 @@ set -euo pipefail root="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" llvm="${root}/function_item_call.ll" +ptx="${root}/function_item_call.ptx" test -s "${llvm}" @@ -21,22 +22,47 @@ extract_llvm_function() { kernel="$(extract_llvm_function '@diverging_callable(')" test -n "${kernel}" -call_count="$( - grep -Ec 'call void @[^ (]*write_then_diverge\(' <<<"${kernel}" || true -)" -if [[ "${call_count}" -ne 1 ]]; then - echo "diverging_callable must contain exactly one direct call to write_then_diverge" >&2 +# A kept diverging call must be followed by LLVM `unreachable` on the next +# line. A `--device-debug` build appends `, !dbg !N` to that line, so the +# match stops at the keyword. Monomorphized callees carry a mangled suffix +# after the Rust name (`apply_neverNvB2_12spin_foreverEB2_`), so the pattern +# allows symbol text on both sides of the name. +expect_kept_diverging_calls() { + local callee="$1" + local expected="$2" + local pattern="call void @[^ (]*${callee}[^ (]*\\(" + local found + found="$(grep -Ec "${pattern}" <<<"${kernel}" || true)" + if [[ "${found}" -ne "${expected}" ]]; then + echo "diverging_callable must keep ${expected} direct call(s) to ${callee}, found ${found}" >&2 + exit 1 + fi + local line next + while IFS= read -r line; do + next="$(sed -n "$((line + 1))p" <<<"${kernel}")" + if ! grep -Eq '^[[:space:]]*unreachable(,|[[:space:]]*$)' <<<"${next}"; then + echo "the direct diverging call to ${callee} at kernel line ${line} is not followed by LLVM unreachable" >&2 + exit 1 + fi + done < <(grep -nE "${pattern}" <<<"${kernel}" | cut -d: -f1) +} + +# The three call sites the importer used to replace with a trap: two through +# the generic `apply_never` (a function item and a closure) and the ordinary +# direct call to the store-then-loop helper. +expect_kept_diverging_calls 'apply_never' 2 +expect_kept_diverging_calls 'write_then_diverge' 1 + +# Retained-trap control. `panic_only` has a body that is nothing but a panic, +# which the collector does not emit, so its call site must still lower to the +# device trap and the callee must not be called. +trap_count="$(grep -Ec 'asm sideeffect "trap;"' <<<"${kernel}" || true)" +if [[ "${trap_count}" -ne 1 ]]; then + echo "diverging_callable must keep one trap for the panic-only callee, found ${trap_count}" >&2 exit 1 fi - -call_line="$( - grep -nE 'call void @[^ (]*write_then_diverge\(' <<<"${kernel}" | - head -n1 | - cut -d: -f1 -)" -next_line="$(sed -n "$((call_line + 1))p" <<<"${kernel}")" -if ! grep -Eq '^[[:space:]]*unreachable[[:space:]]*$' <<<"${next_line}"; then - echo "direct diverging call must be immediately followed by LLVM unreachable" >&2 +if grep -Eq 'call void @[^ (]*panic_only[^ (]*\(' <<<"${kernel}"; then + echo "the panic-only callee must not be called: the collector does not emit its body" >&2 exit 1 fi @@ -47,4 +73,32 @@ if ! grep -Eq 'store i32 286331153, ptr ' <<<"${helper}"; then exit 1 fi +# The direct llc route also writes PTX; the NVVM route writes only `.ll` and +# `.target`. Check the PTX only when it exists and is at least as new as the +# IR, so a leftover from an older build is never checked. In the kernel entry +# the sentinel store must survive `opt` and `llc` (the helper is normally +# inlined) or the helper must still be called, and the panic-only control +# must keep its trap. +if [[ -s "${ptx}" && ! "${ptx}" -ot "${llvm}" ]]; then + extract_ptx_function() { + local pattern="$1" + awk -v pattern="${pattern}" ' + /^\.(visible )?\.(entry|func) / && $0 ~ pattern && $0 ~ /\($/ { in_function = 1 } + in_function { print } + in_function && /^}/ { exit } + ' "${ptx}" + } + entry="$(extract_ptx_function '[.]entry diverging_callable[(]')" + test -n "${entry}" + if ! grep -Eq 'st\.global\.b32[[:space:]]+\[[^]]+\], 286331153|call\.uni[[:space:]].*write_then_diverge' <<<"${entry}"; then + echo "the diverging_callable PTX entry neither stores the sentinel nor calls write_then_diverge" >&2 + exit 1 + fi + ptx_trap_count="$(grep -Ec '^[[:space:]]*trap;' <<<"${entry}" || true)" + if [[ "${ptx_trap_count}" -ne 1 ]]; then + echo "the diverging_callable PTX entry must keep one trap for the panic-only callee, found ${ptx_trap_count}" >&2 + exit 1 + fi +fi + echo "function_item_call code shape: PASS" From ec738f8f1304d52c160f45f8b72db399a05fe4c5 Mon Sep 17 00:00:00 2001 From: nihalpasham Date: Mon, 19 Oct 2026 15:38:37 +0530 Subject: [PATCH 4/5] docs(book): describe how diverging device functions lower A device function declared `-> !` with a real body is called normally and the call is followed by `unreachable`, so its work before diverging is kept. Panic entry points, `-> !` bodies that are nothing but a panic, and calls through function pointers still lower to the device trap. State this in the panic-behavior section and in the Panic / Unwinding row of the supported-features table. Signed-off-by: nihalpasham --- cuda-oxide/cuda-oxide-book/appendix/supported-features.md | 2 +- .../gpu-programming/kernels-and-device-functions.md | 8 ++++++++ 2 files changed, 9 insertions(+), 1 deletion(-) diff --git a/cuda-oxide/cuda-oxide-book/appendix/supported-features.md b/cuda-oxide/cuda-oxide-book/appendix/supported-features.md index c038225a85..c8a77a3e00 100644 --- a/cuda-oxide/cuda-oxide-book/appendix/supported-features.md +++ b/cuda-oxide/cuda-oxide-book/appendix/supported-features.md @@ -304,6 +304,6 @@ Anonymous promoted allocations remain unsupported. | Dynamic Dispatch (`dyn Trait`) | **N/A** | Use generics with static dispatch. Haven't found a real need for this. | | Heap Allocation (`Box`, `Vec`) | **N/A** | CUDA has a device-side heap (`malloc`/`free` in kernels), and the compiler allows the `alloc` crate through -- but no device-side `#[global_allocator]` is wired up today. Even if it were, device `malloc` is extremely slow (serialized, fragmented, uncoalesced). Use slices and `SharedArray`. | | `String` / `format_args!` | **N/A** | Use `gpu_printf!` for formatted output. | -| Panic / Unwinding | **N/A** | Panic paths exist in MIR but the compiler strips `core::panicking::*` and all unwind edges. The GPU hardware *can* support unwinding (absolute branches + per-thread call stack tracking post-Volta), but the CUDA toolchain (nvcc/ptxas) doesn't expose it today -- no landing pads survive to PTX. If a panic path is reached at runtime the GPU traps (same as `panic=abort`). NVIDIA has an active project to add C++ exception support to CUDA for automotive safety; the current cuda-oxide design is forward-compatible with that work. Use `gpu_assert!()` + `trap()` for explicit runtime checks today. | +| Panic / Unwinding | **N/A** | Panic paths exist in MIR but the compiler strips `core::panicking::*` and all unwind edges. The GPU hardware *can* support unwinding (absolute branches + per-thread call stack tracking post-Volta), but the CUDA toolchain (nvcc/ptxas) doesn't expose it today -- no landing pads survive to PTX. If a panic path is reached at runtime the GPU traps (same as `panic=abort`). NVIDIA has an active project to add C++ exception support to CUDA for automotive safety; the current cuda-oxide design is forward-compatible with that work. Use `gpu_assert!()` + `trap()` for explicit runtime checks today. A device function declared `-> !` with a real body, such as an error store followed by `loop {}`, is called normally and the call is followed by `unreachable`; only panic entry points, `-> !` bodies that are nothing but a panic, and calls through function pointers lower to the trap. | | Standard Library (`std`/`alloc`) | **N/A** | `std` is forbidden. `alloc` is allowed by the collector but has no backing allocator. Only `core` is fully functional. `Option`, `Result`, iterators all work. | | Texture Memory | **N/A** | Lower priority given TMA availability on Hopper+. | diff --git a/cuda-oxide/cuda-oxide-book/gpu-programming/kernels-and-device-functions.md b/cuda-oxide/cuda-oxide-book/gpu-programming/kernels-and-device-functions.md index 7fee6612fa..30f6965a82 100644 --- a/cuda-oxide/cuda-oxide-book/gpu-programming/kernels-and-device-functions.md +++ b/cuda-oxide/cuda-oxide-book/gpu-programming/kernels-and-device-functions.md @@ -513,6 +513,14 @@ In practice this means: once the call is dropped. Reach for `gpu_printf!` before the check when you need to see the values, or `gpu_assert!` for an explicit check. +A device function declared `-> !` with a real body is different. It never +returns to its caller, for example because it writes an error code and then +runs `loop {}`. The compiler calls it like any other function and places +`unreachable` after the call, so the work it does before diverging is kept. +Three cases still lower to the device trap instead: the panic entry points in +`core::panicking`, a `-> !` body that is nothing but a panic, and a call +through a function pointer. + :::{seealso} The [Error Handling and Debugging](error-handling-and-debugging.md) chapter covers `gpu_printf!`, `gpu_assert!`, and `cargo oxide debug` for diagnosing From 4f11e505e2b3ac6d9e0497daa92b74a8e0cd8415 Mon Sep 17 00:00:00 2001 From: nihalpasham Date: Mon, 19 Oct 2026 15:39:29 +0530 Subject: [PATCH 5/5] docs(book): show how a diverging device call lowers A one-page figure for the book and the pull request, in the style of the tile-libraries direct-linkage figure. Signed-off-by: nihalpasham --- .../images/diverging-call-lowering.svg | 112 ++++++++++++++++++ .../kernels-and-device-functions.md | 11 ++ 2 files changed, 123 insertions(+) create mode 100644 cuda-oxide/cuda-oxide-book/gpu-programming/images/diverging-call-lowering.svg diff --git a/cuda-oxide/cuda-oxide-book/gpu-programming/images/diverging-call-lowering.svg b/cuda-oxide/cuda-oxide-book/gpu-programming/images/diverging-call-lowering.svg new file mode 100644 index 0000000000..9da0b71355 --- /dev/null +++ b/cuda-oxide/cuda-oxide-book/gpu-programming/images/diverging-call-lowering.svg @@ -0,0 +1,112 @@ + + + + + +A device function that never returns is now called, not replaced by a trap +Such a function may store an error code and then loop. The importer replaced the call with trap; and lost the store. Now it emits the call, then unreachable, as rustc's own backend does. +RUST SOURCE (the function_item_call example) +BEFORE, ON MAIN (the call became a trap) +AFTER, WITH #1364 (the call is kept) + +write_then_diverge (src/main.rs) +#[device] +unsafe fn write_then_diverge(ptr: *mut u32) -> ! { + unsafe { + ptr.write(0x1111_1111); + } + loop {} +} +Returns `!`: it never comes back. The store is the +work that has to survive. + +in the #[kernel] diverging_callable +if flag == 3 { + unsafe { + write_then_diverge(out.as_mut_ptr()); + } +} +An ordinary direct call. Nothing follows it: the +callee's type promises it does not return. +Both backends get the same MIR from rustc: a Call with +no return edge (target: None), because the callee +never returns. The middle and right columns differ +only in what the importer emits for that call. + +What the programmer expects +out[0] = 0x1111_1111, then the thread spins forever. +An ordinary call, as on the host, and as nvcc treats +a __noreturn__ device function. + +LLVM IR before opt (@diverging_callable) +bb6: + %v14 = extractvalue { ptr, i64 } %v6, 0 + call void asm sideeffect "trap;", ""() + unreachable +The call to write_then_diverge is gone. bb2 and bb4, +the other two diverging call sites, look the same. + +opt -O2, then llc -march=nvptx64 + +PTX (.entry diverging_callable) +$L__BB1_5: // %bb6 + // begin inline asm + trap; + // end inline asm +trap; appears three times in this kernel, once per +diverging call site. .func write_then_diverge is in +the module, marked .noreturn, and nothing calls it. + +On the GPU, flag == 3: no store. The thread hits trap; +and the launch fails. The value never lands in memory. + +LLVM IR before opt (same block) +bb6: + %v14 = extractvalue { ptr, i64 } %v6, 0 + call void @write_then_diverge(ptr %v14) #0 + unreachable +The call stays. unreachable only records that control +never comes back to this block. + +opt -O2, then llc -march=nvptx64 + +PTX (.entry diverging_callable) +$L__BB0_7: // %bb6 + st.global.b32 [%rd1], 286331153; +$L__BB0_8: // %bb1.i + bra.uni $L__BB0_8; +opt inlined the callee: the store (286331153 is +0x1111_1111), then the self-loop. No trap; on this path. +.func write_then_diverge is kept too, marked .noreturn. + +On the GPU, flag == 3: out[0] = 0x1111_1111, then the +thread spins. Seen on an RTX 5090 with a scratch kernel. + +Still lowered to trap; (by design) +1 panic entry points in core::panicking: + panic!(..), x.unwrap() on None, a[i] out of range +2 a -> ! body that is only a panic; not collected: + fn panic_only() -> ! { panic!("not collected") } +3 a call through a function pointer, not resolved: + let f: unsafe fn(*mut u32) -> ! = helper; f(p) + + + +How the importer decides (translate_call, for a MIR Call with no return edge) +run_pipeline records the mangled name of every function it emits for this build. translate_call resolves the callee and checks that set: +in the set: emit the call, then unreachable, which is what rustc's own LLVM backend emits for a call with no destination. +not in the set: trap; then unreachable, as before. That blanket trap was written for panic entry points, which have no body on the GPU; it swallowed user functions as well. + diff --git a/cuda-oxide/cuda-oxide-book/gpu-programming/kernels-and-device-functions.md b/cuda-oxide/cuda-oxide-book/gpu-programming/kernels-and-device-functions.md index 30f6965a82..01eecd45df 100644 --- a/cuda-oxide/cuda-oxide-book/gpu-programming/kernels-and-device-functions.md +++ b/cuda-oxide/cuda-oxide-book/gpu-programming/kernels-and-device-functions.md @@ -521,6 +521,17 @@ Three cases still lower to the device trap instead: the panic entry points in `core::panicking`, a `-> !` body that is nothing but a panic, and a call through a function pointer. +```{figure} images/diverging-call-lowering.svg +:align: center +:width: 100% + +How a call to a `-> !` device function lowers. Left, the Rust source: `write_then_diverge` stores a +value and then loops. Middle, what the importer used to emit: the call replaced by `trap;`, so the +store never happened. Right, what it emits now: the call, then `unreachable`, so the PTX stores the +value and spins. Panic entry points, panic-only `-> !` bodies and calls through function pointers +still lower to the trap. +``` + :::{seealso} The [Error Handling and Debugging](error-handling-and-debugging.md) chapter covers `gpu_printf!`, `gpu_assert!`, and `cargo oxide debug` for diagnosing