feat(cuda-device): derive an SM floor cfg and use redux.sync in warp_reduce

cuda-oxide/crates/cuda-device/src/cooperative_groups.rs

@@ -909,6 +909,44 @@ 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 = true;
+
+            // Keep the intrinsic call as the return expression of the fast arm
+            // and fall back to the shuffle butterfly when the resolved target
+            // floor is below sm_80.
+            #[inline(always)]
+            fn reduce_full_warp($mask: u32, $value: $ty) -> $ty {
+                if crate::warp::__target_has_integer_redux_sync() {
+                    return $call;
+                }
+                let mut acc = $value;
+                let mut delta: u32 = 16;
+                while delta > 0 {
+                    let other = crate::warp::shuffle_xor_sync($mask, acc as u32, delta) as $ty;
+                    acc = <Self as ReduceOp<$ty>>::combine(acc, other);
+                    delta >>= 1;
+                }
+                acc
+            }
+        };
     }
 
     /// Sum: wrapping integer addition; IEEE-754 add for floats.
@@ -945,6 +983,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)]
@@ -955,6 +994,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)]
@@ -965,6 +1005,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)]
@@ -975,6 +1016,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)]
@@ -985,6 +1027,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)]
@@ -995,6 +1038,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 ---
@@ -1007,6 +1051,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)]
@@ -1017,6 +1062,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)]
@@ -1027,6 +1073,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 ---
@@ -1256,13 +1303,13 @@ impl WarpShuffle for i64 {
 /// 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.
@@ -1290,6 +1337,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 {
@@ -1748,6 +1800,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() {
@@ -1774,4 +1827,51 @@ 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.
+    #[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");
+        }
+    }
+
+    /// 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);
+    }
 }

cuda-oxide/crates/cuda-device/src/warp.rs

@@ -691,6 +691,14 @@ pub fn popc(predicate: bool) -> u32 {
     ballot(predicate).count_ones()
 }
 
+/// Compiler-resolved predicate returning whether the target floor supports
+/// integer `redux.sync` (`sm_80` or newer).
+#[doc(hidden)]
+#[inline(never)]
+pub fn __target_has_integer_redux_sync() -> bool {
+    unreachable!("__target_has_integer_redux_sync called outside CUDA kernel context")
+}
+
 // =============================================================================
 // Warp Match Operations (sm_70+)
 // =============================================================================

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);
     }
 

cuda-oxide/crates/cuda-oxide-codegen/src/lib.rs

@@ -136,6 +136,17 @@ pub mod __private {
     #[doc(hidden)]
     pub use mir_transforms::unroll::{HintLocation, HintWarning};
 
+    /// Whether the resolved target floor supports Ampere integer `redux.sync` (`sm_80+`).
+    #[doc(hidden)]
+    pub fn target_has_integer_redux_sync(
+        target_arch: Option<&str>,
+        device_arch_hint: Option<&DeviceArchHint>,
+    ) -> bool {
+        let device_arch = device_arch_hint.and_then(|hint| hint.as_device_arch().ok());
+        cuda_target_spec::resolve_sm_floor(target_arch, device_arch.map(AsRef::as_ref))
+            .is_ok_and(|floor| floor >= 80)
+    }
+
     /// Compiler-only attribute used to carry the Rust generated-intrinsic ABI marker.
     #[doc(hidden)]
     pub const GENERATED_INTRINSIC_MARKER_ATTR: &str =

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

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,
@@ -208,8 +235,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"))
 }