fix(mir-importer): preserve emitted diverging device calls

[YunusMutlu]

What this adds

A direct call to a #[device] fn ... -> ! was lowered to trap; by the importer, even when the callee's body had been compiled, so whatever the function did before diverging (an error store, say) was lost. The importer now does what rustc's own backend does: emit the call, then unreachable. Only panic entry points, -> ! bodies that are nothing but a panic, and calls through function pointers still lower to the trap.

#[device] unsafe fn write_then_diverge(p: *mut u32) -> ! { p.write(0x1111_1111); loop {} }

before: trap;                                           // the store never happened
after:  st.global.b32 [%rd1], 286331153; bra.uni <self>; // the store lands, the thread spins

What we fixed

• Every emitted diverging callee is called. The trap stays only where there is nothing to call.
• The function_item_call check covers all three diverging call sites, reads the PTX entry when one exists, accepts --device-debug output, and has a retained-trap control (panic_only).
• facts.rs has a unit test for the emitted-function set (install, check, replace, gone).
• The book and the supported-features row say which -> ! calls still trap.

Implementation

• translate_call: for a MIR Call with no return edge, resolve the callee and check the pipeline's emitted set. In the set means call then unreachable; otherwise trap; then unreachable, as before.
• is_emitted_function delegates to a name-based is_emitted_name.
• The example gains a flag == 4 arm calling a panic-only -> ! helper, which the collector skips, so its call site must still trap.

Verification

Checked at 4f11e505e (nightly-2026-10-11, RTX 5090):

• ✅ cargo test -p mir-importer --all-targets, clippy and rustdoc with warnings as errors, fmt, the example's fmt and clippy gates, the 18 guards and the lock check: clean.
• ✅ Fails before, passes after: the example built with main's backend traps all three sites and the check exits 1; this branch keeps three call + unreachable pairs, and the PTX stores the sentinel before its self-loop with no trap; on that path.
• ✅ cargo oxide run function_item_call: PASS with memcheck clean, through the llc route and the libNVVM route. device_panic_trap still traps out of range.
• ✅ The diverging path ran for real in a scratch kernel: the callee's store became visible to the host while the thread spun.

Caveats

• A -> ! body that stores and then panics within one or two blocks is still skipped by the collector's placeholder heuristic, so that store is lost before the trap. Separate issue.
• Calls through function pointers to diverging functions still trap silently. An up-front diagnostic is a separate change.

Closes #1363.

[copy-pr-bot]

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

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

[nihalpasham]

/ok to test 4f11e50

[nihalpasham]

Thanks, this is the right fix and it matches what rustc's own backend does: emit the call, then unreachable when there is no destination. On top:

• The code-shape check now covers all three call sites (on main the two apply_never calls trapped too), reads the PTX entry, accepts --device-debug output, and has a retained-trap control: a -> ! helper whose body is only a panic is not collected, and its call site must still lower to trap;.
• A unit test for the emitted-function set, and one book paragraph plus a figure naming the cases that still trap.

At 4f11e505e the example passes on sm_120a through both routes, the same source built with main's backend fails the check, and all CI checks passed. The diverging path also ran for real: the callee's store became visible to the host while the thread spun. That takes care of #1363.