diff --git a/cuda-oxide/crates/cuda-device/src/cooperative_groups.rs b/cuda-oxide/crates/cuda-device/src/cooperative_groups.rs index c6c1de84a3..d59a9148d0 100644 --- a/cuda-oxide/crates/cuda-device/src/cooperative_groups.rs +++ b/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 = >::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 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 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 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 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 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 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 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, { + // 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", >::HAS_FULL_WARP_FORM), + ("u32 Min", >::HAS_FULL_WARP_FORM), + ("u32 Max", >::HAS_FULL_WARP_FORM), + ("u32 BitAnd", >::HAS_FULL_WARP_FORM), + ("u32 BitOr", >::HAS_FULL_WARP_FORM), + ("u32 BitXor", >::HAS_FULL_WARP_FORM), + ("i32 Sum", >::HAS_FULL_WARP_FORM), + ("i32 Min", >::HAS_FULL_WARP_FORM), + ("i32 Max", >::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!(!>::HAS_FULL_WARP_FORM); + assert!(!>::HAS_FULL_WARP_FORM); + assert!(!>::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!(>::reduce_full_warp(u32::MAX, 3.5), 3.5); + } } diff --git a/cuda-oxide/crates/cuda-device/src/warp.rs b/cuda-oxide/crates/cuda-device/src/warp.rs index f4043e128c..08b11ba89a 100644 --- a/cuda-oxide/crates/cuda-device/src/warp.rs +++ b/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+) // ============================================================================= 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 { + 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/lib.rs b/cuda-oxide/crates/cuda-oxide-codegen/src/lib.rs index 3d7c596273..dca437110a 100644 --- a/cuda-oxide/crates/cuda-oxide-codegen/src/lib.rs +++ b/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 = 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 d065224532..7f1c9ebc74 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 { + 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")) } 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 f3536259c7..ade95b8944 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 @@ -540,3 +540,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::().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::().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::().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::().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::().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::().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::().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 a5467be911..806461028a 100644 --- a/cuda-oxide/crates/cuda-target-spec/src/lib.rs +++ b/cuda-oxide/crates/cuda-target-spec/src/lib.rs @@ -566,9 +566,146 @@ pub fn spelling_at_least(floor: u16) -> Option { 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 { + if let Some(target) = target.map(str::trim).filter(|target| !target.is_empty()) { + return Ok(target.parse::()?.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 { + 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::().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![70, 75, 80, 86] + ); + assert_eq!( + sm_floors_at_most(120).collect::>(), + SM_FLOOR_LADDER.to_vec() + ); + assert_eq!(sm_floors_at_most(75).collect::>(), 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![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/mir-importer/src/pipeline.rs b/cuda-oxide/crates/mir-importer/src/pipeline.rs index be77e619b5..14a3c8251b 100644 --- a/cuda-oxide/crates/mir-importer/src/pipeline.rs +++ b/cuda-oxide/crates/mir-importer/src/pipeline.rs @@ -534,6 +534,10 @@ pub fn run_pipeline( let mut legaliser = Legaliser::default(); let mut kernel_launch_bounds = BTreeMap::new(); let mut expected_kernels = Vec::new(); + let has_integer_redux_sync = cuda_oxide_codegen::__private::target_has_integer_redux_sync( + config.target_arch.as_deref(), + config.device_arch_hint.as_ref(), + ); // Step 3: Translate all functions for func in functions { @@ -553,13 +557,18 @@ pub fn run_pipeline( .instance .body() .ok_or_else(|| PipelineError::NoBody(func.export_name.clone()))?; + let mono_successors = crate::translator::body::narrow_target_predicate_switches( + &body, + &func.rustc_mono_successors, + has_integer_redux_sync, + ); let func_op_ptr = crate::translator::body::translate_body( &mut ctx, &body, &func.instance, func.rustc_mir_block_count, - &func.rustc_mono_successors, + &mono_successors, func.is_kernel, func.inline, Some(&func.export_name), diff --git a/cuda-oxide/crates/mir-importer/src/translator/body.rs b/cuda-oxide/crates/mir-importer/src/translator/body.rs index 7149aac141..c5b97ef8a5 100644 --- a/cuda-oxide/crates/mir-importer/src/translator/body.rs +++ b/cuda-oxide/crates/mir-importer/src/translator/body.rs @@ -604,6 +604,81 @@ fn validate_monomorphized_successors( Ok(()) } +/// Narrow `SwitchInt` terminators controlled by `__target_has_integer_redux_sync` +/// to the single arm selected by the resolved target floor (`sm_80+` vs below). +pub fn narrow_target_predicate_switches( + body: &mir::Body, + mono_successors: &[Vec], + has_integer_redux_sync: bool, +) -> Vec> { + use rustc_public::mir::{Operand, StatementKind, TerminatorKind}; + use rustc_public::ty::{RigidTy, TyKind}; + + let mut narrowed = mono_successors.to_vec(); + for block in &body.blocks { + let TerminatorKind::Call { + func, + destination, + target: Some(next_bb), + .. + } = &block.terminator.kind + else { + continue; + }; + if !destination.projection.is_empty() || *next_bb >= body.blocks.len() { + continue; + } + let is_redux_predicate = func.ty(body.locals()).ok().is_some_and(|ty| { + matches!( + ty.kind(), + TyKind::RigidTy(RigidTy::FnDef(def, _)) + if def.name() == "cuda_device::warp::__target_has_integer_redux_sync" + || def.name().ends_with("::__target_has_integer_redux_sync") + ) + }); + if !is_redux_predicate { + continue; + } + let next_block = &body.blocks[*next_bb]; + let overwritten = next_block.statements.iter().any(|stmt| { + matches!( + &stmt.kind, + StatementKind::Assign(place, _) if place.local == destination.local + ) + }); + if overwritten { + continue; + } + let TerminatorKind::SwitchInt { + discr: Operand::Move(place) | Operand::Copy(place), + targets, + } = &next_block.terminator.kind + else { + continue; + }; + if !place.projection.is_empty() || place.local != destination.local { + continue; + } + let chosen = if has_integer_redux_sync { + targets + .branches() + .find(|(val, _)| *val == 1) + .map_or_else(|| targets.otherwise(), |(_, target)| target) + } else { + targets + .branches() + .find(|(val, _)| *val == 0) + .map_or_else(|| targets.otherwise(), |(_, target)| target) + }; + if let Some(edges) = narrowed.get_mut(*next_bb) + && edges.contains(&chosen) + { + *edges = vec![chosen]; + } + } + narrowed +} + /// BFS from the entry block following rustc's exact per-block monomorphized /// successors, intersected with the importer's existing non-unwind policy. /// diff --git a/cuda-oxide/crates/mir-importer/src/translator/terminator/mod.rs b/cuda-oxide/crates/mir-importer/src/translator/terminator/mod.rs index 1e39ce7733..fc74e6b48b 100644 --- a/cuda-oxide/crates/mir-importer/src/translator/terminator/mod.rs +++ b/cuda-oxide/crates/mir-importer/src/translator/terminator/mod.rs @@ -3021,7 +3021,8 @@ fn try_dispatch_intrinsic( // integer `checked_pow`/`overflowing_pow` and `checked_ilog`, so // rejecting it also rejects integer `pow` and `ilog`, which call them. "core::intrinsics::is_val_statically_known" - | "std::intrinsics::is_val_statically_known" => { + | "std::intrinsics::is_val_statically_known" + | "cuda_device::warp::__target_has_integer_redux_sync" => { Ok(Some(helpers::emit_bool_constant_intrinsic( ctx, body, 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, @@ -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) -> 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, 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, 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, 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 { 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, 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/warp_reduce_redux/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/warp_reduce_redux/Cargo.lock new file mode 100644 index 0000000000..cd716f616b --- /dev/null +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/warp_reduce_redux/Cargo.lock @@ -0,0 +1,630 @@ +# 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.72.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "993776b509cfb49c750f11b8f07a46fa23e0a1386ffc01fb1e7d343efc387895" +dependencies = [ + "bitflags", + "cexpr", + "clang-sys", + "itertools", + "log", + "prettyplease", + "proc-macro2", + "quote", + "regex", + "rustc-hash", + "shlex", + "syn", +] + +[[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 = "bytemuck" +version = "1.25.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c8efb64bd706a16a1bdde310ae86b351e4d21550d98d056f22f8a7f7a2183fec" +dependencies = [ + "bytemuck_derive", +] + +[[package]] +name = "bytemuck_derive" +version = "1.10.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f9abbd1bc6865053c427f7198e6af43bfdedc55ab791faed4fbd361d789575ff" +dependencies = [ + "proc-macro2", + "quote", + "syn", +] + +[[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.3.0" +dependencies = [ + "libnvvm-sys", + "nvjitlink-sys", + "serde", + "sha2", + "thiserror", +] + +[[package]] +name = "cuda-bindings" +version = "0.5.0" +dependencies = [ + "bindgen", + "libloading", + "prettyplease", + "proc-macro2", + "quote", + "syn", +] + +[[package]] +name = "cuda-core" +version = "0.5.0" +dependencies = [ + "anyhow", + "bytemuck", + "cuda-bindings", + "cuda-core-derive", + "cuda-repr", + "half", + "oxide-artifacts", +] + +[[package]] +name = "cuda-core-derive" +version = "0.5.0" +dependencies = [ + "proc-macro2", + "quote", + "syn", +] + +[[package]] +name = "cuda-device" +version = "0.3.0" +dependencies = [ + "bytemuck", + "cuda-macros", + "cuda-repr", +] + +[[package]] +name = "cuda-host" +version = "0.3.0" +dependencies = [ + "cuda-artifact-finalizer", + "cuda-core", + "cuda-device", + "cuda-macros", + "half", + "libc", + "oxide-artifacts", + "ptx-parse", + "sha2", + "thiserror", +] + +[[package]] +name = "cuda-macros" +version = "0.3.0" +dependencies = [ + "oxide-tile-export", + "proc-macro2", + "quote", + "reserved-oxide-symbols", + "syn", +] + +[[package]] +name = "cuda-repr" +version = "0.5.0" + +[[package]] +name = "cuda-target-spec" +version = "0.3.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 = "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 = [ + "bytemuck", + "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 = "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 = "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.3.0" +dependencies = [ + "cuda-target-spec", + "libloading", + "thiserror", +] + +[[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.3.0" +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 = "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.3.0" +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.3.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.3.0" + +[[package]] +name = "rustc-hash" +version = "2.1.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6b1e7f9a428571be2dc5bc0505c13fb6bf936822b894ec87abf8a08a4e51742d" + +[[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 = "windows-link" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f0805222e57f7521d6a62e36fa9163bc891acd422f971defe97d64e70d0a4fe5" + +[[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..e32f0c9d88 --- /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) { + 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::(&warp, lane); + let r_min = warp_reduce::(&warp, lane); + let r_max = warp_reduce::(&warp, lane); + let r_and = warp_reduce::(&warp, v_inv); + let r_or = warp_reduce::(&warp, v_one); + let r_xor = warp_reduce::(&warp, v_one); + let r_min_high = warp_reduce::(&warp, v_high); + let r_max_high = warp_reduce::(&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) { + let warp = this_thread_block().tiled_partition::<32>(); + let lane = warp.thread_rank(); + let v = lane as i32 - 16; + + let r_sum = warp_reduce::(&warp, v); + let r_min = warp_reduce::(&warp, v); + let r_max = warp_reduce::(&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) { + let warp = this_thread_block().tiled_partition::<32>(); + let lane = warp.thread_rank(); + let v = lane as f32; + + let r_sum = warp_reduce::(&warp, v); + let r_min = warp_reduce::(&warp, v); + let r_max = warp_reduce::(&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) { + 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::(&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::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::::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::::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::::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::::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 : 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 +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/cuda-oxide-book/advanced/warp-level-programming.md b/cuda-oxide/cuda-oxide-book/advanced/warp-level-programming.md index 4df7a75d2a..c21149fff2 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. The compiler resolves that floor 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. The branch is folded during MIR import before feature detection runs, +so a build below the floor drops the `redux.sync` arm 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 668013f72a..c47e9c12b0 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 +the floor `cuda-device` primitives were lowered for. 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..27be36ebc5 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 the floor `cuda-device` primitives +were lowered for. ```{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 994c02084f..3ee2fbb7aa 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 @@ -264,7 +264,8 @@ 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`'s own `warp_reduce` selects `redux.sync` vs shuffles +through an internal compiler-resolved target floor predicate). `--device-cfg NAME` is how you supply one yourself. It is repeatable, and each occurrence becomes a `--cfg NAME` in the build's rustflags: