[nihalpasham]

What this adds

The drop_glue example's compile-only probes used #[inline(never)] on their destructors and empty callees to keep each destructor's own MIR in front of the importer's drop no-op proof, and its verify-code-shape.sh expects the optimized kernels to show the pre-drop store followed by a self-loop. That only holds while LLVM inlines those destructors anyway, which is true today because the attribute is dropped on the way to LLVM. #1511 carries it as noinline, so the destructors become real PTX functions and the check fails. The package now switches MIR inlining off instead, the same route disjoint_slice_len and kernel_reference_validity take, and the attributes go.

before (with #1511 merged):  st.global.b32 [%rd1], 286331153;
                             call.uni <LoopingDrop as Drop>::drop, (...);     <- verify-code-shape.sh: FAIL
after:                       st.global.b32 [%rd1], 286331153;
                             $L__BB1_2: bra.uni $L__BB1_2;                    <- PASS, same PTX as main today

# drop_glue/Cargo.toml
cargo-features = ["profile-rustflags"]
[profile.release.package.drop_glue]
rustflags = ["-Zinline-mir=no"]

What we fixed

• The probes no longer depend on an inline attribute being ignored. The importer sees the same destructor MIR as before (the loop, the recursive call, the diamond), and LLVM inlines the destructors because nothing tells it not to.
• The script's and the README's statement that #[inline(never)] is not propagated as noinline is replaced by why the destructors carry no attribute.

Verification

Checked at ae2116b6c (nightly-2026-10-11, RTX 5090, CUDA 13.4):

• ✅ cargo oxide run drop_glue on main with this change: SUCCESS: drop glue wrote sentinel in all 256 elements, all nine PASS lines, and verify-code-shape.sh passes for the two diverging probes and the reconverging control.
• ✅ Same directory, same backend, patch reversed and re-applied: the optimized PTX is instruction-identical for all five kernels; only block-name comments move.
• ✅ With feat(cuda-oxide): carry inline intent through the device pipeline #1511 merged on top of main plus fix(cuda-oxide): keep host-written constant memory through LTO #1509 and fix(cargo-oxide): pin release-like codegen in Cargo profiles #1510 (the tree the full smoketest ran on): drop_glue fails its check without this change (pre-drop store is not followed by an unconditional self-loop, the kernel calls the destructor instead) and passes with it; the cluster, inline_intent and kernel_reference_validity lanes pass alongside it.
• ✅ The example's cargo fmt --check and cargo clippy --all-targets -- -D warnings, scripts/sync-example-locks.sh --check (239 locks agree) and the SPDX guard pass.

Caveats

• The package-scoped flag is the documented form for an example that needs a special rustc flag; a global RUSTFLAGS entry would fork the shared dependency cache in CI.
• Test-only change: no compiler code moves.

[nihalpasham]

/ok to test ae2116b