Skip to content

fix(reborrow): recursive implementation - #159103

Merged
rust-bors[bot] merged 1 commit into
rust-lang:mainfrom
aapoalas:reborrow/fix-borrowed-place
Aug 28, 2026
Merged

fix(reborrow): recursive implementation#159103
rust-bors[bot] merged 1 commit into
rust-lang:mainfrom
aapoalas:reborrow/fix-borrowed-place

Conversation

@aapoalas

@aapoalas aapoalas commented Jul 10, 2026

Copy link
Copy Markdown
Contributor

View all comments

If Reborrow finds '&'a mut T' fields where 'a is the Reborrowed type's first lifetime parameter (currently the only lifetime that is allowed to reborrow) then it inserts a Deref and borrow of the T, and likewise if it finds a 'T: Reborrow' field then the field type is recursed into.

This makes Reborrow always produce the correct borrow checking logic at the cost of most probably being inconsiderately expensive. The thinking is that performance will be a followup consideration.

r? @oli-obk

@rustbot

rustbot commented Jul 10, 2026

Copy link
Copy Markdown
Collaborator

Some changes occurred to MIR optimizations

cc @rust-lang/wg-mir-opt

This PR changes MIR

cc @oli-obk, @RalfJung, @JakobDegen, @vakaras

Some changes occurred to the CTFE machinery

cc @RalfJung, @oli-obk, @lcnr

Some changes occurred to the CTFE / Miri interpreter

cc @rust-lang/miri

Some changes occurred to constck

cc @fee1-dead

This PR changes rustc_public

cc @oli-obk, @celinval, @ouz-a, @makai410

@rustbot rustbot added S-waiting-on-review Status: Awaiting review from the assignee but also interested parties. T-compiler Relevant to the compiler team, which will review and decide on the PR/issue. labels Jul 10, 2026
@rust-log-analyzer

This comment has been minimized.

@rust-log-analyzer

This comment has been minimized.

@rustbot

rustbot commented Jul 10, 2026

Copy link
Copy Markdown
Collaborator

rustc_codegen_cranelift is developed in its own repository. If possible, consider making this change to rust-lang/rustc_codegen_cranelift instead.

cc @bjorn3

clippy is developed in its own repository. If possible, consider making this change to rust-lang/rust-clippy instead.

cc @rust-lang/clippy

@rustbot rustbot added the T-clippy Relevant to the Clippy team. label Jul 10, 2026
@rust-log-analyzer

This comment has been minimized.

@aapoalas
aapoalas force-pushed the reborrow/fix-borrowed-place branch from c74824e to 4ac1ece Compare July 19, 2026 14:32
@rustbot

This comment has been minimized.

@aapoalas
aapoalas force-pushed the reborrow/fix-borrowed-place branch 3 times, most recently from aebdf5d to b845ed2 Compare July 20, 2026 09:37
@rust-bors

This comment has been minimized.

@aapoalas
aapoalas force-pushed the reborrow/fix-borrowed-place branch from b845ed2 to 549d21b Compare July 23, 2026 11:59
@rustbot

This comment has been minimized.

@oli-obk oli-obk left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Started a review a few days ago, but other things came in between, so nothing architectural yet, but docs are always useful

View changes since this review

Comment thread compiler/rustc_middle/src/mir/syntax.rs
Comment thread tests/ui/reborrow/custom_marker_mut_field_borrow.rs
Comment thread tests/ui/reborrow/custom_marker_place_conflict.rs Outdated
Comment thread tests/ui/reborrow/custom_marker_place_conflict_deref.rs Outdated
@rustbot rustbot added S-waiting-on-author Status: This is awaiting some action (such as code changes or more information) from the author. and removed S-waiting-on-review Status: Awaiting review from the assignee but also interested parties. labels Jul 23, 2026
@rustbot

rustbot commented Jul 23, 2026

Copy link
Copy Markdown
Collaborator

Reminder, once the PR becomes ready for a review, use @rustbot ready.

@aapoalas
aapoalas force-pushed the reborrow/fix-borrowed-place branch from 44ad71e to 2c3a3fa Compare July 27, 2026 18:10
@rustbot

This comment has been minimized.

@rust-bors

This comment has been minimized.

@aapoalas
aapoalas force-pushed the reborrow/fix-borrowed-place branch from 2c3a3fa to d815250 Compare August 7, 2026 16:45
@rustbot

This comment has been minimized.

@rust-bors

This comment has been minimized.

@aapoalas
aapoalas force-pushed the reborrow/fix-borrowed-place branch from d815250 to c289b88 Compare August 13, 2026 16:00
@oli-obk oli-obk added S-waiting-on-review Status: Awaiting review from the assignee but also interested parties. and removed S-waiting-on-author Status: This is awaiting some action (such as code changes or more information) from the author. labels Aug 14, 2026

@oli-obk oli-obk left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't fully grok all this yet, but I have enough of an understanding that I'm confident with landing this. It does seem like a coherent thing to me and it will fully "click" in my mind at some point

View changes since this review

Comment thread compiler/rustc_middle/src/mir/syntax.rs Outdated
Comment thread tests/ui/reborrow/custom_marker_place_conflict.rs Outdated
Comment thread compiler/rustc_borrowck/src/borrow_set.rs Outdated
Comment thread compiler/rustc_borrowck/src/type_check/mod.rs Outdated
@rustbot rustbot added S-waiting-on-author Status: This is awaiting some action (such as code changes or more information) from the author. and removed S-waiting-on-review Status: Awaiting review from the assignee but also interested parties. labels Aug 18, 2026
@oli-obk

oli-obk commented Aug 18, 2026

Copy link
Copy Markdown
Contributor

@bors delegate+

r=me with remaining things resolved

@rust-bors

rust-bors Bot commented Aug 18, 2026

Copy link
Copy Markdown
Contributor

✌️ @aapoalas, you can now approve this pull request!

If @oli-obk told you to "r=me" after making some further change, then please make that change and post @bors r=oli-obk.

View changes since this delegation.

@aapoalas

Copy link
Copy Markdown
Contributor Author

@bors squash

@rust-bors

This comment has been minimized.

* Implement Reborrow as a recursive operation

If Reborrow finds '&'a mut T' fields then it inserts
a Deref and borrow of the T, and likewise if it finds a
'T: Reborrow' field then the field type is recursed into.

This makes Reborrow always produce the correct borrow checking logic at
the cost of most probably being inconsiderately expensive. The thinking
is that performance will be a followup consideration.
* PhantomDeref
* Simpler deref test

* Add more PhantomDeref unreachability assertions
* Write out lifetime omission
* Document ProjectionElem::PhantomDeref
* Comment half of reborrow tests
* fix PhantomDeref conflicting with AccessDepth::Shallow
* Recheck CoerceShared in borrowck TypeChecker to ensure its lifetimes make sense
* Fix rebase
* Changes... but where to?
* Typo fix

* Improve comment
* Use fully_perform_op to evaluate CoerceShared trait in borrowck
* More CoerceShared comment
* Comment rest of Reborrow tests

Co-authored-by: Oli Scherer <github35764891676564198441@oli-obk.de>
@rust-bors

rust-bors Bot commented Aug 27, 2026

Copy link
Copy Markdown
Contributor

🔨 16 commits were squashed into 5686e56.

@rust-bors
rust-bors Bot force-pushed the reborrow/fix-borrowed-place branch from 7b76c2d to 5686e56 Compare August 27, 2026 20:04
@aapoalas

Copy link
Copy Markdown
Contributor Author

All of the comments should now be handled, and I've also decided on the next steps:

  1. CoerceShared will lose the type argument and go back to an associated type.
  2. Despite me already deciding the other way around, CoerceShared will not be written as impl<'a: 'b, 'b> CoerceShared for Source<'a> { type Target = Target<'b>; }. This isn't even allowed as 'b is unconstrained. What we'll do instead is that in borrowck we'll project from Source<'a> into type Target = Target<'a> and then do a normal relate_types from the Target<'a> into our assignment target's Target<'b>.
  3. I'll have to look into whether we'll still need a custom TypeRelation for CoerceShared so that silly impls like the following won't cause havoc:
struct Source<'a, 'b> {
  a: &'a mut A,
  b: &'b mut B,
}

struct Target<'a, 'b> {
  a: &'a A,
  b: &'b B,
}

impl<'a, 'b> CoerceShared for Source<'a, 'b> {
  type Target = Target<'b, 'a>; // Note: flipped lifetimes.
}

And if we do need that, do we really still need the projection? I guess we do to catch implementations like:

impl<'a, 'b> CoerceShared for Source<'a, 'b> {
  type Target = Target<'static, 'static>;
}
  1. I believe I'll eventually be adding the following blanket impls:
impl<T: Copy> Reborrow for T {}
impl<T: Copy> CoerceShared for T {
  type Target = T;
}

If the compiler can reason about reborrows of generic T: Reborrow then this'd mean someone could write code that is generic over &mut T and &T where a particular trait (doing actual useful work) is implemented for both. That seems potentially quite useful.

@bors r=oli-obk

@rust-bors

rust-bors Bot commented Aug 27, 2026

Copy link
Copy Markdown
Contributor

📌 Commit 5686e56 has been approved by oli-obk

It is now in the queue for this repository.

@rust-bors rust-bors Bot added S-waiting-on-bors Status: Waiting on bors to run and complete tests. Bors will change the label on completion. and removed S-waiting-on-author Status: This is awaiting some action (such as code changes or more information) from the author. labels Aug 27, 2026
GuillaumeGomez added a commit to GuillaumeGomez/rust that referenced this pull request Aug 27, 2026
…ce, r=oli-obk

fix(reborrow): recursive implementation

If Reborrow finds '&'a mut T' fields where 'a is the Reborrowed type's first lifetime parameter (currently the only lifetime that is allowed to reborrow) then it inserts a Deref and borrow of the T, and likewise if it finds a 'T: Reborrow' field then the field type is recursed into.

This makes Reborrow always produce the correct borrow checking logic at the cost of most probably being inconsiderately expensive. The thinking is that performance will be a followup consideration.

r? @oli-obk
rust-bors Bot pushed a commit that referenced this pull request Aug 28, 2026
Rollup of 14 pull requests

Successful merges:

 - #150075 (Implement clamp_to)
 - #159103 (fix(reborrow): recursive implementation)
 - #160848 (std: avoid aliasing violations when wrapping opaque C types)
 - #161421 (Include startup crt objects on WASI for more outputs)
 - #161805 (Prefer ambiguous candidates when deduplicating traits in scope, so `ambiguous_glob_imported_traits` doesn't depend on import order)
 - #161862 (Put data segment in specified section with link_section on wasm)
 - #161866 (delegation: add tests fixating behavior of delegating to default trait implementations)
 - #161456 (reduce perf impact of scalar size checks)
 - #161666 (Print vendor instructions in `x vendor`)
 - #161730 (Improve type mismatch annotation for lets with block-wrapped initializers)
 - #161828 (Never type after-stabilization cleanup)
 - #161860 (atomicptr.rs test: remove unused import)
 - #161870 (bind to [::1] instead of 127.0.0.1 in documentation examples for v6 UDP methods)
 - #161876 (rustdoc: Correctly handle when a macro generates multiple items in `--generate-macro-expansion`)
JonathanBrouwer added a commit to JonathanBrouwer/rust that referenced this pull request Aug 28, 2026
…ce, r=oli-obk

fix(reborrow): recursive implementation

If Reborrow finds '&'a mut T' fields where 'a is the Reborrowed type's first lifetime parameter (currently the only lifetime that is allowed to reborrow) then it inserts a Deref and borrow of the T, and likewise if it finds a 'T: Reborrow' field then the field type is recursed into.

This makes Reborrow always produce the correct borrow checking logic at the cost of most probably being inconsiderately expensive. The thinking is that performance will be a followup consideration.

r? @oli-obk
JonathanBrouwer added a commit to JonathanBrouwer/rust that referenced this pull request Aug 28, 2026
…ce, r=oli-obk

fix(reborrow): recursive implementation

If Reborrow finds '&'a mut T' fields where 'a is the Reborrowed type's first lifetime parameter (currently the only lifetime that is allowed to reborrow) then it inserts a Deref and borrow of the T, and likewise if it finds a 'T: Reborrow' field then the field type is recursed into.

This makes Reborrow always produce the correct borrow checking logic at the cost of most probably being inconsiderately expensive. The thinking is that performance will be a followup consideration.

r? @oli-obk
rust-bors Bot pushed a commit that referenced this pull request Aug 28, 2026
…uwer

Rollup of 21 pull requests

Successful merges:

 - #158609 (Update sccache to 0.16.0)
 - #150075 (Implement clamp_to)
 - #159103 (fix(reborrow): recursive implementation)
 - #160562 (add target feature ABI checks for SPARC)
 - #160848 (std: avoid aliasing violations when wrapping opaque C types)
 - #161421 (Include startup crt objects on WASI for more outputs)
 - #161805 (Prefer ambiguous candidates when deduplicating traits in scope, so `ambiguous_glob_imported_traits` doesn't depend on import order)
 - #161862 (Put data segment in specified section with link_section on wasm)
 - #161866 (delegation: add tests fixating behavior of delegating to default trait implementations)
 - #161456 (reduce perf impact of scalar size checks)
 - #161528 (Add regression test to ensure optimal compilation)
 - #161666 (Print vendor instructions in `x vendor`)
 - #161730 (Improve type mismatch annotation for lets with block-wrapped initializers)
 - #161828 (Never type after-stabilization cleanup)
 - #161859 (Do not optimize MIR for comptime ConstFns)
 - #161860 (atomicptr.rs test: remove unused import)
 - #161870 (bind to [::1] instead of 127.0.0.1 in documentation examples for v6 UDP methods)
 - #161876 (rustdoc: Correctly handle when a macro generates multiple items in `--generate-macro-expansion`)
 - #161889 (Add link to ownership section in ptr::read docs)
 - #161890 (rustdoc: some clarifying comments)
 - #161891 (Mark `extern_item_impls` feature as incomplete)

Failed merges:

 - #161702 (Use `drop_guard` in some places in {core,alloc,std})
rust-bors Bot pushed a commit that referenced this pull request Aug 28, 2026
…uwer

Rollup of 21 pull requests

Successful merges:

 - #150075 (Implement clamp_to)
 - #159103 (fix(reborrow): recursive implementation)
 - #160562 (add target feature ABI checks for SPARC)
 - #160848 (std: avoid aliasing violations when wrapping opaque C types)
 - #161421 (Include startup crt objects on WASI for more outputs)
 - #161805 (Prefer ambiguous candidates when deduplicating traits in scope, so `ambiguous_glob_imported_traits` doesn't depend on import order)
 - #161862 (Put data segment in specified section with link_section on wasm)
 - #161866 (delegation: add tests fixating behavior of delegating to default trait implementations)
 - #157218 (Track items behind `cfg_select` in the same way we do for `cfg`)
 - #161456 (reduce perf impact of scalar size checks)
 - #161528 (Add regression test to ensure optimal compilation)
 - #161666 (Print vendor instructions in `x vendor`)
 - #161730 (Improve type mismatch annotation for lets with block-wrapped initializers)
 - #161828 (Never type after-stabilization cleanup)
 - #161859 (Do not optimize MIR for comptime ConstFns)
 - #161860 (atomicptr.rs test: remove unused import)
 - #161870 (bind to [::1] instead of 127.0.0.1 in documentation examples for v6 UDP methods)
 - #161876 (rustdoc: Correctly handle when a macro generates multiple items in `--generate-macro-expansion`)
 - #161889 (Add link to ownership section in ptr::read docs)
 - #161890 (rustdoc: some clarifying comments)
 - #161891 (Mark `extern_item_impls` feature as incomplete)

Failed merges:

 - #161702 (Use `drop_guard` in some places in {core,alloc,std})
@rust-bors
rust-bors Bot merged commit 8e3d555 into rust-lang:main Aug 28, 2026
13 checks passed
@rustbot rustbot added this to the 1.100.0 milestone Aug 28, 2026
rust-bors Bot pushed a commit that referenced this pull request Aug 28, 2026
Rollup merge of #159103 - aapoalas:reborrow/fix-borrowed-place, r=oli-obk

fix(reborrow): recursive implementation

If Reborrow finds '&'a mut T' fields where 'a is the Reborrowed type's first lifetime parameter (currently the only lifetime that is allowed to reborrow) then it inserts a Deref and borrow of the T, and likewise if it finds a 'T: Reborrow' field then the field type is recursed into.

This makes Reborrow always produce the correct borrow checking logic at the cost of most probably being inconsiderately expensive. The thinking is that performance will be a followup consideration.

r? @oli-obk
nihalpasham added a commit to NVlabs/cuda-oxide that referenced this pull request Aug 29, 2026
* chore: bump pinned nightly to 2026-08-28; port mir-importer + rustc-codegen-cuda

Move the pinned toolchain from nightly-2026-04-03 to nightly-2026-08-28
(rustc 1.100.0-nightly e457a7b0d, bundled LLVM 23.1.0). The pin is updated
in all seven documented copies (both rust-toolchain.toml files, the
`cargo oxide new` scaffold const, the devcontainer, README/book install
one-liners); scripts/check-toolchain-parity.sh passes.

Mechanical port (rustc_private / rustc_public churn):
- `rustc_middle::mir::mono` -> `rustc_middle::mono` (CodegenUnit, MonoItem)
- `InstanceKind::DropGlue(..)` -> `InstanceKind::Shim(ShimKind::DropGlue(..))`;
  `Instance::resolve_drop_in_place` -> `Instance::resolve_drop_glue`
- `EarlyBinder::bind` now takes `tcx`
- `TyKind::FnDef` args sit behind a `Binder`; both collector sites use
  `no_bound_vars().expect(..)`, matching rustc_monomorphize's own discipline
  (late-bound vars are instantiated before FnDef types reach built MIR)
- `join_codegen` gained `incr_comp_session`/`crate_info` params and returns
  `WorkProductMap`; `codegen_crate` lost `crate_info`; `CompiledModule`
  gained `global_asm_object`
- `fn_sig.abi`/`c_variadic` are methods; `type_of(..).instantiate_identity()`
  returns `Unnormalized<Ty>`, resolved via `normalize_erasing_regions`
  following `TyCtxt::static_ptr_ty`'s discipline
- rustc_public: `Statement.span` -> `.source_info.span`; `ty_with_args`
  moved to the `CrateDefType` trait; `StatementKind::Retag` removed
  (retags are now the `WithRetag` flag on `Rvalue::Use`, still a device
  no-op); `CastKind::BoxDerefTransmute` maps to the Transmute cast kind,
  exactly as rustc_codegen_ssa treats it

Behavioral ports (compiler gaps the new nightly exposed; each verified by
the full on-GPU smoketest):
- drop shims are now built from `core::ptr::drop_glue` (`&mut T` parameter,
  with `drop_in_place` an ordinary forwarding fn): accept the `&mut`
  parameter in emit_drop_glue and elide provably-no-op `drop_glue::<T>`
  calls with the same shared predicate as `drop_in_place`
- collect `ShimKind::FnPtrAsPtr` shim bodies (they back `FnPtr::addr`,
  have no HIR body, but instance_mir synthesises one)
- translate `RigidTy::Foreign` (extern types, e.g. `core::ops::Code`) as an
  opaque empty struct pointee, like FnPtrTarget
- intercept `std::sys::cmath::{asinh,acosh,atanh}{,f}` (pure-Rust formulas
  before this nightly) and lower to `__nv_asinh(f)`/... libdevice calls
- core atomic load/store gained a `const VOLATILE: bool` generic: parse it
  (bool-typed const) separately from orderings and fail closed on
  `VOLATILE = true`
- re-pack `llvm.*.with.overflow` results into the type converter's padded
  `(T, bool)` tuple form; the new RangeInclusive iterator MIR forwards the
  checked-op tuple through a block argument where the two forms met
- atomics example: pass the new `VOLATILE` argument at its raw-intrinsic
  call sites

Also fixes the new-toolchain lint fallout so `clippy -D warnings` stays
green (chunks_exact->as_chunks, ?-operator, useless format!, and an
`allow(suspicious_runtime_symbol_definitions)` for bindgen's c_ulong
malloc/realloc declarations in cuda-bindings).

* feat(codegen): lower Rvalue::Reborrow; add reborrow regression example

nightly-2026-08-28 introduces `Rvalue::Reborrow(Ty, Mutability, Place)`
(rust-lang/rust#159103, `feature(reborrow)`): reborrowing a user ADT that
implements the `Reborrow`/`CoerceShared` marker traits. Semantically it is
a bitwise copy of the place: same type for `Mut`, the same-layout
`CoerceShared` target ADT for `Not`. rustc_codegen_ssa lowers it as
`codegen_operand(Operand::Copy(place))`; the importer mirrors that with a
place read, plus the existing same-layout Transmute cast path when the
translated dialect types diverge.

The release MIR pipeline folds these into plain copies before the importer
sees them, so the new `reborrow` example runs with `--device-debug`
(-Zmir-opt-level=0) in the smoketest, where the rvalue reaches the
translator intact (verified by stubbing the arm: the stub fails the
example under --device-debug and not under release).

* test: re-bless trybuild goldens for nightly-2026-08-28 diagnostics

Rendering-only drift: rustc 1.100 trims more type paths in E0277/E0308
notes ("cuda_device::__LaunchContractDisjointSlice" now prints through its
module path, `LaunchContext<...>` args elide as `_` instead of `...`, and
suggestion placeholders print trimmed types). Re-blessed with
TRYBUILD=overwrite and eyeballed: no test changes semantics, every case
still fails to compile for the same reason as before.

* fix(cargo-oxide): fingerprint the toolchain that built the backend

write_toolchain_fingerprint ran `rustc -vV` in the USER's cwd, but the
backend builds with `current_dir = <backend source clone>`, whose nested
rust-toolchain.toml can pin a different nightly. A bare cargo-oxide
invocation could therefore record a fingerprint that never matches the
`.so`: the StaleVsToolchain guard compares the user-cwd toolchain against
the recorded user-cwd toolchain, matches forever, and every application
build loops on "couldn't load codegen backend" with no self-heal.

Resolve the recorded fingerprint from the build directory (same cwd and
inherited env as the build command, so rustup's proxy resolves the same
rust-toolchain.toml). The mismatch CHECK still resolves from the user's
cwd, which is the toolchain that will load the `.so`.

Tests: command-construction test asserting the fingerprint command's cwd
is the build dir, and a behavioral test proving a build-dir pin (an
uninstallable channel) defeats fingerprinting even though the process cwd
resolves fine.

* chore(intrinsics): refresh recorded LLVM identity to 23.1.0 / 16696adc

The new nightly bundles LLVM 23.1.0 built from rust-lang/llvm-project
16696adcd119e6ba9cc175207d984d7021211acb. Refresh every recorded identity:

- intrinsics/upstream.lock: new LLVM revision, an llvm-tblgen built from
  that revision (sha256 199fe46b...), new dump hashes. Ubuntu
  llvm-tblgen-22 byte-for-byte reproduced the Intrinsics.td dump; the
  NVPTX.td dump uses LLVM-23-only TableGen operators (!sort) that pre-23
  binaries reject, recorded as such in the comparison-tools provenance.
- intrinsics/imported.json: re-extracted (2569 -> 2633 NVVM declarations;
  the LLVM 22-era unscoped atomics are gone, fp16 sat-arith / e2m1 / e2m3
  conversions arrived).
- 58 rust-llvm evidence files renamed to rust-llvm-23.1.0-16696adc* with
  llvm_revision / backend_version ("LLVM version
  23.1.0-rust-1.100.0-nightly") / backend_sha256 (2e5bc688..., the
  rustup llvm-tools llc) headers updated; overlay profile references moved
  with them.
- LLVM 23 upstream drift absorbed in the resolver's recorded facts:
  IntrNoCreateUndefOrPoison joined most pure math intrinsic property sets
  (not ex2/lg2), many declarations moved from NVVMPureIntrinsic to the
  new target-generic PureIntrinsic class, tcgen05 ld/st data types became
  overloaded per-register-count type variables (anonymous_9933..9961,
  previously concrete vNi32; evidence signature records updated in both
  backends' files), the tcgen05 MMA TableGen selection strings moved
  `.ashift` before the collector qualifier (declaration-side spelling
  split from the ptxas-validated emission spelling), tensormap.replace
  dim/stride declarations gained Range<arg1,0,5>, and the fabs/ex2
  overload type variable renumbered anonymous_14 -> anonymous_8.
- `generate` refreshed catalog.json (1016 intrinsics, unchanged set),
  generated-reference.md, probes/*.ll, and the 9 managed generated roots
  (identity banners only); the book stamp moved 90f843cc -> c8007834.
- cuda-target-spec: the LLVM-derived test gate moves 22 -> 23 (it
  self-skips on any other major); all 28 recorded PTX floors and 11 ISA
  spellings verified against llc-23, no table changes.
- llc/opt/llvm-link PATH fallbacks probe the -23 names first; docs and
  the atomics README follow (llc-23 honors atomicrmw orderings + scopes,
  PR llvm/llvm-project#176015, verified: `atomicrmw ... syncscope("block")
  acq_rel` -> `atom.acq_rel.cta.global.add.u32`; fence-splitting stays
  while LLVM 21/22 remain supported).

Gates: `cuda-intrinsics-gen check` clean, `probe --all --skip-terminal
--per-target` passes 2365 target configurations across all 1016 generated
routes against the pinned llc-23, `check-abi-history` appends nothing
(the ABI ledger records no llc identity).

* fix: address nightly-bump review findings (reborrow CoerceShared coverage, fingerprint heal guard, evidence re-attestation)

Four verified review fixes on top of the nightly-2026-08-28 bump:

- reborrow example: cover the Rvalue::Reborrow transmute half. The arm's
  Mutability::Not (CoerceShared) path had zero coverage; the example now
  coerces its MutView into a shared View at two read() call sites and
  writes the sum back (0 -> 2.0 -> 4.0). Proven by stubbing the transmute
  branch: the stub fails the example in BOTH release and --device-debug
  device builds. The write-back happens outside v's scope because
  CoerceShared's coherence rules force the same lifetime argument on
  source and target, so a write through v after read(v) is E0506.

- comment truthfulness (smoketest + example header): GVN folds only the
  Mutability::Mut reborrow variant at mir-opt-level>0; the CoerceShared
  variant survives release MIR and reaches the importer there too. The
  smoketest now runs reborrow in both its release and --device-debug
  invocations.

- cargo-oxide: one-attempt toolchain heal guard. When the user's project
  pin and the backend source's nested pin genuinely differ, the
  StaleVsToolchain arm re-cloned and cold-rebuilt on every invocation
  forever. A heal attempt now records the (active, recorded) fingerprint
  pair in toolchain-heal-attempt.txt; a repeated identical pair prints
  both toolchain identities plus alignment guidance and exits instead of
  rebuilding. A passing fingerprint check deletes the marker. First-time
  mismatches still self-heal via re-clone.

- intrinsics evidence: four families (tcgen05-mma, cluster-sreg,
  tcgen05-control, tcgen05-ldst-offset) carried llc stage records still
  naming LLVM 22.1.2 / ff1dcdc1 while their headers claim llc-23.
  Re-attested all 73 catalog intrinsics of those families against the
  pinned llc-23 (probe green) and updated the stage tool_version /
  tool_sha256 to 23.1.0-rust-1.100.0 / 2e5bc688; regenerated the
  catalog-stamp banners. Also fixed the asinh/acosh/atanh comments:
  only asinh/acosh became std::sys::cmath shims on rustc e457a7b0d;
  atanh stays pure-Rust ln_1p, its entries are defensive.

* feat(host-apis): diagnose host/device TypeId naming divergence on generic kernel misses

Generic kernels are looked up by "<base>_TID_<hex32>", where the backend
computes the hash with tcx.type_id_hash and the host extracts the same
value from core::any::TypeId (cuda-host/src/type_id.rs). The values
cannot drift while TypeId's runtime layout is the raw hash, but that
layout is a std internal: if a future nightly changes it, the transmute
still compiles and every generic launch fails with an opaque
DriverError(500, "named symbol not found").

Make that failure self-explaining, at zero happy-path cost:

- cuda-host entry_registry (new): load_all_ptx_bundles_merged retains
  the merged module's .entry names (parsed with ptx-parse). Retention is
  Weak-keyed by module allocation identity, prunes dropped modules on
  each registration, and costs a few dozen bytes per kernel. Best-effort
  by design: a parse failure retains nothing and never fails a load the
  driver accepted.
- On a failed _TID_ load_function, the macro-generated launch paths now
  consult the diagnosis. If the module holds the same base kernel under
  a different 32-hex hash, it panics with both entry names, the cause
  (host core::intrinsics::type_id vs device tcx.type_id_hash disagreed,
  so the TypeId-to-u128 extraction in cuda_host::type_id is no longer
  value-correct on this toolchain), and the remedy (build host and
  device with the same pinned nightly; re-verify the type_id.rs contract
  and the backend's compute_kernel_export_name together). An ordinary
  miss keeps today's error and panic text unchanged.
- Wired sites: the #[cuda_module] generic function binding routes
  through diagnose_generic_kernel_load_error (returns the DriverError
  untouched for plain misses), and the cuda_launch! /
  cuda_launch_async! generic and closure branches route through
  panic_generic_kernel_load_failed. Non-generic lookups are untouched
  and stay duck-typed over the module expression.
- The load-time canary kernel (layer 2 of the decision memo) is
  deliberately not implemented; the entry_registry module rustdoc
  records why so future readers see it was considered.

Tests: cuda-host unit tests drive the real ptx-parse -> match -> message
path (divergence fires with both names; different-base miss, non-_TID_
name, .func-only same-base entries, and malformed hash suffixes all keep
the ordinary path; registry identity and pruning). cuda-macros expansion
tests assert the generic and closure paths call the helpers and the
non-generic paths do not.

* chore(intrinsics): re-attest scalar-math appended fixture under the pinned llc-23

The scalar_math_inline_sm80_ptx70_appended fixture in
rust-llvm-23.1.0-16696adc-scalar-math.json still carried the llc-22
identity (LLVM 22.1.2-rust-1.96.0 / ff1dcdc1) in its backend_codegen
detail, plus an llc-22-vs-21 behavioral claim that had not been checked
since. Re-attested instead of reworded:

- Re-ran the evidence probes for all 40 scalar_math catalog intrinsics
  against the pinned sysroot llc-23 (LLVM 23.1.0-rust-1.100.0-nightly,
  SHA-256 2e5bc688...), terminal ptxas (CUDA 13.3 V13.3.33)
  revalidation included: 40/40 green.
- The three appended per-op PTX SHA-256s (ex2_approx_f32,
  ex2_approx_ftz_f32, tanh_approx_f32) are byte-identical under llc-23
  because inline-PTX probe bodies pass through llc verbatim; the detail
  now states that and keeps the original llc-22 identity as the
  historical reference.
- The typed llvm.nvvm.tanh.approx.f32 route was re-checked with current
  binaries: pinned llc-23 and Ubuntu LLVM 22.1.7 select it via
  NVVMIntrinsic-class matching, while Ubuntu LLVM 21.1.8 still lowers it
  to an extern funcall, so the inline route remains required while
  llc-21 remains a supported fallback (doctor probes llc-23/22/21).
- Added the structured llc tool identity fields to the stage, matching
  the convention of the four families re-attested in 0aec1b7.
- `generate` refreshed catalog.json, generated-reference.md,
  probes/*.ll, and the generated-root banners (catalog SHA-256
  478f564b -> 49bcd626); `check` reports outputs current.

* fix: clippy-clean examples under nightly-2026-08-28; re-stamp book catalog page

The bumped nightly's clippy trips two examples out of 222 lint targets:

- field_array_assign: needless_range_loop on the runtime-index loop. The
  indexed loop is the (Field, Index) MIR projection under test (issue
  #235), so keep it and allow the lint on the loop with a reason comment.
- function_item_call: stable_features on #![feature(never_type)], which
  stabilized in 1.100.0-nightly. Drop the dead gate; the never-type
  usage under test is unchanged.

Also carries the cuda-oxide-book catalog-stamp refresh (c8007834 ->
49bcd626) matching the regenerated dialect-nvvm ops/generated headers.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

S-waiting-on-bors Status: Waiting on bors to run and complete tests. Bors will change the label on completion. T-clippy Relevant to the Clippy team. T-compiler Relevant to the compiler team, which will review and decide on the PR/issue.

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants