From 770c4e8f3217c3b1e0a65cb5291637bb2b43d1b8 Mon Sep 17 00:00:00 2001
From: Zihua Wu <13583761+lucifer1004@users.noreply.github.com>
Date: Tue, 29 Sep 2026 17:39:15 -0700
Subject: [PATCH] feat(finalizer): fail a link that drops an expected kernel

nvJitLink can report success while dropping a module's kernels: an
unresolved device-runtime symbol makes it discard the module, and the
link-time optimizer strips a kernel it finds unreachable. The result is a
well-formed but empty image, which is_valid_cubin accepts.

Every finalizer link now takes the kernels its output must define and
checks the output's entry points after the link: cubin entries from the
ELF symbol table (FUNC symbols with STO_CUDA_ENTRY), PTX entries from
ptx-parse. A missing kernel is FinalizerError::MissingKernels, and an
image whose entries cannot be read is UnreadableEntryInventory. Rooting
stays with the @llvm.used the exporter already emits; nvJitLink's
-kernels-used is not passed, since it also lets the link drop kernels it
does not list.

The expected kernels reach every finalization route:
- build-time materialization: the collector's kernel export names;
- embedded bundles finalized at load: the bundle's Kernel entries;
- sidecar files (cargo oxide interop, cuda-host's file loader): a new
  <module>.kernels sidecar the compiler writes next to the .ll, before the
  completing .target. An artifact the compiler emitted must have it; a
  manual .ll is checked against one only when present.

Cached cubins are checked on a hit too, since a stored image may predate
the check. The cuda-host build/link functions and the finalizer entry
points take the expected kernels as a new argument.

Tests: synthetic empty and partial cubins, PTX with only device
functions, unreadable inventories, the sidecar format, and a live link
whose unrooted kernel nvJitLink drops (cubin and PTX); the live file-cache
test now roots its kernel, which it previously linked away.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com>
---
 .gitignore                                    |   1 +
 Cargo.lock                                    |   2 +
 crates/cargo-oxide/src/commands/artifacts.rs  |   1 +
 crates/cargo-oxide/src/commands/interop.rs    |  23 +-
 crates/cuda-artifact-finalizer/Cargo.toml     |   1 +
 crates/cuda-artifact-finalizer/src/entries.rs | 291 ++++++++++++++++++
 crates/cuda-artifact-finalizer/src/lib.rs     |  85 ++++-
 crates/cuda-artifact-finalizer/src/link.rs    |  34 +-
 crates/cuda-host/examples/finalize_nvvm_ir.rs |   3 +-
 crates/cuda-host/src/embedded.rs              |  12 +
 crates/cuda-host/src/ltoir.rs                 | 148 ++++++++-
 crates/mir-importer/Cargo.toml                |   1 +
 crates/mir-importer/src/pipeline.rs           |  25 ++
 crates/rustc-codegen-cuda/Cargo.lock          |   2 +
 .../examples/abi_hmm/Cargo.lock               |   1 +
 .../examples/addressof_sharedarray/Cargo.lock |   1 +
 .../examples/aligned_field_loads/Cargo.lock   |   1 +
 .../examples/aligned_field_stores/Cargo.lock  |   1 +
 .../examples/array_constants/Cargo.lock       |   1 +
 .../examples/array_for_loop/Cargo.lock        |   1 +
 .../examples/array_index/Cargo.lock           |   1 +
 .../examples/array_init/Cargo.lock            |   1 +
 .../examples/asin_acos_smoke/Cargo.lock       |   1 +
 .../assume_init_uninhabited_trap/Cargo.lock   |   1 +
 .../examples/async_mlp/Cargo.lock             |   1 +
 .../examples/async_vecadd/Cargo.lock          |   1 +
 .../examples/atomic_f16/Cargo.lock            |   1 +
 .../examples/atomics/Cargo.lock               |   1 +
 .../examples/barrier/Cargo.lock               |   1 +
 .../examples/bf16x2_arith/Cargo.lock          |   1 +
 .../examples/bf16x2_fma/Cargo.lock            |   1 +
 .../examples/bool_phi_cmp/Cargo.lock          |   1 +
 .../examples/carrying_mul_add/Cargo.lock      |   1 +
 .../examples/cast_chunks/Cargo.lock           |   1 +
 .../examples/cast_tests/Cargo.lock            |   1 +
 .../examples/cbrt_smoke/Cargo.lock            |   1 +
 .../examples/checked_arith/Cargo.lock         |   1 +
 .../examples/clc/Cargo.lock                   |   1 +
 .../examples/cluster/Cargo.lock               |   1 +
 .../examples/compiler_features/Cargo.lock     |   1 +
 .../examples/complex_mul/Cargo.lock           |   1 +
 .../examples/const_aggregate/Cargo.lock       |   1 +
 .../const_bool_dead_branch/Cargo.lock         |   1 +
 .../examples/const_generic/Cargo.lock         |   1 +
 .../examples/const_table_lookup/Cargo.lock    |   1 +
 .../const_tuple_table_lookup/Cargo.lock       |   1 +
 .../constant_index_from_end/Cargo.lock        |   1 +
 .../examples/constant_memory/Cargo.lock       |   1 +
 .../constant_memory_coeffs/Cargo.lock         |   1 +
 .../constant_memory_simple/Cargo.lock         |   1 +
 .../examples/constant_memory_table/Cargo.lock |   1 +
 .../examples/coop_groups_demo/Cargo.lock      |   1 +
 .../examples/copy_aggregate_borrow/Cargo.lock |   1 +
 .../examples/copy_nonoverlapping/Cargo.lock   |   1 +
 .../examples/counted_barrier/Cargo.lock       |   1 +
 .../examples/cp_async_small/Cargo.lock        |   1 +
 .../examples/cp_async_zfill/Cargo.lock        |   1 +
 .../cpp_consumes_rust_device/Cargo.lock       |   1 +
 .../examples/cross_crate_embedded/Cargo.lock  |   1 +
 .../examples/cross_crate_kernel/Cargo.lock    |   1 +
 .../cross_crate_merged_symbols/Cargo.lock     |   1 +
 .../examples/cross_crate_recursion/Cargo.lock |   1 +
 .../examples/cuda_module_contract/Cargo.lock  |   1 +
 .../examples/cuda_module_in_lib/Cargo.lock    |   1 +
 .../examples/cuda_module_nested/Cargo.lock    |   1 +
 .../cutile_inter_kernel/simt/Cargo.lock       |   1 +
 .../examples/cvt_f16x2/Cargo.lock             |   1 +
 .../examples/cvt_packed/Cargo.lock            |   1 +
 .../examples/debug-tests/Cargo.lock           |   1 +
 .../examples/debug/Cargo.lock                 |   1 +
 .../examples/debug_pointer_locals/Cargo.lock  |   1 +
 .../examples/device_closures/Cargo.lock       |   1 +
 .../examples/device_ffi_test/Cargo.lock       |   1 +
 .../examples/device_global/Cargo.lock         |   1 +
 .../examples/device_panic_trap/Cargo.lock     |   1 +
 .../disjoint_from_raw_parts/Cargo.lock        |   1 +
 .../examples/disjoint_slice_len/Cargo.lock    |   1 +
 .../examples/dotprod/Cargo.lock               |   1 +
 .../examples/dpx_minmax_chains/Cargo.lock     |   1 +
 .../examples/drop_glue/Cargo.lock             |   1 +
 .../examples/dst_metadata/Cargo.lock          |   1 +
 .../examples/dynamic_smem/Cargo.lock          |   1 +
 .../examples/elect_leader/Cargo.lock          |   1 +
 .../enum_abi_multi_payload/Cargo.lock         |   1 +
 .../examples/enum_array_match/Cargo.lock      |   1 +
 .../enum_constant_provenance/Cargo.lock       |   1 +
 .../examples/enum_payload_addr/Cargo.lock     |   1 +
 .../enum_state_transitions/Cargo.lock         |   1 +
 .../examples/enum_table_lookup/Cargo.lock     |   1 +
 .../examples/error/Cargo.lock                 |   1 +
 .../error_enum_bool_payload_addr/Cargo.lock   |   1 +
 .../error_enum_pointer_overlap/Cargo.lock     |   1 +
 .../Cargo.lock                                |   1 +
 .../error_generated_intrinsic_abi/Cargo.lock  |   1 +
 .../Cargo.lock                                |   1 +
 .../Cargo.lock                                |   1 +
 .../Cargo.lock                                |   1 +
 .../examples/error_heap_alloc/Cargo.lock      |   1 +
 .../error_host_arch_intrinsic/Cargo.lock      |   1 +
 .../error_host_target_feature/Cargo.lock      |   1 +
 .../error_kernel_shared_param/Cargo.lock      |   1 +
 .../error_missing_device_attr/Cargo.lock      |   1 +
 .../Cargo.lock                                |   1 +
 .../examples/ex2_approx_f16/Cargo.lock        |   1 +
 .../examples/exact_div_chunks/Cargo.lock      |   1 +
 .../examples/export_name_policy/Cargo.lock    |   1 +
 .../examples/extern_crate_closure/Cargo.lock  |   1 +
 .../examples/extern_libdevice/Cargo.lock      |   1 +
 .../examples/f16_stress/Cargo.lock            |   1 +
 .../examples/f16x2_arith/Cargo.lock           |   1 +
 .../examples/f32x2_arith/Cargo.lock           |   1 +
 .../examples/fcmp_ne_nan_smoke/Cargo.lock     |   1 +
 .../examples/field_array_assign/Cargo.lock    |   1 +
 .../examples/float_rounding/Cargo.lock        |   1 +
 .../examples/fmaxmin_smoke/Cargo.lock         |   1 +
 .../examples/fn_ptr_naming/Cargo.lock         |   1 +
 .../fp_special_literal_smoke/Cargo.lock       |   1 +
 .../examples/function_item_call/Cargo.lock    |   1 +
 .../examples/future_apis/Cargo.lock           |   1 +
 .../examples/gemm/Cargo.lock                  |   1 +
 .../examples/gemm_sol/Cargo.lock              |   1 +
 .../examples/gemm_sol_final/Cargo.lock        |   1 +
 .../examples/gemm_views/Cargo.lock            |   1 +
 .../examples/generated_intrinsics/Cargo.lock  |   1 +
 .../generated_intrinsics_blackwell/Cargo.lock |   1 +
 .../examples/generated_ldmatrix/Cargo.lock    |   1 +
 .../examples/generic/Cargo.lock               |   1 +
 .../examples/hashmap/Cargo.lock               |   1 +
 .../examples/hashmap_v2/Cargo.lock            |   1 +
 .../examples/hashmap_v3/Cargo.lock            |   1 +
 .../examples/hello_constant/Cargo.lock        |   1 +
 .../examples/helper_fn/Cargo.lock             |   1 +
 .../examples/host_closure/Cargo.lock          |   1 +
 .../host_closure_target_feature/Cargo.lock    |   1 +
 .../examples/iket_trace/Cargo.lock            |   1 +
 .../examples/image_convolution/Cargo.lock     |   1 +
 .../examples/index2d_const/Cargo.lock         |   1 +
 .../examples/index_field_assign/Cargo.lock    |   1 +
 .../examples/inline_ptx/Cargo.lock            |   1 +
 .../examples/integer_minmax/Cargo.lock        |   1 +
 .../examples/interior_mutability/Cargo.lock   |   1 +
 .../interop_cubin_identity/device/Cargo.lock  |   1 +
 .../kernel_reference_validity/Cargo.lock      |   1 +
 .../examples/lanemask_scan/Cargo.lock         |   1 +
 .../examples/legacy_atomic_fadd/Cargo.lock    |   1 +
 .../examples/legacy_atomic_rmw_cas/Cargo.lock |   1 +
 .../legacy_nvvm_pointer_shapes/Cargo.lock     |   1 +
 .../examples/libdevice_math/Cargo.lock        |   1 +
 .../examples/libm_math/Cargo.lock             |   1 +
 .../examples/manual_launch_generic/Cargo.lock |   1 +
 .../manual_launch_libdevice/Cargo.lock        |   1 +
 .../examples/map_sum/Cargo.lock               |   1 +
 .../examples/math_atan/Cargo.lock             |   1 +
 .../examples/math_hyperbolic/Cargo.lock       |   1 +
 .../examples/math_no_mir_inline/Cargo.lock    |   1 +
 .../examples/math_tan/Cargo.lock              |   1 +
 .../examples/mathdx_ffi_test/Cargo.lock       |   1 +
 .../examples/mcast_barrier_test/Cargo.lock    |   1 +
 .../examples/mem_swap/Cargo.lock              |   1 +
 .../mir_projected_call_destination/Cargo.lock |   1 +
 .../examples/mma_mxf8f6f4/Cargo.lock          |   1 +
 .../examples/ne_bytes_transmute/Cargo.lock    |   1 +
 .../nested_index_assignment/Cargo.lock        |   1 +
 .../examples/numeric_stress/Cargo.lock        |   1 +
 .../examples/nvls_all_reduce/Cargo.lock       |   1 +
 .../examples/option_ref_transmute/Cargo.lock  |   1 +
 .../examples/option_ref_unwrap_or/Cargo.lock  |   1 +
 .../examples/option_unwrap_trap/Cargo.lock    |   1 +
 .../examples/ord_cmp/Cargo.lock               |   1 +
 .../examples/packed_aggregate_abi/Cargo.lock  |   1 +
 .../examples/packed_atomic_add/Cargo.lock     |   1 +
 .../examples/padded_constants/Cargo.lock      |   1 +
 .../examples/partial_warp_reduce/Cargo.lock   |   1 +
 .../examples/pinned_overlap/Cargo.lock        |   1 +
 .../examples/place_read_fallbacks/Cargo.lock  |   1 +
 .../examples/policy_config/Cargo.lock         |   1 +
 .../examples/primitive_stress/Cargo.lock      |   1 +
 .../examples/printf/Cargo.lock                |   1 +
 .../examples/projected_place_read/Cargo.lock  |   1 +
 .../examples/proof_carrying_views/Cargo.lock  |   1 +
 .../examples/ptr_copy/Cargo.lock              |   1 +
 .../examples/ptr_offset_from/Cargo.lock       |   1 +
 .../ptr_without_provenance/Cargo.lock         |   1 +
 .../examples/reborrow/Cargo.lock              |   1 +
 .../examples/redux_f32/Cargo.lock             |   1 +
 .../examples/redux_minmax/Cargo.lock          |   1 +
 .../examples/redux_sum/Cargo.lock             |   1 +
 .../examples/ref_index_projections/Cargo.lock |   1 +
 .../examples/ref_operand_mul/Cargo.lock       |   1 +
 .../examples/repr_align_block_args/Cargo.lock |   1 +
 .../examples/repr_transparent_abi/Cargo.lock  |   1 +
 .../examples/repr_u32_enum_stride/Cargo.lock  |   1 +
 .../examples/row_width_slice/Cargo.lock       |   1 +
 .../examples/rustlantis-smoke/Cargo.lock      |   1 +
 .../scoped_atomic_load_store/Cargo.lock       |   1 +
 .../examples/select_unpredictable/Cargo.lock  |   1 +
 .../examples/set_discriminant/Cargo.lock      |   1 +
 .../set_discriminant_niche/Cargo.lock         |   1 +
 .../examples/shared_debug/Cargo.lock          |   1 +
 .../examples/shared_memory_helper/Cargo.lock  |   1 +
 .../examples/shared_ptr_transmute/Cargo.lock  |   1 +
 .../examples/shared_slice_unsize/Cargo.lock   |   1 +
 .../examples/shared_stage_select/Cargo.lock   |   1 +
 .../examples/sharedmem/Cargo.lock             |   1 +
 .../examples/shuffle_64/Cargo.lock            |   1 +
 .../slice_from_raw_parts_cast/Cargo.lock      |   1 +
 .../examples/slice_get_mut/Cargo.lock         |   1 +
 .../slice_index_bounds_check/Cargo.lock       |   1 +
 .../examples/slice_iterators/Cargo.lock       |   1 +
 .../examples/slice_reslice/Cargo.lock         |   1 +
 .../examples/small_type_ffi_test/Cargo.lock   |   1 +
 .../examples/sparse_mma_fp8/Cargo.lock        |   1 +
 .../sparse_mma_ordered_float/Cargo.lock       |   1 +
 .../examples/split_at_mut/Cargo.lock          |   1 +
 .../examples/standalone_device_fn/Cargo.lock  |   1 +
 .../examples/static_slice_addend/Cargo.lock   |   1 +
 .../examples/static_slice_unsize/Cargo.lock   |   1 +
 .../examples/step_by/Cargo.lock               |   1 +
 .../struct_constant_provenance/Cargo.lock     |   1 +
 .../examples/struct_field_layout/Cargo.lock   |   1 +
 .../examples/subslice_projection/Cargo.lock   |   1 +
 .../examples/switchint_128bit/Cargo.lock      |   1 +
 .../examples/swizzle_smem/Cargo.lock          |   1 +
 .../examples/tcgen05/Cargo.lock               |   1 +
 .../examples/tcgen05_matmul/Cargo.lock        |   1 +
 .../examples/thread_runs/Cargo.lock           |   1 +
 .../examples/tiled_gemm/Cargo.lock            |   1 +
 .../examples/tma_copy/Cargo.lock              |   1 +
 .../examples/tma_multicast/Cargo.lock         |   1 +
 .../toolchain_pairing_smoke/Cargo.lock        |   1 +
 .../transmute_scalar_struct/Cargo.lock        |   1 +
 .../tuple_array_provenance/Cargo.lock         |   1 +
 .../tuple_constant_provenance/Cargo.lock      |   1 +
 .../examples/tuple_field_lookup/Cargo.lock    |   1 +
 .../examples/tuple_return/Cargo.lock          |   1 +
 .../examples/unaligned_memory/Cargo.lock      |   1 +
 .../examples/unchecked_indexing/Cargo.lock    |   1 +
 .../examples/uninit_const/Cargo.lock          |   1 +
 .../examples/union_aggregate/Cargo.lock       |   1 +
 .../examples/unroll_bounds_check/Cargo.lock   |   1 +
 .../examples/unroll_smoke/Cargo.lock          |   1 +
 .../examples/vecadd/Cargo.lock                |   1 +
 .../examples/vectorization/Cargo.lock         |   1 +
 .../examples/volatile_memory/Cargo.lock       |   1 +
 .../examples/warp_reduce/Cargo.lock           |   1 +
 .../examples/warp_sums/Cargo.lock             |   1 +
 .../examples/wgmma/Cargo.lock                 |   1 +
 .../examples/wgmma_mma_bf16/Cargo.lock        |   1 +
 crates/rustc-codegen-cuda/src/lib.rs          |   9 +
 crates/rustc-codegen-cuda/src/materialize.rs  |  10 +-
 .../error-handling-and-debugging.md           |  10 +-
 251 files changed, 859 insertions(+), 33 deletions(-)
 create mode 100644 crates/cuda-artifact-finalizer/src/entries.rs

diff --git a/.gitignore b/.gitignore
index d87b0336a5..5386f6f41c 100644
--- a/.gitignore
+++ b/.gitignore
@@ -7,6 +7,7 @@ target/
 *.ltoir
 *.target
 *.options
+*.kernels
 .ltoir_arch
 *.ll
 *.bc
diff --git a/Cargo.lock b/Cargo.lock
index b3bdf7fc4b..714ff4443b 100644
--- a/Cargo.lock
+++ b/Cargo.lock
@@ -364,6 +364,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "serde_json",
  "sha2",
@@ -994,6 +995,7 @@ name = "mir-importer"
 version = "0.2.1"
 dependencies = [
  "combine",
+ "cuda-artifact-finalizer",
  "cuda-oxide-codegen",
  "dialect-iket",
  "dialect-mir",
diff --git a/crates/cargo-oxide/src/commands/artifacts.rs b/crates/cargo-oxide/src/commands/artifacts.rs
index 85ddd2eaf1..df5b2eeb66 100644
--- a/crates/cargo-oxide/src/commands/artifacts.rs
+++ b/crates/cargo-oxide/src/commands/artifacts.rs
@@ -111,6 +111,7 @@ pub(super) const GENERATED_ARTIFACT_SUFFIXES: &[&str] = &[
     "ptx.identity",
     "target",
     "options",
+    "kernels",
     "cubin.target",
 ];
 
diff --git a/crates/cargo-oxide/src/commands/interop.rs b/crates/cargo-oxide/src/commands/interop.rs
index 13d8c532ec..d48392de1a 100644
--- a/crates/cargo-oxide/src/commands/interop.rs
+++ b/crates/cargo-oxide/src/commands/interop.rs
@@ -477,6 +477,27 @@ fn finalize_interop_device_artifact(
                         );
                         std::process::exit(1);
                     });
+            let kernels_path = ir_path.with_extension("kernels");
+            let kernels_text = std::fs::read_to_string(&kernels_path).unwrap_or_else(|error| {
+                eprintln!(
+                    "Error: could not read emitted kernel list at {}: {error}",
+                    kernels_path.display()
+                );
+                std::process::exit(1);
+            });
+            let expected_kernels =
+                cuda_artifact_finalizer::parse_expected_kernels_sidecar(&kernels_text)
+                    .unwrap_or_else(|| {
+                        eprintln!(
+                            "Error: invalid emitted kernel list at {}",
+                            kernels_path.display()
+                        );
+                        std::process::exit(1);
+                    });
+            let expected_kernels = expected_kernels
+                .iter()
+                .map(String::as_str)
+                .collect::<Vec<_>>();
             let finalizer = cuda_artifact_finalizer::Finalizer::discover().unwrap_or_else(|error| {
                 eprintln!("Error: could not initialize the CUDA artifact finalizer: {error}");
                 eprintln!(
@@ -486,7 +507,7 @@ fn finalize_interop_device_artifact(
             });
             let options = finalization_options_from_artifact(&target, compile_options);
             let cubin = finalizer
-                .materialize_nvvm_ir(artifact_name, &ir, &options)
+                .materialize_nvvm_ir(artifact_name, &ir, &options, &expected_kernels)
                 .unwrap_or_else(|error| {
                     eprintln!(
                         "Error: could not finalize {} for {}: {error}",
diff --git a/crates/cuda-artifact-finalizer/Cargo.toml b/crates/cuda-artifact-finalizer/Cargo.toml
index 8049890dd1..838f431489 100644
--- a/crates/cuda-artifact-finalizer/Cargo.toml
+++ b/crates/cuda-artifact-finalizer/Cargo.toml
@@ -11,6 +11,7 @@ repository.workspace = true
 [dependencies]
 libnvvm-sys = { workspace = true }
 nvjitlink-sys = { workspace = true }
+ptx-parse = { workspace = true }
 serde = { workspace = true }
 sha2 = { workspace = true }
 thiserror = { workspace = true }
diff --git a/crates/cuda-artifact-finalizer/src/entries.rs b/crates/cuda-artifact-finalizer/src/entries.rs
new file mode 100644
index 0000000000..1320db3ed2
--- /dev/null
+++ b/crates/cuda-artifact-finalizer/src/entries.rs
@@ -0,0 +1,291 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+//! Launchable entry inventory of a finalized image.
+//!
+//! nvJitLink can succeed while dropping a module's kernels (an unresolved
+//! device-runtime symbol, or a kernel the link-time optimizer found
+//! unreachable), returning a well-formed but empty image. The finalizer
+//! therefore checks that every kernel its caller expects is an entry point of
+//! the output, and fails otherwise.
+
+use crate::FinalizerError;
+use crate::options::FinalizerOutput;
+
+const SECTION_TYPE_SYMBOL_TABLE: u32 = 2;
+const ELF64_SECTION_HEADER_LENGTH: usize = 64;
+const ELF64_SYMBOL_LENGTH: usize = 24;
+const SYMBOL_TYPE_FUNCTION: u8 = 2;
+/// `st_other` flag CUDA sets on the symbol of a launchable kernel.
+const STO_CUDA_ENTRY: u8 = 0x10;
+
+/// Fail unless every name in `expected` is an entry point of `image`.
+///
+/// Every finalizer link applies this check. Callers that reuse a stored image
+/// (a cache hit) apply it themselves. An empty `expected` checks nothing: the
+/// module declares no kernels.
+pub fn require_expected_kernels(
+    image: &[u8],
+    output: FinalizerOutput,
+    expected: &[&str],
+) -> Result<(), FinalizerError> {
+    if expected.is_empty() {
+        return Ok(());
+    }
+    let entries = match output {
+        FinalizerOutput::Cubin => cubin_entry_names(image),
+        FinalizerOutput::Ptx => ptx_entry_names(image),
+    }
+    .ok_or(FinalizerError::UnreadableEntryInventory { output })?;
+    let missing = expected
+        .iter()
+        .filter(|name| !entries.iter().any(|entry| entry == *name))
+        .map(|name| (*name).to_string())
+        .collect::<Vec<_>>();
+    if missing.is_empty() {
+        Ok(())
+    } else {
+        Err(FinalizerError::MissingKernels {
+            output,
+            missing,
+            present: entries.len(),
+        })
+    }
+}
+
+/// First line of a `<module>.kernels` sidecar. The sidecar lists, one export
+/// name per line, the kernels a module's finalized image must define; the
+/// compiler writes it next to the NVVM IR so that finalizing from files checks
+/// the same kernels as finalizing an embedded bundle.
+pub const EXPECTED_KERNELS_SIDECAR_HEADER: &str = "cuda-oxide expected-kernels v1";
+
+/// Text of a `<module>.kernels` sidecar listing `kernels`.
+pub fn expected_kernels_sidecar_text<'a>(kernels: impl IntoIterator<Item = &'a str>) -> String {
+    let mut text = format!("{EXPECTED_KERNELS_SIDECAR_HEADER}\n");
+    for kernel in kernels {
+        debug_assert!(!kernel.is_empty() && !kernel.contains(char::is_whitespace));
+        text.push_str(kernel);
+        text.push('\n');
+    }
+    text
+}
+
+/// Kernels listed by a `<module>.kernels` sidecar, or `None` when it has
+/// another header or a malformed name.
+pub fn parse_expected_kernels_sidecar(text: &str) -> Option<Vec<String>> {
+    let mut lines = text.lines();
+    if lines.next()? != EXPECTED_KERNELS_SIDECAR_HEADER {
+        return None;
+    }
+    lines
+        .map(|name| {
+            (!name.is_empty() && !name.contains(char::is_whitespace)).then(|| name.to_string())
+        })
+        .collect()
+}
+
+/// Names of the entry-point function symbols of a CUDA ELF, or `None` when
+/// its symbol table cannot be read. An image without a symbol table defines
+/// no entries.
+pub fn cubin_entry_names(bytes: &[u8]) -> Option<Vec<String>> {
+    let section_offset = usize::try_from(read_u64(bytes, 40)?).ok()?;
+    let section_count = usize::from(read_u16(bytes, 60)?);
+    let section = |index: usize| -> Option<usize> {
+        (index < section_count).then(|| section_offset + index * ELF64_SECTION_HEADER_LENGTH)
+    };
+    let mut names = Vec::new();
+    for index in 0..section_count {
+        let header = section(index)?;
+        if read_u32(bytes, header + 4)? != SECTION_TYPE_SYMBOL_TABLE {
+            continue;
+        }
+        let symbols = file_range(bytes, header)?;
+        let strings = file_range(
+            bytes,
+            section(usize::try_from(read_u32(bytes, header + 40)?).ok()?)?,
+        )?;
+        let (symbols, rest) = symbols.as_chunks::<ELF64_SYMBOL_LENGTH>();
+        if !rest.is_empty() {
+            return None;
+        }
+        for symbol in symbols {
+            let info = symbol[4];
+            let other = symbol[5];
+            if info & 0xf != SYMBOL_TYPE_FUNCTION || other & STO_CUDA_ENTRY == 0 {
+                continue;
+            }
+            let name_offset = usize::try_from(read_u32(symbol, 0)?).ok()?;
+            let name = strings.get(name_offset..)?;
+            let end = name.iter().position(|&byte| byte == 0)?;
+            names.push(std::str::from_utf8(&name[..end]).ok()?.to_string());
+        }
+    }
+    Some(names)
+}
+
+/// Names of the `.entry` kernels of a PTX module, or `None` when it cannot
+/// be parsed.
+pub fn ptx_entry_names(bytes: &[u8]) -> Option<Vec<String>> {
+    let text = std::str::from_utf8(bytes.strip_suffix(b"\0").unwrap_or(bytes)).ok()?;
+    let document = ptx_parse::Document::parse(text).ok()?;
+    Some(
+        document
+            .callables()
+            .iter()
+            .filter(|callable| callable.kind() == ptx_parse::CallableKind::Entry)
+            .map(|callable| callable.name().to_string())
+            .collect(),
+    )
+}
+
+/// File bytes of the section whose header starts at `header`.
+fn file_range(bytes: &[u8], header: usize) -> Option<&[u8]> {
+    let offset = usize::try_from(read_u64(bytes, header + 24)?).ok()?;
+    let length = usize::try_from(read_u64(bytes, header + 32)?).ok()?;
+    bytes.get(offset..offset.checked_add(length)?)
+}
+
+fn read_u16(bytes: &[u8], offset: usize) -> Option<u16> {
+    Some(u16::from_le_bytes(
+        bytes.get(offset..offset + 2)?.try_into().ok()?,
+    ))
+}
+
+fn read_u32(bytes: &[u8], offset: usize) -> Option<u32> {
+    Some(u32::from_le_bytes(
+        bytes.get(offset..offset + 4)?.try_into().ok()?,
+    ))
+}
+
+fn read_u64(bytes: &[u8], offset: usize) -> Option<u64> {
+    Some(u64::from_le_bytes(
+        bytes.get(offset..offset + 8)?.try_into().ok()?,
+    ))
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    /// A CUDA ELF with a symbol table holding `symbols` as
+    /// `(name, st_info, st_other)`, a string table, and nothing else.
+    fn cubin_with_symbols(symbols: &[(&str, u8, u8)]) -> Vec<u8> {
+        let mut strings = vec![0_u8];
+        let mut table = vec![0_u8; ELF64_SYMBOL_LENGTH];
+        for &(name, info, other) in symbols {
+            let mut symbol = vec![0_u8; ELF64_SYMBOL_LENGTH];
+            symbol[0..4].copy_from_slice(&(strings.len() as u32).to_le_bytes());
+            symbol[4] = info;
+            symbol[5] = other;
+            table.extend_from_slice(&symbol);
+            strings.extend_from_slice(name.as_bytes());
+            strings.push(0);
+        }
+        let headers = 64;
+        let symbols_offset = headers + 3 * ELF64_SECTION_HEADER_LENGTH;
+        let strings_offset = symbols_offset + table.len();
+        let mut bytes = vec![0_u8; strings_offset + strings.len()];
+        bytes[..4].copy_from_slice(b"\x7fELF");
+        bytes[4] = 2;
+        bytes[5] = 1;
+        bytes[6] = 1;
+        bytes[16..18].copy_from_slice(&2_u16.to_le_bytes());
+        bytes[18..20].copy_from_slice(&190_u16.to_le_bytes());
+        bytes[20..24].copy_from_slice(&1_u32.to_le_bytes());
+        bytes[40..48].copy_from_slice(&(headers as u64).to_le_bytes());
+        bytes[52..54].copy_from_slice(&64_u16.to_le_bytes());
+        bytes[58..60].copy_from_slice(&(ELF64_SECTION_HEADER_LENGTH as u16).to_le_bytes());
+        bytes[60..62].copy_from_slice(&3_u16.to_le_bytes());
+        let symtab = headers + ELF64_SECTION_HEADER_LENGTH;
+        bytes[symtab + 4..symtab + 8].copy_from_slice(&SECTION_TYPE_SYMBOL_TABLE.to_le_bytes());
+        bytes[symtab + 24..symtab + 32].copy_from_slice(&(symbols_offset as u64).to_le_bytes());
+        bytes[symtab + 32..symtab + 40].copy_from_slice(&(table.len() as u64).to_le_bytes());
+        bytes[symtab + 40..symtab + 44].copy_from_slice(&2_u32.to_le_bytes());
+        let strtab = symtab + ELF64_SECTION_HEADER_LENGTH;
+        bytes[strtab + 4..strtab + 8].copy_from_slice(&3_u32.to_le_bytes());
+        bytes[strtab + 24..strtab + 32].copy_from_slice(&(strings_offset as u64).to_le_bytes());
+        bytes[strtab + 32..strtab + 40].copy_from_slice(&(strings.len() as u64).to_le_bytes());
+        bytes[symbols_offset..strings_offset].copy_from_slice(&table);
+        bytes[strings_offset..].copy_from_slice(&strings);
+        bytes
+    }
+
+    const GLOBAL_FUNCTION: u8 = 0x12;
+    const LOCAL_FUNCTION: u8 = 0x02;
+
+    #[test]
+    fn cubin_inventory_lists_only_entry_flagged_functions() {
+        let cubin = cubin_with_symbols(&[
+            ("kernel", GLOBAL_FUNCTION, STO_CUDA_ENTRY),
+            ("helper", LOCAL_FUNCTION, 0),
+            ("device_global", GLOBAL_FUNCTION, 0),
+        ]);
+        assert_eq!(cubin_entry_names(&cubin).unwrap(), ["kernel"]);
+    }
+
+    #[test]
+    fn a_dropped_kernel_fails_the_link_even_in_a_well_formed_image() {
+        // The shape nvJitLink returns when it drops a module's kernels: a
+        // valid ELF whose symbol table has no entries.
+        let empty = cubin_with_symbols(&[]);
+        assert!(crate::is_valid_cubin(&empty));
+        assert!(matches!(
+            require_expected_kernels(&empty, FinalizerOutput::Cubin, &["kernel"]),
+            Err(FinalizerError::MissingKernels { ref missing, present: 0, .. })
+                if missing == &["kernel"]
+        ));
+
+        let partial = cubin_with_symbols(&[("first", GLOBAL_FUNCTION, STO_CUDA_ENTRY)]);
+        assert!(matches!(
+            require_expected_kernels(&partial, FinalizerOutput::Cubin, &["first", "second"]),
+            Err(FinalizerError::MissingKernels { ref missing, present: 1, .. })
+                if missing == &["second"]
+        ));
+        assert!(require_expected_kernels(&partial, FinalizerOutput::Cubin, &["first"]).is_ok());
+    }
+
+    #[test]
+    fn ptx_inventory_requires_entries_not_device_functions() {
+        let ptx = b".version 8.0\n.target sm_80\n.address_size 64\n\
+            .func helper()\n{\n\tret;\n}\n\
+            .visible .entry kernel()\n{\n\tret;\n}\n";
+        assert_eq!(ptx_entry_names(ptx).unwrap(), ["kernel"]);
+        assert!(matches!(
+            require_expected_kernels(ptx, FinalizerOutput::Ptx, &["helper"]),
+            Err(FinalizerError::MissingKernels { ref missing, .. }) if missing == &["helper"]
+        ));
+    }
+
+    #[test]
+    fn the_kernels_sidecar_round_trips_and_rejects_other_text() {
+        let text = expected_kernels_sidecar_text(["first", "second_TID_0123"]);
+        assert_eq!(
+            parse_expected_kernels_sidecar(&text).unwrap(),
+            ["first", "second_TID_0123"]
+        );
+        assert_eq!(
+            parse_expected_kernels_sidecar(&expected_kernels_sidecar_text([])).unwrap(),
+            Vec::<String>::new()
+        );
+        assert!(parse_expected_kernels_sidecar("first\n").is_none());
+        assert!(parse_expected_kernels_sidecar("").is_none());
+        let blank = format!("{EXPECTED_KERNELS_SIDECAR_HEADER}\nfirst\n\n");
+        assert!(parse_expected_kernels_sidecar(&blank).is_none());
+    }
+
+    #[test]
+    fn an_unreadable_inventory_fails_closed_and_no_expectation_checks_nothing() {
+        let not_text = b"\xff\xfe";
+        assert!(matches!(
+            require_expected_kernels(not_text, FinalizerOutput::Ptx, &["kernel"]),
+            Err(FinalizerError::UnreadableEntryInventory { .. })
+        ));
+        assert!(matches!(
+            require_expected_kernels(b"\x7fELF", FinalizerOutput::Cubin, &["kernel"]),
+            Err(FinalizerError::UnreadableEntryInventory { .. })
+        ));
+        assert!(require_expected_kernels(not_text, FinalizerOutput::Ptx, &[]).is_ok());
+    }
+}
diff --git a/crates/cuda-artifact-finalizer/src/lib.rs b/crates/cuda-artifact-finalizer/src/lib.rs
index 6a014422de..7d4b9ac4ec 100644
--- a/crates/cuda-artifact-finalizer/src/lib.rs
+++ b/crates/cuda-artifact-finalizer/src/lib.rs
@@ -11,6 +11,7 @@
 //! FMA, debug, input-order, validation, and provenance rules.
 
 mod diagnostics;
+mod entries;
 mod link;
 mod nvvm;
 mod options;
@@ -19,6 +20,10 @@ mod ptx;
 mod validation;
 
 pub use diagnostics::KernelResourceUsage;
+pub use entries::{
+    EXPECTED_KERNELS_SIDECAR_HEADER, cubin_entry_names, expected_kernels_sidecar_text,
+    parse_expected_kernels_sidecar, ptx_entry_names, require_expected_kernels,
+};
 pub use libnvvm_sys::{CudaArch, CudaArchParseError, LibdeviceNotFound, NvvmError, find_libdevice};
 pub use link::{LinkReport, LtoLinker};
 pub use nvjitlink_sys::NvJitLinkError;
@@ -130,6 +135,29 @@ pub enum FinalizerError {
     #[error("nvJitLink returned an empty PTX artifact")]
     EmptyPtx,
 
+    /// The linked image lacks kernels its inputs define. nvJitLink reports
+    /// success when it drops a module's kernels, so the finalizer checks the
+    /// output's entry points against the caller's expected kernels.
+    #[error(
+        "the linked {output:?} is missing expected kernel(s) {missing:?} ({present} entry point(s) present)"
+    )]
+    MissingKernels {
+        /// Requested output format.
+        output: FinalizerOutput,
+        /// Expected kernels that are not entry points of the output.
+        missing: Vec<String>,
+        /// Number of entry points the output does define.
+        present: usize,
+    },
+
+    /// The entry points of a linked image could not be read, so the expected
+    /// kernels could not be checked.
+    #[error("could not read the entry points of the linked {output:?}")]
+    UnreadableEntryInventory {
+        /// Requested output format.
+        output: FinalizerOutput,
+    },
+
     /// A pinned CUDA compilation tool changed around an operation. Its output can
     /// no longer be attributed to the provenance used by Cargo or a cache key.
     #[error(
@@ -188,12 +216,14 @@ impl Finalizer {
         ))
     }
 
-    /// Compile one NVVM IR module and return a validated target-specific cubin.
+    /// Compile one NVVM IR module and return a validated target-specific cubin
+    /// that defines every kernel in `expected_kernels`.
     pub fn materialize_nvvm_ir(
         &self,
         module_name: &str,
         nvvm_ir: &[u8],
         options: &FinalizationOptions,
+        expected_kernels: &[&str],
     ) -> Result<Vec<u8>, FinalizerError> {
         let ltoir = self
             .compiler
@@ -203,6 +233,7 @@ impl Finalizer {
             &[NamedInput::new(&ltoir_name, &ltoir)],
             options,
             FinalizerOutput::Cubin,
+            expected_kernels,
         )
     }
 
@@ -217,6 +248,7 @@ impl Finalizer {
         module_name: &str,
         nvvm_ir: &[u8],
         options: &FinalizationOptions,
+        expected_kernels: &[&str],
     ) -> Result<LinkReport, FinalizerError> {
         let ltoir = self
             .compiler
@@ -226,17 +258,21 @@ impl Finalizer {
             &[NamedInput::new(&ltoir_name, &ltoir)],
             options,
             FinalizerOutput::Cubin,
+            expected_kernels,
         )
     }
 
-    /// Link ordered LTOIR modules to cubin or PTX.
+    /// Link ordered LTOIR modules to cubin or PTX that defines every kernel in
+    /// `expected_kernels`.
     pub fn link_ltoir(
         &self,
         inputs: &[NamedInput<'_>],
         options: &FinalizationOptions,
         output: FinalizerOutput,
+        expected_kernels: &[&str],
     ) -> Result<Vec<u8>, FinalizerError> {
-        self.linker.link_ltoir(inputs, options, output)
+        self.linker
+            .link_ltoir(inputs, options, output, expected_kernels)
     }
 
     /// Link ordered LTOIR modules while collecting ptxas resource diagnostics.
@@ -250,8 +286,10 @@ impl Finalizer {
         inputs: &[NamedInput<'_>],
         options: &FinalizationOptions,
         output: FinalizerOutput,
+        expected_kernels: &[&str],
     ) -> Result<LinkReport, FinalizerError> {
-        self.linker.link_ltoir_with_report(inputs, options, output)
+        self.linker
+            .link_ltoir_with_report(inputs, options, output, expected_kernels)
     }
 
     /// Compiler component, including exact libdevice bytes and provenance.
@@ -408,7 +446,7 @@ entry:
             assert!(!ltoir.is_empty());
             let input = [NamedInput::new("kernel.ltoir", &ltoir)];
             let cubin = finalizer
-                .link_ltoir(&input, &options, FinalizerOutput::Cubin)
+                .link_ltoir(&input, &options, FinalizerOutput::Cubin, &["kernel"])
                 .unwrap();
             assert!(is_valid_cubin(&cubin));
             // A cubin whose kernel was stripped is still a well-formed ELF, so
@@ -421,7 +459,7 @@ entry:
                 cubin.len()
             );
             let ptx = finalizer
-                .link_ltoir(&input, &options, FinalizerOutput::Ptx)
+                .link_ltoir(&input, &options, FinalizerOutput::Ptx, &["kernel"])
                 .unwrap();
             assert!(
                 ptx.windows(b".version".len())
@@ -439,6 +477,39 @@ entry:
         }
     }
 
+    /// The same kernel without `@llvm.used`: nvJitLink's link-time optimizer
+    /// strips it and still reports success with a well-formed, empty image.
+    /// The finalizer must refuse that image for both outputs.
+    #[test]
+    #[ignore = "requires discoverable CUDA Toolkit libNVVM, nvJitLink, and libdevice"]
+    fn live_link_fails_when_nvjitlink_drops_an_expected_kernel() {
+        let unrooted = String::from_utf8_lossy(LEGACY_NVVM_IR)
+            .lines()
+            .filter(|line| !line.starts_with("@llvm.used"))
+            .collect::<Vec<_>>()
+            .join("\n");
+        let finalizer = Finalizer::discover().unwrap();
+        let options = FinalizationOptions::new("sm_86".parse().unwrap());
+        let ltoir = finalizer
+            .compiler()
+            .compile_nvvm_ir_to_ltoir("dropped.ll", unrooted.as_bytes(), &options)
+            .unwrap();
+        let input = [NamedInput::new("dropped.ltoir", &ltoir)];
+        for output in [FinalizerOutput::Cubin, FinalizerOutput::Ptx] {
+            let error = finalizer
+                .link_ltoir(&input, &options, output, &["kernel"])
+                .unwrap_err();
+            assert!(
+                matches!(
+                    error,
+                    FinalizerError::MissingKernels { ref missing, present: 0, .. }
+                        if missing == &["kernel"]
+                ),
+                "{output:?}: {error}"
+            );
+        }
+    }
+
     /// Legacy-dialect NVVM IR for a kernel that ptxas must spill: `values`
     /// floats stay live across `rounds` of all-to-all mixing while a
     /// `maxnreg` annotation caps the kernel far below that live set.
@@ -515,7 +586,7 @@ entry:
         let input = [NamedInput::new("spill.ltoir", &ltoir)];
 
         let report = finalizer
-            .link_ltoir_with_report(&input, &options, FinalizerOutput::Cubin)
+            .link_ltoir_with_report(&input, &options, FinalizerOutput::Cubin, &["spill_kernel"])
             .unwrap();
         assert!(is_valid_cubin(&report.image));
 
diff --git a/crates/cuda-artifact-finalizer/src/link.rs b/crates/cuda-artifact-finalizer/src/link.rs
index 2cb0c1c86f..ecabb9e962 100644
--- a/crates/cuda-artifact-finalizer/src/link.rs
+++ b/crates/cuda-artifact-finalizer/src/link.rs
@@ -4,6 +4,7 @@
  */
 
 use crate::diagnostics::{KernelResourceUsage, parse_ptxas_resource_usage};
+use crate::entries::require_expected_kernels;
 use crate::nvvm::{loaded_tool_digest_with_expected, report_changed_tool};
 use crate::options::{FinalizationOptions, FinalizerOutput, NamedInput};
 use crate::provenance::{
@@ -80,14 +81,18 @@ impl LtoLinker {
             .map(|digest| linker_provenance_digest(&digest))
     }
 
-    /// Link one or more LTOIR modules in the exact supplied order.
+    /// Link one or more LTOIR modules in the exact supplied order. The output
+    /// must define every kernel in `expected_kernels` as an entry point.
     pub fn link_ltoir(
         &self,
         inputs: &[NamedInput<'_>],
         options: &FinalizationOptions,
         output: FinalizerOutput,
+        expected_kernels: &[&str],
     ) -> Result<Vec<u8>, FinalizerError> {
-        Ok(self.link_ltoir_impl(inputs, options, output, false)?.image)
+        Ok(self
+            .link_ltoir_impl(inputs, options, output, expected_kernels, false)?
+            .image)
     }
 
     /// Link LTOIR while collecting non-semantic ptxas resource diagnostics.
@@ -112,8 +117,9 @@ impl LtoLinker {
         inputs: &[NamedInput<'_>],
         options: &FinalizationOptions,
         output: FinalizerOutput,
+        expected_kernels: &[&str],
     ) -> Result<LinkReport, FinalizerError> {
-        self.link_ltoir_impl(inputs, options, output, true)
+        self.link_ltoir_impl(inputs, options, output, expected_kernels, true)
     }
 
     fn link_ltoir_impl(
@@ -121,6 +127,7 @@ impl LtoLinker {
         inputs: &[NamedInput<'_>],
         options: &FinalizationOptions,
         output: FinalizerOutput,
+        expected_kernels: &[&str],
         collect_resource_usage: bool,
     ) -> Result<LinkReport, FinalizerError> {
         validate_inputs(inputs)?;
@@ -129,15 +136,18 @@ impl LtoLinker {
             InputType::Ltoir,
             options,
             output,
+            expected_kernels,
             collect_resource_usage,
         )
     }
 
-    /// Compile and link one PTX module to a validated target-specific cubin.
+    /// Compile and link one PTX module to a validated target-specific cubin
+    /// that defines every kernel in `expected_kernels`.
     pub fn link_ptx_to_cubin(
         &self,
         input: NamedInput<'_>,
         options: &FinalizationOptions,
+        expected_kernels: &[&str],
     ) -> Result<Vec<u8>, FinalizerError> {
         validate_inputs(std::slice::from_ref(&input))?;
         // Enforce the PTX C-string rule up front so this route rejects
@@ -151,6 +161,7 @@ impl LtoLinker {
                 InputType::Ptx,
                 options,
                 FinalizerOutput::Cubin,
+                expected_kernels,
                 false,
             )?
             .image)
@@ -162,6 +173,7 @@ impl LtoLinker {
         input_type: InputType,
         options: &FinalizationOptions,
         output: FinalizerOutput,
+        expected_kernels: &[&str],
         collect_resource_usage: bool,
     ) -> Result<LinkReport, FinalizerError> {
         with_revalidated_tool_identity(
@@ -169,7 +181,14 @@ impl LtoLinker {
             self.tool.digest,
             || current_linker_tool_digest(&self.tool),
             || {
-                match self.run_link(inputs, input_type, options, output, collect_resource_usage) {
+                match self.run_link(
+                    inputs,
+                    input_type,
+                    options,
+                    output,
+                    expected_kernels,
+                    collect_resource_usage,
+                ) {
                     // Older nvJitLink versions reject the diagnostic-only
                     // reporting options with NVJITLINK_ERROR_UNRECOGNIZED_OPTION.
                     // The caller asked for the same program plus a best-effort
@@ -178,7 +197,7 @@ impl LtoLinker {
                     Err(FinalizerError::NvJitLink(error))
                         if collect_resource_usage && error.is_unrecognized_option() =>
                     {
-                        self.run_link(inputs, input_type, options, output, false)
+                        self.run_link(inputs, input_type, options, output, expected_kernels, false)
                     }
                     result => result,
                 }
@@ -192,6 +211,7 @@ impl LtoLinker {
         input_type: InputType,
         options: &FinalizationOptions,
         output: FinalizerOutput,
+        expected_kernels: &[&str],
         collect_resource_usage: bool,
     ) -> Result<LinkReport, FinalizerError> {
         let mut option_storage = if input_type == InputType::Ptx {
@@ -221,6 +241,7 @@ impl LtoLinker {
         if output == FinalizerOutput::Ptx && image.is_empty() {
             return Err(FinalizerError::EmptyPtx);
         }
+        require_expected_kernels(&image, output, expected_kernels)?;
 
         let resource_usage = if collect_resource_usage {
             info_log
@@ -550,6 +571,7 @@ mod live_tests {
             .link_ptx_to_cubin(
                 NamedInput::new("acquire-load.ptx", ACQUIRE_LOAD_PTX),
                 &options,
+                &["acquire_load"],
             )
             .unwrap();
         assert!(is_valid_cubin(&cubin));
diff --git a/crates/cuda-host/examples/finalize_nvvm_ir.rs b/crates/cuda-host/examples/finalize_nvvm_ir.rs
index 2a383aec0c..c13e5f9d3f 100644
--- a/crates/cuda-host/examples/finalize_nvvm_ir.rs
+++ b/crates/cuda-host/examples/finalize_nvvm_ir.rs
@@ -68,7 +68,8 @@ fn main() -> Result<(), Box<dyn Error>> {
     let options = FinalizationOptions::new(target.clone());
     let finalizer = Finalizer::discover()?;
     let provenance = finalizer.provenance();
-    let cubin = finalizer.materialize_nvvm_ir("example.ll", EXAMPLE_NVVM_IR, &options)?;
+    let cubin =
+        finalizer.materialize_nvvm_ir("example.ll", EXAMPLE_NVVM_IR, &options, &["kernel"])?;
     std::fs::write(&output_path, &cubin)?;
 
     let e_ident_version = cubin[6];
diff --git a/crates/cuda-host/src/embedded.rs b/crates/cuda-host/src/embedded.rs
index 48776448f2..478b2e5523 100644
--- a/crates/cuda-host/src/embedded.rs
+++ b/crates/cuda-host/src/embedded.rs
@@ -12,6 +12,7 @@ pub use cuda_core::embedded::{
     embedded_modules_from_current_exe,
 };
 use cuda_core::{CudaContext, CudaModule, DriverError};
+use oxide_artifacts::ArtifactEntryKind;
 use std::sync::Arc;
 use thiserror::Error;
 
@@ -305,6 +306,13 @@ fn load_bundle(
         return Ok(ctx.load_module_from_image(ptx)?);
     }
 
+    let expected_kernels = bundle
+        .entries
+        .iter()
+        .filter(|entry| entry.kind == ArtifactEntryKind::Kernel)
+        .map(|entry| entry.symbol.as_str())
+        .collect::<Vec<_>>();
+
     if let Some(nvvm_ir) = bundle.payload(ArtifactPayloadKind::NvvmIr) {
         let emitted = target_arch_for_bundle(bundle)?;
         let execution = ltoir::execution_arch_for_context(ctx)?;
@@ -314,12 +322,14 @@ fn load_bundle(
                 &bundle.name,
                 &emitted.sm(),
                 bundle.compile_options,
+                &expected_kernels,
             )?,
             ltoir::ExecutionRoute::PtxBridge => ltoir::build_ptx_from_nvvm_ir_with_compile_options(
                 nvvm_ir,
                 &bundle.name,
                 &emitted.sm(),
                 bundle.compile_options,
+                &expected_kernels,
             )?,
         };
         return Ok(ctx.load_module_from_image(&image)?);
@@ -334,12 +344,14 @@ fn load_bundle(
                 &bundle.name,
                 &emitted.sm(),
                 bundle.compile_options,
+                &expected_kernels,
             )?,
             ltoir::ExecutionRoute::PtxBridge => ltoir::link_ltoir_to_ptx_with_compile_options(
                 ltoir,
                 &bundle.name,
                 &emitted.sm(),
                 bundle.compile_options,
+                &expected_kernels,
             )?,
         };
         return Ok(ctx.load_module_from_image(&image)?);
diff --git a/crates/cuda-host/src/ltoir.rs b/crates/cuda-host/src/ltoir.rs
index 810c56a396..ba4069abe7 100644
--- a/crates/cuda-host/src/ltoir.rs
+++ b/crates/cuda-host/src/ltoir.rs
@@ -77,7 +77,8 @@
 use crate::ltoir_cache::{BuiltArtifacts, CacheResult, cache_or_build};
 use cuda_artifact_finalizer::{
     CudaArch, CudaArchParseError, FinalizationOptions, Finalizer, FinalizerError, FinalizerOutput,
-    LtoLinker, NamedInput, NvJitLinkError, NvvmError,
+    LtoLinker, NamedInput, NvJitLinkError, NvvmError, parse_expected_kernels_sidecar,
+    require_expected_kernels,
 };
 #[cfg(test)]
 use cuda_artifact_finalizer::{
@@ -273,6 +274,11 @@ fn build_cubin_from_ll_file(ll_path: &Path, arch: &CudaArch) -> Result<FileCubin
         source,
     })?;
     let compile_options = read_compile_options(ll_path)?;
+    let expected_kernels = read_expected_kernels(ll_path)?;
+    let expected_kernels = expected_kernels
+        .iter()
+        .map(String::as_str)
+        .collect::<Vec<_>>();
     validate_ir_target_sidecar(ll_path, arch)?;
     // Record the supplied target for older or manually created `.ll` files.
     // The sibling `.ltoir` can then be loaded after the `.ll` is removed.
@@ -293,6 +299,7 @@ fn build_cubin_from_ll_file(ll_path: &Path, arch: &CudaArch) -> Result<FileCubin
         &ltoir_path.display().to_string(),
         arch,
         compile_options,
+        &expected_kernels,
     )?;
     let ltoir = cached
         .ltoir
@@ -327,8 +334,9 @@ pub fn build_cubin_from_nvvm_ir(
     nvvm_ir: &[u8],
     module_name: &str,
     arch: &str,
+    expected_kernels: &[&str],
 ) -> Result<Vec<u8>, LtoirError> {
-    build_cubin_from_nvvm_ir_with_options(nvvm_ir, module_name, arch, true)
+    build_cubin_from_nvvm_ir_with_options(nvvm_ir, module_name, arch, true, expected_kernels)
 }
 
 /// Compile NVVM IR bytes to a loadable cubin with an explicit FMA policy.
@@ -341,12 +349,14 @@ pub fn build_cubin_from_nvvm_ir_with_options(
     module_name: &str,
     arch: &str,
     allow_fma_contraction: bool,
+    expected_kernels: &[&str],
 ) -> Result<Vec<u8>, LtoirError> {
     build_cubin_from_nvvm_ir_with_compile_options(
         nvvm_ir,
         module_name,
         arch,
         ArtifactCompileOptions::new().with_fma_contraction(allow_fma_contraction),
+        expected_kernels,
     )
 }
 
@@ -357,10 +367,16 @@ pub fn build_cubin_from_nvvm_ir_with_compile_options(
     module_name: &str,
     arch: &str,
     compile_options: ArtifactCompileOptions,
+    expected_kernels: &[&str],
 ) -> Result<Vec<u8>, LtoirError> {
     let arch: CudaArch = arch.parse()?;
     let options = finalization_options(&arch, compile_options);
-    Ok(Finalizer::discover()?.materialize_nvvm_ir(module_name, nvvm_ir, &options)?)
+    Ok(Finalizer::discover()?.materialize_nvvm_ir(
+        module_name,
+        nvvm_ir,
+        &options,
+        expected_kernels,
+    )?)
 }
 
 /// Compile NVVM IR bytes to forward-compatible PTX.
@@ -372,8 +388,9 @@ pub fn build_ptx_from_nvvm_ir(
     nvvm_ir: &[u8],
     module_name: &str,
     arch: &str,
+    expected_kernels: &[&str],
 ) -> Result<Vec<u8>, LtoirError> {
-    build_ptx_from_nvvm_ir_with_options(nvvm_ir, module_name, arch, true)
+    build_ptx_from_nvvm_ir_with_options(nvvm_ir, module_name, arch, true, expected_kernels)
 }
 
 /// Compile NVVM IR bytes to forward-compatible PTX with an explicit FMA
@@ -387,12 +404,14 @@ pub fn build_ptx_from_nvvm_ir_with_options(
     module_name: &str,
     arch: &str,
     allow_fma_contraction: bool,
+    expected_kernels: &[&str],
 ) -> Result<Vec<u8>, LtoirError> {
     build_ptx_from_nvvm_ir_with_compile_options(
         nvvm_ir,
         module_name,
         arch,
         ArtifactCompileOptions::new().with_fma_contraction(allow_fma_contraction),
+        expected_kernels,
     )
 }
 
@@ -403,6 +422,7 @@ pub fn build_ptx_from_nvvm_ir_with_compile_options(
     module_name: &str,
     arch: &str,
     compile_options: ArtifactCompileOptions,
+    expected_kernels: &[&str],
 ) -> Result<Vec<u8>, LtoirError> {
     let arch: CudaArch = arch.parse()?;
     let finalizer = Finalizer::discover()?;
@@ -415,6 +435,7 @@ pub fn build_ptx_from_nvvm_ir_with_compile_options(
         &[NamedInput::new(&ltoir_name, &ltoir)],
         &options,
         FinalizerOutput::Ptx,
+        expected_kernels,
     )?)
 }
 
@@ -423,8 +444,9 @@ pub fn link_ltoir_to_cubin(
     ltoir: &[u8],
     module_name: &str,
     arch: &str,
+    expected_kernels: &[&str],
 ) -> Result<Vec<u8>, LtoirError> {
-    link_ltoir_to_cubin_with_options(ltoir, module_name, arch, true)
+    link_ltoir_to_cubin_with_options(ltoir, module_name, arch, true, expected_kernels)
 }
 
 /// Link a single LTOIR payload to a cubin with an explicit FMA policy.
@@ -437,12 +459,14 @@ pub fn link_ltoir_to_cubin_with_options(
     module_name: &str,
     arch: &str,
     allow_fma_contraction: bool,
+    expected_kernels: &[&str],
 ) -> Result<Vec<u8>, LtoirError> {
     link_ltoir_to_cubin_with_compile_options(
         ltoir,
         module_name,
         arch,
         ArtifactCompileOptions::new().with_fma_contraction(allow_fma_contraction),
+        expected_kernels,
     )
 }
 
@@ -453,9 +477,16 @@ pub fn link_ltoir_to_cubin_with_compile_options(
     module_name: &str,
     arch: &str,
     compile_options: ArtifactCompileOptions,
+    expected_kernels: &[&str],
 ) -> Result<Vec<u8>, LtoirError> {
     let arch: CudaArch = arch.parse()?;
-    link_ltoir_to_cubin_parsed_with_compile_options(ltoir, module_name, &arch, compile_options)
+    link_ltoir_to_cubin_parsed_with_compile_options(
+        ltoir,
+        module_name,
+        &arch,
+        compile_options,
+        expected_kernels,
+    )
 }
 
 /// Link one LTOIR payload to forward-compatible PTX.
@@ -466,8 +497,9 @@ pub fn link_ltoir_to_ptx(
     ltoir: &[u8],
     module_name: &str,
     arch: &str,
+    expected_kernels: &[&str],
 ) -> Result<Vec<u8>, LtoirError> {
-    link_ltoir_to_ptx_with_options(ltoir, module_name, arch, true)
+    link_ltoir_to_ptx_with_options(ltoir, module_name, arch, true, expected_kernels)
 }
 
 /// Link one LTOIR payload to forward-compatible PTX with an explicit FMA
@@ -481,12 +513,14 @@ pub fn link_ltoir_to_ptx_with_options(
     module_name: &str,
     arch: &str,
     allow_fma_contraction: bool,
+    expected_kernels: &[&str],
 ) -> Result<Vec<u8>, LtoirError> {
     link_ltoir_to_ptx_with_compile_options(
         ltoir,
         module_name,
         arch,
         ArtifactCompileOptions::new().with_fma_contraction(allow_fma_contraction),
+        expected_kernels,
     )
 }
 
@@ -497,9 +531,16 @@ pub fn link_ltoir_to_ptx_with_compile_options(
     module_name: &str,
     arch: &str,
     compile_options: ArtifactCompileOptions,
+    expected_kernels: &[&str],
 ) -> Result<Vec<u8>, LtoirError> {
     let arch: CudaArch = arch.parse()?;
-    link_ltoir_to_ptx_parsed_with_compile_options(ltoir, module_name, &arch, compile_options)
+    link_ltoir_to_ptx_parsed_with_compile_options(
+        ltoir,
+        module_name,
+        &arch,
+        compile_options,
+        expected_kernels,
+    )
 }
 
 fn link_ltoir_to_cubin_parsed_with_compile_options(
@@ -507,6 +548,7 @@ fn link_ltoir_to_cubin_parsed_with_compile_options(
     module_name: &str,
     arch: &CudaArch,
     compile_options: ArtifactCompileOptions,
+    expected_kernels: &[&str],
 ) -> Result<Vec<u8>, LtoirError> {
     link_ltoir_with_shared_finalizer(
         ltoir,
@@ -514,6 +556,7 @@ fn link_ltoir_to_cubin_parsed_with_compile_options(
         arch,
         compile_options,
         FinalizerOutput::Cubin,
+        expected_kernels,
     )
 }
 
@@ -522,6 +565,7 @@ fn link_ltoir_to_ptx_parsed_with_compile_options(
     module_name: &str,
     arch: &CudaArch,
     compile_options: ArtifactCompileOptions,
+    expected_kernels: &[&str],
 ) -> Result<Vec<u8>, LtoirError> {
     link_ltoir_with_shared_finalizer(
         ltoir,
@@ -529,6 +573,7 @@ fn link_ltoir_to_ptx_parsed_with_compile_options(
         arch,
         compile_options,
         FinalizerOutput::Ptx,
+        expected_kernels,
     )
 }
 
@@ -538,12 +583,14 @@ fn link_ltoir_with_shared_finalizer(
     arch: &CudaArch,
     compile_options: ArtifactCompileOptions,
     output: FinalizerOutput,
+    expected_kernels: &[&str],
 ) -> Result<Vec<u8>, LtoirError> {
     let options = finalization_options(arch, compile_options);
     Ok(LtoLinker::discover()?.link_ltoir(
         &[NamedInput::new(module_name, ltoir)],
         &options,
         output,
+        expected_kernels,
     )?)
 }
 
@@ -568,6 +615,7 @@ fn cached_nvvm_ir_to_cubin(
     nvvm_module_name: &str,
     ltoir_module_name: &str,
     arch: &CudaArch,
+    expected_kernels: &[&str],
 ) -> Result<CacheResult, LtoirError> {
     cached_nvvm_ir_to_cubin_with_compile_options(
         source_dir,
@@ -576,6 +624,7 @@ fn cached_nvvm_ir_to_cubin(
         ltoir_module_name,
         arch,
         ArtifactCompileOptions::new(),
+        expected_kernels,
     )
 }
 
@@ -586,6 +635,7 @@ fn cached_nvvm_ir_to_cubin_with_compile_options(
     ltoir_module_name: &str,
     arch: &CudaArch,
     compile_options: ArtifactCompileOptions,
+    expected_kernels: &[&str],
 ) -> Result<CacheResult, LtoirError> {
     let finalizer = Finalizer::discover()?;
     let options = finalization_options(arch, compile_options);
@@ -606,6 +656,7 @@ fn cached_nvvm_ir_to_cubin_with_compile_options(
             &[NamedInput::new(ltoir_module_name, &ltoir)],
             &options,
             FinalizerOutput::Cubin,
+            expected_kernels,
         )?;
         Ok(BuiltArtifacts::new(cubin, Some(ltoir)))
     };
@@ -614,6 +665,8 @@ fn cached_nvvm_ir_to_cubin_with_compile_options(
         Some(key) => cache_or_build(source_dir, &key, build)?,
         None => uncached_result(build()?),
     };
+    // A fresh build was checked by the link; a hit may predate the check.
+    require_expected_kernels(&result.cubin, FinalizerOutput::Cubin, expected_kernels)?;
     report_cache_result(&result);
     Ok(result)
 }
@@ -624,6 +677,7 @@ fn cached_ltoir_to_cubin(
     ltoir: &[u8],
     module_name: &str,
     arch: &CudaArch,
+    expected_kernels: &[&str],
 ) -> Result<CacheResult, LtoirError> {
     cached_ltoir_to_cubin_with_compile_options(
         source_dir,
@@ -631,6 +685,7 @@ fn cached_ltoir_to_cubin(
         module_name,
         arch,
         ArtifactCompileOptions::new(),
+        expected_kernels,
     )
 }
 
@@ -640,6 +695,7 @@ fn cached_ltoir_to_cubin_with_compile_options(
     module_name: &str,
     arch: &CudaArch,
     compile_options: ArtifactCompileOptions,
+    expected_kernels: &[&str],
 ) -> Result<CacheResult, LtoirError> {
     let linker = LtoLinker::discover()?;
     let options = finalization_options(arch, compile_options);
@@ -647,7 +703,7 @@ fn cached_ltoir_to_cubin_with_compile_options(
     let key = linker.artifact_digest(&inputs, &options, FinalizerOutput::Cubin);
     let build = || -> Result<BuiltArtifacts, LtoirError> {
         Ok(BuiltArtifacts::new(
-            linker.link_ltoir(&inputs, &options, FinalizerOutput::Cubin)?,
+            linker.link_ltoir(&inputs, &options, FinalizerOutput::Cubin, expected_kernels)?,
             None,
         ))
     };
@@ -656,6 +712,8 @@ fn cached_ltoir_to_cubin_with_compile_options(
         Some(key) => cache_or_build(source_dir, &key, build)?,
         None => uncached_result(build()?),
     };
+    // A fresh build was checked by the link; a hit may predate the check.
+    require_expected_kernels(&result.cubin, FinalizerOutput::Cubin, expected_kernels)?;
     report_cache_result(&result);
     Ok(result)
 }
@@ -864,11 +922,16 @@ pub fn load_kernel_module(
                 ExecutionRoute::PtxBridge => {
                     let nvvm_ir = read_artifact(&ll)?;
                     let compile_options = read_compile_options(&ll)?;
+                    let expected_kernels = read_expected_kernels(&ll)?;
                     let ptx = build_ptx_from_nvvm_ir_with_compile_options(
                         &nvvm_ir,
                         &ll.display().to_string(),
                         &emitted.sm(),
                         compile_options,
+                        &expected_kernels
+                            .iter()
+                            .map(String::as_str)
+                            .collect::<Vec<_>>(),
                     )?;
                     Ok(ctx.load_module_from_image(&ptx)?)
                 }
@@ -879,6 +942,11 @@ pub fn load_kernel_module(
             let execution = execution_arch_for_context(ctx)?;
             let bytes = read_artifact(&ltoir)?;
             let compile_options = read_compile_options(&ltoir)?;
+            let expected_kernels = read_expected_kernels(&ltoir)?;
+            let expected_kernels = expected_kernels
+                .iter()
+                .map(String::as_str)
+                .collect::<Vec<_>>();
             let image = match execution_route(&emitted, &execution)? {
                 ExecutionRoute::Cubin => {
                     cached_ltoir_to_cubin_with_compile_options(
@@ -887,6 +955,7 @@ pub fn load_kernel_module(
                         &ltoir.display().to_string(),
                         &emitted,
                         compile_options,
+                        &expected_kernels,
                     )?
                     .cubin
                 }
@@ -895,6 +964,7 @@ pub fn load_kernel_module(
                     &ltoir.display().to_string(),
                     &emitted,
                     compile_options,
+                    &expected_kernels,
                 )?,
             };
             Ok(ctx.load_module_from_image(&image)?)
@@ -1052,6 +1122,46 @@ fn emitted_compile_options_path(ll_path: &Path) -> PathBuf {
     ll_path.with_extension("options")
 }
 
+/// Kernels the finalized image of the NVVM IR or LTOIR at `artifact_path`
+/// must define, from its `<name>.kernels` sidecar.
+///
+/// An artifact the compiler emitted (its `.target` carries the compile-options
+/// marker the compiler always writes) must have the list. A manual artifact is
+/// checked only against a list that is present.
+fn read_expected_kernels(artifact_path: &Path) -> Result<Vec<String>, LtoirError> {
+    let path = artifact_path.with_extension("kernels");
+    let text = match std::fs::read_to_string(&path) {
+        Ok(text) => text,
+        Err(source) if source.kind() == std::io::ErrorKind::NotFound => {
+            let target_path = emitted_target_path(artifact_path);
+            let emitted_by_compiler = match std::fs::read_to_string(&target_path) {
+                Ok(target) => target.lines().nth(1) == Some(COMPILE_OPTIONS_TARGET_MARKER),
+                Err(source) if source.kind() == std::io::ErrorKind::NotFound => false,
+                Err(source) => {
+                    return Err(LtoirError::Io {
+                        path: target_path,
+                        source,
+                    });
+                }
+            };
+            return if emitted_by_compiler {
+                Err(LtoirError::InvalidCompileOptions {
+                    path,
+                    value: "required expected-kernels sidecar is missing; rebuild the device crate"
+                        .to_string(),
+                })
+            } else {
+                Ok(Vec::new())
+            };
+        }
+        Err(source) => return Err(LtoirError::Io { path, source }),
+    };
+    parse_expected_kernels_sidecar(&text).ok_or_else(|| LtoirError::InvalidCompileOptions {
+        path,
+        value: "not a cuda-oxide expected-kernels list".to_string(),
+    })
+}
+
 fn read_compile_options(ll_path: &Path) -> Result<ArtifactCompileOptions, LtoirError> {
     let path = emitted_compile_options_path(ll_path);
     match std::fs::read_to_string(&path) {
@@ -1433,10 +1543,10 @@ mod tests {
 
     #[test]
     fn invalid_link_target_fails_before_loading_nvidia_libraries() {
-        let error = link_ltoir_to_cubin(&[], "empty", "nvvm-ir").unwrap_err();
+        let error = link_ltoir_to_cubin(&[], "empty", "nvvm-ir", &[]).unwrap_err();
         assert!(matches!(error, LtoirError::InvalidTarget(_)));
 
-        let error = link_ltoir_to_ptx(&[], "empty", "nvvm-ir").unwrap_err();
+        let error = link_ltoir_to_ptx(&[], "empty", "nvvm-ir", &[]).unwrap_err();
         assert!(matches!(error, LtoirError::InvalidTarget(_)));
     }
 
@@ -1750,6 +1860,8 @@ mod tests {
 target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-i128:128:128-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 target triple = "nvptx64-nvidia-cuda"
 
+@llvm.used = appending global [1 x i8*] [i8* bitcast (void ()* @kernel to i8*)], section "llvm.metadata"
+
 define void @kernel() {
 entry:
   ret void
@@ -1765,6 +1877,11 @@ entry:
         std::fs::create_dir_all(&dir).unwrap();
         let ll = dir.join("kernel.ll");
         std::fs::write(&ll, LEGACY_NVVM_IR).unwrap();
+        std::fs::write(
+            dir.join("kernel.kernels"),
+            cuda_artifact_finalizer::expected_kernels_sidecar_text(["kernel"]),
+        )
+        .unwrap();
 
         let first = build_cubin_from_ll(&ll, "sm_86").unwrap();
         let first_bytes = std::fs::read(&first).unwrap();
@@ -1787,6 +1904,7 @@ entry:
             &ll.display().to_string(),
             &ltoir_path.display().to_string(),
             &arch,
+            &["kernel"],
         )
         .unwrap();
         assert!(native_hit.cache_hit);
@@ -1801,6 +1919,7 @@ entry:
             &ll.display().to_string(),
             &ltoir_path.display().to_string(),
             &arch,
+            &["kernel"],
         )
         .unwrap();
         assert!(native_hit_again.cache_hit);
@@ -1821,11 +1940,13 @@ entry:
         let standalone_dir = dir.join("standalone");
         std::fs::create_dir(&standalone_dir).unwrap();
         let standalone_first =
-            cached_ltoir_to_cubin(&standalone_dir, &ltoir, "kernel.ltoir", &arch).unwrap();
+            cached_ltoir_to_cubin(&standalone_dir, &ltoir, "kernel.ltoir", &arch, &["kernel"])
+                .unwrap();
         assert!(!standalone_first.cache_hit);
         assert_eq!(standalone_first.ltoir, None);
         let standalone_second =
-            cached_ltoir_to_cubin(&standalone_dir, &ltoir, "kernel.ltoir", &arch).unwrap();
+            cached_ltoir_to_cubin(&standalone_dir, &ltoir, "kernel.ltoir", &arch, &["kernel"])
+                .unwrap();
         assert!(standalone_second.cache_hit);
         assert_eq!(standalone_second.cubin, standalone_first.cubin);
         assert_eq!(
@@ -1845,6 +1966,7 @@ entry:
             &ll.display().to_string(),
             &ltoir_path.display().to_string(),
             &arch,
+            &["kernel"],
         )
         .unwrap();
         assert!(changed_hit.cache_hit);
diff --git a/crates/mir-importer/Cargo.toml b/crates/mir-importer/Cargo.toml
index 3f32e242a4..5d89c7f190 100644
--- a/crates/mir-importer/Cargo.toml
+++ b/crates/mir-importer/Cargo.toml
@@ -17,6 +17,7 @@ dialect-mir = { workspace = true }
 dialect-iket = { workspace = true }
 dialect-nvvm = { workspace = true }
 mir-lower = { workspace = true }
+cuda-artifact-finalizer = { workspace = true }
 oxide-artifacts = { workspace = true }
 cuda-oxide-codegen = { workspace = true }
 reserved-oxide-symbols = { workspace = true }
diff --git a/crates/mir-importer/src/pipeline.rs b/crates/mir-importer/src/pipeline.rs
index 5d88f4a234..ce5a054796 100644
--- a/crates/mir-importer/src/pipeline.rs
+++ b/crates/mir-importer/src/pipeline.rs
@@ -620,6 +620,7 @@ pub fn run_pipeline(
                 config.allow_fma_contraction,
                 config.debug_kind,
             )?;
+            write_expected_kernels_sidecar(&config.output_dir, &config.output_name, functions)?;
             // Publish the target last: its version marker is the completion record
             // that says the sibling options file is required.
             write_nvvm_target_sidecar(&config.output_dir, &config.output_name, &generated.target)?;
@@ -735,6 +736,29 @@ fn write_nvvm_compile_options_sidecar(
     })
 }
 
+/// Records, next to the emitted `.ll`, the kernels its finalized image must
+/// define, so that finalizing from files checks them as the embedded route
+/// does.
+fn write_expected_kernels_sidecar(
+    output_dir: &Path,
+    output_name: &str,
+    functions: &[CollectedFunction],
+) -> Result<(), PipelineError> {
+    let path = output_dir.join(format!("{output_name}.kernels"));
+    let text = cuda_artifact_finalizer::expected_kernels_sidecar_text(
+        functions
+            .iter()
+            .filter(|function| function.is_kernel)
+            .map(|function| function.export_name.as_str()),
+    );
+    std::fs::write(&path, text).map_err(|error| {
+        PipelineError::Export(format!(
+            "failed to record expected kernels in {}: {error}",
+            path.display()
+        ))
+    })
+}
+
 fn stale_compilation_artifact_paths(
     output_dir: &Path,
     output_name: &str,
@@ -746,6 +770,7 @@ fn stale_compilation_artifact_paths(
         "ptx",
         "target",
         "options",
+        "kernels",
         "ltoir",
         "cubin",
         "cubin.target",
diff --git a/crates/rustc-codegen-cuda/Cargo.lock b/crates/rustc-codegen-cuda/Cargo.lock
index 0c08699d3f..b4e0d3da14 100644
--- a/crates/rustc-codegen-cuda/Cargo.lock
+++ b/crates/rustc-codegen-cuda/Cargo.lock
@@ -196,6 +196,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
@@ -497,6 +498,7 @@ name = "mir-importer"
 version = "0.2.1"
 dependencies = [
  "combine",
+ "cuda-artifact-finalizer",
  "cuda-oxide-codegen",
  "dialect-iket",
  "dialect-mir",
diff --git a/crates/rustc-codegen-cuda/examples/abi_hmm/Cargo.lock b/crates/rustc-codegen-cuda/examples/abi_hmm/Cargo.lock
index 9461537672..c210884b75 100644
--- a/crates/rustc-codegen-cuda/examples/abi_hmm/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/abi_hmm/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/addressof_sharedarray/Cargo.lock b/crates/rustc-codegen-cuda/examples/addressof_sharedarray/Cargo.lock
index 25817b78ad..7c6a9194f4 100644
--- a/crates/rustc-codegen-cuda/examples/addressof_sharedarray/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/addressof_sharedarray/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/aligned_field_loads/Cargo.lock b/crates/rustc-codegen-cuda/examples/aligned_field_loads/Cargo.lock
index 2f039583c6..582f997435 100644
--- a/crates/rustc-codegen-cuda/examples/aligned_field_loads/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/aligned_field_loads/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/aligned_field_stores/Cargo.lock b/crates/rustc-codegen-cuda/examples/aligned_field_stores/Cargo.lock
index fef3a581ef..a6419a2147 100644
--- a/crates/rustc-codegen-cuda/examples/aligned_field_stores/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/aligned_field_stores/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/array_constants/Cargo.lock b/crates/rustc-codegen-cuda/examples/array_constants/Cargo.lock
index 05682c563a..0c79a67e97 100644
--- a/crates/rustc-codegen-cuda/examples/array_constants/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/array_constants/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/array_for_loop/Cargo.lock b/crates/rustc-codegen-cuda/examples/array_for_loop/Cargo.lock
index b9284b62f2..a5f55439d2 100644
--- a/crates/rustc-codegen-cuda/examples/array_for_loop/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/array_for_loop/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/array_index/Cargo.lock b/crates/rustc-codegen-cuda/examples/array_index/Cargo.lock
index e1f7a4e21f..ae2d301d8c 100644
--- a/crates/rustc-codegen-cuda/examples/array_index/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/array_index/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/array_init/Cargo.lock b/crates/rustc-codegen-cuda/examples/array_init/Cargo.lock
index a7dba9ef07..b64daaddab 100644
--- a/crates/rustc-codegen-cuda/examples/array_init/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/array_init/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/asin_acos_smoke/Cargo.lock b/crates/rustc-codegen-cuda/examples/asin_acos_smoke/Cargo.lock
index 88bcd5d473..fc2a1e9b1d 100644
--- a/crates/rustc-codegen-cuda/examples/asin_acos_smoke/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/asin_acos_smoke/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/assume_init_uninhabited_trap/Cargo.lock b/crates/rustc-codegen-cuda/examples/assume_init_uninhabited_trap/Cargo.lock
index 9c9b67d2bc..6ce71d46aa 100644
--- a/crates/rustc-codegen-cuda/examples/assume_init_uninhabited_trap/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/assume_init_uninhabited_trap/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/async_mlp/Cargo.lock b/crates/rustc-codegen-cuda/examples/async_mlp/Cargo.lock
index 4745715f3e..b2780c9b98 100644
--- a/crates/rustc-codegen-cuda/examples/async_mlp/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/async_mlp/Cargo.lock
@@ -138,6 +138,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror 2.0.18",
diff --git a/crates/rustc-codegen-cuda/examples/async_vecadd/Cargo.lock b/crates/rustc-codegen-cuda/examples/async_vecadd/Cargo.lock
index bd31eaf12c..f9c61c43b6 100644
--- a/crates/rustc-codegen-cuda/examples/async_vecadd/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/async_vecadd/Cargo.lock
@@ -138,6 +138,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror 2.0.18",
diff --git a/crates/rustc-codegen-cuda/examples/atomic_f16/Cargo.lock b/crates/rustc-codegen-cuda/examples/atomic_f16/Cargo.lock
index 8297803ec6..d4bee3d031 100644
--- a/crates/rustc-codegen-cuda/examples/atomic_f16/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/atomic_f16/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/atomics/Cargo.lock b/crates/rustc-codegen-cuda/examples/atomics/Cargo.lock
index 2e790372f5..e75c6a1b9f 100644
--- a/crates/rustc-codegen-cuda/examples/atomics/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/atomics/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/barrier/Cargo.lock b/crates/rustc-codegen-cuda/examples/barrier/Cargo.lock
index d522fbb17e..7fd50f5b11 100644
--- a/crates/rustc-codegen-cuda/examples/barrier/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/barrier/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/bf16x2_arith/Cargo.lock b/crates/rustc-codegen-cuda/examples/bf16x2_arith/Cargo.lock
index 421675603d..421a4a4f79 100644
--- a/crates/rustc-codegen-cuda/examples/bf16x2_arith/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/bf16x2_arith/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/bf16x2_fma/Cargo.lock b/crates/rustc-codegen-cuda/examples/bf16x2_fma/Cargo.lock
index aa23e0adaf..8f864a9109 100644
--- a/crates/rustc-codegen-cuda/examples/bf16x2_fma/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/bf16x2_fma/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/bool_phi_cmp/Cargo.lock b/crates/rustc-codegen-cuda/examples/bool_phi_cmp/Cargo.lock
index 3dd701229d..cc2ff65cdc 100644
--- a/crates/rustc-codegen-cuda/examples/bool_phi_cmp/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/bool_phi_cmp/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/carrying_mul_add/Cargo.lock b/crates/rustc-codegen-cuda/examples/carrying_mul_add/Cargo.lock
index 66ae2238d9..d5b6555982 100644
--- a/crates/rustc-codegen-cuda/examples/carrying_mul_add/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/carrying_mul_add/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/cast_chunks/Cargo.lock b/crates/rustc-codegen-cuda/examples/cast_chunks/Cargo.lock
index 73f02f09d0..15c8da278a 100644
--- a/crates/rustc-codegen-cuda/examples/cast_chunks/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/cast_chunks/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/cast_tests/Cargo.lock b/crates/rustc-codegen-cuda/examples/cast_tests/Cargo.lock
index f6bf9a8c90..539fd2b080 100644
--- a/crates/rustc-codegen-cuda/examples/cast_tests/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/cast_tests/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/cbrt_smoke/Cargo.lock b/crates/rustc-codegen-cuda/examples/cbrt_smoke/Cargo.lock
index 6d0981f8e1..ca0afc91a3 100644
--- a/crates/rustc-codegen-cuda/examples/cbrt_smoke/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/cbrt_smoke/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/checked_arith/Cargo.lock b/crates/rustc-codegen-cuda/examples/checked_arith/Cargo.lock
index f2ef233ba5..3922586ce9 100644
--- a/crates/rustc-codegen-cuda/examples/checked_arith/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/checked_arith/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/clc/Cargo.lock b/crates/rustc-codegen-cuda/examples/clc/Cargo.lock
index 2392e70a74..a78c6c8ba1 100644
--- a/crates/rustc-codegen-cuda/examples/clc/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/clc/Cargo.lock
@@ -131,6 +131,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/cluster/Cargo.lock b/crates/rustc-codegen-cuda/examples/cluster/Cargo.lock
index 30d7d03df2..116b481eb9 100644
--- a/crates/rustc-codegen-cuda/examples/cluster/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/cluster/Cargo.lock
@@ -131,6 +131,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/compiler_features/Cargo.lock b/crates/rustc-codegen-cuda/examples/compiler_features/Cargo.lock
index 0b587243ce..03ee1675ed 100644
--- a/crates/rustc-codegen-cuda/examples/compiler_features/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/compiler_features/Cargo.lock
@@ -131,6 +131,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/complex_mul/Cargo.lock b/crates/rustc-codegen-cuda/examples/complex_mul/Cargo.lock
index 5638e9cc7c..f60022cf5c 100644
--- a/crates/rustc-codegen-cuda/examples/complex_mul/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/complex_mul/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/const_aggregate/Cargo.lock b/crates/rustc-codegen-cuda/examples/const_aggregate/Cargo.lock
index 9ee2dce022..ebc601f8eb 100644
--- a/crates/rustc-codegen-cuda/examples/const_aggregate/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/const_aggregate/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/const_bool_dead_branch/Cargo.lock b/crates/rustc-codegen-cuda/examples/const_bool_dead_branch/Cargo.lock
index d5f59c979d..3f31b7a600 100644
--- a/crates/rustc-codegen-cuda/examples/const_bool_dead_branch/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/const_bool_dead_branch/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/const_generic/Cargo.lock b/crates/rustc-codegen-cuda/examples/const_generic/Cargo.lock
index 844ba64d3f..b21f2a6b06 100644
--- a/crates/rustc-codegen-cuda/examples/const_generic/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/const_generic/Cargo.lock
@@ -131,6 +131,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/const_table_lookup/Cargo.lock b/crates/rustc-codegen-cuda/examples/const_table_lookup/Cargo.lock
index 3249cb7164..c880b4ffd0 100644
--- a/crates/rustc-codegen-cuda/examples/const_table_lookup/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/const_table_lookup/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/const_tuple_table_lookup/Cargo.lock b/crates/rustc-codegen-cuda/examples/const_tuple_table_lookup/Cargo.lock
index dbf96fa6da..e1ec53e605 100644
--- a/crates/rustc-codegen-cuda/examples/const_tuple_table_lookup/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/const_tuple_table_lookup/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/constant_index_from_end/Cargo.lock b/crates/rustc-codegen-cuda/examples/constant_index_from_end/Cargo.lock
index bcf6f614aa..8ab7798f89 100644
--- a/crates/rustc-codegen-cuda/examples/constant_index_from_end/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/constant_index_from_end/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/constant_memory/Cargo.lock b/crates/rustc-codegen-cuda/examples/constant_memory/Cargo.lock
index e71b118b0f..a8f16ec0a1 100644
--- a/crates/rustc-codegen-cuda/examples/constant_memory/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/constant_memory/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/constant_memory_coeffs/Cargo.lock b/crates/rustc-codegen-cuda/examples/constant_memory_coeffs/Cargo.lock
index f9740dddb6..cf465b7d88 100644
--- a/crates/rustc-codegen-cuda/examples/constant_memory_coeffs/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/constant_memory_coeffs/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/constant_memory_simple/Cargo.lock b/crates/rustc-codegen-cuda/examples/constant_memory_simple/Cargo.lock
index b44db16544..26e90488d5 100644
--- a/crates/rustc-codegen-cuda/examples/constant_memory_simple/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/constant_memory_simple/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/constant_memory_table/Cargo.lock b/crates/rustc-codegen-cuda/examples/constant_memory_table/Cargo.lock
index 5163734956..4df0b2811b 100644
--- a/crates/rustc-codegen-cuda/examples/constant_memory_table/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/constant_memory_table/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/coop_groups_demo/Cargo.lock b/crates/rustc-codegen-cuda/examples/coop_groups_demo/Cargo.lock
index 11a8bc8689..07a4095dfe 100644
--- a/crates/rustc-codegen-cuda/examples/coop_groups_demo/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/coop_groups_demo/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/copy_aggregate_borrow/Cargo.lock b/crates/rustc-codegen-cuda/examples/copy_aggregate_borrow/Cargo.lock
index d402dce173..1c5c030459 100644
--- a/crates/rustc-codegen-cuda/examples/copy_aggregate_borrow/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/copy_aggregate_borrow/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/copy_nonoverlapping/Cargo.lock b/crates/rustc-codegen-cuda/examples/copy_nonoverlapping/Cargo.lock
index 0d87e1c364..6f16f20353 100644
--- a/crates/rustc-codegen-cuda/examples/copy_nonoverlapping/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/copy_nonoverlapping/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/counted_barrier/Cargo.lock b/crates/rustc-codegen-cuda/examples/counted_barrier/Cargo.lock
index 8ac4602cbe..4087ac86fb 100644
--- a/crates/rustc-codegen-cuda/examples/counted_barrier/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/counted_barrier/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/cp_async_small/Cargo.lock b/crates/rustc-codegen-cuda/examples/cp_async_small/Cargo.lock
index 275f6ce89b..ee41cb930d 100644
--- a/crates/rustc-codegen-cuda/examples/cp_async_small/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/cp_async_small/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/cp_async_zfill/Cargo.lock b/crates/rustc-codegen-cuda/examples/cp_async_zfill/Cargo.lock
index e6ab94100b..5bb903bd49 100644
--- a/crates/rustc-codegen-cuda/examples/cp_async_zfill/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/cp_async_zfill/Cargo.lock
@@ -130,6 +130,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/cpp_consumes_rust_device/Cargo.lock b/crates/rustc-codegen-cuda/examples/cpp_consumes_rust_device/Cargo.lock
index 7a3141ac8a..298f6aa00b 100644
--- a/crates/rustc-codegen-cuda/examples/cpp_consumes_rust_device/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/cpp_consumes_rust_device/Cargo.lock
@@ -129,6 +129,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/cross_crate_embedded/Cargo.lock b/crates/rustc-codegen-cuda/examples/cross_crate_embedded/Cargo.lock
index d07a96890c..9af251dd1b 100644
--- a/crates/rustc-codegen-cuda/examples/cross_crate_embedded/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/cross_crate_embedded/Cargo.lock
@@ -131,6 +131,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/cross_crate_kernel/Cargo.lock b/crates/rustc-codegen-cuda/examples/cross_crate_kernel/Cargo.lock
index b7aa9b1d52..e7d5e17a0d 100644
--- a/crates/rustc-codegen-cuda/examples/cross_crate_kernel/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/cross_crate_kernel/Cargo.lock
@@ -132,6 +132,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/cross_crate_merged_symbols/Cargo.lock b/crates/rustc-codegen-cuda/examples/cross_crate_merged_symbols/Cargo.lock
index b8576eb953..09df3776b5 100644
--- a/crates/rustc-codegen-cuda/examples/cross_crate_merged_symbols/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/cross_crate_merged_symbols/Cargo.lock
@@ -131,6 +131,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/cross_crate_recursion/Cargo.lock b/crates/rustc-codegen-cuda/examples/cross_crate_recursion/Cargo.lock
index 0334ded1a0..31839b0ed5 100644
--- a/crates/rustc-codegen-cuda/examples/cross_crate_recursion/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/cross_crate_recursion/Cargo.lock
@@ -131,6 +131,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/cuda_module_contract/Cargo.lock b/crates/rustc-codegen-cuda/examples/cuda_module_contract/Cargo.lock
index 01eeff6910..ffc053d86e 100644
--- a/crates/rustc-codegen-cuda/examples/cuda_module_contract/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/cuda_module_contract/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/cuda_module_in_lib/Cargo.lock b/crates/rustc-codegen-cuda/examples/cuda_module_in_lib/Cargo.lock
index b899c33e0a..82dd546f16 100644
--- a/crates/rustc-codegen-cuda/examples/cuda_module_in_lib/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/cuda_module_in_lib/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/cuda_module_nested/Cargo.lock b/crates/rustc-codegen-cuda/examples/cuda_module_nested/Cargo.lock
index 8b64dec455..a5cd5def2c 100644
--- a/crates/rustc-codegen-cuda/examples/cuda_module_nested/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/cuda_module_nested/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/cutile_inter_kernel/simt/Cargo.lock b/crates/rustc-codegen-cuda/examples/cutile_inter_kernel/simt/Cargo.lock
index 364b61cab0..bfc69d52df 100644
--- a/crates/rustc-codegen-cuda/examples/cutile_inter_kernel/simt/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/cutile_inter_kernel/simt/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/cvt_f16x2/Cargo.lock b/crates/rustc-codegen-cuda/examples/cvt_f16x2/Cargo.lock
index 28a7db4a95..3478b57012 100644
--- a/crates/rustc-codegen-cuda/examples/cvt_f16x2/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/cvt_f16x2/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/cvt_packed/Cargo.lock b/crates/rustc-codegen-cuda/examples/cvt_packed/Cargo.lock
index 692dfc7f4a..4188aa2cdc 100644
--- a/crates/rustc-codegen-cuda/examples/cvt_packed/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/cvt_packed/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/debug-tests/Cargo.lock b/crates/rustc-codegen-cuda/examples/debug-tests/Cargo.lock
index f47883ba38..87ef980eac 100644
--- a/crates/rustc-codegen-cuda/examples/debug-tests/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/debug-tests/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/debug/Cargo.lock b/crates/rustc-codegen-cuda/examples/debug/Cargo.lock
index 2ba759945b..1e4f171ba9 100644
--- a/crates/rustc-codegen-cuda/examples/debug/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/debug/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/debug_pointer_locals/Cargo.lock b/crates/rustc-codegen-cuda/examples/debug_pointer_locals/Cargo.lock
index d67b648691..8d03ded50b 100644
--- a/crates/rustc-codegen-cuda/examples/debug_pointer_locals/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/debug_pointer_locals/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/device_closures/Cargo.lock b/crates/rustc-codegen-cuda/examples/device_closures/Cargo.lock
index 4c6e46a1c0..6fd22c7787 100644
--- a/crates/rustc-codegen-cuda/examples/device_closures/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/device_closures/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/device_ffi_test/Cargo.lock b/crates/rustc-codegen-cuda/examples/device_ffi_test/Cargo.lock
index c85eb38160..4edb8237f5 100644
--- a/crates/rustc-codegen-cuda/examples/device_ffi_test/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/device_ffi_test/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/device_global/Cargo.lock b/crates/rustc-codegen-cuda/examples/device_global/Cargo.lock
index 383f6b0335..16b6558c88 100644
--- a/crates/rustc-codegen-cuda/examples/device_global/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/device_global/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/device_panic_trap/Cargo.lock b/crates/rustc-codegen-cuda/examples/device_panic_trap/Cargo.lock
index f529d5fd6b..49a3251b0a 100644
--- a/crates/rustc-codegen-cuda/examples/device_panic_trap/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/device_panic_trap/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/disjoint_from_raw_parts/Cargo.lock b/crates/rustc-codegen-cuda/examples/disjoint_from_raw_parts/Cargo.lock
index 70cd490a8e..d5dec4e3f2 100644
--- a/crates/rustc-codegen-cuda/examples/disjoint_from_raw_parts/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/disjoint_from_raw_parts/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/disjoint_slice_len/Cargo.lock b/crates/rustc-codegen-cuda/examples/disjoint_slice_len/Cargo.lock
index 5f68f80a5f..00982aebb8 100644
--- a/crates/rustc-codegen-cuda/examples/disjoint_slice_len/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/disjoint_slice_len/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/dotprod/Cargo.lock b/crates/rustc-codegen-cuda/examples/dotprod/Cargo.lock
index d57979e28c..5ac7fc4b9c 100644
--- a/crates/rustc-codegen-cuda/examples/dotprod/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/dotprod/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/dpx_minmax_chains/Cargo.lock b/crates/rustc-codegen-cuda/examples/dpx_minmax_chains/Cargo.lock
index f8d469d2b7..594c207a8a 100644
--- a/crates/rustc-codegen-cuda/examples/dpx_minmax_chains/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/dpx_minmax_chains/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/drop_glue/Cargo.lock b/crates/rustc-codegen-cuda/examples/drop_glue/Cargo.lock
index 5dba90ff46..673ee9cedd 100644
--- a/crates/rustc-codegen-cuda/examples/drop_glue/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/drop_glue/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/dst_metadata/Cargo.lock b/crates/rustc-codegen-cuda/examples/dst_metadata/Cargo.lock
index 52bd8569aa..22afd62a9c 100644
--- a/crates/rustc-codegen-cuda/examples/dst_metadata/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/dst_metadata/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/dynamic_smem/Cargo.lock b/crates/rustc-codegen-cuda/examples/dynamic_smem/Cargo.lock
index 8d5be55af3..6f380232fc 100644
--- a/crates/rustc-codegen-cuda/examples/dynamic_smem/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/dynamic_smem/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/elect_leader/Cargo.lock b/crates/rustc-codegen-cuda/examples/elect_leader/Cargo.lock
index 81a7f81994..e267f7ac79 100644
--- a/crates/rustc-codegen-cuda/examples/elect_leader/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/elect_leader/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/enum_abi_multi_payload/Cargo.lock b/crates/rustc-codegen-cuda/examples/enum_abi_multi_payload/Cargo.lock
index b1d74e24d3..1e625fdbf0 100644
--- a/crates/rustc-codegen-cuda/examples/enum_abi_multi_payload/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/enum_abi_multi_payload/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/enum_array_match/Cargo.lock b/crates/rustc-codegen-cuda/examples/enum_array_match/Cargo.lock
index 09c7177119..09ab90fc67 100644
--- a/crates/rustc-codegen-cuda/examples/enum_array_match/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/enum_array_match/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/enum_constant_provenance/Cargo.lock b/crates/rustc-codegen-cuda/examples/enum_constant_provenance/Cargo.lock
index 262dfc9fc8..575d0d80a1 100644
--- a/crates/rustc-codegen-cuda/examples/enum_constant_provenance/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/enum_constant_provenance/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/enum_payload_addr/Cargo.lock b/crates/rustc-codegen-cuda/examples/enum_payload_addr/Cargo.lock
index 3468d7b47d..dd4f9af180 100644
--- a/crates/rustc-codegen-cuda/examples/enum_payload_addr/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/enum_payload_addr/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/enum_state_transitions/Cargo.lock b/crates/rustc-codegen-cuda/examples/enum_state_transitions/Cargo.lock
index ed89d4cc7e..1b268bf201 100644
--- a/crates/rustc-codegen-cuda/examples/enum_state_transitions/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/enum_state_transitions/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/enum_table_lookup/Cargo.lock b/crates/rustc-codegen-cuda/examples/enum_table_lookup/Cargo.lock
index 5c216116d5..cbb553ebfe 100644
--- a/crates/rustc-codegen-cuda/examples/enum_table_lookup/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/enum_table_lookup/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/error/Cargo.lock b/crates/rustc-codegen-cuda/examples/error/Cargo.lock
index b536acb602..07ed977786 100644
--- a/crates/rustc-codegen-cuda/examples/error/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/error/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/error_enum_bool_payload_addr/Cargo.lock b/crates/rustc-codegen-cuda/examples/error_enum_bool_payload_addr/Cargo.lock
index 7622915028..4d580f930c 100644
--- a/crates/rustc-codegen-cuda/examples/error_enum_bool_payload_addr/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/error_enum_bool_payload_addr/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/error_enum_pointer_overlap/Cargo.lock b/crates/rustc-codegen-cuda/examples/error_enum_pointer_overlap/Cargo.lock
index ad48361fc0..1d8dd721ba 100644
--- a/crates/rustc-codegen-cuda/examples/error_enum_pointer_overlap/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/error_enum_pointer_overlap/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/error_enum_shared_pointer_layout/Cargo.lock b/crates/rustc-codegen-cuda/examples/error_enum_shared_pointer_layout/Cargo.lock
index 6af2cf3dde..3ee82c8bb3 100644
--- a/crates/rustc-codegen-cuda/examples/error_enum_shared_pointer_layout/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/error_enum_shared_pointer_layout/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_abi/Cargo.lock b/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_abi/Cargo.lock
index e506b4ab6c..58d0b609dc 100644
--- a/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_abi/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_abi/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_callable/Cargo.lock b/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_callable/Cargo.lock
index eafbb76d6c..6ebbb9dcd7 100644
--- a/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_callable/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_callable/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_fn_pointer/Cargo.lock b/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_fn_pointer/Cargo.lock
index 99c15f3011..6018b01e9e 100644
--- a/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_fn_pointer/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_fn_pointer/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_unknown_id/Cargo.lock b/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_unknown_id/Cargo.lock
index f78ef8f5d7..def93b1eb4 100644
--- a/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_unknown_id/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_unknown_id/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/error_heap_alloc/Cargo.lock b/crates/rustc-codegen-cuda/examples/error_heap_alloc/Cargo.lock
index 787e9ce175..01024e310b 100644
--- a/crates/rustc-codegen-cuda/examples/error_heap_alloc/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/error_heap_alloc/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/error_host_arch_intrinsic/Cargo.lock b/crates/rustc-codegen-cuda/examples/error_host_arch_intrinsic/Cargo.lock
index 19bd1ee85e..7306745e57 100644
--- a/crates/rustc-codegen-cuda/examples/error_host_arch_intrinsic/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/error_host_arch_intrinsic/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/error_host_target_feature/Cargo.lock b/crates/rustc-codegen-cuda/examples/error_host_target_feature/Cargo.lock
index 9acf06018d..c00e0441e0 100644
--- a/crates/rustc-codegen-cuda/examples/error_host_target_feature/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/error_host_target_feature/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/error_kernel_shared_param/Cargo.lock b/crates/rustc-codegen-cuda/examples/error_kernel_shared_param/Cargo.lock
index 314bb7e525..697c5ec60a 100644
--- a/crates/rustc-codegen-cuda/examples/error_kernel_shared_param/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/error_kernel_shared_param/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/error_missing_device_attr/Cargo.lock b/crates/rustc-codegen-cuda/examples/error_missing_device_attr/Cargo.lock
index ffc20635a8..76dff18749 100644
--- a/crates/rustc-codegen-cuda/examples/error_missing_device_attr/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/error_missing_device_attr/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/error_set_discriminant_uninhabited/Cargo.lock b/crates/rustc-codegen-cuda/examples/error_set_discriminant_uninhabited/Cargo.lock
index 7bfbd566a6..bb89587640 100644
--- a/crates/rustc-codegen-cuda/examples/error_set_discriminant_uninhabited/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/error_set_discriminant_uninhabited/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/ex2_approx_f16/Cargo.lock b/crates/rustc-codegen-cuda/examples/ex2_approx_f16/Cargo.lock
index c57932810b..dfe05208cf 100644
--- a/crates/rustc-codegen-cuda/examples/ex2_approx_f16/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/ex2_approx_f16/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/exact_div_chunks/Cargo.lock b/crates/rustc-codegen-cuda/examples/exact_div_chunks/Cargo.lock
index 6ccca22a38..fb7959f860 100644
--- a/crates/rustc-codegen-cuda/examples/exact_div_chunks/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/exact_div_chunks/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/export_name_policy/Cargo.lock b/crates/rustc-codegen-cuda/examples/export_name_policy/Cargo.lock
index 2df3450bc2..c63dbdb49e 100644
--- a/crates/rustc-codegen-cuda/examples/export_name_policy/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/export_name_policy/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/extern_crate_closure/Cargo.lock b/crates/rustc-codegen-cuda/examples/extern_crate_closure/Cargo.lock
index 980afae2a9..eb78da03d9 100644
--- a/crates/rustc-codegen-cuda/examples/extern_crate_closure/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/extern_crate_closure/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/extern_libdevice/Cargo.lock b/crates/rustc-codegen-cuda/examples/extern_libdevice/Cargo.lock
index adba8e86fb..78fd28ea65 100644
--- a/crates/rustc-codegen-cuda/examples/extern_libdevice/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/extern_libdevice/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/f16_stress/Cargo.lock b/crates/rustc-codegen-cuda/examples/f16_stress/Cargo.lock
index 98b1d3585a..7edf1c7dd1 100644
--- a/crates/rustc-codegen-cuda/examples/f16_stress/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/f16_stress/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/f16x2_arith/Cargo.lock b/crates/rustc-codegen-cuda/examples/f16x2_arith/Cargo.lock
index 851a560cb9..f7220be946 100644
--- a/crates/rustc-codegen-cuda/examples/f16x2_arith/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/f16x2_arith/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/f32x2_arith/Cargo.lock b/crates/rustc-codegen-cuda/examples/f32x2_arith/Cargo.lock
index d04ddd1b48..299028d27b 100644
--- a/crates/rustc-codegen-cuda/examples/f32x2_arith/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/f32x2_arith/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/fcmp_ne_nan_smoke/Cargo.lock b/crates/rustc-codegen-cuda/examples/fcmp_ne_nan_smoke/Cargo.lock
index 7e5bd24388..415a412a2d 100644
--- a/crates/rustc-codegen-cuda/examples/fcmp_ne_nan_smoke/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/fcmp_ne_nan_smoke/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/field_array_assign/Cargo.lock b/crates/rustc-codegen-cuda/examples/field_array_assign/Cargo.lock
index 23a109223b..28ff52c123 100644
--- a/crates/rustc-codegen-cuda/examples/field_array_assign/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/field_array_assign/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/float_rounding/Cargo.lock b/crates/rustc-codegen-cuda/examples/float_rounding/Cargo.lock
index 3c86f49492..d0e8df5130 100644
--- a/crates/rustc-codegen-cuda/examples/float_rounding/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/float_rounding/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/fmaxmin_smoke/Cargo.lock b/crates/rustc-codegen-cuda/examples/fmaxmin_smoke/Cargo.lock
index 3aea5b89c8..3dfbf56f91 100644
--- a/crates/rustc-codegen-cuda/examples/fmaxmin_smoke/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/fmaxmin_smoke/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/fn_ptr_naming/Cargo.lock b/crates/rustc-codegen-cuda/examples/fn_ptr_naming/Cargo.lock
index cfe577c030..80eb25a765 100644
--- a/crates/rustc-codegen-cuda/examples/fn_ptr_naming/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/fn_ptr_naming/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/fp_special_literal_smoke/Cargo.lock b/crates/rustc-codegen-cuda/examples/fp_special_literal_smoke/Cargo.lock
index 15935cdf61..f4607fc473 100644
--- a/crates/rustc-codegen-cuda/examples/fp_special_literal_smoke/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/fp_special_literal_smoke/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/function_item_call/Cargo.lock b/crates/rustc-codegen-cuda/examples/function_item_call/Cargo.lock
index 16f54126d7..fb16cce5c6 100644
--- a/crates/rustc-codegen-cuda/examples/function_item_call/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/function_item_call/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/future_apis/Cargo.lock b/crates/rustc-codegen-cuda/examples/future_apis/Cargo.lock
index 456c8bc76b..c9b389940f 100644
--- a/crates/rustc-codegen-cuda/examples/future_apis/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/future_apis/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/gemm/Cargo.lock b/crates/rustc-codegen-cuda/examples/gemm/Cargo.lock
index 42a247fbd7..94dc5a3701 100644
--- a/crates/rustc-codegen-cuda/examples/gemm/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/gemm/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/gemm_sol/Cargo.lock b/crates/rustc-codegen-cuda/examples/gemm_sol/Cargo.lock
index e0ec81c8bb..9152ca1550 100644
--- a/crates/rustc-codegen-cuda/examples/gemm_sol/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/gemm_sol/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/gemm_sol_final/Cargo.lock b/crates/rustc-codegen-cuda/examples/gemm_sol_final/Cargo.lock
index d804b0e2d3..87a3160b13 100644
--- a/crates/rustc-codegen-cuda/examples/gemm_sol_final/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/gemm_sol_final/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/gemm_views/Cargo.lock b/crates/rustc-codegen-cuda/examples/gemm_views/Cargo.lock
index d80fa49aa9..4ce3dc733d 100644
--- a/crates/rustc-codegen-cuda/examples/gemm_views/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/gemm_views/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/generated_intrinsics/Cargo.lock b/crates/rustc-codegen-cuda/examples/generated_intrinsics/Cargo.lock
index 39ddf5df63..13339c7d7e 100644
--- a/crates/rustc-codegen-cuda/examples/generated_intrinsics/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/generated_intrinsics/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/generated_intrinsics_blackwell/Cargo.lock b/crates/rustc-codegen-cuda/examples/generated_intrinsics_blackwell/Cargo.lock
index e0f4d81a3d..09f13b8389 100644
--- a/crates/rustc-codegen-cuda/examples/generated_intrinsics_blackwell/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/generated_intrinsics_blackwell/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/generated_ldmatrix/Cargo.lock b/crates/rustc-codegen-cuda/examples/generated_ldmatrix/Cargo.lock
index a6a401cbae..2a00938468 100644
--- a/crates/rustc-codegen-cuda/examples/generated_ldmatrix/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/generated_ldmatrix/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/generic/Cargo.lock b/crates/rustc-codegen-cuda/examples/generic/Cargo.lock
index 4c5adc73fe..aea72b2c47 100644
--- a/crates/rustc-codegen-cuda/examples/generic/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/generic/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/hashmap/Cargo.lock b/crates/rustc-codegen-cuda/examples/hashmap/Cargo.lock
index 2319208577..ae7534f007 100644
--- a/crates/rustc-codegen-cuda/examples/hashmap/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/hashmap/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/hashmap_v2/Cargo.lock b/crates/rustc-codegen-cuda/examples/hashmap_v2/Cargo.lock
index 25fac89f43..1ebae5aad1 100644
--- a/crates/rustc-codegen-cuda/examples/hashmap_v2/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/hashmap_v2/Cargo.lock
@@ -152,6 +152,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/hashmap_v3/Cargo.lock b/crates/rustc-codegen-cuda/examples/hashmap_v3/Cargo.lock
index a783fa3efa..26e3007ea5 100644
--- a/crates/rustc-codegen-cuda/examples/hashmap_v3/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/hashmap_v3/Cargo.lock
@@ -152,6 +152,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/hello_constant/Cargo.lock b/crates/rustc-codegen-cuda/examples/hello_constant/Cargo.lock
index f75ea10174..322f7a1110 100644
--- a/crates/rustc-codegen-cuda/examples/hello_constant/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/hello_constant/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/helper_fn/Cargo.lock b/crates/rustc-codegen-cuda/examples/helper_fn/Cargo.lock
index 2d58ec8e57..3a146dda1b 100644
--- a/crates/rustc-codegen-cuda/examples/helper_fn/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/helper_fn/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/host_closure/Cargo.lock b/crates/rustc-codegen-cuda/examples/host_closure/Cargo.lock
index d1c83eea7c..c775bae918 100644
--- a/crates/rustc-codegen-cuda/examples/host_closure/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/host_closure/Cargo.lock
@@ -127,6 +127,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror 2.0.18",
diff --git a/crates/rustc-codegen-cuda/examples/host_closure_target_feature/Cargo.lock b/crates/rustc-codegen-cuda/examples/host_closure_target_feature/Cargo.lock
index 7ffc4317fc..0f77a67ca3 100644
--- a/crates/rustc-codegen-cuda/examples/host_closure_target_feature/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/host_closure_target_feature/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/iket_trace/Cargo.lock b/crates/rustc-codegen-cuda/examples/iket_trace/Cargo.lock
index 244a74a2a3..d2cfa56856 100644
--- a/crates/rustc-codegen-cuda/examples/iket_trace/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/iket_trace/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/image_convolution/Cargo.lock b/crates/rustc-codegen-cuda/examples/image_convolution/Cargo.lock
index 75d597190c..4759201061 100644
--- a/crates/rustc-codegen-cuda/examples/image_convolution/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/image_convolution/Cargo.lock
@@ -133,6 +133,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/index2d_const/Cargo.lock b/crates/rustc-codegen-cuda/examples/index2d_const/Cargo.lock
index 61c4488138..bf5f79b880 100644
--- a/crates/rustc-codegen-cuda/examples/index2d_const/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/index2d_const/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/index_field_assign/Cargo.lock b/crates/rustc-codegen-cuda/examples/index_field_assign/Cargo.lock
index 46eb8d72a2..0b0f883f5d 100644
--- a/crates/rustc-codegen-cuda/examples/index_field_assign/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/index_field_assign/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/inline_ptx/Cargo.lock b/crates/rustc-codegen-cuda/examples/inline_ptx/Cargo.lock
index 3f727fe69a..7e98f13c9f 100644
--- a/crates/rustc-codegen-cuda/examples/inline_ptx/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/inline_ptx/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/integer_minmax/Cargo.lock b/crates/rustc-codegen-cuda/examples/integer_minmax/Cargo.lock
index a295a5530d..450789bef8 100644
--- a/crates/rustc-codegen-cuda/examples/integer_minmax/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/integer_minmax/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/interior_mutability/Cargo.lock b/crates/rustc-codegen-cuda/examples/interior_mutability/Cargo.lock
index fc4d32f5e3..f26ecbaf86 100644
--- a/crates/rustc-codegen-cuda/examples/interior_mutability/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/interior_mutability/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/interop_cubin_identity/device/Cargo.lock b/crates/rustc-codegen-cuda/examples/interop_cubin_identity/device/Cargo.lock
index 207a1fb6d3..207ca1a2a6 100644
--- a/crates/rustc-codegen-cuda/examples/interop_cubin_identity/device/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/interop_cubin_identity/device/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/kernel_reference_validity/Cargo.lock b/crates/rustc-codegen-cuda/examples/kernel_reference_validity/Cargo.lock
index 0935248304..dde25e4d2d 100644
--- a/crates/rustc-codegen-cuda/examples/kernel_reference_validity/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/kernel_reference_validity/Cargo.lock
@@ -127,6 +127,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror 2.0.18",
diff --git a/crates/rustc-codegen-cuda/examples/lanemask_scan/Cargo.lock b/crates/rustc-codegen-cuda/examples/lanemask_scan/Cargo.lock
index 6587b07b31..285e6b3c14 100644
--- a/crates/rustc-codegen-cuda/examples/lanemask_scan/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/lanemask_scan/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/legacy_atomic_fadd/Cargo.lock b/crates/rustc-codegen-cuda/examples/legacy_atomic_fadd/Cargo.lock
index 8a1ca2cc18..e81b7f335b 100644
--- a/crates/rustc-codegen-cuda/examples/legacy_atomic_fadd/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/legacy_atomic_fadd/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/legacy_atomic_rmw_cas/Cargo.lock b/crates/rustc-codegen-cuda/examples/legacy_atomic_rmw_cas/Cargo.lock
index 46c0130603..39bdba697a 100644
--- a/crates/rustc-codegen-cuda/examples/legacy_atomic_rmw_cas/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/legacy_atomic_rmw_cas/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/legacy_nvvm_pointer_shapes/Cargo.lock b/crates/rustc-codegen-cuda/examples/legacy_nvvm_pointer_shapes/Cargo.lock
index 34e88a1108..b63d45985f 100644
--- a/crates/rustc-codegen-cuda/examples/legacy_nvvm_pointer_shapes/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/legacy_nvvm_pointer_shapes/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/libdevice_math/Cargo.lock b/crates/rustc-codegen-cuda/examples/libdevice_math/Cargo.lock
index cfa69d8993..d44f2dc21a 100644
--- a/crates/rustc-codegen-cuda/examples/libdevice_math/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/libdevice_math/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/libm_math/Cargo.lock b/crates/rustc-codegen-cuda/examples/libm_math/Cargo.lock
index c8d4e98cf7..525b4412a2 100644
--- a/crates/rustc-codegen-cuda/examples/libm_math/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/libm_math/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/manual_launch_generic/Cargo.lock b/crates/rustc-codegen-cuda/examples/manual_launch_generic/Cargo.lock
index 19666b0c06..9653dcf92d 100644
--- a/crates/rustc-codegen-cuda/examples/manual_launch_generic/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/manual_launch_generic/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/manual_launch_libdevice/Cargo.lock b/crates/rustc-codegen-cuda/examples/manual_launch_libdevice/Cargo.lock
index 33ab1d54c3..5bdc2fe112 100644
--- a/crates/rustc-codegen-cuda/examples/manual_launch_libdevice/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/manual_launch_libdevice/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/map_sum/Cargo.lock b/crates/rustc-codegen-cuda/examples/map_sum/Cargo.lock
index 601c2a54e8..4c50267a38 100644
--- a/crates/rustc-codegen-cuda/examples/map_sum/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/map_sum/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/math_atan/Cargo.lock b/crates/rustc-codegen-cuda/examples/math_atan/Cargo.lock
index 92da87cf7c..9a90f59a88 100644
--- a/crates/rustc-codegen-cuda/examples/math_atan/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/math_atan/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/math_hyperbolic/Cargo.lock b/crates/rustc-codegen-cuda/examples/math_hyperbolic/Cargo.lock
index e1e2b9218e..5533ab786e 100644
--- a/crates/rustc-codegen-cuda/examples/math_hyperbolic/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/math_hyperbolic/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/math_no_mir_inline/Cargo.lock b/crates/rustc-codegen-cuda/examples/math_no_mir_inline/Cargo.lock
index d81fee134f..99ccb0367b 100644
--- a/crates/rustc-codegen-cuda/examples/math_no_mir_inline/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/math_no_mir_inline/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/math_tan/Cargo.lock b/crates/rustc-codegen-cuda/examples/math_tan/Cargo.lock
index 2cfef54516..900864879c 100644
--- a/crates/rustc-codegen-cuda/examples/math_tan/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/math_tan/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/mathdx_ffi_test/Cargo.lock b/crates/rustc-codegen-cuda/examples/mathdx_ffi_test/Cargo.lock
index d1b06fc902..cad426c527 100644
--- a/crates/rustc-codegen-cuda/examples/mathdx_ffi_test/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/mathdx_ffi_test/Cargo.lock
@@ -141,6 +141,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/mcast_barrier_test/Cargo.lock b/crates/rustc-codegen-cuda/examples/mcast_barrier_test/Cargo.lock
index d5502a8d42..7a7714d71d 100644
--- a/crates/rustc-codegen-cuda/examples/mcast_barrier_test/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/mcast_barrier_test/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/mem_swap/Cargo.lock b/crates/rustc-codegen-cuda/examples/mem_swap/Cargo.lock
index 01de7b85de..25a218d700 100644
--- a/crates/rustc-codegen-cuda/examples/mem_swap/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/mem_swap/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/mir_projected_call_destination/Cargo.lock b/crates/rustc-codegen-cuda/examples/mir_projected_call_destination/Cargo.lock
index 863e54db07..faca3b82a9 100644
--- a/crates/rustc-codegen-cuda/examples/mir_projected_call_destination/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/mir_projected_call_destination/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/mma_mxf8f6f4/Cargo.lock b/crates/rustc-codegen-cuda/examples/mma_mxf8f6f4/Cargo.lock
index 0922e1f8e0..d8a001fa38 100644
--- a/crates/rustc-codegen-cuda/examples/mma_mxf8f6f4/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/mma_mxf8f6f4/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/ne_bytes_transmute/Cargo.lock b/crates/rustc-codegen-cuda/examples/ne_bytes_transmute/Cargo.lock
index b22fe5ef2d..a01209fa1f 100644
--- a/crates/rustc-codegen-cuda/examples/ne_bytes_transmute/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/ne_bytes_transmute/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/nested_index_assignment/Cargo.lock b/crates/rustc-codegen-cuda/examples/nested_index_assignment/Cargo.lock
index a8e8178c7f..a439abedd7 100644
--- a/crates/rustc-codegen-cuda/examples/nested_index_assignment/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/nested_index_assignment/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/numeric_stress/Cargo.lock b/crates/rustc-codegen-cuda/examples/numeric_stress/Cargo.lock
index e00ca11700..b329acd0f7 100644
--- a/crates/rustc-codegen-cuda/examples/numeric_stress/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/numeric_stress/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/nvls_all_reduce/Cargo.lock b/crates/rustc-codegen-cuda/examples/nvls_all_reduce/Cargo.lock
index 057fa2bbbc..8d820373af 100644
--- a/crates/rustc-codegen-cuda/examples/nvls_all_reduce/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/nvls_all_reduce/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/option_ref_transmute/Cargo.lock b/crates/rustc-codegen-cuda/examples/option_ref_transmute/Cargo.lock
index bdeafb01bd..10cad3c8ab 100644
--- a/crates/rustc-codegen-cuda/examples/option_ref_transmute/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/option_ref_transmute/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/option_ref_unwrap_or/Cargo.lock b/crates/rustc-codegen-cuda/examples/option_ref_unwrap_or/Cargo.lock
index 28420dff9a..930c59a511 100644
--- a/crates/rustc-codegen-cuda/examples/option_ref_unwrap_or/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/option_ref_unwrap_or/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/option_unwrap_trap/Cargo.lock b/crates/rustc-codegen-cuda/examples/option_unwrap_trap/Cargo.lock
index 92603a6e8c..a76e6875e4 100644
--- a/crates/rustc-codegen-cuda/examples/option_unwrap_trap/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/option_unwrap_trap/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/ord_cmp/Cargo.lock b/crates/rustc-codegen-cuda/examples/ord_cmp/Cargo.lock
index dfe385a365..e1133a3886 100644
--- a/crates/rustc-codegen-cuda/examples/ord_cmp/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/ord_cmp/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/packed_aggregate_abi/Cargo.lock b/crates/rustc-codegen-cuda/examples/packed_aggregate_abi/Cargo.lock
index 790d39cce9..27dc331bcd 100644
--- a/crates/rustc-codegen-cuda/examples/packed_aggregate_abi/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/packed_aggregate_abi/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/packed_atomic_add/Cargo.lock b/crates/rustc-codegen-cuda/examples/packed_atomic_add/Cargo.lock
index dad75fc12e..e4de1bd416 100644
--- a/crates/rustc-codegen-cuda/examples/packed_atomic_add/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/packed_atomic_add/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/padded_constants/Cargo.lock b/crates/rustc-codegen-cuda/examples/padded_constants/Cargo.lock
index 13eb6c1c41..f46d58214f 100644
--- a/crates/rustc-codegen-cuda/examples/padded_constants/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/padded_constants/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/partial_warp_reduce/Cargo.lock b/crates/rustc-codegen-cuda/examples/partial_warp_reduce/Cargo.lock
index 59cc4cf060..ae3f8f3068 100644
--- a/crates/rustc-codegen-cuda/examples/partial_warp_reduce/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/partial_warp_reduce/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/pinned_overlap/Cargo.lock b/crates/rustc-codegen-cuda/examples/pinned_overlap/Cargo.lock
index 45d4bc755e..7c7fad20cd 100644
--- a/crates/rustc-codegen-cuda/examples/pinned_overlap/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/pinned_overlap/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/place_read_fallbacks/Cargo.lock b/crates/rustc-codegen-cuda/examples/place_read_fallbacks/Cargo.lock
index 836460af70..d9b515455d 100644
--- a/crates/rustc-codegen-cuda/examples/place_read_fallbacks/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/place_read_fallbacks/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/policy_config/Cargo.lock b/crates/rustc-codegen-cuda/examples/policy_config/Cargo.lock
index 7a7d466aed..6a22d261e5 100644
--- a/crates/rustc-codegen-cuda/examples/policy_config/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/policy_config/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/primitive_stress/Cargo.lock b/crates/rustc-codegen-cuda/examples/primitive_stress/Cargo.lock
index 0f780e0c03..1253738e02 100644
--- a/crates/rustc-codegen-cuda/examples/primitive_stress/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/primitive_stress/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/printf/Cargo.lock b/crates/rustc-codegen-cuda/examples/printf/Cargo.lock
index 6ba7b5a341..c6936e6458 100644
--- a/crates/rustc-codegen-cuda/examples/printf/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/printf/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/projected_place_read/Cargo.lock b/crates/rustc-codegen-cuda/examples/projected_place_read/Cargo.lock
index a711f72df0..e83633302d 100644
--- a/crates/rustc-codegen-cuda/examples/projected_place_read/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/projected_place_read/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/proof_carrying_views/Cargo.lock b/crates/rustc-codegen-cuda/examples/proof_carrying_views/Cargo.lock
index 6c51df57bb..da29c7a6e1 100644
--- a/crates/rustc-codegen-cuda/examples/proof_carrying_views/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/proof_carrying_views/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/ptr_copy/Cargo.lock b/crates/rustc-codegen-cuda/examples/ptr_copy/Cargo.lock
index 5fc42300c5..8b4fd05a6f 100644
--- a/crates/rustc-codegen-cuda/examples/ptr_copy/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/ptr_copy/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/ptr_offset_from/Cargo.lock b/crates/rustc-codegen-cuda/examples/ptr_offset_from/Cargo.lock
index 39ec99a09a..47b347c759 100644
--- a/crates/rustc-codegen-cuda/examples/ptr_offset_from/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/ptr_offset_from/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/ptr_without_provenance/Cargo.lock b/crates/rustc-codegen-cuda/examples/ptr_without_provenance/Cargo.lock
index 309b04e6ad..2d38a7471f 100644
--- a/crates/rustc-codegen-cuda/examples/ptr_without_provenance/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/ptr_without_provenance/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/reborrow/Cargo.lock b/crates/rustc-codegen-cuda/examples/reborrow/Cargo.lock
index 62442662a7..8e678f4cab 100644
--- a/crates/rustc-codegen-cuda/examples/reborrow/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/reborrow/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/redux_f32/Cargo.lock b/crates/rustc-codegen-cuda/examples/redux_f32/Cargo.lock
index 6aa124be71..7e41709ce9 100644
--- a/crates/rustc-codegen-cuda/examples/redux_f32/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/redux_f32/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/redux_minmax/Cargo.lock b/crates/rustc-codegen-cuda/examples/redux_minmax/Cargo.lock
index f3b9bee812..7bc91cde87 100644
--- a/crates/rustc-codegen-cuda/examples/redux_minmax/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/redux_minmax/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/redux_sum/Cargo.lock b/crates/rustc-codegen-cuda/examples/redux_sum/Cargo.lock
index 9dc1c7039e..766c66771c 100644
--- a/crates/rustc-codegen-cuda/examples/redux_sum/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/redux_sum/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/ref_index_projections/Cargo.lock b/crates/rustc-codegen-cuda/examples/ref_index_projections/Cargo.lock
index f5f0bc3392..d0b39890a0 100644
--- a/crates/rustc-codegen-cuda/examples/ref_index_projections/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/ref_index_projections/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/ref_operand_mul/Cargo.lock b/crates/rustc-codegen-cuda/examples/ref_operand_mul/Cargo.lock
index 02fa2424b0..c4b8422924 100644
--- a/crates/rustc-codegen-cuda/examples/ref_operand_mul/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/ref_operand_mul/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/repr_align_block_args/Cargo.lock b/crates/rustc-codegen-cuda/examples/repr_align_block_args/Cargo.lock
index 46399531a1..cf54cc0590 100644
--- a/crates/rustc-codegen-cuda/examples/repr_align_block_args/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/repr_align_block_args/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/repr_transparent_abi/Cargo.lock b/crates/rustc-codegen-cuda/examples/repr_transparent_abi/Cargo.lock
index 7074df0e96..efc1cb7ef5 100644
--- a/crates/rustc-codegen-cuda/examples/repr_transparent_abi/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/repr_transparent_abi/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/repr_u32_enum_stride/Cargo.lock b/crates/rustc-codegen-cuda/examples/repr_u32_enum_stride/Cargo.lock
index d451463481..8af3dc77c9 100644
--- a/crates/rustc-codegen-cuda/examples/repr_u32_enum_stride/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/repr_u32_enum_stride/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/row_width_slice/Cargo.lock b/crates/rustc-codegen-cuda/examples/row_width_slice/Cargo.lock
index ae37cf93a1..32fc7e09fb 100644
--- a/crates/rustc-codegen-cuda/examples/row_width_slice/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/row_width_slice/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/rustlantis-smoke/Cargo.lock b/crates/rustc-codegen-cuda/examples/rustlantis-smoke/Cargo.lock
index 4d673e5b07..099612c17a 100644
--- a/crates/rustc-codegen-cuda/examples/rustlantis-smoke/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/rustlantis-smoke/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/scoped_atomic_load_store/Cargo.lock b/crates/rustc-codegen-cuda/examples/scoped_atomic_load_store/Cargo.lock
index 92e2a934ce..ceb0e92cb0 100644
--- a/crates/rustc-codegen-cuda/examples/scoped_atomic_load_store/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/scoped_atomic_load_store/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/select_unpredictable/Cargo.lock b/crates/rustc-codegen-cuda/examples/select_unpredictable/Cargo.lock
index 35519c6b51..2154ad05f8 100644
--- a/crates/rustc-codegen-cuda/examples/select_unpredictable/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/select_unpredictable/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/set_discriminant/Cargo.lock b/crates/rustc-codegen-cuda/examples/set_discriminant/Cargo.lock
index 6f3c3a9b3f..19669085b6 100644
--- a/crates/rustc-codegen-cuda/examples/set_discriminant/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/set_discriminant/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/set_discriminant_niche/Cargo.lock b/crates/rustc-codegen-cuda/examples/set_discriminant_niche/Cargo.lock
index aecbfcee2b..ebc1ac4a4b 100644
--- a/crates/rustc-codegen-cuda/examples/set_discriminant_niche/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/set_discriminant_niche/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/shared_debug/Cargo.lock b/crates/rustc-codegen-cuda/examples/shared_debug/Cargo.lock
index 77bc354e7a..ac0e135aff 100644
--- a/crates/rustc-codegen-cuda/examples/shared_debug/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/shared_debug/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/shared_memory_helper/Cargo.lock b/crates/rustc-codegen-cuda/examples/shared_memory_helper/Cargo.lock
index 4ad967108d..c6cf197270 100644
--- a/crates/rustc-codegen-cuda/examples/shared_memory_helper/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/shared_memory_helper/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/shared_ptr_transmute/Cargo.lock b/crates/rustc-codegen-cuda/examples/shared_ptr_transmute/Cargo.lock
index d02cf251c1..e77989e20c 100644
--- a/crates/rustc-codegen-cuda/examples/shared_ptr_transmute/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/shared_ptr_transmute/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/shared_slice_unsize/Cargo.lock b/crates/rustc-codegen-cuda/examples/shared_slice_unsize/Cargo.lock
index bb356c6539..96634d6a5c 100644
--- a/crates/rustc-codegen-cuda/examples/shared_slice_unsize/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/shared_slice_unsize/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/shared_stage_select/Cargo.lock b/crates/rustc-codegen-cuda/examples/shared_stage_select/Cargo.lock
index 44f7af58e3..2d9a1fd2d9 100644
--- a/crates/rustc-codegen-cuda/examples/shared_stage_select/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/shared_stage_select/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/sharedmem/Cargo.lock b/crates/rustc-codegen-cuda/examples/sharedmem/Cargo.lock
index a216018f98..ae25ba8ed5 100644
--- a/crates/rustc-codegen-cuda/examples/sharedmem/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/sharedmem/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/shuffle_64/Cargo.lock b/crates/rustc-codegen-cuda/examples/shuffle_64/Cargo.lock
index 30945aa733..4c9c1327ed 100644
--- a/crates/rustc-codegen-cuda/examples/shuffle_64/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/shuffle_64/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/slice_from_raw_parts_cast/Cargo.lock b/crates/rustc-codegen-cuda/examples/slice_from_raw_parts_cast/Cargo.lock
index b68ba2a55c..dad51c9991 100644
--- a/crates/rustc-codegen-cuda/examples/slice_from_raw_parts_cast/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/slice_from_raw_parts_cast/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/slice_get_mut/Cargo.lock b/crates/rustc-codegen-cuda/examples/slice_get_mut/Cargo.lock
index bb298f071e..262cc7322c 100644
--- a/crates/rustc-codegen-cuda/examples/slice_get_mut/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/slice_get_mut/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/slice_index_bounds_check/Cargo.lock b/crates/rustc-codegen-cuda/examples/slice_index_bounds_check/Cargo.lock
index b3f901c9d1..cc2ec54bd1 100644
--- a/crates/rustc-codegen-cuda/examples/slice_index_bounds_check/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/slice_index_bounds_check/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/slice_iterators/Cargo.lock b/crates/rustc-codegen-cuda/examples/slice_iterators/Cargo.lock
index 980220cb65..a92a898d6b 100644
--- a/crates/rustc-codegen-cuda/examples/slice_iterators/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/slice_iterators/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/slice_reslice/Cargo.lock b/crates/rustc-codegen-cuda/examples/slice_reslice/Cargo.lock
index cce31d92a6..ca92447445 100644
--- a/crates/rustc-codegen-cuda/examples/slice_reslice/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/slice_reslice/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/small_type_ffi_test/Cargo.lock b/crates/rustc-codegen-cuda/examples/small_type_ffi_test/Cargo.lock
index e824c71ea7..8319aeb2e8 100644
--- a/crates/rustc-codegen-cuda/examples/small_type_ffi_test/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/small_type_ffi_test/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/sparse_mma_fp8/Cargo.lock b/crates/rustc-codegen-cuda/examples/sparse_mma_fp8/Cargo.lock
index 8fd79030f8..ce07e8fd03 100644
--- a/crates/rustc-codegen-cuda/examples/sparse_mma_fp8/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/sparse_mma_fp8/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/sparse_mma_ordered_float/Cargo.lock b/crates/rustc-codegen-cuda/examples/sparse_mma_ordered_float/Cargo.lock
index f158136638..8fbe85e968 100644
--- a/crates/rustc-codegen-cuda/examples/sparse_mma_ordered_float/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/sparse_mma_ordered_float/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/split_at_mut/Cargo.lock b/crates/rustc-codegen-cuda/examples/split_at_mut/Cargo.lock
index c0ad4d412a..e0ecb12ca0 100644
--- a/crates/rustc-codegen-cuda/examples/split_at_mut/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/split_at_mut/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/standalone_device_fn/Cargo.lock b/crates/rustc-codegen-cuda/examples/standalone_device_fn/Cargo.lock
index 38b458f4ba..e84056bc80 100644
--- a/crates/rustc-codegen-cuda/examples/standalone_device_fn/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/standalone_device_fn/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/static_slice_addend/Cargo.lock b/crates/rustc-codegen-cuda/examples/static_slice_addend/Cargo.lock
index e437e5f7b9..df6a319513 100644
--- a/crates/rustc-codegen-cuda/examples/static_slice_addend/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/static_slice_addend/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/static_slice_unsize/Cargo.lock b/crates/rustc-codegen-cuda/examples/static_slice_unsize/Cargo.lock
index 45724fc295..82eea5f095 100644
--- a/crates/rustc-codegen-cuda/examples/static_slice_unsize/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/static_slice_unsize/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/step_by/Cargo.lock b/crates/rustc-codegen-cuda/examples/step_by/Cargo.lock
index 4afbc84252..c84a13be9e 100644
--- a/crates/rustc-codegen-cuda/examples/step_by/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/step_by/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/struct_constant_provenance/Cargo.lock b/crates/rustc-codegen-cuda/examples/struct_constant_provenance/Cargo.lock
index 49b3d30f38..96e8923049 100644
--- a/crates/rustc-codegen-cuda/examples/struct_constant_provenance/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/struct_constant_provenance/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/struct_field_layout/Cargo.lock b/crates/rustc-codegen-cuda/examples/struct_field_layout/Cargo.lock
index 65c9d79bb7..311ed272de 100644
--- a/crates/rustc-codegen-cuda/examples/struct_field_layout/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/struct_field_layout/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/subslice_projection/Cargo.lock b/crates/rustc-codegen-cuda/examples/subslice_projection/Cargo.lock
index 1611b94d24..fd05e78901 100644
--- a/crates/rustc-codegen-cuda/examples/subslice_projection/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/subslice_projection/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/switchint_128bit/Cargo.lock b/crates/rustc-codegen-cuda/examples/switchint_128bit/Cargo.lock
index 89afb020f2..a9cceda1a5 100644
--- a/crates/rustc-codegen-cuda/examples/switchint_128bit/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/switchint_128bit/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/swizzle_smem/Cargo.lock b/crates/rustc-codegen-cuda/examples/swizzle_smem/Cargo.lock
index 9ccf92a128..a6fe4b6981 100644
--- a/crates/rustc-codegen-cuda/examples/swizzle_smem/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/swizzle_smem/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/tcgen05/Cargo.lock b/crates/rustc-codegen-cuda/examples/tcgen05/Cargo.lock
index 03471ce7f3..491bfe4da6 100644
--- a/crates/rustc-codegen-cuda/examples/tcgen05/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/tcgen05/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/tcgen05_matmul/Cargo.lock b/crates/rustc-codegen-cuda/examples/tcgen05_matmul/Cargo.lock
index a674d94714..982d1e0373 100644
--- a/crates/rustc-codegen-cuda/examples/tcgen05_matmul/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/tcgen05_matmul/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/thread_runs/Cargo.lock b/crates/rustc-codegen-cuda/examples/thread_runs/Cargo.lock
index 74af911dd3..f4e2a6738f 100644
--- a/crates/rustc-codegen-cuda/examples/thread_runs/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/thread_runs/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/tiled_gemm/Cargo.lock b/crates/rustc-codegen-cuda/examples/tiled_gemm/Cargo.lock
index 09bd17266f..78ad468a46 100644
--- a/crates/rustc-codegen-cuda/examples/tiled_gemm/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/tiled_gemm/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/tma_copy/Cargo.lock b/crates/rustc-codegen-cuda/examples/tma_copy/Cargo.lock
index ce15b4216f..dc26535205 100644
--- a/crates/rustc-codegen-cuda/examples/tma_copy/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/tma_copy/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/tma_multicast/Cargo.lock b/crates/rustc-codegen-cuda/examples/tma_multicast/Cargo.lock
index 973fd6f0e9..9f7dddf840 100644
--- a/crates/rustc-codegen-cuda/examples/tma_multicast/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/tma_multicast/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/toolchain_pairing_smoke/Cargo.lock b/crates/rustc-codegen-cuda/examples/toolchain_pairing_smoke/Cargo.lock
index f3c613dc02..01f7af5618 100644
--- a/crates/rustc-codegen-cuda/examples/toolchain_pairing_smoke/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/toolchain_pairing_smoke/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/transmute_scalar_struct/Cargo.lock b/crates/rustc-codegen-cuda/examples/transmute_scalar_struct/Cargo.lock
index 6e55450a3b..e78115db18 100644
--- a/crates/rustc-codegen-cuda/examples/transmute_scalar_struct/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/transmute_scalar_struct/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/tuple_array_provenance/Cargo.lock b/crates/rustc-codegen-cuda/examples/tuple_array_provenance/Cargo.lock
index fb0420658c..7e9c724cc2 100644
--- a/crates/rustc-codegen-cuda/examples/tuple_array_provenance/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/tuple_array_provenance/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/tuple_constant_provenance/Cargo.lock b/crates/rustc-codegen-cuda/examples/tuple_constant_provenance/Cargo.lock
index 06d48f94d2..d4b7e8be2e 100644
--- a/crates/rustc-codegen-cuda/examples/tuple_constant_provenance/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/tuple_constant_provenance/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/tuple_field_lookup/Cargo.lock b/crates/rustc-codegen-cuda/examples/tuple_field_lookup/Cargo.lock
index c9ce6633a5..1141b5e5bc 100644
--- a/crates/rustc-codegen-cuda/examples/tuple_field_lookup/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/tuple_field_lookup/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/tuple_return/Cargo.lock b/crates/rustc-codegen-cuda/examples/tuple_return/Cargo.lock
index 3a9dc29402..3cc9d0a399 100644
--- a/crates/rustc-codegen-cuda/examples/tuple_return/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/tuple_return/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/unaligned_memory/Cargo.lock b/crates/rustc-codegen-cuda/examples/unaligned_memory/Cargo.lock
index 88930f17b5..713644f115 100644
--- a/crates/rustc-codegen-cuda/examples/unaligned_memory/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/unaligned_memory/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/unchecked_indexing/Cargo.lock b/crates/rustc-codegen-cuda/examples/unchecked_indexing/Cargo.lock
index bfc405a3cd..9191210df7 100644
--- a/crates/rustc-codegen-cuda/examples/unchecked_indexing/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/unchecked_indexing/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/uninit_const/Cargo.lock b/crates/rustc-codegen-cuda/examples/uninit_const/Cargo.lock
index 38cbad6214..48a00250c6 100644
--- a/crates/rustc-codegen-cuda/examples/uninit_const/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/uninit_const/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/union_aggregate/Cargo.lock b/crates/rustc-codegen-cuda/examples/union_aggregate/Cargo.lock
index 0927d182dc..0a84a1a14e 100644
--- a/crates/rustc-codegen-cuda/examples/union_aggregate/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/union_aggregate/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/unroll_bounds_check/Cargo.lock b/crates/rustc-codegen-cuda/examples/unroll_bounds_check/Cargo.lock
index 1080afd240..9d874e6bfd 100644
--- a/crates/rustc-codegen-cuda/examples/unroll_bounds_check/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/unroll_bounds_check/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/unroll_smoke/Cargo.lock b/crates/rustc-codegen-cuda/examples/unroll_smoke/Cargo.lock
index 3fd89f5d34..c470936692 100644
--- a/crates/rustc-codegen-cuda/examples/unroll_smoke/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/unroll_smoke/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/vecadd/Cargo.lock b/crates/rustc-codegen-cuda/examples/vecadd/Cargo.lock
index dde1a3c108..986cfc7a4e 100644
--- a/crates/rustc-codegen-cuda/examples/vecadd/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/vecadd/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/vectorization/Cargo.lock b/crates/rustc-codegen-cuda/examples/vectorization/Cargo.lock
index 2a9c1e4734..56d670d22c 100644
--- a/crates/rustc-codegen-cuda/examples/vectorization/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/vectorization/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/volatile_memory/Cargo.lock b/crates/rustc-codegen-cuda/examples/volatile_memory/Cargo.lock
index c6dc438147..e02f469f55 100644
--- a/crates/rustc-codegen-cuda/examples/volatile_memory/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/volatile_memory/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/warp_reduce/Cargo.lock b/crates/rustc-codegen-cuda/examples/warp_reduce/Cargo.lock
index ac76620b90..2f7694b676 100644
--- a/crates/rustc-codegen-cuda/examples/warp_reduce/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/warp_reduce/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/warp_sums/Cargo.lock b/crates/rustc-codegen-cuda/examples/warp_sums/Cargo.lock
index f299324612..23e2ca3361 100644
--- a/crates/rustc-codegen-cuda/examples/warp_sums/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/warp_sums/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/wgmma/Cargo.lock b/crates/rustc-codegen-cuda/examples/wgmma/Cargo.lock
index 24394dceac..4beb6cb947 100644
--- a/crates/rustc-codegen-cuda/examples/wgmma/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/wgmma/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/examples/wgmma_mma_bf16/Cargo.lock b/crates/rustc-codegen-cuda/examples/wgmma_mma_bf16/Cargo.lock
index c15f1983ba..aea18e8a88 100644
--- a/crates/rustc-codegen-cuda/examples/wgmma_mma_bf16/Cargo.lock
+++ b/crates/rustc-codegen-cuda/examples/wgmma_mma_bf16/Cargo.lock
@@ -121,6 +121,7 @@ version = "0.2.1"
 dependencies = [
  "libnvvm-sys",
  "nvjitlink-sys",
+ "ptx-parse",
  "serde",
  "sha2",
  "thiserror",
diff --git a/crates/rustc-codegen-cuda/src/lib.rs b/crates/rustc-codegen-cuda/src/lib.rs
index 8deaf0cf0f..bec95eb467 100644
--- a/crates/rustc-codegen-cuda/src/lib.rs
+++ b/crates/rustc-codegen-cuda/src/lib.rs
@@ -863,11 +863,17 @@ fn write_device_artifact_object(
     materialization_request: Option<materialize::MaterializationRequest>,
 ) -> Result<PathBuf, Box<dyn std::error::Error>> {
     let bundle_name = std::env::var("CARGO_PKG_NAME").unwrap_or_else(|_| output_name.to_string());
+    let expected_kernels = functions
+        .iter()
+        .filter(|function| function.is_kernel)
+        .map(|function| function.export_name.as_str())
+        .collect::<Vec<_>>();
     let materialized_artifact = materialize_artifact_for_embedding(
         materialization_request,
         &bundle_name,
         result,
         artifact,
+        &expected_kernels,
     )?;
     if let Some(materialized) = materialized_artifact.as_ref() {
         emit_launch_bounds_spill_warnings(tcx, result, functions, &materialized.resource_usage);
@@ -1000,6 +1006,7 @@ fn materialize_artifact_for_embedding(
     bundle_name: &str,
     result: &device_codegen::DeviceCodegenResult,
     artifact: &device_codegen::DeviceCodegenArtifact,
+    expected_kernels: &[&str],
 ) -> Result<Option<MaterializedDeviceArtifact>, Box<dyn std::error::Error>> {
     let Some(request) = request else {
         return Ok(None);
@@ -1019,6 +1026,7 @@ fn materialize_artifact_for_embedding(
             &result.target,
             result.allow_fma_contraction,
             debug_policy,
+            expected_kernels,
         )?,
         device_codegen::DeviceCodegenArtifactKind::Ltoir => materialize::ltoir_to_cubin(
             request,
@@ -1027,6 +1035,7 @@ fn materialize_artifact_for_embedding(
             &result.target,
             result.allow_fma_contraction,
             debug_policy,
+            expected_kernels,
         )?,
         device_codegen::DeviceCodegenArtifactKind::Ptx => {
             return Err(Box::new(materialize::MaterializeError::PtxInput));
diff --git a/crates/rustc-codegen-cuda/src/materialize.rs b/crates/rustc-codegen-cuda/src/materialize.rs
index 387c4ce4f6..52cef5ea7d 100644
--- a/crates/rustc-codegen-cuda/src/materialize.rs
+++ b/crates/rustc-codegen-cuda/src/materialize.rs
@@ -197,10 +197,16 @@ pub(crate) fn nvvm_ir_to_cubin(
     target: &str,
     allow_fma_contraction: bool,
     debug_policy: DebugPolicy,
+    expected_kernels: &[&str],
 ) -> Result<MaterializedCubin, MaterializeError> {
     let options = options(target, allow_fma_contraction, debug_policy)?;
     let finalizer = checked_finalizer(request)?;
-    let report = finalizer.materialize_nvvm_ir_with_report(module_name, nvvm_ir, &options)?;
+    let report = finalizer.materialize_nvvm_ir_with_report(
+        module_name,
+        nvvm_ir,
+        &options,
+        expected_kernels,
+    )?;
     Ok(MaterializedCubin {
         bytes: report.image,
         resource_usage: report.resource_usage,
@@ -214,6 +220,7 @@ pub(crate) fn ltoir_to_cubin(
     target: &str,
     allow_fma_contraction: bool,
     debug_policy: DebugPolicy,
+    expected_kernels: &[&str],
 ) -> Result<MaterializedCubin, MaterializeError> {
     let options = options(target, allow_fma_contraction, debug_policy)?;
     let finalizer = checked_finalizer(request)?;
@@ -221,6 +228,7 @@ pub(crate) fn ltoir_to_cubin(
         &[NamedInput::new(module_name, ltoir)],
         &options,
         FinalizerOutput::Cubin,
+        expected_kernels,
     )?;
     Ok(MaterializedCubin {
         bytes: report.image,
diff --git a/cuda-oxide-book/gpu-programming/error-handling-and-debugging.md b/cuda-oxide-book/gpu-programming/error-handling-and-debugging.md
index bed8e6ccd2..ac2d15019f 100644
--- a/cuda-oxide-book/gpu-programming/error-handling-and-debugging.md
+++ b/cuda-oxide-book/gpu-programming/error-handling-and-debugging.md
@@ -467,9 +467,11 @@ It keeps ordinary multiply and add/subtract operations separate, with one
 rounding per operation. Explicit fused operations such as `f32::mul_add`
 remain fused.
 
-NVVM IR and LTOIR builds also produce `.options` and versioned `.target`
-sidecars. Keep both files with the artifact so later libNVVM and nvJitLink
-steps preserve the same FMA policy.
+NVVM IR and LTOIR builds also produce `.options`, `.kernels` and versioned
+`.target` sidecars. Keep them with the artifact: later libNVVM and nvJitLink
+steps preserve the same FMA policy, and check that the finalized image still
+defines every kernel listed in `.kernels` (nvJitLink can report success while
+dropping a module's kernels).
 
 Run `memcheck` first when investigating memory safety. The other tools are
 complementary and do not perform full memory-access checking.
@@ -531,7 +533,7 @@ output mode is active in the environment (`CUDA_OXIDE_MATERIALIZE_CUBIN` or
 
 `clean` deletes project-local Cargo `target/` directories and generated
 cuda-oxide device artifacts (`.ptx`, `.ll`, `.opt.ll`, `.ltoir`, `.cubin`,
-plus the `.target` / `.options` / `.cubin.target` sidecars) under the current
+plus the `.target` / `.options` / `.kernels` / `.cubin.target` sidecars) under the current
 workspace or standalone project. It refuses (with an error) to remove a
 symlinked `target/` directory or artifact, and it does **not** wipe the
 shared codegen backend cache at `~/.cargo/cuda-oxide/`.