diff --git a/cuda-oxide/Cargo.lock b/cuda-oxide/Cargo.lock
index cdef05af07..c4eb19c779 100644
--- a/cuda-oxide/Cargo.lock
+++ b/cuda-oxide/Cargo.lock
@@ -423,6 +423,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
  "trybuild",
 ]
 
diff --git a/cuda-oxide/crates/cuda-device/Cargo.toml b/cuda-oxide/crates/cuda-device/Cargo.toml
index 00c2a1b1f0..14145ee9a4 100644
--- a/cuda-oxide/crates/cuda-device/Cargo.toml
+++ b/cuda-oxide/crates/cuda-device/Cargo.toml
@@ -16,6 +16,12 @@ readme = "README.md"
 # `host` explicitly, so any build that loads or launches a module still has it.
 cuda-macros = { workspace = true }
 
+[build-dependencies]
+# Shares one arch parser and one floor rule with the backend, so the cfg
+# ladder cannot drift from the target the PTX is actually built for. This
+# crate has no dependencies of its own, so it costs the kernel graph nothing.
+cuda-target-spec = { workspace = true }
+
 [dev-dependencies]
 trybuild = "1"
 
diff --git a/cuda-oxide/crates/cuda-device/build.rs b/cuda-oxide/crates/cuda-device/build.rs
new file mode 100644
index 0000000000..8b4579bb99
--- /dev/null
+++ b/cuda-oxide/crates/cuda-device/build.rs
@@ -0,0 +1,45 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+//! Expose the target's minimum compute capability as device cfgs.
+//!
+//! An explicit target sets the floor; an advisory device hint can only lower
+//! the default. The shared resolver keeps these cfgs consistent with backend
+//! selection. Unsupported intrinsic branches disappear before MIR lowering.
+
+fn main() {
+    // The two variables the backend already reads to pick a target.
+    println!("cargo::rerun-if-env-changed=CUDA_OXIDE_TARGET");
+    println!("cargo::rerun-if-env-changed=CUDA_OXIDE_DEVICE_ARCH");
+
+    // Declare the whole ladder, not just the rungs set for this build, so an
+    // unset rung is a false `cfg` rather than an `unexpected_cfgs` warning.
+    let values = cuda_target_spec::SM_FLOOR_LADDER
+        .iter()
+        .map(|floor| format!("\"{floor}\""))
+        .collect::<Vec<_>>()
+        .join(", ");
+    println!("cargo::rustc-check-cfg=cfg(cuda_oxide_sm_at_least, values({values}))");
+
+    let target = std::env::var("CUDA_OXIDE_TARGET").ok();
+    // Read the hint with the backend's grammar. One that does not parse is
+    // left out: the backend rejects it, so nothing is built against it.
+    let device_arch = std::env::var("CUDA_OXIDE_DEVICE_ARCH")
+        .ok()
+        .and_then(|value| value.parse::<cuda_target_spec::DeviceArch>().ok());
+    match cuda_target_spec::resolve_sm_floor(
+        target.as_deref(),
+        device_arch.as_ref().map(AsRef::as_ref),
+    ) {
+        Ok(floor) => {
+            for rung in cuda_target_spec::sm_floors_at_most(floor) {
+                println!("cargo::rustc-cfg=cuda_oxide_sm_at_least=\"{rung}\"");
+            }
+        }
+        // Fail the build rather than guess. Guessing high emits instructions
+        // the target cannot run; guessing low silently drops the faster form.
+        Err(error) => println!("cargo::error={error}"),
+    }
+}
diff --git a/cuda-oxide/crates/cuda-device/src/cooperative_groups.rs b/cuda-oxide/crates/cuda-device/src/cooperative_groups.rs
index 07d50d592f..41e31baea1 100644
--- a/cuda-oxide/crates/cuda-device/src/cooperative_groups.rs
+++ b/cuda-oxide/crates/cuda-device/src/cooperative_groups.rs
@@ -822,6 +822,39 @@ pub mod ops {
         fn identity() -> T;
         /// Binary combiner. Must be associative and commutative.
         fn combine(a: T, b: T) -> T;
+
+        /// Whether this operation has a full-warp instruction at the target floor.
+        /// Other operation/type pairs use the shuffle implementation.
+        const HAS_FULL_WARP_FORM: bool = false;
+
+        /// Reduce the participating warp. Called only when `HAS_FULL_WARP_FORM`
+        /// is true and the tile contains all 32 lanes. An implementation that
+        /// sets `HAS_FULL_WARP_FORM` must override it: this default returns
+        /// `value` unreduced.
+        #[inline(always)]
+        fn reduce_full_warp(_mask: u32, value: T) -> T {
+            value
+        }
+    }
+
+    // Exclude unsupported intrinsics before MIR and feature detection run.
+    macro_rules! redux_full_warp {
+        ($ty:ty, |$mask:ident, $value:ident| $call:expr) => {
+            const HAS_FULL_WARP_FORM: bool = cfg!(cuda_oxide_sm_at_least = "80");
+
+            // Keep the intrinsic call as the return expression for device lowering.
+            #[cfg(cuda_oxide_sm_at_least = "80")]
+            #[inline(always)]
+            fn reduce_full_warp($mask: u32, $value: $ty) -> $ty {
+                $call
+            }
+
+            #[cfg(not(cuda_oxide_sm_at_least = "80"))]
+            #[inline(always)]
+            fn reduce_full_warp(_mask: u32, value: $ty) -> $ty {
+                value
+            }
+        };
     }
 
     /// Sum: wrapping integer addition; IEEE-754 add for floats.
@@ -858,6 +891,7 @@ pub mod ops {
         fn combine(a: u32, b: u32) -> u32 {
             a.wrapping_add(b)
         }
+        redux_full_warp!(u32, |m, v| crate::warp::redux_sync_add(m, v));
     }
     impl ReduceOp<u32> for Min {
         #[inline(always)]
@@ -868,6 +902,7 @@ pub mod ops {
         fn combine(a: u32, b: u32) -> u32 {
             if a < b { a } else { b }
         }
+        redux_full_warp!(u32, |m, v| crate::warp::redux_sync_min_u32(m, v));
     }
     impl ReduceOp<u32> for Max {
         #[inline(always)]
@@ -878,6 +913,7 @@ pub mod ops {
         fn combine(a: u32, b: u32) -> u32 {
             if a > b { a } else { b }
         }
+        redux_full_warp!(u32, |m, v| crate::warp::redux_sync_max_u32(m, v));
     }
     impl ReduceOp<u32> for BitAnd {
         #[inline(always)]
@@ -888,6 +924,7 @@ pub mod ops {
         fn combine(a: u32, b: u32) -> u32 {
             a & b
         }
+        redux_full_warp!(u32, |m, v| crate::warp::redux_sync_and(m, v));
     }
     impl ReduceOp<u32> for BitOr {
         #[inline(always)]
@@ -898,6 +935,7 @@ pub mod ops {
         fn combine(a: u32, b: u32) -> u32 {
             a | b
         }
+        redux_full_warp!(u32, |m, v| crate::warp::redux_sync_or(m, v));
     }
     impl ReduceOp<u32> for BitXor {
         #[inline(always)]
@@ -908,6 +946,7 @@ pub mod ops {
         fn combine(a: u32, b: u32) -> u32 {
             a ^ b
         }
+        redux_full_warp!(u32, |m, v| crate::warp::redux_sync_xor(m, v));
     }
 
     // --- i32 ---
@@ -920,6 +959,7 @@ pub mod ops {
         fn combine(a: i32, b: i32) -> i32 {
             a.wrapping_add(b)
         }
+        redux_full_warp!(i32, |m, v| crate::warp::redux_sync_add(m, v as u32) as i32);
     }
     impl ReduceOp<i32> for Min {
         #[inline(always)]
@@ -930,6 +970,7 @@ pub mod ops {
         fn combine(a: i32, b: i32) -> i32 {
             if a < b { a } else { b }
         }
+        redux_full_warp!(i32, |m, v| crate::warp::redux_sync_min_i32(m, v));
     }
     impl ReduceOp<i32> for Max {
         #[inline(always)]
@@ -940,6 +981,7 @@ pub mod ops {
         fn combine(a: i32, b: i32) -> i32 {
             if a > b { a } else { b }
         }
+        redux_full_warp!(i32, |m, v| crate::warp::redux_sync_max_i32(m, v));
     }
 
     // --- f32 ---
@@ -1048,13 +1090,13 @@ impl WarpShuffle for f32 {
 /// bit-identical results. That is the primitive in isolation — a kernel that
 /// reduces once after a memory-bound pass will see far less.
 ///
-/// This function does not select that form for you: `redux.sync` will not
-/// assemble below `sm_80`, and device code has no way to ask what target it is
-/// being compiled for (see
-/// [#811](https://github.com/NVlabs/cuda-oxide/issues/811)). Call it directly
-/// when you know the target is Ampere+, gating on
-/// `CudaContext::compute_capability` as the `redux_sum` example does. Floats
-/// keep this butterfly — there is no `f32` `redux` before `sm_100`.
+/// Full-warp integer reductions use `redux.sync` when the target floor is
+/// `sm_80` or newer. Floats, smaller tiles and lower floors use shuffles.
+///
+/// `CUDA_OXIDE_TARGET` sets the floor. An advisory `CUDA_OXIDE_DEVICE_ARCH`
+/// can lower the default `sm_80` floor, but cannot raise it: incompatible
+/// kernels may target another architecture. The build script and backend
+/// share this rule. Use `--arch sm_75` to select the shuffle implementation.
 ///
 /// `N` is the tile size (1, 2, 4, 8, 16, or 32) — already validated by
 /// [`ThreadBlock::tiled_partition`] at construction time.
@@ -1082,6 +1124,11 @@ where
     T: WarpShuffle,
     Op: ops::ReduceOp<T>,
 {
+    // The const conditions remove this branch for floats and smaller tiles.
+    if N == 32 && Op::HAS_FULL_WARP_FORM {
+        return Op::reduce_full_warp(tile.mask(), value);
+    }
+
     let mut acc = value;
     let mut delta: u32 = N >> 1;
     while delta > 0 {
@@ -1520,6 +1567,7 @@ where
 #[cfg(test)]
 mod tests {
     use super::CoalescedThreads;
+    use super::ops::{BitAnd, BitOr, BitXor, Max, Min, ReduceOp, Sum};
 
     #[test]
     fn pack_lanes_uses_group_relative_bit_positions() {
@@ -1546,4 +1594,84 @@ mod tests {
             );
         }
     }
+
+    /// Every integer pair `warp_reduce` can be asked for, and whether the
+    /// Ampere rung should give it a single-instruction form.
+    ///
+    /// A `const` hook makes this checkable from the host with no device call,
+    /// which matters because the crate is `no_std` and the `redux_sync_*`
+    /// stubs are device functions.
+    const INTEGER_PAIRS: [(&str, bool); 9] = [
+        ("u32 Sum", <Sum as ReduceOp<u32>>::HAS_FULL_WARP_FORM),
+        ("u32 Min", <Min as ReduceOp<u32>>::HAS_FULL_WARP_FORM),
+        ("u32 Max", <Max as ReduceOp<u32>>::HAS_FULL_WARP_FORM),
+        ("u32 BitAnd", <BitAnd as ReduceOp<u32>>::HAS_FULL_WARP_FORM),
+        ("u32 BitOr", <BitOr as ReduceOp<u32>>::HAS_FULL_WARP_FORM),
+        ("u32 BitXor", <BitXor as ReduceOp<u32>>::HAS_FULL_WARP_FORM),
+        ("i32 Sum", <Sum as ReduceOp<i32>>::HAS_FULL_WARP_FORM),
+        ("i32 Min", <Min as ReduceOp<i32>>::HAS_FULL_WARP_FORM),
+        ("i32 Max", <Max as ReduceOp<i32>>::HAS_FULL_WARP_FORM),
+    ];
+
+    /// All nine are wired, not just the one that happened to get tested.
+    #[cfg(cuda_oxide_sm_at_least = "80")]
+    #[test]
+    fn every_integer_reduction_has_the_single_instruction_form_at_the_ampere_floor() {
+        for (name, wired) in INTEGER_PAIRS {
+            assert!(wired, "{name} lost its redux.sync form");
+        }
+    }
+
+    /// Below the floor every one of them goes dark, so nothing can emit
+    /// `redux.sync` into a module that cannot assemble it.
+    #[cfg(not(cuda_oxide_sm_at_least = "80"))]
+    #[test]
+    fn no_integer_reduction_claims_the_form_below_the_ampere_floor() {
+        for (name, wired) in INTEGER_PAIRS {
+            assert!(!wired, "{name} claims redux.sync below sm_80");
+        }
+    }
+
+    /// There is no `f32` `redux` before `sm_100`, so floats keep the butterfly
+    /// at every floor. The control a broadened hook would break: claiming the
+    /// form here would name an instruction that does not exist for the type.
+    #[test]
+    fn floats_never_claim_the_single_instruction_form() {
+        // `const` blocks: these are compile-time facts, so a regression is a
+        // build failure rather than a test failure.
+        const {
+            assert!(!<Sum as ReduceOp<f32>>::HAS_FULL_WARP_FORM);
+            assert!(!<Min as ReduceOp<f32>>::HAS_FULL_WARP_FORM);
+            assert!(!<Max as ReduceOp<f32>>::HAS_FULL_WARP_FORM);
+        }
+    }
+
+    /// The default hook must stay inert: it is reached only when the `const`
+    /// is false, and must not be mistaken for a reduction.
+    #[test]
+    fn the_default_hook_is_the_identity_on_its_input() {
+        assert_eq!(<Sum as ReduceOp<f32>>::reduce_full_warp(u32::MAX, 3.5), 3.5);
+    }
+
+    /// The build script sets every rung at or below the resolved capability,
+    /// and device code relies on that: an `sm_86` build must still answer yes
+    /// to "at least sm_80".
+    #[test]
+    fn the_floor_ladder_reaching_this_crate_is_cumulative() {
+        let rungs = [
+            cfg!(cuda_oxide_sm_at_least = "70"),
+            cfg!(cuda_oxide_sm_at_least = "75"),
+            cfg!(cuda_oxide_sm_at_least = "80"),
+            cfg!(cuda_oxide_sm_at_least = "86"),
+            cfg!(cuda_oxide_sm_at_least = "90"),
+            cfg!(cuda_oxide_sm_at_least = "100"),
+            cfg!(cuda_oxide_sm_at_least = "120"),
+        ];
+        // Once false, every higher rung is false: no gaps in the ladder.
+        let first_unset = rungs.iter().position(|set| !set).unwrap_or(rungs.len());
+        assert!(
+            rungs[first_unset..].iter().all(|set| !set),
+            "floor ladder has a gap: {rungs:?}"
+        );
+    }
 }
diff --git a/cuda-oxide/crates/cuda-macros/tests/device-only/Cargo.lock b/cuda-oxide/crates/cuda-macros/tests/device-only/Cargo.lock
index 4a78f429b9..ed8b75b3fc 100644
--- a/cuda-oxide/crates/cuda-macros/tests/device-only/Cargo.lock
+++ b/cuda-oxide/crates/cuda-macros/tests/device-only/Cargo.lock
@@ -7,6 +7,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
@@ -20,6 +21,10 @@ dependencies = [
  "syn 2.0.119",
 ]
 
+[[package]]
+name = "cuda-target-spec"
+version = "0.1.0"
+
 [[package]]
 name = "device-only-kernels"
 version = "0.0.0"
diff --git a/cuda-oxide/crates/cuda-oxide-codegen/src/export.rs b/cuda-oxide/crates/cuda-oxide-codegen/src/export.rs
index 12fe5ecda2..40eef02b00 100644
--- a/cuda-oxide/crates/cuda-oxide-codegen/src/export.rs
+++ b/cuda-oxide/crates/cuda-oxide-codegen/src/export.rs
@@ -7,8 +7,10 @@ use crate::error::PipelineError;
 use crate::generated::GeneratedModuleRequirements;
 use crate::target::{
     DetectedFeatures, arch_satisfies, generated_ptx_isa_requirement, generated_target_satisfied,
-    select_target_with_generated, validate_generated_target, validate_target_features,
+    honour_sm_floor, select_target_with_generated, validate_generated_target,
+    validate_target_features,
 };
+use cuda_target_spec::resolve_sm_floor;
 use libnvvm_sys::CudaArch;
 use llvm_export::export::{
     DebugKind, DeviceExternType, ExportBackendConfig, FunctionLocalStaticPlacement, NvvmIrDialect,
@@ -82,6 +84,16 @@ impl llvm_export::export::AsDeviceExtern for DeviceExternDecl {
     }
 }
 
+/// Resolve the shared device cfg floor when no target was explicitly selected.
+fn sm_floor_for_selection(device_arch_hint: Option<&CudaArch>) -> Result<u32, PipelineError> {
+    resolve_sm_floor(None, device_arch_hint).map_err(|error| PipelineError::TargetSelection {
+        target: String::new(),
+        reason: format!(
+            "cannot resolve the compute-capability floor device code is compiled against: {error}"
+        ),
+    })
+}
+
 // mir-importer pipeline plumbing; not part of the frontend contract.
 #[doc(hidden)]
 #[cfg(test)]
@@ -153,6 +165,12 @@ pub(crate) fn resolve_nvvm_target_with_generated(
         }
         let target =
             select_target_with_generated(features, generated).map_err(PipelineError::Export)?;
+        // No explicit target reaches here, so this is the shared default
+        // floor; it is resolved rather than restated so the ladder device
+        // code was built against and the target selected here cannot drift.
+        let floor = sm_floor_for_selection(device_arch_hint)?;
+        let target =
+            honour_sm_floor(target, floor, features, generated).map_err(PipelineError::Export)?;
         return Ok(target);
     }
 
@@ -166,6 +184,9 @@ pub(crate) fn resolve_nvvm_target_with_generated(
     if !generated.is_empty() {
         let target = select_target_with_generated(DetectedFeatures::Basic, generated)
             .map_err(PipelineError::Export)?;
+        let floor = sm_floor_for_selection(device_arch_hint)?;
+        let target = honour_sm_floor(target, floor, DetectedFeatures::Basic, generated)
+            .map_err(PipelineError::Export)?;
         return Ok(target);
     }
 
diff --git a/cuda-oxide/crates/cuda-oxide-codegen/src/target/detect.rs b/cuda-oxide/crates/cuda-oxide-codegen/src/target/detect.rs
index 20ce82e87a..2bfff0af70 100644
--- a/cuda-oxide/crates/cuda-oxide-codegen/src/target/detect.rs
+++ b/cuda-oxide/crates/cuda-oxide-codegen/src/target/detect.rs
@@ -5,6 +5,7 @@
 
 use super::features::{DetectedFeatures, ModuleRequirements, PtxIsaRequirement};
 use crate::error::PipelineError;
+use std::borrow::Cow;
 use std::path::Path;
 
 pub(super) fn contains_wgmma_features(contents: &str) -> bool {
@@ -95,11 +96,266 @@ fn contains_multimem_blackwell_features(contents: &str) -> bool {
     })
 }
 
-/// Checks the PTX 8.6 floating-point extension to `redux.sync`.
+/// Checks the PTX 8.6 floating-point extension to `redux.sync`: a call to a
+/// float NVVM redux intrinsic, or a `redux.sync ... .f32` instruction in
+/// inline PTX. Decided on tokens, so text that only mentions either -- in a
+/// comment, a string or another symbol's name -- does not count.
 fn contains_redux_f32_features(contents: &str) -> bool {
-    contents
-        .split(';')
-        .any(|statement| statement.contains("redux.sync") && statement.contains(".f32"))
+    let tokens = lex_llvm_ir(contents);
+    tokens.iter().enumerate().any(|(index, token)| match token {
+        IrToken::Global(name) => is_float_redux_callee(name),
+        IrToken::Word(word) if *word == "asm" => inline_asm_template(&tokens[index + 1..])
+            .is_some_and(|template| ptx_contains_float_redux(&unescape_llvm_string(template))),
+        // A raw PTX fragment rather than IR, which is how tests pass PTX.
+        IrToken::Word(word) => ptx_contains_float_redux(word),
+        IrToken::Label(_) | IrToken::String(_) => false,
+    })
+}
+
+/// `llvm.nvvm.redux.sync.` followed by a float operation. Those are the ones
+/// beginning with `f` (`fmin`, `fmax`, each with optional `.abs` and `.NaN`);
+/// no integer operation does.
+fn is_float_redux_callee(name: &str) -> bool {
+    name.strip_prefix("llvm.nvvm.redux.sync.")
+        .is_some_and(|operation| operation.starts_with('f'))
+}
+
+/// Does this PTX contain a `redux` instruction with an `.f32` modifier?
+///
+/// An opcode's modifiers are the run of `.name` tokens after it, in any
+/// order, so `redux.sync.min.f32`, `redux.f32.sync.min` and
+/// `redux .sync .min .f32` are the same instruction. ptxas requires `.sync`
+/// on every `redux`, so it is not checked.
+fn ptx_contains_float_redux(ptx: &str) -> bool {
+    if !ptx.contains("redux") {
+        return false;
+    }
+    let tokens = lex_ptx(ptx);
+    tokens.iter().enumerate().any(|(index, token)| {
+        *token == PtxToken::Identifier("redux")
+            && tokens[index + 1..]
+                .iter()
+                .map_while(|token| match token {
+                    PtxToken::Modifier(modifier) => Some(*modifier),
+                    _ => None,
+                })
+                .any(|modifier| modifier == ".f32")
+    })
+}
+
+/// A PTX token. Whitespace, comments and string literals produce none.
+#[derive(Debug, PartialEq)]
+enum PtxToken<'a> {
+    /// An opcode or other identifier: `redux`, `%r1`, `$0`.
+    Identifier(&'a str),
+    /// A `.name` token: an instruction modifier or a directive.
+    Modifier(&'a str),
+    /// Punctuation or a number.
+    Other,
+}
+
+/// Split PTX into [`PtxToken`]s, dropping whitespace, `//` and `/* */`
+/// comments, and string literals such as a `.file` path.
+fn lex_ptx(ptx: &str) -> Vec<PtxToken<'_>> {
+    fn is_identifier(byte: u8) -> bool {
+        byte.is_ascii_alphanumeric() || matches!(byte, b'_' | b'$')
+    }
+
+    let bytes = ptx.as_bytes();
+    // In an asm template a bare `$` starts an operand and `$$` is a literal
+    // dollar, so `.f32$0` is the modifier `.f32` followed by `$0`.
+    let run_end = |from: usize| {
+        let mut end = from;
+        while end < bytes.len() && is_identifier(bytes[end]) {
+            if bytes[end] == b'$' {
+                if bytes.get(end + 1) != Some(&b'$') {
+                    break;
+                }
+                end += 1;
+            }
+            end += 1;
+        }
+        end
+    };
+
+    let mut tokens = Vec::new();
+    let mut index = 0;
+    while index < bytes.len() {
+        let rest = &bytes[index..];
+        if rest.starts_with(b"//") {
+            index += rest
+                .iter()
+                .position(|&byte| byte == b'\n')
+                .unwrap_or(rest.len());
+        } else if rest.starts_with(b"/*") {
+            index += rest[2..]
+                .windows(2)
+                .position(|pair| pair == b"*/")
+                .map_or(rest.len(), |offset| offset + 4);
+        } else if rest[0] == b'"' {
+            // A backslash is literal in a PTX string -- ptxas accepts
+            // `.file 1 "C:\dir\"` -- so the next quote always closes it.
+            index += rest[1..]
+                .iter()
+                .position(|&byte| byte == b'"')
+                .map_or(rest.len(), |offset| offset + 2);
+        } else if rest[0] == b'.' && rest.get(1).is_some_and(|&byte| is_identifier(byte)) {
+            let end = run_end(index + 1);
+            tokens.push(PtxToken::Modifier(&ptx[index..end]));
+            index = end;
+        } else if rest[0].is_ascii_digit() {
+            // A number consumes its whole run -- `0f3F800000`, `0x1F`, `123U` --
+            // so an identifier is never read out of its tail.
+            index = run_end(index + 1);
+            tokens.push(PtxToken::Other);
+        } else if rest[0].is_ascii_alphabetic() || matches!(rest[0], b'_' | b'$' | b'%') {
+            let end = run_end(index + 1);
+            tokens.push(PtxToken::Identifier(&ptx[index..end]));
+            index = end;
+        } else {
+            if !rest[0].is_ascii_whitespace() {
+                tokens.push(PtxToken::Other);
+            }
+            index += 1;
+        }
+    }
+    tokens
+}
+
+/// The template of an inline-asm expression, given the tokens after `asm`:
+/// the string after `[sideeffect] [alignstack] [inteldialect] [unwind]`,
+/// each optional and in that order, as `LLParser` accepts them. Also covers
+/// `module asm`.
+fn inline_asm_template<'a>(after_asm: &[IrToken<'a>]) -> Option<&'a str> {
+    const MODIFIERS: [&str; 4] = ["sideeffect", "alignstack", "inteldialect", "unwind"];
+    let mut tokens = after_asm.iter();
+    let mut next = tokens.next();
+    for modifier in MODIFIERS {
+        if matches!(next, Some(IrToken::Word(word)) if *word == modifier) {
+            next = tokens.next();
+        }
+    }
+    match next {
+        Some(IrToken::String(template)) => Some(template),
+        _ => None,
+    }
+}
+
+/// An LLVM IR token. Comments, punctuation, local and metadata names and
+/// attribute groups produce none.
+#[derive(Debug, PartialEq)]
+enum IrToken<'a> {
+    /// A bare keyword, type or identifier.
+    Word(&'a str),
+    /// A basic-block label, `name:`.
+    Label(&'a str),
+    /// A global symbol with the sigil removed and quoting decoded, so
+    /// `@name` and `@"quoted name"` both yield the name.
+    Global(Cow<'a, str>),
+    /// A string literal between its quotes, still escaped.
+    String(&'a str),
+}
+
+/// Split LLVM IR into [`IrToken`]s. Quoted text is consumed whole before
+/// anything inside it is read, so a `;` in a string is not a comment and a
+/// `@` in a quoted name does not start another symbol.
+fn lex_llvm_ir(contents: &str) -> Vec<IrToken<'_>> {
+    fn is_identifier(byte: u8) -> bool {
+        byte.is_ascii_alphanumeric() || matches!(byte, b'-' | b'$' | b'.' | b'_')
+    }
+
+    let bytes = contents.as_bytes();
+    // LLVM writes a quote inside quoted text as `\22`, so the next `"` closes it.
+    let closing_quote = |from: usize| {
+        bytes[from..]
+            .iter()
+            .position(|&byte| byte == b'"')
+            .map_or(bytes.len(), |offset| from + offset)
+    };
+    let identifier_end = |from: usize| {
+        bytes[from..]
+            .iter()
+            .position(|&byte| !is_identifier(byte))
+            .map_or(bytes.len(), |offset| from + offset)
+    };
+
+    let mut tokens = Vec::new();
+    let mut index = 0;
+    while index < bytes.len() {
+        match bytes[index] {
+            b';' => {
+                index = bytes[index..]
+                    .iter()
+                    .position(|&byte| byte == b'\n')
+                    .map_or(bytes.len(), |offset| index + offset);
+            }
+            b'"' => {
+                let close = closing_quote(index + 1);
+                tokens.push(IrToken::String(&contents[index + 1..close]));
+                index = bytes.len().min(close + 1);
+            }
+            sigil @ (b'@' | b'%' | b'!' | b'#') => {
+                let (name, next) = if bytes.get(index + 1) == Some(&b'"') {
+                    let close = closing_quote(index + 2);
+                    (
+                        unescape_llvm_string(&contents[index + 2..close]),
+                        bytes.len().min(close + 1),
+                    )
+                } else {
+                    let end = identifier_end(index + 1);
+                    (Cow::Borrowed(&contents[index + 1..end]), end)
+                };
+                if sigil == b'@' {
+                    tokens.push(IrToken::Global(name));
+                }
+                index = next;
+            }
+            byte if is_identifier(byte) => {
+                let end = identifier_end(index);
+                let word = &contents[index..end];
+                if bytes.get(end) == Some(&b':') {
+                    tokens.push(IrToken::Label(word));
+                    index = end + 1;
+                } else {
+                    tokens.push(IrToken::Word(word));
+                    index = end;
+                }
+            }
+            _ => index += 1,
+        }
+    }
+    tokens
+}
+
+/// Decode LLVM quoted text as `llvm::UnEscapeLexed` does: `\\` is one
+/// backslash, `\XX` is the byte with that hex value, and any other backslash
+/// stands for itself.
+fn unescape_llvm_string(text: &str) -> Cow<'_, str> {
+    if !text.contains('\\') {
+        return Cow::Borrowed(text);
+    }
+    let hex = |byte: Option<&u8>| byte.and_then(|&byte| (byte as char).to_digit(16));
+    let bytes = text.as_bytes();
+    let mut decoded = Vec::with_capacity(bytes.len());
+    let mut index = 0;
+    while index < bytes.len() {
+        if bytes[index] == b'\\' {
+            if bytes.get(index + 1) == Some(&b'\\') {
+                decoded.push(b'\\');
+                index += 2;
+                continue;
+            }
+            if let (Some(high), Some(low)) = (hex(bytes.get(index + 1)), hex(bytes.get(index + 2)))
+            {
+                decoded.push((high * 16 + low) as u8);
+                index += 3;
+                continue;
+            }
+        }
+        decoded.push(bytes[index]);
+        index += 1;
+    }
+    Cow::Owned(String::from_utf8_lossy(&decoded).into_owned())
 }
 
 /// Checks for forward-compatible instructions whose minimum target is sm_90.
diff --git a/cuda-oxide/crates/cuda-oxide-codegen/src/target/mod.rs b/cuda-oxide/crates/cuda-oxide-codegen/src/target/mod.rs
index 97ce92d883..d77ba264fb 100644
--- a/cuda-oxide/crates/cuda-oxide-codegen/src/target/mod.rs
+++ b/cuda-oxide/crates/cuda-oxide-codegen/src/target/mod.rs
@@ -25,11 +25,11 @@ pub(crate) use generated_requirements::{
     generated_ptx_isa_requirement, generated_target_satisfied, merge_generated_module_requirements,
     merge_generated_module_requirements_for_target, validate_generated_target,
 };
-pub use select::{required_ptx_feature, validate_target_features, validate_target_for_llvm_major};
 pub(crate) use select::{
-    resolve_ptx_target_with_generated, select_target_with_generated,
+    honour_sm_floor, resolve_ptx_target_with_generated, select_target_with_generated,
     validate_ptx_isa_for_llvm_major,
 };
+pub use select::{required_ptx_feature, validate_target_features, validate_target_for_llvm_major};
 
 // The re-exports below preserve the module's pre-split surface. In-crate
 // consumers currently reach these names only from test code (or straight
diff --git a/cuda-oxide/crates/cuda-oxide-codegen/src/target/select.rs b/cuda-oxide/crates/cuda-oxide-codegen/src/target/select.rs
index efee412907..766ca4437c 100644
--- a/cuda-oxide/crates/cuda-oxide-codegen/src/target/select.rs
+++ b/cuda-oxide/crates/cuda-oxide-codegen/src/target/select.rs
@@ -11,7 +11,7 @@ use super::generated_requirements::{
 };
 use crate::error::PipelineError;
 use crate::generated::GeneratedModuleRequirements;
-use cuda_target_spec::{CudaArch, RECORDED_PTX_FLOORS, recorded_ptx_floor};
+use cuda_target_spec::{CudaArch, RECORDED_PTX_FLOORS, recorded_ptx_floor, resolve_sm_floor};
 
 /// Select a concrete architecture that satisfies every detected feature.
 ///
@@ -167,6 +167,33 @@ pub fn resolve_ptx_target(
     )
 }
 
+/// Raise an automatically selected target to the device cfg floor.
+/// Reject the combination if that floor cannot run the module's features.
+pub(crate) fn honour_sm_floor(
+    selected: CudaArch,
+    floor: u32,
+    features: DetectedFeatures,
+    generated: &GeneratedModuleRequirements,
+) -> Result<CudaArch, String> {
+    if selected.capability() >= floor {
+        return Ok(selected);
+    }
+
+    let at_floor = CudaArch::new(floor, None).map_err(|error| {
+        format!("compute capability {floor} is not a CUDA target this compiler knows: {error}")
+    })?;
+    if arch_satisfies(&at_floor, features) && generated_target_satisfied(&at_floor, generated) {
+        return Ok(at_floor);
+    }
+
+    Err(format!(
+        "device code is compiled against compute capability {floor} or newer, but this module's \
+         features run only on {}, which that floor cannot host. Pin the architecture you mean to \
+         build for with `--arch` (or CUDA_OXIDE_TARGET), so device code and the emitted PTX agree.",
+        selected.sm()
+    ))
+}
+
 pub(crate) fn resolve_ptx_target_with_generated(
     explicit_override: Option<&str>,
     explicit_override_source: &'static str,
@@ -206,8 +233,21 @@ pub(crate) fn resolve_ptx_target_with_generated(
         return Ok((device.clone(), "detected GPU"));
     }
 
+    // Reached only with no explicit target, so only the detected-device hint
+    // can move the floor -- downwards. Resolved through the same function the
+    // build script uses rather than restated here.
+    let floor = resolve_sm_floor(explicit_override, device_hint).map_err(|error| {
+        PipelineError::TargetSelection {
+            target: String::new(),
+            reason: format!(
+                "cannot resolve the compute-capability floor device code is compiled against: {error}"
+            ),
+        }
+    })?;
     let target =
         select_target_with_generated(detected, generated).map_err(PipelineError::PtxGeneration)?;
+    let target = honour_sm_floor(target, floor, detected, generated)
+        .map_err(PipelineError::PtxGeneration)?;
     Ok((target, "feature requirement"))
 }
 
diff --git a/cuda-oxide/crates/cuda-oxide-codegen/src/target/tests/detection_async_memory.rs b/cuda-oxide/crates/cuda-oxide-codegen/src/target/tests/detection_async_memory.rs
index 9c644810bb..9f24089ec6 100644
--- a/cuda-oxide/crates/cuda-oxide-codegen/src/target/tests/detection_async_memory.rs
+++ b/cuda-oxide/crates/cuda-oxide-codegen/src/target/tests/detection_async_memory.rs
@@ -586,3 +586,209 @@ fn test_tma_multicast_detection_requires_cta_mask() {
     );
     assert!(!contains_tma_cta_group_features(unrelated_i32));
 }
+
+/// `redux.sync` is f32 by its own token, never by an `.f32` sitting nearby
+/// (#1303).
+#[test]
+fn test_redux_f32_detection_is_scoped_to_the_instruction() {
+    // The shape that found this: integer reductions in one kernel and
+    // ordinary float work in another. In the module that was rejected, the
+    // two matches were 547 lines apart.
+    let integer_redux_and_distant_f32 = r#"
+        declare i32 @llvm.nvvm.redux.sync.add(i32, i32) #0
+        declare float @llvm.nvvm.shfl.sync.bfly.f32(i32, float, i32, i32) #0
+    "#;
+    // LLVM does not require a newline between declarations, so scoping to a
+    // line is not enough either.
+    let same_line = "declare i32 @llvm.nvvm.redux.sync.add(i32, i32) declare float @llvm.nvvm.shfl.sync.bfly.f32(i32, float, i32, i32)";
+    // An SSA name may end in `.f32` and says nothing about the callee.
+    let unrelated_ssa_name = "%acc.f32 = call i32 @llvm.nvvm.redux.sync.umax(i32 -1, i32 %v)";
+
+    for text in [integer_redux_and_distant_f32, same_line, unrelated_ssa_name] {
+        assert!(
+            !detect_features_in_llvm_text(text).contains(DetectedFeatures::ReduxF32),
+            "an integer reduction must not select the f32 extension: {text}"
+        );
+    }
+
+    // The intrinsic spelling has no type suffix -- the `f`-prefixed operation
+    // is what makes it float. These were detected before only when some
+    // unrelated `.f32` happened to sit nearby.
+    let float_intrinsic = "%r = call float @llvm.nvvm.redux.sync.fmin(float %v, i32 -1)";
+    let float_intrinsic_qualified =
+        "%r = call float @llvm.nvvm.redux.sync.fmax.abs.NaN(float %v, i32 -1)";
+    // Inline PTX carries the type as a modifier of the opcode token.
+    let float_asm = r#"call float asm sideeffect "redux.sync.min.abs.NaN.f32 $0, $1, $2;", "=f,f,r"(float %v, i32 %m)"#;
+
+    for text in [float_intrinsic, float_intrinsic_qualified, float_asm] {
+        assert!(
+            detect_features_in_llvm_text(text).contains(DetectedFeatures::ReduxF32),
+            "a float reduction must select the f32 extension: {text}"
+        );
+    }
+}
+
+/// Mentioning an intrinsic is not using one, and a real call may be spelled
+/// with a quoted identifier (#1303).
+#[test]
+fn test_redux_f32_detection_reads_code_and_not_prose() {
+    // Text that names the callee without calling it.
+    let in_a_comment =
+        "%r = call i32 @llvm.nvvm.redux.sync.add(i32 -1, i32 %v) ; @llvm.nvvm.redux.sync.fmin";
+    let in_a_data_string = r#"@.msg = private constant [27 x i8] c"@llvm.nvvm.redux.sync.fmin\00""#;
+    let in_metadata = r#"!0 = !{!"@llvm.nvvm.redux.sync.fmin"}"#;
+    // The same for the inline-PTX spelling: a comment and a data constant are
+    // not instructions.
+    let ptx_in_a_comment = "  ; redux.sync.min.abs.NaN.f32 $0, $1, $2;";
+    let ptx_in_a_data_string = r#"@.msg = private constant [25 x i8] c"redux.sync.min.f32 a,b,c;""#;
+
+    for text in [
+        in_a_comment,
+        in_a_data_string,
+        in_metadata,
+        ptx_in_a_comment,
+        ptx_in_a_data_string,
+    ] {
+        assert!(
+            !detect_features_in_llvm_text(text).contains(DetectedFeatures::ReduxF32),
+            "text that only mentions the f32 reduction must not select it: {text}"
+        );
+    }
+
+    // A real call, spelled with a quoted identifier and with one of its
+    // characters escaped -- `\66` is `f`, so both name `fmin`.
+    let quoted = r#"%r = call float @"llvm.nvvm.redux.sync.fmin"(float %v, i32 -1)"#;
+    let quoted_escaped = r#"%r = call float @"llvm.nvvm.redux.sync.\66min"(float %v, i32 -1)"#;
+    // And the constraint string beside an asm template is data, while the
+    // template itself is code.
+    let asm_template = r#"call float asm sideeffect "redux.sync.min.abs.NaN.f32 $0, $1, $2;", "=f,f,r"(float %v, i32 %m)"#;
+
+    for text in [quoted, quoted_escaped, asm_template] {
+        assert!(
+            detect_features_in_llvm_text(text).contains(DetectedFeatures::ReduxF32),
+            "a float reduction must select the f32 extension: {text}"
+        );
+    }
+}
+
+/// Token identity at the edges the flattened search lost (#1303).
+///
+/// Each input is a line from a module `llvm-as` from the pinned toolchain
+/// accepts, so none of these is a shape valid IR cannot produce. The escape
+/// cases were also round-tripped through `llvm-dis` to confirm what LLVM
+/// itself decodes them to.
+#[test]
+fn test_redux_f32_detection_keeps_token_identity() {
+    // `unwind` is the last of the four modifiers `asm` may carry, in order:
+    // sideeffect, alignstack, inteldialect, unwind.
+    let unwind = r#"%r = call float asm sideeffect unwind "redux.sync.min.f32 $0, $1, $2;", "=f,f,r"(float %v, i32 %m)"#;
+    // `\72` is `r`: LLVM decodes the template before the backend sees it.
+    let escaped_opcode = r#"%r = call float asm sideeffect "\72edux.sync.min.f32 $0, $1, $2;", "=f,f,r"(float %v, i32 %m)"#;
+    // A PTX comment ended by an *escaped* newline, then a real instruction.
+    // Only decoding before stripping comments finds it: the raw template is
+    // one line, so a `//` there would appear to run to its end.
+    let after_escaped_newline = r#"%r = call float asm sideeffect "// scale first\0Aredux.sync.min.f32 $0, $1, $2;", "=f,f,r"(float %v, i32 %m)"#;
+    // Module-level inline asm has the same grammar and reaches the same PTX.
+    let module_asm = r#"module asm "redux.sync.min.f32 %f1, %f2, %r3;""#;
+
+    for text in [unwind, escaped_opcode, after_escaped_newline, module_asm] {
+        assert!(
+            detect_features_in_llvm_text(text).contains(DetectedFeatures::ReduxF32),
+            "a float reduction must select the f32 extension: {text}"
+        );
+    }
+
+    // PTX comments inside a template are prose, in either comment form.
+    let ptx_line_comment = r#"%r = call float asm sideeffect "// redux.sync.min.f32 $0, $1, $2;\0Amov.f32 $0, $1;", "=f,f,r"(float %v, i32 %m)"#;
+    let ptx_block_comment = r#"%r = call float asm sideeffect "/* redux.sync.min.f32 $0, $1, $2; */ mov.f32 $0, $1;", "=f,f,r"(float %v, i32 %m)"#;
+    // `\\` is one backslash and the hex digits after it are plain text: `llc`
+    // emits this template as `\20redux.sync.min.f32 ...`, which is not an
+    // instruction. A decoder that skipped the `\\` rule would read `\20` as a
+    // space instead and find a `redux.sync.min.f32` that is not there.
+    let escaped_backslash = r#"%r = call float asm sideeffect "\\20redux.sync.min.f32 $0, $1, $2;", "=f,f,r"(float %v, i32 %m)"#;
+    // One symbol whose quoted name happens to contain another's spelling.
+    let embedded_in_a_name = r#"@"unused @llvm.nvvm.redux.sync.fmin" = global i32 0"#;
+    // A basic block may be named like an instruction without being one.
+    let label = "redux.sync.min.f32:\n  ret void";
+
+    for text in [
+        ptx_line_comment,
+        ptx_block_comment,
+        escaped_backslash,
+        embedded_in_a_name,
+        label,
+    ] {
+        assert!(
+            !detect_features_in_llvm_text(text).contains(DetectedFeatures::ReduxF32),
+            "no float reduction is executed here: {text}"
+        );
+    }
+}
+
+/// PTX lexical rules inside an asm template (#1303).
+///
+/// Every input is a complete module accepted by `llvm-as` and `llc` from the
+/// pinned toolchain and by ptxas at sm_100a. The float cases are the ones
+/// ptxas then refuses at sm_80 as `redux.f32`; the others assemble there.
+#[test]
+fn test_redux_f32_detection_reads_ptx_as_ptx() {
+    // Whitespace, a newline or a comment may separate an opcode from its
+    // modifiers; it is still one instruction.
+    let spaced_type = r#"%r = call float asm sideeffect "redux.sync.min .f32 $0, $1, $2;", "=f,f,r"(float %v, i32 %m)"#;
+    let comment_between = r#"%r = call float asm sideeffect "redux.sync.min/* qualifier */.f32 $0, $1, $2;", "=f,f,r"(float %v, i32 %m)"#;
+    let fully_spaced = r#"%r = call float asm sideeffect "redux .sync .min .f32 $0, $1, $2;", "=f,f,r"(float %v, i32 %m)"#;
+    let newline_between = r#"%r = call float asm sideeffect "redux.sync.min\0A.f32 $0, $1, $2;", "=f,f,r"(float %v, i32 %m)"#;
+    // A backslash is literal inside a PTX string, so `"nounroll\"` closes at
+    // its last quote and the instruction after it is real.
+    let after_trailing_backslash = r#"%r = call float asm sideeffect ".pragma \22nounroll\5C\22;\0Aredux.sync.min.f32 $0, $1, $2;", "=f,f,r"(float %v, i32 %m)"#;
+    // A bare `$` starts an operand, so it ends the modifier before it. llc
+    // emits `redux.sync.min.f32%r1, %r2, %r3;`, which ptxas assembles at
+    // sm_100a and refuses at sm_80 with `Instruction 'redux.f32' not supported`.
+    let glued_operand = r#"%r = call float asm sideeffect "redux.sync.min.f32$0, $1, $2;", "=f,f,r"(float %v, i32 %m)"#;
+    // Modifiers may come in any order; ptxas gives both of these the same
+    // two results as the glued form.
+    let sync_after_operation = r#"%r = call float asm sideeffect "redux.min.sync.f32 $0, $1, $2;", "=f,f,r"(float %v, i32 %m)"#;
+    let type_first = r#"%r = call float asm sideeffect "redux.f32.sync.min $0, $1, $2;", "=f,f,r"(float %v, i32 %m)"#;
+
+    for text in [
+        spaced_type,
+        comment_between,
+        fully_spaced,
+        newline_between,
+        after_trailing_backslash,
+        glued_operand,
+        sync_after_operation,
+        type_first,
+    ] {
+        assert!(
+            detect_features_in_llvm_text(text).contains(DetectedFeatures::ReduxF32),
+            "a float reduction must select the f32 extension: {text}"
+        );
+    }
+
+    // A directive's quoted data is not an instruction.
+    let file_path = r#"module asm ".file 1 \22name redux.sync.min.f32 label.cu\22""#;
+    // Spacing does not make an integer reduction float.
+    let spaced_integer = r#"%r = call i32 asm sideeffect "redux .sync .min .u32 $0, $1, $2;", "=r,r,r"(i32 %v, i32 %m)"#;
+    // Nor does float work beside it in the same template: `.f32` counts only
+    // as a modifier of the `redux` opcode itself.
+    let integer_beside_f32 = r#"%r = call i32 asm sideeffect "{ .reg .f32 t; redux.sync.min.u32 $0, $1, $2; mov.f32 t, 0f3F800000; }", "=r,r,r"(i32 %v, i32 %m)"#;
+    // The integer forms of the glued and reordered inputs above; ptxas
+    // assembles both at sm_80.
+    let glued_integer =
+        r#"%r = call i32 asm sideeffect "redux.sync.min.u32$0, $1, $2;", "=r,r,r"(i32 %v, i32 %m)"#;
+    let reordered_integer = r#"%r = call i32 asm sideeffect "redux.min.sync.u32 $0, $1, $2;", "=r,r,r"(i32 %v, i32 %m)"#;
+
+    for text in [
+        file_path,
+        spaced_integer,
+        integer_beside_f32,
+        glued_integer,
+        reordered_integer,
+    ] {
+        assert!(
+            !detect_features_in_llvm_text(text).contains(DetectedFeatures::ReduxF32),
+            "no float reduction is executed here: {text}"
+        );
+    }
+}
diff --git a/cuda-oxide/crates/cuda-oxide-codegen/src/target/tests/detection_matrix.rs b/cuda-oxide/crates/cuda-oxide-codegen/src/target/tests/detection_matrix.rs
index 5ee2a970f1..036169666c 100644
--- a/cuda-oxide/crates/cuda-oxide-codegen/src/target/tests/detection_matrix.rs
+++ b/cuda-oxide/crates/cuda-oxide-codegen/src/target/tests/detection_matrix.rs
@@ -1112,9 +1112,27 @@ fn m8n8k16_int8_mma_detection_applies_exact_sm75_and_ptx65_floors() {
         "{$0, $1}, {$4}, {$5}, {$2, $3};"
     );
     let requirements = detect_module_requirements_in_llvm_text(representative);
+    // The feature's own floor is sm_75 and stays there.
+    assert_eq!(
+        select_target(requirements.features)
+            .expect("feature floor")
+            .sm(),
+        "sm_75"
+    );
+    // Selection does not land below the capability device code was compiled
+    // against (#811). With nothing pinned that is the sm_80 default, which
+    // runs everything sm_75 does; pinning `--arch sm_75` selects sm_75.
     let (target, _) = resolve_ptx_target(None, "CUDA_OXIDE_TARGET", None, requirements.features)
         .expect("auto-resolve");
-    assert_eq!(target.sm(), "sm_75");
+    assert_eq!(target.sm(), "sm_80");
+    let (pinned, _) = resolve_ptx_target(
+        Some("sm_75"),
+        "CUDA_OXIDE_TARGET",
+        None,
+        requirements.features,
+    )
+    .expect("pinned");
+    assert_eq!(pinned.sm(), "sm_75");
     assert_eq!(
         required_ptx_feature(&"sm_75".parse().unwrap(), requirements.ptx_isa).unwrap(),
         Some("+ptx65")
diff --git a/cuda-oxide/crates/cuda-oxide-codegen/src/target/tests/selection.rs b/cuda-oxide/crates/cuda-oxide-codegen/src/target/tests/selection.rs
index 54d0cc026f..b65b21837b 100644
--- a/cuda-oxide/crates/cuda-oxide-codegen/src/target/tests/selection.rs
+++ b/cuda-oxide/crates/cuda-oxide-codegen/src/target/tests/selection.rs
@@ -518,3 +518,207 @@ fn malformed_low_width_targets_keep_override_errors() {
         assert_eq!(explicit.to_string(), expected_reason);
     }
 }
+
+/// The property `cuda-device`'s `cuda_oxide_sm_at_least` ladder depends on:
+/// selection never lands below the floor device code was compiled against
+/// (#811).
+///
+/// The build script fixes that ladder before rustc runs, from the same
+/// [`resolve_sm_floor`] both resolvers below consult, so this asserts the
+/// property end to end rather than re-checking that one helper's table. Both
+/// stages are covered because a module reaches PTX through one or the other.
+#[test]
+fn selection_never_lands_below_the_floor_device_code_assumes() {
+    let generated = crate::generated::GeneratedModuleRequirements::default();
+    let [sm_75, sm_86, sm_120a] = ["sm_75", "sm_86", "sm_120a"]
+        .map(|raw| raw.parse::<cuda_target_spec::DeviceArch>().unwrap());
+    for explicit in [None, Some("sm_75"), Some("sm_86"), Some("sm_90a")] {
+        for hint in [None, Some(&sm_75), Some(&sm_86), Some(&sm_120a)] {
+            let floor = cuda_target_spec::resolve_sm_floor(explicit, hint.map(AsRef::as_ref))
+                .expect("floor");
+            for features in [
+                DetectedFeatures::Basic,
+                DetectedFeatures::Ldmatrix,
+                DetectedFeatures::Movmatrix,
+                DetectedFeatures::Sm75,
+                DetectedFeatures::Sm80,
+                DetectedFeatures::Cluster,
+                DetectedFeatures::Wgmma,
+            ] {
+                // An explicit target that cannot host the module is rejected
+                // outright; that is the other half of the contract, asserted
+                // separately below. Only successful selections are bound.
+                if let Ok((target, source)) =
+                    resolve_ptx_target(explicit, "CUDA_OXIDE_TARGET", hint, features)
+                {
+                    assert!(
+                        target.capability() >= floor,
+                        "PTX: {features:?} with explicit={explicit:?} hint={hint:?} \
+                         selected {} (from {source}) below floor {floor}",
+                        target.sm()
+                    );
+                }
+                if let Ok(target) = crate::export::resolve_nvvm_target_with_generated(
+                    explicit,
+                    hint,
+                    Some(features),
+                    &generated,
+                ) {
+                    assert!(
+                        target.capability() >= floor,
+                        "NVVM: {features:?} with explicit={explicit:?} hint={hint:?} \
+                         selected {} below floor {floor}",
+                        target.sm()
+                    );
+                }
+            }
+        }
+    }
+}
+
+/// The two cases that motivated the floor, named so a regression says which.
+#[test]
+fn the_floor_is_honoured_without_overriding_a_runnable_device() {
+    let generated = crate::generated::GeneratedModuleRequirements::default();
+
+    // Nothing pinned: the ladder advertises the sm_80 default, so a module
+    // whose features only need sm_75 must not be built for sm_75.
+    let (ptx, _) =
+        resolve_ptx_target(None, "CUDA_OXIDE_TARGET", None, DetectedFeatures::Ldmatrix).unwrap();
+    assert_eq!(ptx.sm(), "sm_80");
+    let nvvm = crate::export::resolve_nvvm_target_with_generated(
+        None,
+        None,
+        Some(DetectedFeatures::Ldmatrix),
+        &generated,
+    )
+    .unwrap();
+    assert_eq!(nvvm.sm(), "sm_80");
+
+    // A detected Blackwell with a WGMMA module still selects sm_90a, because
+    // the hint cannot raise the floor: it is advisory, and this is exactly the
+    // case the backend is documented to build elsewhere for.
+    let blackwell = "sm_120a".parse::<cuda_target_spec::DeviceArch>().unwrap();
+    let (wgmma, _) = resolve_ptx_target(
+        None,
+        "CUDA_OXIDE_TARGET",
+        Some(&blackwell),
+        DetectedFeatures::Wgmma,
+    )
+    .unwrap();
+    assert_eq!(wgmma.sm(), "sm_90a");
+    assert_eq!(
+        cuda_target_spec::resolve_sm_floor(None, Some(blackwell.as_ref())).unwrap(),
+        80
+    );
+
+    // A detected device *below* the default pulls the floor down with it, so
+    // running on Turing still produces sm_75 rather than PTX it cannot load.
+    let sm_75 = "sm_75".parse::<cuda_target_spec::DeviceArch>().unwrap();
+    let (turing, source) = resolve_ptx_target(
+        None,
+        "CUDA_OXIDE_TARGET",
+        Some(&sm_75),
+        DetectedFeatures::Basic,
+    )
+    .unwrap();
+    assert_eq!((turing.sm().as_str(), source), ("sm_75", "detected GPU"));
+    assert_eq!(
+        cuda_target_spec::resolve_sm_floor(None, Some(sm_75.as_ref())).unwrap(),
+        75
+    );
+}
+
+/// An architecture family that cannot host the module is refused where it is
+/// stated, rather than quietly resolved to something else.
+#[test]
+fn an_explicit_target_that_cannot_host_the_module_is_refused() {
+    let error = resolve_ptx_target(
+        Some("sm_120a"),
+        "CUDA_OXIDE_TARGET",
+        None,
+        DetectedFeatures::Wgmma,
+    )
+    .unwrap_err()
+    .to_string();
+    assert!(error.contains("Wgmma"), "{error}");
+
+    // A blank target, which the two stages read differently.
+    //
+    // `resolve_sm_floor` treats a blank `CUDA_OXIDE_TARGET` as unset and falls
+    // back to the default floor, while the backend reads the variable with
+    // `env::var().ok()` and hands the empty string to the target parser, which
+    // refuses it. The strict side is the one that decides whether anything is
+    // built, so the ladder can never be the thing that overstates: the build
+    // fails before a module exists.
+    assert_eq!(
+        cuda_target_spec::resolve_sm_floor(Some("  "), None).unwrap(),
+        80
+    );
+    assert!(
+        resolve_ptx_target(
+            Some("  "),
+            "CUDA_OXIDE_TARGET",
+            None,
+            DetectedFeatures::Basic
+        )
+        .is_err()
+    );
+}
+
+/// The floor's reject arm. A feature-selected target below the floor is
+/// raised only when the floor can host the module: WGMMA exists only on
+/// sm_90a, so a floor of 100 cannot, and that is reported instead of resolved
+/// to another architecture family.
+#[test]
+fn a_floor_that_cannot_host_the_module_is_reported() {
+    let generated = crate::generated::GeneratedModuleRequirements::default();
+    let sm_90a = "sm_90a".parse::<cuda_target_spec::CudaArch>().unwrap();
+
+    let error =
+        honour_sm_floor(sm_90a.clone(), 100, DetectedFeatures::Wgmma, &generated).unwrap_err();
+    assert!(error.contains("compute capability 100"), "{error}");
+    assert!(error.contains("run only on sm_90a"), "{error}");
+
+    // Controls: a floor the selection already meets leaves it alone, and one
+    // that can host the module raises it.
+    let kept = honour_sm_floor(sm_90a, 90, DetectedFeatures::Wgmma, &generated).unwrap();
+    assert_eq!(kept.sm(), "sm_90a");
+    let sm_75 = "sm_75".parse::<cuda_target_spec::CudaArch>().unwrap();
+    let raised = honour_sm_floor(sm_75, 80, DetectedFeatures::Ldmatrix, &generated).unwrap();
+    assert_eq!(raised.sm(), "sm_80");
+}
+
+/// `cuda-device`'s build script reads `CUDA_OXIDE_DEVICE_ARCH` as
+/// `value.parse::<DeviceArch>().ok()`. Its floor matches selection only while
+/// the backend's input boundary accepts exactly those values and rejects the
+/// rest before any target is chosen.
+#[test]
+fn the_build_script_reads_the_device_hint_as_the_backend_does() {
+    for raw in [
+        "sm_75",
+        "sm_86",
+        "sm_120a",
+        "",
+        "  ",
+        " sm_86",
+        "sm_86 ",
+        "86",
+        "sm_9",
+        "compute_86",
+        "sm_100f",
+    ] {
+        let build_script = raw.parse::<cuda_target_spec::DeviceArch>().ok();
+        let backend = crate::options::DeviceArchHint::from_env_value(Ok(raw.to_string()))
+            .expect("a set variable is a hint")
+            .as_device_arch()
+            .ok()
+            .cloned();
+        assert_eq!(build_script, backend, "CUDA_OXIDE_DEVICE_ARCH={raw:?}");
+    }
+    assert!(
+        crate::options::DeviceArchHint::from_env_value(Err(std::env::VarError::NotPresent))
+            .is_none(),
+        "an unset variable is no hint on either side"
+    );
+}
diff --git a/cuda-oxide/crates/cuda-target-spec/src/lib.rs b/cuda-oxide/crates/cuda-target-spec/src/lib.rs
index 62b2f65d99..09a7908db8 100644
--- a/cuda-oxide/crates/cuda-target-spec/src/lib.rs
+++ b/cuda-oxide/crates/cuda-target-spec/src/lib.rs
@@ -556,9 +556,146 @@ pub fn spelling_at_least(floor: u16) -> Option<u16> {
     PtxSpelling::round_up(floor).map(PtxSpelling::get)
 }
 
+/// Compute-capability floors device code can test with the
+/// `cuda_oxide_sm_at_least` cfg, ascending.
+///
+/// `cfg` compares strings, not numbers, so a single "the floor is 86" cfg
+/// could not answer "is this at least sm_80". `cuda-device`'s build script
+/// instead sets every rung at or below the resolved capability, and device
+/// code tests the one rung its instruction needs.
+pub const SM_FLOOR_LADDER: &[u32] = &[70, 75, 80, 86, 90, 100, 120];
+
+/// Capability floor assumed when nothing names a target.
+///
+/// Matches the backend's own "Basic CUDA / sm_80 / Ampere+ (max compat)"
+/// default, so a cross-compile with no environment set agrees with the PTX it
+/// will actually produce.
+pub const DEFAULT_SM_FLOOR: u32 = 80;
+
+/// Resolve the minimum compute capability assumed by device code.
+///
+/// An explicit target sets the floor. A device hint can only lower the default
+/// `sm_80` floor because backend selection may override an incompatible hint.
+/// Both device cfg generation and backend selection use this rule.
+///
+/// A blank target is unset and a malformed one is an error; the hint arrives
+/// parsed. Architecture suffixes do not affect this numeric floor, which does
+/// not imply support for architecture-specific instruction families.
+pub fn resolve_sm_floor(
+    target: Option<&str>,
+    device_arch: Option<&CudaArch>,
+) -> Result<u32, CudaArchParseError> {
+    if let Some(target) = target.map(str::trim).filter(|target| !target.is_empty()) {
+        return Ok(target.parse::<CudaArch>()?.capability());
+    }
+    Ok(device_arch.map_or(DEFAULT_SM_FLOOR, |device| {
+        DEFAULT_SM_FLOOR.min(device.capability())
+    }))
+}
+
+/// The ladder rungs at or below `capability`, ascending.
+pub fn sm_floors_at_most(capability: u32) -> impl Iterator<Item = u32> {
+    SM_FLOOR_LADDER
+        .iter()
+        .copied()
+        .filter(move |floor| *floor <= capability)
+}
+
 #[cfg(test)]
 mod tests {
     use super::*;
+
+    /// A detected device, as the backend and the build script both carry it.
+    fn device(raw: &str) -> CudaArch {
+        raw.parse::<DeviceArch>().unwrap().into()
+    }
+
+    /// An explicit target sets the floor exactly; a detected device can only
+    /// lower it, because the backend may build for that device as-is but is
+    /// documented to build elsewhere when the kernel needs it.
+    #[test]
+    fn a_detected_device_lowers_the_floor_and_never_raises_it() {
+        assert_eq!(resolve_sm_floor(Some("sm_86"), None).unwrap(), 86);
+        assert_eq!(resolve_sm_floor(None, None).unwrap(), DEFAULT_SM_FLOOR);
+
+        // Below the default: running on Turing really does produce sm_75, so
+        // device code must not be told it has Ampere rungs.
+        assert_eq!(resolve_sm_floor(None, Some(&device("sm_75"))).unwrap(), 75);
+
+        // Above the default: selection may still land on sm_90a for a WGMMA
+        // module, so the hint promises nothing higher.
+        assert_eq!(
+            resolve_sm_floor(None, Some(&device("sm_120a"))).unwrap(),
+            DEFAULT_SM_FLOOR
+        );
+
+        // An explicit target wins over the hint in both directions.
+        assert_eq!(
+            resolve_sm_floor(Some("sm_75"), Some(&device("sm_120"))).unwrap(),
+            75
+        );
+        assert_eq!(
+            resolve_sm_floor(Some("sm_120"), Some(&device("sm_75"))).unwrap(),
+            120
+        );
+    }
+
+    /// An env var set to the empty string is how a shell spells "unset", and
+    /// must not be read as a target that failed to parse.
+    #[test]
+    fn blank_env_values_are_not_targets() {
+        assert_eq!(resolve_sm_floor(Some(""), None).unwrap(), DEFAULT_SM_FLOOR);
+        assert_eq!(
+            resolve_sm_floor(Some("  "), Some(&device("sm_75"))).unwrap(),
+            75
+        );
+        assert_eq!(resolve_sm_floor(Some(" sm_86 "), None).unwrap(), 86);
+    }
+
+    /// Family suffixes carry no capability of their own.
+    #[test]
+    fn architecture_family_suffixes_do_not_change_the_floor() {
+        assert_eq!(resolve_sm_floor(Some("sm_90a"), None).unwrap(), 90);
+        assert_eq!(resolve_sm_floor(Some("sm_100f"), None).unwrap(), 100);
+    }
+
+    /// Guessing high emits instructions the target cannot run and guessing low
+    /// silently drops the faster form, so a named value that does not parse
+    /// fails instead of falling back.
+    #[test]
+    fn a_named_target_that_does_not_parse_is_an_error() {
+        assert!(resolve_sm_floor(Some("ampere"), None).is_err());
+        assert!(resolve_sm_floor(Some("sm_9"), None).is_err());
+    }
+
+    #[test]
+    fn the_floor_ladder_is_ascending_and_carries_the_default() {
+        assert!(SM_FLOOR_LADDER.windows(2).all(|pair| pair[0] < pair[1]));
+        assert!(SM_FLOOR_LADDER.contains(&DEFAULT_SM_FLOOR));
+    }
+
+    /// The cumulative rungs are what `cfg` needs: every floor at or below the
+    /// capability, so device code can ask for the one it needs.
+    #[test]
+    fn floors_at_most_is_cumulative_and_excludes_higher_rungs() {
+        assert_eq!(
+            sm_floors_at_most(86).collect::<Vec<_>>(),
+            vec![70, 75, 80, 86]
+        );
+        assert_eq!(
+            sm_floors_at_most(120).collect::<Vec<_>>(),
+            SM_FLOOR_LADDER.to_vec()
+        );
+        assert_eq!(sm_floors_at_most(75).collect::<Vec<_>>(), vec![70, 75]);
+        // Between rungs: sm_89 is Ada, which is Ampere-or-newer for our
+        // purposes but is not itself a rung.
+        assert_eq!(
+            sm_floors_at_most(89).collect::<Vec<_>>(),
+            vec![70, 75, 80, 86]
+        );
+        assert!(sm_floors_at_most(50).next().is_none());
+    }
+
     #[test]
     fn device_arch_requires_sm_spelling_and_reuses_cuda_grammar() {
         for entry in RECORDED_PTX_FLOORS {
diff --git a/cuda-oxide/crates/nvvm-transforms/src/legalize.rs b/cuda-oxide/crates/nvvm-transforms/src/legalize.rs
index a3fb9f3e79..54e0d777bb 100644
--- a/cuda-oxide/crates/nvvm-transforms/src/legalize.rs
+++ b/cuda-oxide/crates/nvvm-transforms/src/legalize.rs
@@ -55,7 +55,8 @@ const NNEG_ATTR: &str = "llvm_nneg_flag";
 /// and compare-exchange operations with stronger failure ordering are rewritten
 /// to inline PTX instead of weakening scope or ordering. Other atomic and fence
 /// operations return an error instead of being emitted with unverified
-/// semantics.
+/// semantics. Integer `redux.sync`, which LLVM 7 has no intrinsic for, becomes
+/// convergent inline PTX.
 pub(crate) fn legalize_for_legacy_nvvm(
     ctx: &mut Context,
     module: Ptr<Operation>,
@@ -137,6 +138,9 @@ pub(crate) fn legalize_for_legacy_nvvm(
         } else if vote_mode(&name).is_some() {
             rewrite_vote(ctx, op, &name)?;
             obsolete_declarations.push(name);
+        } else if integer_redux_ptx(&name).is_some() {
+            rewrite_integer_redux(ctx, op, &name)?;
+            obsolete_declarations.push(name);
         } else if let Some(legacy_name) = legacy_memory_intrinsic_name(&name) {
             rewrite_renamed_call(ctx, op, &legacy_name)?;
             obsolete_declarations.push(name);
@@ -1203,6 +1207,8 @@ fn validate_rewrite_candidate(ctx: &Context, op: Ptr<Operation>) -> Result<()> {
         validate_shuffle_call(ctx, op, &name)?;
     } else if vote_mode(&name).is_some() {
         validate_vote_call(ctx, op, &name)?;
+    } else if integer_redux_ptx(&name).is_some() {
+        validate_integer_redux_call(ctx, op, &name)?;
     }
     Ok(())
 }
@@ -2011,6 +2017,57 @@ fn rewrite_vote(ctx: &mut Context, op: Ptr<Operation>, name: &str) -> Result<()>
     replace_one_result(ctx, op, result)
 }
 
+/// The PTX instruction for an integer `llvm.nvvm.redux.sync.*` call.
+fn integer_redux_ptx(name: &str) -> Option<&'static str> {
+    match name {
+        "llvm_nvvm_redux_sync_add" => Some("redux.sync.add.s32"),
+        "llvm_nvvm_redux_sync_min" => Some("redux.sync.min.s32"),
+        "llvm_nvvm_redux_sync_max" => Some("redux.sync.max.s32"),
+        "llvm_nvvm_redux_sync_umin" => Some("redux.sync.min.u32"),
+        "llvm_nvvm_redux_sync_umax" => Some("redux.sync.max.u32"),
+        "llvm_nvvm_redux_sync_and" => Some("redux.sync.and.b32"),
+        "llvm_nvvm_redux_sync_or" => Some("redux.sync.or.b32"),
+        "llvm_nvvm_redux_sync_xor" => Some("redux.sync.xor.b32"),
+        _ => None,
+    }
+}
+
+fn validate_integer_redux_call(ctx: &Context, op: Ptr<Operation>, name: &str) -> Result<()> {
+    let operands: Vec<_> = op.deref(ctx).operands().collect();
+    let i32_ty: TypeHandle = IntegerType::get(ctx, 32, Signedness::Signless).into();
+    if operands.len() != 2
+        || op.deref(ctx).get_num_results() != 1
+        || operands
+            .iter()
+            .any(|operand| operand.get_type(ctx) != i32_ty)
+        || op.deref(ctx).get_result(0).get_type(ctx) != i32_ty
+    {
+        return pliron::input_err!(
+            op.deref(ctx).loc(),
+            "{name} must take (i32 value, i32 member mask) and return i32"
+        );
+    }
+    Ok(())
+}
+
+/// Replace an integer `redux.sync` call with the same PTX instruction. The
+/// call's operands are already in PTX order: value, then member mask.
+fn rewrite_integer_redux(ctx: &mut Context, op: Ptr<Operation>, name: &str) -> Result<()> {
+    let opcode = integer_redux_ptx(name).expect("validated redux name");
+    let operands: Vec<_> = op.deref(ctx).operands().collect();
+    let i32_ty: TypeHandle = IntegerType::get(ctx, 32, Signedness::Signless).into();
+    let asm = llvm::InlineAsmOp::build(
+        ctx,
+        i32_ty,
+        operands,
+        &format!("{opcode} $0, $1, $2;"),
+        "=r,r,r",
+        AsmKind::Convergent,
+    );
+    let result = insert_before(ctx, op, asm.get_operation());
+    replace_one_result(ctx, op, result)
+}
+
 fn legacy_memory_intrinsic_name(name: &str) -> Option<String> {
     for stem in ["memcpy", "memmove"] {
         if let Some(rest) = name.strip_prefix(&format!("llvm_{stem}_p")) {
@@ -2148,6 +2205,7 @@ fn verify_legacy_subset(ctx: &Context, module: Ptr<Operation>, capability: u32)
             || legacy_bit_rewrite(&name).is_some()
             || shuffle_mode(&name).is_some()
             || vote_mode(&name).is_some()
+            || integer_redux_ptx(&name).is_some()
             || legacy_memory_intrinsic_name(&name).is_some()
         {
             return pliron::input_err!(
@@ -2830,6 +2888,67 @@ mod tests {
         module.get_operation().deref(&ctx).verify(&ctx).unwrap();
     }
 
+    /// LLVM 7 has no `llvm.nvvm.redux.sync.*`, so legacy libNVVM would leave
+    /// these calls as undefined external functions that ptxas rejects. Each
+    /// becomes its PTX instruction, convergent, with the operands in the
+    /// intrinsic's order: value, then member mask.
+    #[test]
+    fn integer_redux_calls_become_convergent_inline_ptx() {
+        for (name, opcode) in [
+            ("llvm_nvvm_redux_sync_add", "redux.sync.add.s32"),
+            ("llvm_nvvm_redux_sync_min", "redux.sync.min.s32"),
+            ("llvm_nvvm_redux_sync_max", "redux.sync.max.s32"),
+            ("llvm_nvvm_redux_sync_umin", "redux.sync.min.u32"),
+            ("llvm_nvvm_redux_sync_umax", "redux.sync.max.u32"),
+            ("llvm_nvvm_redux_sync_and", "redux.sync.and.b32"),
+            ("llvm_nvvm_redux_sync_or", "redux.sync.or.b32"),
+            ("llvm_nvvm_redux_sync_xor", "redux.sync.xor.b32"),
+        ] {
+            let mut ctx = Context::new();
+            let module = ModuleOp::new(&mut ctx, name.try_into().unwrap());
+            let i32_ty: TypeHandle = IntegerType::get(&ctx, 32, Signedness::Signless).into();
+            direct_intrinsic_call(&mut ctx, &module, name, i32_ty, vec![i32_ty, i32_ty]);
+
+            legalize_for_legacy_nvvm(&mut ctx, module.get_operation(), 80).unwrap();
+
+            module.get_operation().deref(&ctx).verify(&ctx).unwrap();
+            let ir = export_module_to_string_with_config(
+                &ctx,
+                &module,
+                &NvvmExportConfig::new(NvvmIrDialect::LegacyLlvm7),
+            )
+            .expect("legalized redux exports as LLVM 7 NVVM IR");
+            let asm = format!(
+                r#"call i32 asm sideeffect "{opcode} $0, $1, $2;", "=r,r,r"(i32 %v0, i32 %v1) #0"#
+            );
+            assert_eq!(ir.matches(&asm).count(), 1, "missing `{asm}`:\n{ir}");
+            assert!(ir.contains("attributes #0 = { convergent }"), "{ir}");
+            assert!(!ir.contains("llvm.nvvm.redux"), "{ir}");
+        }
+    }
+
+    /// The float forms need sm_100, which uses the modern dialect, so only
+    /// the eight integer calls are rewritten; a malformed one is an error.
+    #[test]
+    fn integer_redux_call_with_a_wrong_signature_is_rejected() {
+        let mut ctx = Context::new();
+        let module = ModuleOp::new(&mut ctx, "redux_i64".try_into().unwrap());
+        let i64_ty: TypeHandle = IntegerType::get(&ctx, 64, Signedness::Signless).into();
+        let i32_ty: TypeHandle = IntegerType::get(&ctx, 32, Signedness::Signless).into();
+        direct_intrinsic_call(
+            &mut ctx,
+            &module,
+            "llvm_nvvm_redux_sync_add",
+            i64_ty,
+            vec![i64_ty, i32_ty],
+        );
+
+        let error = legalize_for_legacy_nvvm(&mut ctx, module.get_operation(), 80).unwrap_err();
+        let text = error.disp(&ctx).to_string();
+        assert!(text.contains("llvm_nvvm_redux_sync_add"), "{text}");
+        assert!(text.contains("return i32"), "{text}");
+    }
+
     #[test]
     fn f16_function_signature_is_rejected_without_needing_an_fneg() {
         let mut ctx = Context::new();
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/abi_hmm/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/abi_hmm/Cargo.lock
index ce0fdb7a46..7d2c12cde4 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/abi_hmm/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/abi_hmm/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/addressof_sharedarray/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/addressof_sharedarray/Cargo.lock
index 298682ebe7..90b4206abf 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/addressof_sharedarray/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/addressof_sharedarray/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/aligned_field_loads/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/aligned_field_loads/Cargo.lock
index f5335376d4..7b791e38e8 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/aligned_field_loads/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/aligned_field_loads/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/aligned_field_stores/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/aligned_field_stores/Cargo.lock
index 4c5f7cb96a..14f058dbb6 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/aligned_field_stores/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/aligned_field_stores/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/array_constants/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/array_constants/Cargo.lock
index ef1ecaa2a8..ce35a52a27 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/array_constants/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/array_constants/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/array_for_loop/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/array_for_loop/Cargo.lock
index 3d737175ef..9c37ad4746 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/array_for_loop/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/array_for_loop/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/array_index/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/array_index/Cargo.lock
index d3a3501bfa..77fcad0815 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/array_index/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/array_index/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/array_init/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/array_init/Cargo.lock
index 86fba05561..3d601a9385 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/array_init/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/array_init/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/asin_acos_smoke/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/asin_acos_smoke/Cargo.lock
index 802b3b906e..bf2dd09d04 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/asin_acos_smoke/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/asin_acos_smoke/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/assume_init_uninhabited_trap/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/assume_init_uninhabited_trap/Cargo.lock
index 3adf8ffba8..a9ab6e0eaf 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/assume_init_uninhabited_trap/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/assume_init_uninhabited_trap/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/async_mlp/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/async_mlp/Cargo.lock
index 7028e34dbe..a622267f19 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/async_mlp/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/async_mlp/Cargo.lock
@@ -195,6 +195,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/async_vecadd/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/async_vecadd/Cargo.lock
index 2174866807..df129462f8 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/async_vecadd/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/async_vecadd/Cargo.lock
@@ -195,6 +195,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/atomic_f16/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/atomic_f16/Cargo.lock
index b9a5b8237d..322acaf17b 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/atomic_f16/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/atomic_f16/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/atomics/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/atomics/Cargo.lock
index 88b7fea84b..fede44e2db 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/atomics/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/atomics/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/barrier/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/barrier/Cargo.lock
index 8b838a20e4..400eca739f 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/barrier/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/barrier/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/bf16x2_arith/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/bf16x2_arith/Cargo.lock
index de59e51102..174790ea10 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/bf16x2_arith/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/bf16x2_arith/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/bf16x2_fma/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/bf16x2_fma/Cargo.lock
index 251e9a2a67..d24b0270d6 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/bf16x2_fma/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/bf16x2_fma/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/bool_phi_cmp/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/bool_phi_cmp/Cargo.lock
index 670e54de0c..0d03aa9c63 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/bool_phi_cmp/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/bool_phi_cmp/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/carrying_mul_add/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/carrying_mul_add/Cargo.lock
index c07608a163..560a6b26b5 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/carrying_mul_add/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/carrying_mul_add/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/cast_chunks/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/cast_chunks/Cargo.lock
index 7354ec93d3..2f78702a18 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/cast_chunks/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/cast_chunks/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/cast_tests/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/cast_tests/Cargo.lock
index 5a0f4f1248..31b2859cd2 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/cast_tests/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/cast_tests/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/cbrt_smoke/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/cbrt_smoke/Cargo.lock
index 3dac7b1997..4e0b5d92db 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/cbrt_smoke/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/cbrt_smoke/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/checked_arith/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/checked_arith/Cargo.lock
index aef438f58d..6a3a84f463 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/checked_arith/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/checked_arith/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/clc/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/clc/Cargo.lock
index b2358c8f96..ae578c1618 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/clc/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/clc/Cargo.lock
@@ -173,6 +173,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/cluster/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/cluster/Cargo.lock
index 428561d84b..59c6ebcb99 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/cluster/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/cluster/Cargo.lock
@@ -173,6 +173,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/compiler_features/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/compiler_features/Cargo.lock
index 33acd7147e..31e2968aba 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/compiler_features/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/compiler_features/Cargo.lock
@@ -173,6 +173,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/complex_mul/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/complex_mul/Cargo.lock
index c4fa3cf464..2994850293 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/complex_mul/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/complex_mul/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/const_aggregate/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/const_aggregate/Cargo.lock
index 6ea623f850..d9670c2974 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/const_aggregate/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/const_aggregate/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/const_bool_dead_branch/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/const_bool_dead_branch/Cargo.lock
index 91b2876e08..1282da8d82 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/const_bool_dead_branch/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/const_bool_dead_branch/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/const_generic/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/const_generic/Cargo.lock
index b15a635b4c..6f1603f622 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/const_generic/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/const_generic/Cargo.lock
@@ -173,6 +173,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/const_table_lookup/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/const_table_lookup/Cargo.lock
index 8ae82696df..57d0b519c5 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/const_table_lookup/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/const_table_lookup/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/const_tuple_table_lookup/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/const_tuple_table_lookup/Cargo.lock
index 53037c381d..b42f652e25 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/const_tuple_table_lookup/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/const_tuple_table_lookup/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/constant_index_from_end/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/constant_index_from_end/Cargo.lock
index ab06a74fe1..23c60b7fd6 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/constant_index_from_end/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/constant_index_from_end/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/constant_memory/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/constant_memory/Cargo.lock
index 441dc70016..ba9b15877a 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/constant_memory/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/constant_memory/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/constant_memory_coeffs/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/constant_memory_coeffs/Cargo.lock
index cbd2c4a773..c6d99a07bb 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/constant_memory_coeffs/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/constant_memory_coeffs/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/constant_memory_simple/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/constant_memory_simple/Cargo.lock
index 8656b34c60..9f9b183fae 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/constant_memory_simple/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/constant_memory_simple/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/constant_memory_table/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/constant_memory_table/Cargo.lock
index e1b980e909..1e4db21b90 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/constant_memory_table/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/constant_memory_table/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/coop_groups_demo/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/coop_groups_demo/Cargo.lock
index cd1195cbd1..8632e9b398 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/coop_groups_demo/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/coop_groups_demo/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/copy_aggregate_borrow/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/copy_aggregate_borrow/Cargo.lock
index 0ef4472db0..9d0ba925b7 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/copy_aggregate_borrow/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/copy_aggregate_borrow/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/copy_nonoverlapping/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/copy_nonoverlapping/Cargo.lock
index 03fc3fb565..d9a3eb8d82 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/copy_nonoverlapping/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/copy_nonoverlapping/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/counted_barrier/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/counted_barrier/Cargo.lock
index ee38589b3f..8db06bb71d 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/counted_barrier/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/counted_barrier/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/cp_async_small/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/cp_async_small/Cargo.lock
index 8884ed45e2..cbf4847e40 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/cp_async_small/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/cp_async_small/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/cp_async_zfill/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/cp_async_zfill/Cargo.lock
index 4c768d4fe9..3c9da53424 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/cp_async_zfill/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/cp_async_zfill/Cargo.lock
@@ -172,6 +172,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/cpp_consumes_rust_device/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/cpp_consumes_rust_device/Cargo.lock
index 51381eb264..8b304c5544 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/cpp_consumes_rust_device/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/cpp_consumes_rust_device/Cargo.lock
@@ -171,6 +171,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/cross_crate_embedded/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/cross_crate_embedded/Cargo.lock
index 58f9f47452..d47e365398 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/cross_crate_embedded/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/cross_crate_embedded/Cargo.lock
@@ -173,6 +173,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/cross_crate_kernel/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/cross_crate_kernel/Cargo.lock
index 4142ea1aa4..6e62fdf5e4 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/cross_crate_kernel/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/cross_crate_kernel/Cargo.lock
@@ -174,6 +174,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/cross_crate_merged_symbols/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/cross_crate_merged_symbols/Cargo.lock
index 12297db68c..32dd548ccf 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/cross_crate_merged_symbols/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/cross_crate_merged_symbols/Cargo.lock
@@ -173,6 +173,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/cross_crate_recursion/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/cross_crate_recursion/Cargo.lock
index 8f2b797d96..6c082aee9c 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/cross_crate_recursion/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/cross_crate_recursion/Cargo.lock
@@ -173,6 +173,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/cuda_module_contract/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/cuda_module_contract/Cargo.lock
index 08be6237f6..7999275a82 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/cuda_module_contract/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/cuda_module_contract/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/cuda_module_in_lib/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/cuda_module_in_lib/Cargo.lock
index ed02b84ccf..99d34e1097 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/cuda_module_in_lib/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/cuda_module_in_lib/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/cuda_module_nested/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/cuda_module_nested/Cargo.lock
index 971388c6c3..80ad3282d1 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/cuda_module_nested/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/cuda_module_nested/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/cutile_inter_kernel/simt/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/cutile_inter_kernel/simt/Cargo.lock
index 35aaf4c667..e098d8f8a6 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/cutile_inter_kernel/simt/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/cutile_inter_kernel/simt/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/cvt_f16x2/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/cvt_f16x2/Cargo.lock
index 818576913e..a52834b597 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/cvt_f16x2/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/cvt_f16x2/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/cvt_packed/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/cvt_packed/Cargo.lock
index 94d0618251..e55366cca0 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/cvt_packed/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/cvt_packed/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/debug-tests/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/debug-tests/Cargo.lock
index f5df8eec21..22ff9950c7 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/debug-tests/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/debug-tests/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/debug/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/debug/Cargo.lock
index 06c9f62bfd..a267c52e34 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/debug/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/debug/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/debug_pointer_locals/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/debug_pointer_locals/Cargo.lock
index a163c07916..4d81c25f2d 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/debug_pointer_locals/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/debug_pointer_locals/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/device_closures/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/device_closures/Cargo.lock
index a71448a33a..05d2607ef2 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/device_closures/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/device_closures/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/device_ffi_test/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/device_ffi_test/Cargo.lock
index a158e03b8d..797df39f5f 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/device_ffi_test/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/device_ffi_test/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/device_global/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/device_global/Cargo.lock
index 35ee0cd474..b8838aba25 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/device_global/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/device_global/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/device_panic_trap/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/device_panic_trap/Cargo.lock
index 6407c721d4..e44736ad56 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/device_panic_trap/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/device_panic_trap/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/disjoint_from_raw_parts/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/disjoint_from_raw_parts/Cargo.lock
index 16f5eeea64..bf3dc8d9b3 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/disjoint_from_raw_parts/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/disjoint_from_raw_parts/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/disjoint_slice_len/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/disjoint_slice_len/Cargo.lock
index dd634310c2..cebe8bd5bf 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/disjoint_slice_len/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/disjoint_slice_len/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/dotprod/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/dotprod/Cargo.lock
index 6eae6d938e..069e34dcdb 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/dotprod/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/dotprod/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/dpx_minmax_chains/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/dpx_minmax_chains/Cargo.lock
index e089b4e94a..22b6fe4262 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/dpx_minmax_chains/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/dpx_minmax_chains/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/drop_glue/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/drop_glue/Cargo.lock
index 266b7f6694..56f88f26f6 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/drop_glue/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/drop_glue/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/dst_metadata/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/dst_metadata/Cargo.lock
index 9c6fb2bc3e..6fa6923f1c 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/dst_metadata/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/dst_metadata/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/dynamic_smem/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/dynamic_smem/Cargo.lock
index d00f6456ab..2f91b1b80e 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/dynamic_smem/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/dynamic_smem/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/elect_leader/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/elect_leader/Cargo.lock
index 24b95937f2..ac7ed0960c 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/elect_leader/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/elect_leader/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_abi_multi_payload/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_abi_multi_payload/Cargo.lock
index 6139e1050f..b7b1ff11d8 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_abi_multi_payload/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_abi_multi_payload/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_array_match/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_array_match/Cargo.lock
index 4985ca6d9f..729b7498d1 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_array_match/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_array_match/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_constant_provenance/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_constant_provenance/Cargo.lock
index 8af4e5d135..03d5cce642 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_constant_provenance/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_constant_provenance/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_payload_addr/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_payload_addr/Cargo.lock
index c952ef0989..6c2739d99b 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_payload_addr/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_payload_addr/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_state_transitions/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_state_transitions/Cargo.lock
index 3a6c6a60cc..f264165a88 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_state_transitions/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_state_transitions/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_table_lookup/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_table_lookup/Cargo.lock
index 984dd1bb34..fd162b7375 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_table_lookup/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/enum_table_lookup/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/error/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/error/Cargo.lock
index aaa54e5982..4c5de7befb 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/error/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/error/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_enum_bool_payload_addr/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_enum_bool_payload_addr/Cargo.lock
index ce15431ac5..1da6540a30 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_enum_bool_payload_addr/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_enum_bool_payload_addr/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_enum_pointer_overlap/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_enum_pointer_overlap/Cargo.lock
index 0fa3b0a4ee..22ebef9cda 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_enum_pointer_overlap/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_enum_pointer_overlap/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_enum_shared_pointer_layout/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_enum_shared_pointer_layout/Cargo.lock
index f06797a06b..f587eefdbd 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_enum_shared_pointer_layout/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_enum_shared_pointer_layout/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_abi/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_abi/Cargo.lock
index 4c5a1cecc5..7ced6fd3c9 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_abi/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_abi/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_callable/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_callable/Cargo.lock
index edd815e111..6e3908dc5e 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_callable/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_callable/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_fn_pointer/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_fn_pointer/Cargo.lock
index d5446db5d0..768e54f50d 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_fn_pointer/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_fn_pointer/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_unknown_id/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_unknown_id/Cargo.lock
index 9b25129d83..314a8e3a24 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_unknown_id/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_generated_intrinsic_unknown_id/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_heap_alloc/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_heap_alloc/Cargo.lock
index 1a466544e6..cf95245bfa 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_heap_alloc/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_heap_alloc/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_host_arch_intrinsic/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_host_arch_intrinsic/Cargo.lock
index 7191053847..a2281de6c3 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_host_arch_intrinsic/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_host_arch_intrinsic/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_host_target_feature/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_host_target_feature/Cargo.lock
index 3219ee0d34..34236d8eeb 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_host_target_feature/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_host_target_feature/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_kernel_shared_param/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_kernel_shared_param/Cargo.lock
index 44916f1a99..62de228b0e 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_kernel_shared_param/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_kernel_shared_param/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_kernel_shared_ptr_param/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_kernel_shared_ptr_param/Cargo.lock
index c46db5e9e1..452b398de4 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_kernel_shared_ptr_param/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_kernel_shared_ptr_param/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_missing_device_attr/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_missing_device_attr/Cargo.lock
index 649e784194..5f10d797ed 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_missing_device_attr/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_missing_device_attr/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_set_discriminant_uninhabited/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_set_discriminant_uninhabited/Cargo.lock
index 5c81d90910..a09039fb2f 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_set_discriminant_uninhabited/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_set_discriminant_uninhabited/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_union_shared_pointer_addr/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_union_shared_pointer_addr/Cargo.lock
index 746e3e1063..069728fbf3 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_union_shared_pointer_addr/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_union_shared_pointer_addr/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_union_shared_pointer_layout/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_union_shared_pointer_layout/Cargo.lock
index ac804086d6..51114bcd30 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/error_union_shared_pointer_layout/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/error_union_shared_pointer_layout/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/ex2_approx_f16/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/ex2_approx_f16/Cargo.lock
index 33cd0aa21e..9f4b0dbc3d 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/ex2_approx_f16/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/ex2_approx_f16/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/exact_div_chunks/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/exact_div_chunks/Cargo.lock
index 93a2a2860b..b40cde9b3a 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/exact_div_chunks/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/exact_div_chunks/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/export_name_policy/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/export_name_policy/Cargo.lock
index ad50057c7b..c2658e679d 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/export_name_policy/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/export_name_policy/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/extern_crate_closure/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/extern_crate_closure/Cargo.lock
index a3de61cf65..c108e5725c 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/extern_crate_closure/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/extern_crate_closure/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/extern_libdevice/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/extern_libdevice/Cargo.lock
index 676af798d0..70d23cff54 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/extern_libdevice/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/extern_libdevice/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/f16_stress/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/f16_stress/Cargo.lock
index 1981d414f2..b1bb905e4f 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/f16_stress/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/f16_stress/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/f16x2_arith/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/f16x2_arith/Cargo.lock
index 31579841ea..efa08a21e6 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/f16x2_arith/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/f16x2_arith/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/f32x2_arith/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/f32x2_arith/Cargo.lock
index 9bc296863f..ba8266062d 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/f32x2_arith/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/f32x2_arith/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/fcmp_ne_nan_smoke/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/fcmp_ne_nan_smoke/Cargo.lock
index 71b0d21466..69d5da0813 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/fcmp_ne_nan_smoke/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/fcmp_ne_nan_smoke/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/field_array_assign/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/field_array_assign/Cargo.lock
index b198305b8e..8acce50102 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/field_array_assign/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/field_array_assign/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/float_rounding/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/float_rounding/Cargo.lock
index 80dfa026eb..0389eb2121 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/float_rounding/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/float_rounding/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/fmaxmin_smoke/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/fmaxmin_smoke/Cargo.lock
index a3233b3c21..76164b9541 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/fmaxmin_smoke/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/fmaxmin_smoke/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/fn_ptr_naming/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/fn_ptr_naming/Cargo.lock
index 2a909c83ac..c201e9c730 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/fn_ptr_naming/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/fn_ptr_naming/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/fp_special_literal_smoke/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/fp_special_literal_smoke/Cargo.lock
index 86eb3f25e8..bc07a66fed 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/fp_special_literal_smoke/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/fp_special_literal_smoke/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/Cargo.lock
index fc1c163579..a4affef696 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/future_apis/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/future_apis/Cargo.lock
index cbfd3853ea..a6eb1b4cab 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/future_apis/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/future_apis/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/gemm/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/gemm/Cargo.lock
index d83a865a01..3ed2f413bd 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/gemm/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/gemm/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/gemm_sol/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/gemm_sol/Cargo.lock
index 4ce39a1971..dc3781d854 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/gemm_sol/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/gemm_sol/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/gemm_sol_final/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/gemm_sol_final/Cargo.lock
index ad956f3af4..5e57abc8ab 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/gemm_sol_final/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/gemm_sol_final/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/gemm_views/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/gemm_views/Cargo.lock
index 9a3324ea0b..fd1e9a9508 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/gemm_views/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/gemm_views/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/generated_intrinsics/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/generated_intrinsics/Cargo.lock
index cd9f2ba0b2..55330c773e 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/generated_intrinsics/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/generated_intrinsics/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/generated_intrinsics_blackwell/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/generated_intrinsics_blackwell/Cargo.lock
index 04de2c5f8b..e147cf0415 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/generated_intrinsics_blackwell/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/generated_intrinsics_blackwell/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/generated_ldmatrix/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/generated_ldmatrix/Cargo.lock
index bb356b611b..10a9b6e21d 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/generated_ldmatrix/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/generated_ldmatrix/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/generic/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/generic/Cargo.lock
index 6a78ca4f12..0669b206f8 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/generic/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/generic/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/hashmap/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/hashmap/Cargo.lock
index 86b24f7d2a..a044b7648f 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/hashmap/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/hashmap/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/hashmap_v2/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/hashmap_v2/Cargo.lock
index 0fb7f0e99e..3808449ee9 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/hashmap_v2/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/hashmap_v2/Cargo.lock
@@ -194,6 +194,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/hashmap_v3/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/hashmap_v3/Cargo.lock
index b781c994c5..1a4bd36e21 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/hashmap_v3/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/hashmap_v3/Cargo.lock
@@ -194,6 +194,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/hello_constant/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/hello_constant/Cargo.lock
index f3fe99746f..bdaf705a2c 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/hello_constant/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/hello_constant/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/helper_fn/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/helper_fn/Cargo.lock
index 7de74fdb2b..6826769094 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/helper_fn/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/helper_fn/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/host_closure/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/host_closure/Cargo.lock
index 05cdf65074..a68eed40e0 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/host_closure/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/host_closure/Cargo.lock
@@ -184,6 +184,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/host_closure_target_feature/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/host_closure_target_feature/Cargo.lock
index 6c77a5ff67..8b7506e1ca 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/host_closure_target_feature/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/host_closure_target_feature/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/iket_trace/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/iket_trace/Cargo.lock
index 86c75b8ff3..2a92b8cdae 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/iket_trace/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/iket_trace/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/image_convolution/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/image_convolution/Cargo.lock
index d5aa2ae65e..5388dd2a42 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/image_convolution/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/image_convolution/Cargo.lock
@@ -175,6 +175,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/index2d_const/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/index2d_const/Cargo.lock
index 3836acca0b..980e4244ba 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/index2d_const/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/index2d_const/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/index_field_assign/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/index_field_assign/Cargo.lock
index 1a1e93d7b5..80c3f0cf8d 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/index_field_assign/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/index_field_assign/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/inline_intent/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/inline_intent/Cargo.lock
index cb8ef2d6db..d131521698 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/inline_intent/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/inline_intent/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/inline_ptx/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/inline_ptx/Cargo.lock
index 53bc623360..6b030bf31b 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/inline_ptx/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/inline_ptx/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/integer_minmax/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/integer_minmax/Cargo.lock
index fe88f463bb..848bbd372c 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/integer_minmax/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/integer_minmax/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/interior_mutability/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/interior_mutability/Cargo.lock
index a6f716be40..9cbfdc41ce 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/interior_mutability/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/interior_mutability/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/interop_cubin_identity/device/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/interop_cubin_identity/device/Cargo.lock
index f22694c1a4..d8cb91f670 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/interop_cubin_identity/device/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/interop_cubin_identity/device/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/kernel_reference_validity/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/kernel_reference_validity/Cargo.lock
index 6100df50d1..155f6ea15c 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/kernel_reference_validity/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/kernel_reference_validity/Cargo.lock
@@ -184,6 +184,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/lanemask_scan/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/lanemask_scan/Cargo.lock
index 394622c351..e3d368896d 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/lanemask_scan/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/lanemask_scan/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/legacy_atomic_fadd/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/legacy_atomic_fadd/Cargo.lock
index 1e2ee59815..b5fdbce0b1 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/legacy_atomic_fadd/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/legacy_atomic_fadd/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/legacy_atomic_rmw_cas/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/legacy_atomic_rmw_cas/Cargo.lock
index ecfb175d50..c5a4eec96e 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/legacy_atomic_rmw_cas/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/legacy_atomic_rmw_cas/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/legacy_nvvm_pointer_shapes/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/legacy_nvvm_pointer_shapes/Cargo.lock
index 1aa0595d80..59cca1065c 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/legacy_nvvm_pointer_shapes/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/legacy_nvvm_pointer_shapes/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/libdevice_math/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/libdevice_math/Cargo.lock
index d285772669..fe8be5ce9b 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/libdevice_math/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/libdevice_math/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/libm_math/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/libm_math/Cargo.lock
index 460643529b..04cc7b4245 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/libm_math/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/libm_math/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/manual_launch_generic/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/manual_launch_generic/Cargo.lock
index 3f172b6f05..33d112c717 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/manual_launch_generic/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/manual_launch_generic/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/manual_launch_libdevice/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/manual_launch_libdevice/Cargo.lock
index 4118e90b10..ad12554b99 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/manual_launch_libdevice/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/manual_launch_libdevice/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/map_sum/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/map_sum/Cargo.lock
index 9924014ffa..a9578fc261 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/map_sum/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/map_sum/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/math_atan/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/math_atan/Cargo.lock
index 8ce1894329..d65ee5c323 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/math_atan/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/math_atan/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/math_hyperbolic/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/math_hyperbolic/Cargo.lock
index f7f1e6c274..8f675ccc91 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/math_hyperbolic/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/math_hyperbolic/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/math_no_mir_inline/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/math_no_mir_inline/Cargo.lock
index b7718a33a7..195447b60b 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/math_no_mir_inline/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/math_no_mir_inline/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/math_tan/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/math_tan/Cargo.lock
index 552212e5e4..fe7c7e0574 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/math_tan/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/math_tan/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/mathdx_ffi_test/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/mathdx_ffi_test/Cargo.lock
index 76fc9daba6..bdb4f16cb0 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/mathdx_ffi_test/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/mathdx_ffi_test/Cargo.lock
@@ -183,6 +183,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/mcast_barrier_test/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/mcast_barrier_test/Cargo.lock
index 47eec7302e..d6c39cd56d 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/mcast_barrier_test/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/mcast_barrier_test/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/mem_swap/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/mem_swap/Cargo.lock
index 611d7a182b..bc20bae99b 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/mem_swap/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/mem_swap/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/mir_projected_call_destination/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/mir_projected_call_destination/Cargo.lock
index 8791af0505..aa611ce7a7 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/mir_projected_call_destination/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/mir_projected_call_destination/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/mma_mxf8f6f4/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/mma_mxf8f6f4/Cargo.lock
index da7974c0f7..7cb1302c59 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/mma_mxf8f6f4/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/mma_mxf8f6f4/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/ne_bytes_transmute/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/ne_bytes_transmute/Cargo.lock
index 4bd0371e5b..c2ea4c8d18 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/ne_bytes_transmute/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/ne_bytes_transmute/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/nested_index_assignment/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/nested_index_assignment/Cargo.lock
index c52786e005..26698cc3b7 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/nested_index_assignment/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/nested_index_assignment/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/numeric_stress/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/numeric_stress/Cargo.lock
index bfd2fa4115..3eb22b3101 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/numeric_stress/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/numeric_stress/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/nvls_all_reduce/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/nvls_all_reduce/Cargo.lock
index 781a8d3f06..3d59b5eb5b 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/nvls_all_reduce/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/nvls_all_reduce/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/option_ref_transmute/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/option_ref_transmute/Cargo.lock
index b5af8f30b4..241d1f71b5 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/option_ref_transmute/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/option_ref_transmute/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/option_ref_unwrap_or/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/option_ref_unwrap_or/Cargo.lock
index f5893ba8a7..2cfd5247cd 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/option_ref_unwrap_or/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/option_ref_unwrap_or/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/option_unwrap_trap/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/option_unwrap_trap/Cargo.lock
index 85a0cbe346..ab9330cc61 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/option_unwrap_trap/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/option_unwrap_trap/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/ord_cmp/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/ord_cmp/Cargo.lock
index b80b1f9dcb..7c94d80d01 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/ord_cmp/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/ord_cmp/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/packed_aggregate_abi/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/packed_aggregate_abi/Cargo.lock
index f4c8348b24..0fd8145d5b 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/packed_aggregate_abi/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/packed_aggregate_abi/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/packed_atomic_add/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/packed_atomic_add/Cargo.lock
index 5c9aae69d3..e55602beeb 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/packed_atomic_add/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/packed_atomic_add/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/padded_constants/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/padded_constants/Cargo.lock
index 7621003859..850235692e 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/padded_constants/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/padded_constants/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/partial_warp_reduce/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/partial_warp_reduce/Cargo.lock
index 1dd0ec55a8..fcb32c2d00 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/partial_warp_reduce/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/partial_warp_reduce/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/pinned_overlap/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/pinned_overlap/Cargo.lock
index f00ea7f41c..3b17dbf8ef 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/pinned_overlap/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/pinned_overlap/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/place_read_fallbacks/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/place_read_fallbacks/Cargo.lock
index 63206cdbab..3593bc86e3 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/place_read_fallbacks/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/place_read_fallbacks/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/policy_config/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/policy_config/Cargo.lock
index ac860e426a..3a31705ad2 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/policy_config/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/policy_config/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/primitive_stress/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/primitive_stress/Cargo.lock
index f07ca73faa..b1e142be20 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/primitive_stress/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/primitive_stress/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/printf/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/printf/Cargo.lock
index cfaa742b4f..9876e60506 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/printf/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/printf/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/projected_place_read/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/projected_place_read/Cargo.lock
index 30bc83e4e5..cf03858c6c 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/projected_place_read/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/projected_place_read/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/proof_carrying_views/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/proof_carrying_views/Cargo.lock
index 42b08210f7..59d1530a47 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/proof_carrying_views/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/proof_carrying_views/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/ptr_copy/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/ptr_copy/Cargo.lock
index de5ac4ae48..96c32f646c 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/ptr_copy/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/ptr_copy/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/ptr_offset_from/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/ptr_offset_from/Cargo.lock
index 0ec7cfb44f..c944f05998 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/ptr_offset_from/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/ptr_offset_from/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/ptr_without_provenance/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/ptr_without_provenance/Cargo.lock
index fa5df66742..2a87651aa9 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/ptr_without_provenance/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/ptr_without_provenance/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/reborrow/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/reborrow/Cargo.lock
index 8d7de86367..3eb3e7a2f5 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/reborrow/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/reborrow/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/redux_f32/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/redux_f32/Cargo.lock
index a74777042e..0a1510d5e3 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/redux_f32/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/redux_f32/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/redux_minmax/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/redux_minmax/Cargo.lock
index bbb321da29..0ddd107d3e 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/redux_minmax/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/redux_minmax/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/redux_sum/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/redux_sum/Cargo.lock
index 3495a7289a..15e25fdce4 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/redux_sum/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/redux_sum/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/ref_index_projections/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/ref_index_projections/Cargo.lock
index 58a388e07a..d8153f9e8a 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/ref_index_projections/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/ref_index_projections/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/ref_operand_mul/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/ref_operand_mul/Cargo.lock
index ec2404f932..775d9e5614 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/ref_operand_mul/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/ref_operand_mul/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/repr_align_block_args/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/repr_align_block_args/Cargo.lock
index cc0078e643..9b04b546ed 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/repr_align_block_args/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/repr_align_block_args/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/repr_transparent_abi/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/repr_transparent_abi/Cargo.lock
index fa4c02c424..b680236d23 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/repr_transparent_abi/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/repr_transparent_abi/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/repr_u32_enum_stride/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/repr_u32_enum_stride/Cargo.lock
index 6af97378f8..fd57b9440b 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/repr_u32_enum_stride/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/repr_u32_enum_stride/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/row_width_slice/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/row_width_slice/Cargo.lock
index 85b2f5f477..4fc0d1f3a0 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/row_width_slice/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/row_width_slice/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/rustlantis-smoke/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/rustlantis-smoke/Cargo.lock
index 67e55271c0..14df4f7753 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/rustlantis-smoke/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/rustlantis-smoke/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/scoped_atomic_load_store/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/scoped_atomic_load_store/Cargo.lock
index 013a3379fb..4ff7e54f62 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/scoped_atomic_load_store/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/scoped_atomic_load_store/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/select_unpredictable/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/select_unpredictable/Cargo.lock
index 86eafa2b6e..f0c7a80076 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/select_unpredictable/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/select_unpredictable/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/set_discriminant/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/set_discriminant/Cargo.lock
index bbd354e38b..31edee41d8 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/set_discriminant/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/set_discriminant/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/set_discriminant_niche/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/set_discriminant_niche/Cargo.lock
index e157fea2ac..5dcde3d6c2 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/set_discriminant_niche/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/set_discriminant_niche/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_debug/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_debug/Cargo.lock
index 3195c5052d..b5d80744b0 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_debug/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_debug/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_memory_helper/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_memory_helper/Cargo.lock
index e72dcc60f5..184c8aa902 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_memory_helper/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_memory_helper/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_ptr_api/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_ptr_api/Cargo.lock
index 0df03ea524..5668e223c6 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_ptr_api/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_ptr_api/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_ptr_transmute/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_ptr_transmute/Cargo.lock
index 42647acf94..8c9f984130 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_ptr_transmute/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_ptr_transmute/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_slice_unsize/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_slice_unsize/Cargo.lock
index 7517a55603..e289fcb537 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_slice_unsize/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_slice_unsize/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_stage_select/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_stage_select/Cargo.lock
index 98f07a9a47..6fbded2f30 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_stage_select/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/shared_stage_select/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/sharedmem/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/sharedmem/Cargo.lock
index e60e6cf190..9bb0d78e59 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/sharedmem/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/sharedmem/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/shuffle_64/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/shuffle_64/Cargo.lock
index 46b3b08f05..0f8603f039 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/shuffle_64/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/shuffle_64/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/slice_from_raw_parts_cast/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/slice_from_raw_parts_cast/Cargo.lock
index 52b8d52887..e802bc80df 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/slice_from_raw_parts_cast/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/slice_from_raw_parts_cast/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/slice_get_mut/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/slice_get_mut/Cargo.lock
index 20152e80f8..32cf42a160 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/slice_get_mut/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/slice_get_mut/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/slice_index_bounds_check/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/slice_index_bounds_check/Cargo.lock
index e28e9d5ddf..1e31a33c93 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/slice_index_bounds_check/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/slice_index_bounds_check/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/slice_iterators/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/slice_iterators/Cargo.lock
index 0ac8bde6f2..76fc689175 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/slice_iterators/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/slice_iterators/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/slice_reslice/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/slice_reslice/Cargo.lock
index fe2e92d2b2..33da6fbe80 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/slice_reslice/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/slice_reslice/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/small_type_ffi_test/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/small_type_ffi_test/Cargo.lock
index d446e6625f..eeaf12fb8f 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/small_type_ffi_test/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/small_type_ffi_test/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/sparse_mma_fp8/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/sparse_mma_fp8/Cargo.lock
index c781eb7524..bc097a2811 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/sparse_mma_fp8/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/sparse_mma_fp8/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/sparse_mma_ordered_float/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/sparse_mma_ordered_float/Cargo.lock
index 03d67c5e3a..c4857b2ea1 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/sparse_mma_ordered_float/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/sparse_mma_ordered_float/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/split_at_mut/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/split_at_mut/Cargo.lock
index 4324edf483..da31062308 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/split_at_mut/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/split_at_mut/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/standalone_device_fn/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/standalone_device_fn/Cargo.lock
index bbf2681720..d7da353c80 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/standalone_device_fn/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/standalone_device_fn/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/static_slice_addend/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/static_slice_addend/Cargo.lock
index 15f0537070..ad906ed87f 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/static_slice_addend/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/static_slice_addend/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/static_slice_unsize/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/static_slice_unsize/Cargo.lock
index c074858de1..7ce691d222 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/static_slice_unsize/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/static_slice_unsize/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/step_by/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/step_by/Cargo.lock
index 84cba98cab..f8ecd4f3f5 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/step_by/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/step_by/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/struct_constant_provenance/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/struct_constant_provenance/Cargo.lock
index 690c24a327..276cfb0f77 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/struct_constant_provenance/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/struct_constant_provenance/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/struct_field_layout/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/struct_field_layout/Cargo.lock
index fe53acdb76..7c60ab7414 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/struct_field_layout/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/struct_field_layout/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/subslice_projection/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/subslice_projection/Cargo.lock
index 7ef72e0da5..8cffc182b9 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/subslice_projection/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/subslice_projection/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/switchint_128bit/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/switchint_128bit/Cargo.lock
index 03811356e6..ab807d611f 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/switchint_128bit/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/switchint_128bit/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/swizzle_smem/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/swizzle_smem/Cargo.lock
index f152cecc4c..7f2a725dae 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/swizzle_smem/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/swizzle_smem/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/tcgen05/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/tcgen05/Cargo.lock
index 8c6054ce23..6503ef5de1 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/tcgen05/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/tcgen05/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/tcgen05_matmul/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/tcgen05_matmul/Cargo.lock
index 651ffe661e..3566ee9735 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/tcgen05_matmul/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/tcgen05_matmul/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/thread_runs/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/thread_runs/Cargo.lock
index 75c208fc87..f42576446e 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/thread_runs/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/thread_runs/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/tiled_gemm/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/tiled_gemm/Cargo.lock
index cf9576c158..9e5641f5d2 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/tiled_gemm/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/tiled_gemm/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/tilelib_relu/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/tilelib_relu/Cargo.lock
index c86e4ed6e2..2b44ad8dbf 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/tilelib_relu/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/tilelib_relu/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/tma_copy/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/tma_copy/Cargo.lock
index 5a9c81ad8f..a8abc1c058 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/tma_copy/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/tma_copy/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/tma_multicast/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/tma_multicast/Cargo.lock
index 3977ab9ddc..35203ec7c1 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/tma_multicast/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/tma_multicast/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/toolchain_pairing_smoke/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/toolchain_pairing_smoke/Cargo.lock
index 29f8b4dec6..e9490ea086 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/toolchain_pairing_smoke/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/toolchain_pairing_smoke/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/transmute_scalar_struct/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/transmute_scalar_struct/Cargo.lock
index 0e2998f8d5..11f7e71e47 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/transmute_scalar_struct/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/transmute_scalar_struct/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/tuple_array_provenance/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/tuple_array_provenance/Cargo.lock
index de10cae507..1a11848cfb 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/tuple_array_provenance/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/tuple_array_provenance/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/tuple_constant_provenance/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/tuple_constant_provenance/Cargo.lock
index 48212e960c..65b470274d 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/tuple_constant_provenance/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/tuple_constant_provenance/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/tuple_field_lookup/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/tuple_field_lookup/Cargo.lock
index 47aad8a282..2a8dc92e90 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/tuple_field_lookup/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/tuple_field_lookup/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/tuple_return/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/tuple_return/Cargo.lock
index 2ac68321ee..d575390667 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/tuple_return/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/tuple_return/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/unaligned_memory/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/unaligned_memory/Cargo.lock
index 4affc8b321..fb67d22fcd 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/unaligned_memory/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/unaligned_memory/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/unchecked_indexing/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/unchecked_indexing/Cargo.lock
index 53f368594d..a35430399f 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/unchecked_indexing/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/unchecked_indexing/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/uninit_const/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/uninit_const/Cargo.lock
index 1517384a3b..cbc40b040c 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/uninit_const/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/uninit_const/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/union_aggregate/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/union_aggregate/Cargo.lock
index c2a4fec2d7..93d8acc488 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/union_aggregate/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/union_aggregate/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/unroll_bounds_check/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/unroll_bounds_check/Cargo.lock
index 71e089c287..92e3383a18 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/unroll_bounds_check/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/unroll_bounds_check/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/unroll_smoke/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/unroll_smoke/Cargo.lock
index 780f6f3e9d..aaba93322e 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/unroll_smoke/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/unroll_smoke/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/vecadd/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/vecadd/Cargo.lock
index 9d76a571bd..41b1406225 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/vecadd/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/vecadd/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/vectorization/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/vectorization/Cargo.lock
index 6cb5e3d2a9..02a5fc0d01 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/vectorization/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/vectorization/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/volatile_memory/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/volatile_memory/Cargo.lock
index 51689b847a..c6d2894dbb 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/volatile_memory/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/volatile_memory/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/warp_reduce/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/warp_reduce/Cargo.lock
index 49805d2fe3..6b90d6b3a8 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/warp_reduce/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/warp_reduce/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/warp_reduce_redux/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/warp_reduce_redux/Cargo.lock
new file mode 100644
index 0000000000..4d8f92986f
--- /dev/null
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/warp_reduce_redux/Cargo.lock
@@ -0,0 +1,754 @@
+# This file is automatically @generated by Cargo.
+# It is not intended for manual editing.
+version = 4
+
+[[package]]
+name = "aho-corasick"
+version = "1.1.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ddd31a130427c27518df266943a5308ed92d4b226cc639f5a8f1002816174301"
+dependencies = [
+ "memchr",
+]
+
+[[package]]
+name = "anyhow"
+version = "1.0.102"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7f202df86484c868dbad7eaa557ef785d5c66295e41b460ef922eca0723b842c"
+
+[[package]]
+name = "bindgen"
+version = "0.69.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "271383c67ccabffb7381723dea0672a673f292304fcb45c01cc648c7a8d58088"
+dependencies = [
+ "bitflags",
+ "cexpr",
+ "clang-sys",
+ "itertools",
+ "lazy_static",
+ "lazycell",
+ "log",
+ "prettyplease",
+ "proc-macro2",
+ "quote",
+ "regex",
+ "rustc-hash",
+ "shlex",
+ "syn",
+ "which",
+]
+
+[[package]]
+name = "bitflags"
+version = "2.13.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b4388bee8683e3d04af747c73422af53102d2bd24d9eadb6cbc100baef4b43f8"
+
+[[package]]
+name = "block-buffer"
+version = "0.10.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3078c7629b62d3f0439517fa394996acacc5cbc91c5a20d8c658e77abd503a71"
+dependencies = [
+ "generic-array",
+]
+
+[[package]]
+name = "cexpr"
+version = "0.6.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6fac387a98bb7c37292057cffc56d62ecb629900026402633ae9160df93a8766"
+dependencies = [
+ "nom",
+]
+
+[[package]]
+name = "cfg-if"
+version = "1.0.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9330f8b2ff13f34540b44e946ef35111825727b38d33286ef986142615121801"
+
+[[package]]
+name = "clang-sys"
+version = "1.8.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0b023947811758c97c59bf9d1c188fd619ad4718dcaa767947df1cadb14f39f4"
+dependencies = [
+ "glob",
+ "libc",
+ "libloading",
+]
+
+[[package]]
+name = "cpufeatures"
+version = "0.2.17"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "59ed5838eebb26a2bb2e58f6d5b5316989ae9d08bab10e0e6d103e656d1b0280"
+dependencies = [
+ "libc",
+]
+
+[[package]]
+name = "crc32fast"
+version = "1.5.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9481c1c90cbf2ac953f07c8d4a58aa3945c425b7185c9154d67a65e4230da511"
+dependencies = [
+ "cfg-if",
+]
+
+[[package]]
+name = "crunchy"
+version = "0.2.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "460fbee9c2c2f33933d720630a6a0bac33ba7053db5344fac858d4b8952d77d5"
+
+[[package]]
+name = "crypto-common"
+version = "0.1.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "78c8292055d1c1df0cce5d180393dc8cce0abec0a7102adb6c7b1eef6016d60a"
+dependencies = [
+ "generic-array",
+ "typenum",
+]
+
+[[package]]
+name = "cuda-artifact-finalizer"
+version = "0.2.1"
+dependencies = [
+ "libnvvm-sys",
+ "nvjitlink-sys",
+ "serde",
+ "sha2",
+ "thiserror",
+]
+
+[[package]]
+name = "cuda-bindings"
+version = "0.4.0"
+dependencies = [
+ "bindgen",
+ "libloading",
+ "prettyplease",
+ "proc-macro2",
+ "quote",
+ "syn",
+]
+
+[[package]]
+name = "cuda-core"
+version = "0.4.0"
+dependencies = [
+ "anyhow",
+ "cuda-bindings",
+ "cuda-core-derive",
+ "half",
+ "oxide-artifacts",
+]
+
+[[package]]
+name = "cuda-core-derive"
+version = "0.4.0"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn",
+]
+
+[[package]]
+name = "cuda-device"
+version = "0.2.1"
+dependencies = [
+ "cuda-macros",
+ "cuda-target-spec",
+]
+
+[[package]]
+name = "cuda-host"
+version = "0.2.1"
+dependencies = [
+ "cuda-artifact-finalizer",
+ "cuda-core",
+ "cuda-macros",
+ "half",
+ "libc",
+ "oxide-artifacts",
+ "ptx-parse",
+ "sha2",
+ "thiserror",
+]
+
+[[package]]
+name = "cuda-macros"
+version = "0.2.1"
+dependencies = [
+ "oxide-tile-export",
+ "proc-macro2",
+ "quote",
+ "reserved-oxide-symbols",
+ "syn",
+]
+
+[[package]]
+name = "cuda-target-spec"
+version = "0.1.0"
+
+[[package]]
+name = "digest"
+version = "0.10.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9ed9a281f7bc9b7576e61468ba615a66a5c8cfdff42420a70aa82701a3b1e292"
+dependencies = [
+ "block-buffer",
+ "crypto-common",
+]
+
+[[package]]
+name = "either"
+version = "1.16.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "91622ff5e7162018101f2fea40d6ebf4a78bbe5a49736a2020649edf9693679e"
+
+[[package]]
+name = "equivalent"
+version = "1.0.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "877a4ace8713b0bcf2a4e7eec82529c029f1d0619886d18145fea96c3ffe5c0f"
+
+[[package]]
+name = "errno"
+version = "0.3.14"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "39cab71617ae0d63f51a36d69f866391735b51691dbda63cf6f96d042b63efeb"
+dependencies = [
+ "libc",
+ "windows-sys 0.61.2",
+]
+
+[[package]]
+name = "foldhash"
+version = "0.1.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d9c4f5dac5e15c24eb999c26181a6ca40b39fe946cbe4c263c7209467bc83af2"
+
+[[package]]
+name = "generic-array"
+version = "0.14.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "85649ca51fd72272d7821adaf274ad91c288277713d9c18820d8499a7ff69e9a"
+dependencies = [
+ "typenum",
+ "version_check",
+]
+
+[[package]]
+name = "glob"
+version = "0.3.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0cc23270f6e1808e30a928bdc84dea0b9b4136a8bc82338574f23baf47bbd280"
+
+[[package]]
+name = "half"
+version = "2.7.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6ea2d84b969582b4b1864a92dc5d27cd2b77b622a8d79306834f1be5ba20d84b"
+dependencies = [
+ "cfg-if",
+ "crunchy",
+ "zerocopy",
+]
+
+[[package]]
+name = "hashbrown"
+version = "0.15.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9229cfe53dfd69f0609a49f65461bd93001ea1ef889cd5529dd176593f5338a1"
+dependencies = [
+ "foldhash",
+]
+
+[[package]]
+name = "hashbrown"
+version = "0.17.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ed5909b6e89a2db4456e54cd5f673791d7eca6732202bbf2a9cc504fe2f9b84a"
+
+[[package]]
+name = "home"
+version = "0.5.12"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "cc627f471c528ff0c4a49e1d5e60450c8f6461dd6d10ba9dcd3a61d3dff7728d"
+dependencies = [
+ "windows-sys 0.61.2",
+]
+
+[[package]]
+name = "indexmap"
+version = "2.14.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d466e9454f08e4a911e14806c24e16fba1b4c121d1ea474396f396069cf949d9"
+dependencies = [
+ "equivalent",
+ "hashbrown 0.17.1",
+]
+
+[[package]]
+name = "itertools"
+version = "0.12.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ba291022dbbd398a455acf126c1e341954079855bc60dfdda641363bd6922569"
+dependencies = [
+ "either",
+]
+
+[[package]]
+name = "itoa"
+version = "1.0.18"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8f42a60cbdf9a97f5d2305f08a87dc4e09308d1276d28c869c684d7777685682"
+
+[[package]]
+name = "lazy_static"
+version = "1.5.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "bbd2bcb4c963f2ddae06a2efc7e9f3591312473c50c6685e1f298068316e66fe"
+
+[[package]]
+name = "lazycell"
+version = "1.3.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "830d08ce1d1d941e6b30645f1a0eb5643013d835ce3779a5fc208261dbe10f55"
+
+[[package]]
+name = "libc"
+version = "0.2.186"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "68ab91017fe16c622486840e4c83c9a37afeff978bd239b5293d61ece587de66"
+
+[[package]]
+name = "libloading"
+version = "0.8.9"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d7c4b02199fee7c5d21a5ae7d8cfa79a6ef5bb2fc834d6e9058e89c825efdc55"
+dependencies = [
+ "cfg-if",
+ "windows-link",
+]
+
+[[package]]
+name = "libnvvm-sys"
+version = "0.2.1"
+dependencies = [
+ "cuda-target-spec",
+ "libloading",
+ "thiserror",
+]
+
+[[package]]
+name = "linux-raw-sys"
+version = "0.4.15"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d26c52dbd32dccf2d10cac7725f8eae5296885fb5703b261f7d0a0739ec807ab"
+
+[[package]]
+name = "log"
+version = "0.4.29"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "5e5032e24019045c762d3c0f28f5b6b8bbf38563a65908389bf7978758920897"
+
+[[package]]
+name = "memchr"
+version = "2.8.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "cf8baf1c55e62ffcace7a9f06f4bd9cd3f0c4beb022d3b367256b91b87513d98"
+
+[[package]]
+name = "minimal-lexical"
+version = "0.2.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "68354c5c6bd36d73ff3feceb05efa59b6acb7626617f4962be322a825e61f79a"
+
+[[package]]
+name = "nom"
+version = "7.1.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d273983c5a657a70a3e8f2a01329822f3b8c8172b73826411a55751e404a0a4a"
+dependencies = [
+ "memchr",
+ "minimal-lexical",
+]
+
+[[package]]
+name = "nvjitlink-sys"
+version = "0.2.1"
+dependencies = [
+ "libloading",
+ "thiserror",
+]
+
+[[package]]
+name = "object"
+version = "0.36.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "62948e14d923ea95ea2c7c86c71013138b66525b86bdc08d2dcc262bdb497b87"
+dependencies = [
+ "crc32fast",
+ "hashbrown 0.15.5",
+ "indexmap",
+ "memchr",
+]
+
+[[package]]
+name = "once_cell"
+version = "1.21.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9f7c3e4beb33f85d45ae3e3a1792185706c8e16d043238c593331cc7cd313b50"
+
+[[package]]
+name = "oxide-artifacts"
+version = "0.2.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "74d9a941c0c43d2dceaf25f5c9ca4c9f9cceb28dfa5e4e3cfd14e476f003e7c9"
+dependencies = [
+ "object",
+]
+
+[[package]]
+name = "oxide-tile-export"
+version = "0.2.1"
+dependencies = [
+ "serde",
+ "serde_json",
+ "thiserror",
+]
+
+[[package]]
+name = "prettyplease"
+version = "0.2.37"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "479ca8adacdd7ce8f1fb39ce9ecccbfe93a3f1344b3d0d97f20bc0196208f62b"
+dependencies = [
+ "proc-macro2",
+ "syn",
+]
+
+[[package]]
+name = "proc-macro2"
+version = "1.0.106"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8fd00f0bb2e90d81d1044c2b32617f68fcb9fa3bb7640c23e9c748e53fb30934"
+dependencies = [
+ "unicode-ident",
+]
+
+[[package]]
+name = "ptx-parse"
+version = "0.1.0"
+
+[[package]]
+name = "quote"
+version = "1.0.45"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "41f2619966050689382d2b44f664f4bc593e129785a36d6ee376ddf37259b924"
+dependencies = [
+ "proc-macro2",
+]
+
+[[package]]
+name = "regex"
+version = "1.12.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e10754a14b9137dd7b1e3e5b0493cc9171fdd105e0ab477f51b72e7f3ac0e276"
+dependencies = [
+ "aho-corasick",
+ "memchr",
+ "regex-automata",
+ "regex-syntax",
+]
+
+[[package]]
+name = "regex-automata"
+version = "0.4.14"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6e1dd4122fc1595e8162618945476892eefca7b88c52820e74af6262213cae8f"
+dependencies = [
+ "aho-corasick",
+ "memchr",
+ "regex-syntax",
+]
+
+[[package]]
+name = "regex-syntax"
+version = "0.8.11"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d6f6ff9a378485b298a5286656da665ba74413d36db0979633275d2e708145d4"
+
+[[package]]
+name = "reserved-oxide-symbols"
+version = "0.2.1"
+
+[[package]]
+name = "rustc-hash"
+version = "1.1.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "08d43f7aa6b08d49f382cde6a7982047c3426db949b1424bc4b7ec9ae12c6ce2"
+
+[[package]]
+name = "rustix"
+version = "0.38.44"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "fdb5bc1ae2baa591800df16c9ca78619bf65c0488b41b96ccec5d11220d8c154"
+dependencies = [
+ "bitflags",
+ "errno",
+ "libc",
+ "linux-raw-sys",
+ "windows-sys 0.59.0",
+]
+
+[[package]]
+name = "serde"
+version = "1.0.228"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9a8e94ea7f378bd32cbbd37198a4a91436180c5bb472411e48b5ec2e2124ae9e"
+dependencies = [
+ "serde_core",
+ "serde_derive",
+]
+
+[[package]]
+name = "serde_core"
+version = "1.0.228"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "41d385c7d4ca58e59fc732af25c3983b67ac852c1a25000afe1175de458b67ad"
+dependencies = [
+ "serde_derive",
+]
+
+[[package]]
+name = "serde_derive"
+version = "1.0.228"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d540f220d3187173da220f885ab66608367b6574e925011a9353e4badda91d79"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn",
+]
+
+[[package]]
+name = "serde_json"
+version = "1.0.154"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e7e9cc8b1b85264074fbcc02a88680c4096b1e47df8f739dceb03bf482f04bd6"
+dependencies = [
+ "itoa",
+ "memchr",
+ "serde",
+ "serde_core",
+ "zmij",
+]
+
+[[package]]
+name = "sha2"
+version = "0.10.9"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a7507d819769d01a365ab707794a4084392c824f54a7a6a7862f8c3d0892b283"
+dependencies = [
+ "cfg-if",
+ "cpufeatures",
+ "digest",
+]
+
+[[package]]
+name = "shlex"
+version = "1.3.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0fda2ff0d084019ba4d7c6f371c95d8fd75ce3524c3cb8fb653a3023f6323e64"
+
+[[package]]
+name = "syn"
+version = "2.0.117"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e665b8803e7b1d2a727f4023456bbbbe74da67099c585258af0ad9c5013b9b99"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "unicode-ident",
+]
+
+[[package]]
+name = "thiserror"
+version = "2.0.18"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "4288b5bcbc7920c07a1149a35cf9590a2aa808e0bc1eafaade0b80947865fbc4"
+dependencies = [
+ "thiserror-impl",
+]
+
+[[package]]
+name = "thiserror-impl"
+version = "2.0.18"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ebc4ee7f67670e9b64d05fa4253e753e016c6c95ff35b89b7941d6b856dec1d5"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn",
+]
+
+[[package]]
+name = "typenum"
+version = "1.20.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b6f5e870be6c3b371b77fe0ee0bafb859fa4964b4404c27de1d380043c4dda20"
+
+[[package]]
+name = "unicode-ident"
+version = "1.0.24"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e6e4313cd5fcd3dad5cafa179702e2b244f760991f45397d14d4ebf38247da75"
+
+[[package]]
+name = "version_check"
+version = "0.9.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a"
+
+[[package]]
+name = "warp_reduce_redux"
+version = "0.1.0"
+dependencies = [
+ "cuda-core",
+ "cuda-device",
+ "cuda-host",
+]
+
+[[package]]
+name = "which"
+version = "4.4.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "87ba24419a2078cd2b0f2ede2691b6c66d8e47836da3b6db8265ebad47afbfc7"
+dependencies = [
+ "either",
+ "home",
+ "once_cell",
+ "rustix",
+]
+
+[[package]]
+name = "windows-link"
+version = "0.2.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f0805222e57f7521d6a62e36fa9163bc891acd422f971defe97d64e70d0a4fe5"
+
+[[package]]
+name = "windows-sys"
+version = "0.59.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1e38bc4d79ed67fd075bcc251a1c39b32a1776bbe92e5bef1f0bf1f8c531853b"
+dependencies = [
+ "windows-targets",
+]
+
+[[package]]
+name = "windows-sys"
+version = "0.61.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ae137229bcbd6cdf0f7b80a31df61766145077ddf49416a728b02cb3921ff3fc"
+dependencies = [
+ "windows-link",
+]
+
+[[package]]
+name = "windows-targets"
+version = "0.52.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9b724f72796e036ab90c1021d4780d4d3d648aca59e491e6b98e725b84e99973"
+dependencies = [
+ "windows_aarch64_gnullvm",
+ "windows_aarch64_msvc",
+ "windows_i686_gnu",
+ "windows_i686_gnullvm",
+ "windows_i686_msvc",
+ "windows_x86_64_gnu",
+ "windows_x86_64_gnullvm",
+ "windows_x86_64_msvc",
+]
+
+[[package]]
+name = "windows_aarch64_gnullvm"
+version = "0.52.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "32a4622180e7a0ec044bb555404c800bc9fd9ec262ec147edd5989ccd0c02cd3"
+
+[[package]]
+name = "windows_aarch64_msvc"
+version = "0.52.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "09ec2a7bb152e2252b53fa7803150007879548bc709c039df7627cabbd05d469"
+
+[[package]]
+name = "windows_i686_gnu"
+version = "0.52.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8e9b5ad5ab802e97eb8e295ac6720e509ee4c243f69d781394014ebfe8bbfa0b"
+
+[[package]]
+name = "windows_i686_gnullvm"
+version = "0.52.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0eee52d38c090b3caa76c563b86c3a4bd71ef1a819287c19d586d7334ae8ed66"
+
+[[package]]
+name = "windows_i686_msvc"
+version = "0.52.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "240948bc05c5e7c6dabba28bf89d89ffce3e303022809e73deaefe4f6ec56c66"
+
+[[package]]
+name = "windows_x86_64_gnu"
+version = "0.52.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "147a5c80aabfbf0c7d901cb5895d1de30ef2907eb21fbbab29ca94c5b08b1a78"
+
+[[package]]
+name = "windows_x86_64_gnullvm"
+version = "0.52.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "24d5b23dc417412679681396f2b49f3de8c1473deb516bd34410872eff51ed0d"
+
+[[package]]
+name = "windows_x86_64_msvc"
+version = "0.52.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "589f6da84c646204747d1270a2a5661ea66ed1cced2631d546fdfb155959f9ec"
+
+[[package]]
+name = "zerocopy"
+version = "0.8.52"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ce1022995ff5ff5d841ad7d994facc23098cd40152f2c1d11cd607c6f530653f"
+dependencies = [
+ "zerocopy-derive",
+]
+
+[[package]]
+name = "zerocopy-derive"
+version = "0.8.52"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1ae7f38b72ec2a254e2b87ef277cf2cd4fb97cbebf944faa6f33354da0867930"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn",
+]
+
+[[package]]
+name = "zmij"
+version = "1.0.23"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "29666d0abbfad1e3dc4dcf6144730dd3a3ab225bbbdac83319345b1b44ccfc1b"
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/warp_reduce_redux/Cargo.toml b/cuda-oxide/crates/rustc-codegen-cuda/examples/warp_reduce_redux/Cargo.toml
new file mode 100644
index 0000000000..9cabf31879
--- /dev/null
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/warp_reduce_redux/Cargo.toml
@@ -0,0 +1,26 @@
+[package]
+name = "warp_reduce_redux"
+version = "0.1.0"
+edition = "2024"
+license = "Apache-2.0"
+
+# Mark as standalone crate (not part of parent workspace)
+[workspace]
+
+# This example shows `warp_reduce` taking the single-instruction
+# `redux.sync` form at the Ampere floor and the shuffle butterfly below it.
+# The form is chosen by the `cuda_oxide_sm_at_least` cfg, so the same source
+# builds both ways:
+#   cargo oxide build warp_reduce_redux --arch sm_80   # redux.sync
+#   cargo oxide build warp_reduce_redux --arch sm_75   # shfl butterfly
+#   cargo oxide run   warp_reduce_redux
+
+[dependencies]
+# Device-side cooperative groups (warp_reduce, tiled_partition)
+cuda-device = { path = "../../../cuda-device" }
+
+# Host-side utilities (CudaKernel trait, etc.)
+cuda-host = { path = "../../../cuda-host" }
+
+# Host-side CUDA runtime
+cuda-core = { path = "../../../../../cuda-core" }
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/warp_reduce_redux/src/main.rs b/cuda-oxide/crates/rustc-codegen-cuda/examples/warp_reduce_redux/src/main.rs
new file mode 100644
index 0000000000..526345c353
--- /dev/null
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/warp_reduce_redux/src/main.rs
@@ -0,0 +1,284 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+//! `warp_reduce` on the Ampere single-instruction form, and its controls.
+//!
+//! `warp_reduce` reduces a full warp with a `shfl`-based butterfly: five
+//! rounds of shuffle-and-combine, ten instructions. Ampere does the integer
+//! cases in one (`redux.sync`), which #811 measured at 4.74x for the
+//! primitive. The form is not chosen at run time -- there is no branch in the
+//! emitted code -- it is chosen by the `cuda_oxide_sm_at_least` cfg that
+//! `cuda-device`'s build script derives from the target, so the same source
+//! builds both ways:
+//!
+//! ```text
+//!   cargo oxide build warp_reduce_redux --arch sm_80    redux.sync, no shfl
+//!   cargo oxide build warp_reduce_redux --arch sm_75    shfl butterfly
+//! ```
+//!
+//! Three things have to stay true, and each has a kernel below:
+//!
+//!   * the nine integer pairs must take the fast form and still be correct;
+//!   * `f32` must not -- there is no `redux` for floats before sm_100, so a
+//!     hook that claimed one would name an instruction that does not exist;
+//!   * a sub-warp tile must not -- `redux.sync` reduces over the lanes in
+//!     `membermask`, and a `WarpTile<16>`'s contract is not that shape yet.
+//!
+//! Run with:
+//!   cargo oxide run warp_reduce_redux
+
+use cuda_device::cooperative_groups::{
+    ThreadGroup,
+    ops::{BitAnd, BitOr, BitXor, Max, Min, Sum},
+    this_thread_block, warp_reduce,
+};
+use cuda_device::{DisjointSlice, kernel, thread};
+use cuda_host::cuda_module;
+
+// =============================================================================
+// KERNELS
+// =============================================================================
+#[cuda_module]
+mod kernels {
+    use super::*;
+
+    /// Every `u32` op `warp_reduce` supports, over one full warp.
+    ///
+    /// Lane `l` contributes `l` for sum/min/max, and a single distinct bit
+    /// `1 << l` for the bitwise ops, so each answer pins the whole warp:
+    /// an `or` or `xor` of all 32 single-bit values is `0xFFFFFFFF`, and an
+    /// `and` of their complements is `0` -- every bit is cleared by exactly
+    /// one lane. A reduction that silently dropped lanes could not produce
+    /// those.
+    ///
+    /// Min and max run again over `l << 27`, which sets bit 31 on lanes 16..32.
+    /// Unsigned, the answers are `0` and `0xF800_0000`; a signed reduction
+    /// would read those lanes as negative and answer `0x8000_0000` and
+    /// `0x7800_0000`.
+    ///
+    /// Lane 0 writes `[sum, min, max, and, or, xor, min_high, max_high]`.
+    #[kernel]
+    pub fn reduce_u32_all_ops(mut out: DisjointSlice<u32>) {
+        let warp = this_thread_block().tiled_partition::<32>();
+        let lane = warp.thread_rank();
+
+        let v_one = 1u32 << lane;
+        let v_inv = !v_one;
+        let v_high = lane << 27;
+
+        let r_sum = warp_reduce::<u32, Sum, _>(&warp, lane);
+        let r_min = warp_reduce::<u32, Min, _>(&warp, lane);
+        let r_max = warp_reduce::<u32, Max, _>(&warp, lane);
+        let r_and = warp_reduce::<u32, BitAnd, _>(&warp, v_inv);
+        let r_or = warp_reduce::<u32, BitOr, _>(&warp, v_one);
+        let r_xor = warp_reduce::<u32, BitXor, _>(&warp, v_one);
+        let r_min_high = warp_reduce::<u32, Min, _>(&warp, v_high);
+        let r_max_high = warp_reduce::<u32, Max, _>(&warp, v_high);
+
+        if lane == 0 {
+            // SAFETY: the host allocates eight elements for this kernel.
+            unsafe {
+                *out.get_unchecked_mut(0) = r_sum;
+                *out.get_unchecked_mut(1) = r_min;
+                *out.get_unchecked_mut(2) = r_max;
+                *out.get_unchecked_mut(3) = r_and;
+                *out.get_unchecked_mut(4) = r_or;
+                *out.get_unchecked_mut(5) = r_xor;
+                *out.get_unchecked_mut(6) = r_min_high;
+                *out.get_unchecked_mut(7) = r_max_high;
+            }
+        }
+    }
+
+    /// The signed ops, with values that straddle zero (`-16..=15`).
+    ///
+    /// Signedness is the part a single-instruction form can get wrong
+    /// silently: `redux.sync.min.s32` and `.u32` disagree on exactly these
+    /// inputs, and only the signed answer is `-16`.
+    ///
+    /// Lane 0 writes `[sum, min, max]`.
+    #[kernel]
+    pub fn reduce_i32_signed_ops(mut out: DisjointSlice<i32>) {
+        let warp = this_thread_block().tiled_partition::<32>();
+        let lane = warp.thread_rank();
+        let v = lane as i32 - 16;
+
+        let r_sum = warp_reduce::<i32, Sum, _>(&warp, v);
+        let r_min = warp_reduce::<i32, Min, _>(&warp, v);
+        let r_max = warp_reduce::<i32, Max, _>(&warp, v);
+
+        if lane == 0 {
+            // SAFETY: the host allocates three elements for this kernel.
+            unsafe {
+                *out.get_unchecked_mut(0) = r_sum;
+                *out.get_unchecked_mut(1) = r_min;
+                *out.get_unchecked_mut(2) = r_max;
+            }
+        }
+    }
+
+    /// Control: `f32` keeps the butterfly at every floor.
+    ///
+    /// There is no `redux` for floats before sm_100, so this must reduce the
+    /// same way at sm_75 and sm_80 and stay correct.
+    ///
+    /// Lane 0 writes `[sum, min, max]`.
+    #[kernel]
+    pub fn reduce_f32_control(mut out: DisjointSlice<f32>) {
+        let warp = this_thread_block().tiled_partition::<32>();
+        let lane = warp.thread_rank();
+        let v = lane as f32;
+
+        let r_sum = warp_reduce::<f32, Sum, _>(&warp, v);
+        let r_min = warp_reduce::<f32, Min, _>(&warp, v);
+        let r_max = warp_reduce::<f32, Max, _>(&warp, v);
+
+        if lane == 0 {
+            // SAFETY: the host allocates three elements for this kernel.
+            unsafe {
+                *out.get_unchecked_mut(0) = r_sum;
+                *out.get_unchecked_mut(1) = r_min;
+                *out.get_unchecked_mut(2) = r_max;
+            }
+        }
+    }
+
+    /// Control: a `WarpTile<16>` keeps the butterfly, and each half reduces
+    /// only its own lanes.
+    ///
+    /// Lane `l` contributes its rank within the tile, so both tiles answer
+    /// `0 + 1 + ... + 15 = 120`. A fast path that reduced the whole warp
+    /// instead would answer `496` and be caught here.
+    ///
+    /// Tile leader `t` writes `out[t]`.
+    #[kernel]
+    pub fn reduce_subwarp_control(mut out: DisjointSlice<u32>) {
+        let tile = this_thread_block().tiled_partition::<16>();
+        let rank = tile.thread_rank();
+        let tile_index = thread::index_1d().get() / 16;
+
+        let r_sum = warp_reduce::<u32, Sum, _>(&tile, rank);
+
+        if rank == 0 {
+            // SAFETY: the host allocates one element per tile (two).
+            unsafe {
+                *out.get_unchecked_mut(tile_index as usize) = r_sum;
+            }
+        }
+    }
+}
+
+// =============================================================================
+// HOST CODE
+// =============================================================================
+
+fn main() {
+    use cuda_core::simt::LaunchConfig;
+    use cuda_core::{CudaContext, DeviceBuffer};
+
+    println!("=== warp_reduce: single-instruction form and its controls ===\n");
+
+    let ctx = CudaContext::new(0).expect("Failed to create CUDA context");
+    let stream = ctx.default_stream();
+
+    let (major, minor) = ctx.compute_capability().expect("compute capability");
+    println!("GPU Compute Capability: sm_{}{}", major, minor);
+
+    let module = kernels::load(&ctx).expect("Failed to load embedded CUDA module");
+    // One warp is the whole contract; a second would only repeat it.
+    let cfg = LaunchConfig {
+        block_dim: (32, 1, 1),
+        grid_dim: (1, 1, 1),
+        shared_mem_bytes: 0,
+    };
+
+    let mut failed = false;
+
+    // ===== Test 1: every u32 op =====
+    println!("\n--- Test 1: warp_reduce over u32, all six ops ---");
+    let mut u32_dev = DeviceBuffer::<u32>::zeroed(&stream, 8).unwrap();
+    // SAFETY: launch shape/resources match the kernel; the buffer covers its writes.
+    unsafe { module.reduce_u32_all_ops((stream).as_ref(), cfg, &mut u32_dev) }
+        .expect("Kernel launch failed");
+    let got = u32_dev.to_host_vec(&stream).unwrap();
+    let want: [u32; 8] = [496, 0, 31, 0, u32::MAX, u32::MAX, 0, 0xF800_0000];
+    println!(
+        "[sum, min, max, and, or, xor, min_high, max_high] = {:?}",
+        got
+    );
+    println!(
+        "expected                                          = {:?}",
+        want
+    );
+    println!("(a signed min/max would read 2147483648 / 2013265920 for the last two)");
+    if got == want {
+        println!("✓ every u32 reduction correct");
+    } else {
+        println!("✗ u32 reduction mismatch!");
+        failed = true;
+    }
+
+    // ===== Test 2: signed ops across zero =====
+    println!("\n--- Test 2: warp_reduce over i32, values -16..=15 ---");
+    let mut i32_dev = DeviceBuffer::<i32>::zeroed(&stream, 3).unwrap();
+    // SAFETY: launch shape/resources match the kernel; the buffer covers its writes.
+    unsafe { module.reduce_i32_signed_ops((stream).as_ref(), cfg, &mut i32_dev) }
+        .expect("Kernel launch failed");
+    let got = i32_dev.to_host_vec(&stream).unwrap();
+    let want: [i32; 3] = [-16, -16, 15];
+    println!("[sum, min, max] = {:?}", got);
+    println!(
+        "expected        = {:?}   (unsigned min/max would read 0 / 4294967295)",
+        want
+    );
+    if got == want {
+        println!("✓ signed reductions correct");
+    } else {
+        println!("✗ signed reduction mismatch!");
+        failed = true;
+    }
+
+    // ===== Test 3: f32 control =====
+    println!("\n--- Test 3: control, f32 keeps the butterfly ---");
+    let mut f32_dev = DeviceBuffer::<f32>::zeroed(&stream, 3).unwrap();
+    // SAFETY: launch shape/resources match the kernel; the buffer covers its writes.
+    unsafe { module.reduce_f32_control((stream).as_ref(), cfg, &mut f32_dev) }
+        .expect("Kernel launch failed");
+    let got = f32_dev.to_host_vec(&stream).unwrap();
+    let want: [f32; 3] = [496.0, 0.0, 31.0];
+    println!("[sum, min, max] = {:?}", got);
+    println!("expected        = {:?}", want);
+    if got == want {
+        println!("✓ f32 reductions correct");
+    } else {
+        println!("✗ f32 reduction mismatch!");
+        failed = true;
+    }
+
+    // ===== Test 4: sub-warp control =====
+    println!("\n--- Test 4: control, WarpTile<16> reduces only its own lanes ---");
+    let mut tile_dev = DeviceBuffer::<u32>::zeroed(&stream, 2).unwrap();
+    // SAFETY: launch shape/resources match the kernel; the buffer covers its writes.
+    unsafe { module.reduce_subwarp_control((stream).as_ref(), cfg, &mut tile_dev) }
+        .expect("Kernel launch failed");
+    let got = tile_dev.to_host_vec(&stream).unwrap();
+    let want: [u32; 2] = [120, 120];
+    println!("per-tile sums = {:?}", got);
+    println!(
+        "expected      = {:?}   (a whole-warp reduction would read 496)",
+        want
+    );
+    if got == want {
+        println!("✓ sub-warp tiles reduce independently");
+    } else {
+        println!("✗ sub-warp reduction mismatch!");
+        failed = true;
+    }
+
+    if failed {
+        std::process::exit(1);
+    }
+    println!("\nSUCCESS: warp_reduce is correct in both forms");
+}
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/warp_reduce_redux/verify-code-shape.sh b/cuda-oxide/crates/rustc-codegen-cuda/examples/warp_reduce_redux/verify-code-shape.sh
new file mode 100755
index 0000000000..41993a0cb5
--- /dev/null
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/warp_reduce_redux/verify-code-shape.sh
@@ -0,0 +1,80 @@
+#!/usr/bin/env bash
+# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: Apache-2.0
+
+# The runtime PASS gate only proves the reductions are correct, which the
+# shuffle butterfly also achieves. Which form was emitted lives in the PTX:
+# at a floor of sm_80 or newer each integer reduction must be one `redux.sync`
+# and the butterfly must be gone; the f32 and sub-warp controls must never use
+# `redux`; and below sm_80 no `redux` may appear at all.
+
+set -euo pipefail
+
+root="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
+ptx="${root}/warp_reduce_redux.ptx"
+
+test -s "${ptx}"
+
+target="$(grep -m1 -oE '^\.target sm_[0-9]+' "${ptx}" | grep -oE '[0-9]+$')" || true
+if [[ -z "${target}" ]]; then
+    echo "error: no .target sm_NN line in ${ptx}" >&2
+    exit 1
+fi
+
+# Print the body of one PTX entry. Waits for the header's `{` so a forward
+# declaration (which ends in `;`) is skipped.
+entry_body() {
+    local symbol="$1"
+    awk -v marker="${symbol}(" '
+        !emit && !candidate && index($0, marker) && index($0, ".entry") {
+            candidate = 1
+        }
+        candidate && $0 ~ /^[[:space:]]*;[[:space:]]*$/ {
+            candidate = 0
+            next
+        }
+        candidate && $0 ~ /^[[:space:]]*\{[[:space:]]*$/ {
+            emit = 1
+            candidate = 0
+        }
+        emit { print }
+        emit && index($0, "End function") != 0 { exit }
+    ' "${ptx}"
+}
+
+# count <kernel> <opcode-regex>: instructions in the kernel body whose opcode
+# matches, anchored at instruction position so comments never count.
+count() {
+    local symbol="$1" opcode="$2" body
+    body="$(entry_body "${symbol}")"
+    if [[ -z "${body}" ]]; then
+        echo "error: no .entry body found for ${symbol} in ${ptx}" >&2
+        exit 1
+    fi
+    grep -cE "^[[:space:]]*(@%p[0-9]+[[:space:]]+)?${opcode}" <<<"${body}" || true
+}
+
+# require <kernel> <opcode-regex> <expected-count>
+require() {
+    local symbol="$1" opcode="$2" want="$3" got
+    got="$(count "${symbol}" "${opcode}")"
+    if [[ "${got}" -ne "${want}" ]]; then
+        echo "error: sm_${target}: ${symbol} has ${got}x ${opcode}, expected ${want}x in ${ptx}" >&2
+        exit 1
+    fi
+}
+
+if [[ "${target}" -ge 80 ]]; then
+    # Six u32 ops plus the unsigned min/max over values above 2^31.
+    require reduce_u32_all_ops 'redux\.sync\.' 8
+    require reduce_i32_signed_ops 'redux\.sync\.' 3
+    require reduce_u32_all_ops 'shfl\.sync\.' 0
+    require reduce_i32_signed_ops 'shfl\.sync\.' 0
+else
+    require reduce_u32_all_ops 'redux\.' 0
+    require reduce_i32_signed_ops 'redux\.' 0
+fi
+require reduce_f32_control 'redux\.' 0
+require reduce_subwarp_control 'redux\.' 0
+
+echo "warp_reduce_redux code shape OK at sm_${target}"
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/warp_sums/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/warp_sums/Cargo.lock
index 1b446984c3..2b1fe2af5a 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/warp_sums/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/warp_sums/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/wgmma/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/wgmma/Cargo.lock
index 11fbfa1d79..c0668a6f1e 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/wgmma/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/wgmma/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/wgmma_mma_bf16/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/wgmma_mma_bf16/Cargo.lock
index e5e504178f..33b4dee4c1 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/wgmma_mma_bf16/Cargo.lock
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/wgmma_mma_bf16/Cargo.lock
@@ -163,6 +163,7 @@ name = "cuda-device"
 version = "0.2.1"
 dependencies = [
  "cuda-macros",
+ "cuda-target-spec",
 ]
 
 [[package]]
diff --git a/cuda-oxide/cuda-oxide-book/advanced/warp-level-programming.md b/cuda-oxide/cuda-oxide-book/advanced/warp-level-programming.md
index 4df7a75d2a..254b82d7f1 100644
--- a/cuda-oxide/cuda-oxide-book/advanced/warp-level-programming.md
+++ b/cuda-oxide/cuda-oxide-book/advanced/warp-level-programming.md
@@ -161,10 +161,20 @@ Three constraints:
   example does, on `ctx.compute_capability()`, or restrict the build to Ampere
   and newer.
 
-`warp_reduce` does not select this form for you today; making it do so
-automatically needs a way for device code to know its target architecture,
-which is the open question in
-[#811](https://github.com/NVlabs/cuda-oxide/issues/811).
+`warp_reduce` selects this form for you, for integer element types on a full
+32-lane tile, whenever the build's compute-capability floor is `sm_80` or
+newer. That floor comes from the `cuda_oxide_sm_at_least` cfg, which
+`cuda-device`'s build script derives with the same rule the backend uses to
+choose a target. An explicit target (`--arch` or `CUDA_OXIDE_TARGET`) sets the
+floor. Without one it is `sm_80`, or the detected GPU
+(`CUDA_OXIDE_DEVICE_ARCH`) if that is older: a newer GPU never raises it,
+because the backend may still build for an older architecture the kernel
+needs. A cfg is resolved before rustc runs, so a build below the floor drops
+the branch entirely and never emits an instruction it cannot assemble.
+
+Calling `redux_sync_add` and friends directly is still the right move when you
+want the single instruction regardless of the surrounding generic code, and it
+is the only option for a sub-warp tile.
 
 ### Floats on Blackwell: `redux.sync` for `f32`
 
diff --git a/cuda-oxide/cuda-oxide-book/compiler/lowering-pipeline.md b/cuda-oxide/cuda-oxide-book/compiler/lowering-pipeline.md
index e227a9e59c..41e70dc026 100644
--- a/cuda-oxide/cuda-oxide-book/compiler/lowering-pipeline.md
+++ b/cuda-oxide/cuda-oxide-book/compiler/lowering-pipeline.md
@@ -612,7 +612,10 @@ to the backend so the generated module is guaranteed to load on the local GPU.
 The full precedence is `--arch` > `CUDA_OXIDE_TARGET` > host CC (for `run`
 only) > backend feature-based default. `cargo oxide build` and
 `cargo oxide pipeline` deliberately skip the host-CC step so they remain
-usable for cross-compilation.
+usable for cross-compilation. Without an explicit target, a feature-based
+choice below the device-code floor -- `sm_80`, or the detected GPU if that is
+older -- is raised to it, so the module never targets less than
+`cuda-device`'s `cuda_oxide_sm_at_least` cfg assumed.
 
 If the lowered module calls CUDA libdevice, cuda-oxide switches from `llc` to
 the NVVM path automatically. It resolves the target with the same precedence
diff --git a/cuda-oxide/cuda-oxide-book/compiler/rustc-codegen-cuda.md b/cuda-oxide/cuda-oxide-book/compiler/rustc-codegen-cuda.md
index 683605591e..d3ad8e65ea 100644
--- a/cuda-oxide/cuda-oxide-book/compiler/rustc-codegen-cuda.md
+++ b/cuda-oxide/cuda-oxide-book/compiler/rustc-codegen-cuda.md
@@ -411,7 +411,10 @@ without fighting the compiler driver's flag plumbing.
 `CUDA_OXIDE_TARGET` sits in the middle of a precedence chain that
 `cargo oxide` honours: explicit `--arch` wins, then `CUDA_OXIDE_TARGET`,
 then `cargo oxide run`'s host-CC auto-detect (only for `run`, not `build`
-or `pipeline`), and finally the backend's feature-based default.
+or `pipeline`), and finally the backend's feature-based default. That default
+is raised to the device-code floor -- `sm_80`, or the detected GPU if that is
+older -- so the target never falls below what `cuda-device`'s
+`cuda_oxide_sm_at_least` cfg assumed.
 
 ```{note}
 `CUDA_OXIDE_VERBOSE=1 cargo oxide build` is your best friend when debugging
diff --git a/cuda-oxide/cuda-oxide-book/gpu-programming/kernels-and-device-functions.md b/cuda-oxide/cuda-oxide-book/gpu-programming/kernels-and-device-functions.md
index 303334a50e..c66e1ef54e 100644
--- a/cuda-oxide/cuda-oxide-book/gpu-programming/kernels-and-device-functions.md
+++ b/cuda-oxide/cuda-oxide-book/gpu-programming/kernels-and-device-functions.md
@@ -262,7 +262,10 @@ target-derived `cfg`, so a kernel has no way to ask which architecture it is
 being compiled for. Arch requirements live in doc comments and are enforced by
 the caller, which is why an example that needs `redux.sync` checks
 `ctx.compute_capability()` on the host and skips rather than specializing the
-kernel.
+kernel. `cuda-device` does derive a `cuda_oxide_sm_at_least` cfg for its own
+`warp_reduce`, but a build script's `cfg` reaches only its own crate: in a
+kernel crate the name is undeclared, so `#[cfg(cuda_oxide_sm_at_least = "80")]`
+is always false and warns as `unexpected_cfgs`.
 
 `--device-cfg NAME` is how you supply one yourself. It is repeatable, and each
 occurrence becomes a `--cfg NAME` in the build's rustflags: