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..db134cb78d 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,28 @@ 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 { + 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 /// contract using rustc's type system. MIR locals have erased regions, so the /// outer lifetime is read from the declaration before instantiation instead. @@ -1104,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/helpers.rs b/cuda-oxide/crates/mir-importer/src/translator/terminator/helpers.rs index ea76a12bfd..301415d0a4 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 911a3a4c23..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,25 +1254,47 @@ 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. 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. // - // 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. + // 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() { - // 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..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 @@ -7,3 +7,20 @@ 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` 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 68b442f644..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 @@ -10,8 +10,11 @@ //! 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 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. @@ -57,6 +60,30 @@ 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 {} +} + +// 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) @@ -133,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(); @@ -148,6 +176,19 @@ 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 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 new file mode 100755 index 0000000000..a866ee8104 --- /dev/null +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/verify-code-shape.sh @@ -0,0 +1,104 @@ +#!/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" +ptx="${root}/function_item_call.ptx" + +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}" + +# 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 +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 + +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 + +# 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" diff --git a/cuda-oxide/crates/rustc-codegen-cuda/src/collector.rs b/cuda-oxide/crates/rustc-codegen-cuda/src/collector.rs index 4fcbf51b3e..97e5b7c45a 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/src/collector.rs +++ b/cuda-oxide/crates/rustc-codegen-cuda/src/collector.rs @@ -2695,10 +2695,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>, @@ -2778,8 +2777,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; } @@ -2808,7 +2808,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() } 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/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 7fee6612fa..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 @@ -513,6 +513,25 @@ 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. + +```{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