From 61c862c34887706f8479d66ed8cdebf947e92d2f Mon Sep 17 00:00:00 2001 From: Trent Nelson Date: Mon, 19 Oct 2026 23:45:41 -0700 Subject: [PATCH 1/3] Support cluster queries on targets below sm90 Treat each block as one cluster when a known earlier target is selected. Bind a device hint only after this fallback is used, so later feature selection cannot silently change its semantics. Keep raw intrinsic and cluster operation target requirements intact. Signed-off-by: Trent Nelson --- cuda-oxide/crates/cuda-device/src/cluster.rs | 35 +++- .../crates/mir-importer/src/pipeline.rs | 55 +++++- .../mir-importer/src/translator/body.rs | 10 +- .../terminator/intrinsics/cluster_query.rs | 187 ++++++++++++++++++ .../translator/terminator/intrinsics/mod.rs | 1 + .../src/translator/terminator/mod.rs | 16 ++ .../mir-importer/src/translator/values.rs | 6 + 7 files changed, 297 insertions(+), 13 deletions(-) create mode 100644 cuda-oxide/crates/mir-importer/src/translator/terminator/intrinsics/cluster_query.rs diff --git a/cuda-oxide/crates/cuda-device/src/cluster.rs b/cuda-oxide/crates/cuda-device/src/cluster.rs index 07d6c93076..c61d373184 100644 --- a/cuda-oxide/crates/cuda-device/src/cluster.rs +++ b/cuda-oxide/crates/cuda-device/src/cluster.rs @@ -4,7 +4,7 @@ */ #![allow(non_snake_case)] -//! Thread Block Cluster intrinsics for Hopper (sm_90+). +//! Thread block cluster queries and Hopper (sm_90+) cluster operations. //! //! Thread Block Clusters are a new hierarchy level introduced in NVIDIA Hopper //! that groups multiple thread blocks together. Blocks within a cluster can: @@ -25,9 +25,24 @@ //! //! # Hardware Requirements //! -//! - Minimum: sm_90 (Hopper: H100, H200) +//! - Native cluster operations require sm_90 (Hopper: H100, H200) //! - Blackwell (sm_100/sm_120) extends cluster capabilities //! +//! # Queries on earlier targets +//! +//! When compiling for a known target below sm90, the position and dimension +//! queries in this module treat each block as its own cluster. Block coordinates +//! and [`block_rank`] are zero; cluster dimensions and [`cluster_size`] are one. +//! [`cluster_idx`] returns the linear block index, and [`num_clusters`] returns +//! the number of blocks in the grid. +//! +//! An explicit target (for example, `cargo oxide build --arch sm_80`) takes +//! precedence over a detected device. If a pre-sm90 device hint selects these +//! semantics, the compiler binds that target and rejects incompatible operations. +//! Without a target or device hint, queries retain native sm90+ requirements. +//! Raw `cuda_intrinsics` cluster queries, synchronization, and distributed shared +//! memory operations always retain their hardware requirements. +//! //! # Cluster Dimensions //! //! Clusters are defined as a 3D arrangement of thread blocks: @@ -76,7 +91,7 @@ /// /// # PTX /// -/// Lowers to: `mov.u32 %r, %cluster_ctaid.x` +/// On sm90+: `mov.u32 %r, %cluster_ctaid.x`. On earlier targets: zero. #[inline(never)] pub fn cluster_ctaidX() -> u32 { unreachable!("cluster_ctaidX called outside CUDA kernel context") @@ -88,7 +103,7 @@ pub fn cluster_ctaidX() -> u32 { /// /// # PTX /// -/// Lowers to: `mov.u32 %r, %cluster_ctaid.y` +/// On sm90+: `mov.u32 %r, %cluster_ctaid.y`. On earlier targets: zero. #[inline(never)] pub fn cluster_ctaidY() -> u32 { unreachable!("cluster_ctaidY called outside CUDA kernel context") @@ -100,7 +115,7 @@ pub fn cluster_ctaidY() -> u32 { /// /// # PTX /// -/// Lowers to: `mov.u32 %r, %cluster_ctaid.z` +/// On sm90+: `mov.u32 %r, %cluster_ctaid.z`. On earlier targets: zero. #[inline(never)] pub fn cluster_ctaidZ() -> u32 { unreachable!("cluster_ctaidZ called outside CUDA kernel context") @@ -117,7 +132,7 @@ pub fn cluster_ctaidZ() -> u32 { /// /// # PTX /// -/// Lowers to: `mov.u32 %r, %cluster_nctaid.x` +/// On sm90+: `mov.u32 %r, %cluster_nctaid.x`. On earlier targets: one. #[inline(never)] pub fn cluster_nctaidX() -> u32 { unreachable!("cluster_nctaidX called outside CUDA kernel context") @@ -127,7 +142,7 @@ pub fn cluster_nctaidX() -> u32 { /// /// # PTX /// -/// Lowers to: `mov.u32 %r, %cluster_nctaid.y` +/// On sm90+: `mov.u32 %r, %cluster_nctaid.y`. On earlier targets: one. #[inline(never)] pub fn cluster_nctaidY() -> u32 { unreachable!("cluster_nctaidY called outside CUDA kernel context") @@ -137,7 +152,7 @@ pub fn cluster_nctaidY() -> u32 { /// /// # PTX /// -/// Lowers to: `mov.u32 %r, %cluster_nctaid.z` +/// On sm90+: `mov.u32 %r, %cluster_nctaid.z`. On earlier targets: one. #[inline(never)] pub fn cluster_nctaidZ() -> u32 { unreachable!("cluster_nctaidZ called outside CUDA kernel context") @@ -157,7 +172,8 @@ include!("generated/cluster_sreg.rs"); /// # PTX /// /// Lowers to documented `%clusterid.{x,y,z}` / `%nclusterid.{x,y}` reads and -/// computes `x + y * nx + z * nx * ny`. +/// computes `x + y * nx + z * nx * ny`. Below sm90, uses ordinary block +/// coordinates and grid dimensions in the same expression. #[inline(always)] pub fn cluster_idx() -> u32 { let x = __cluster_idxX(); @@ -173,6 +189,7 @@ pub fn cluster_idx() -> u32 { /// # PTX /// /// Lowers to documented `%nclusterid.{x,y,z}` reads and computes `nx * ny * nz`. +/// Below sm90, uses ordinary grid dimensions instead. #[inline(always)] pub fn num_clusters() -> u32 { __cluster_grid_dimX() * __cluster_grid_dimY() * __cluster_grid_dimZ() diff --git a/cuda-oxide/crates/mir-importer/src/pipeline.rs b/cuda-oxide/crates/mir-importer/src/pipeline.rs index 12d74984ef..1d9af3ea3e 100644 --- a/cuda-oxide/crates/mir-importer/src/pipeline.rs +++ b/cuda-oxide/crates/mir-importer/src/pipeline.rs @@ -450,6 +450,21 @@ fn backend_options_for(config: &PipelineConfig) -> BackendOptions { backend_options } +/// Select single-block-cluster semantics only for a known pre-sm90 target. +/// Invalid inputs are still diagnosed by the shared backend validator. +fn cluster_query_fallback_target(options: &BackendOptions) -> Option { + if let Some(target) = &options.target_arch { + let target = cuda_oxide_codegen::experimental::Target::parse(target).ok()?; + return (target.capability() < 90).then(|| target.sm()); + } + match &options.device_arch_hint { + Some(DeviceArchHint::Valid(arch)) if arch.as_ref().capability() < 90 => { + Some(arch.as_ref().sm()) + } + _ => None, + } +} + /// Runs the full compilation pipeline on collected functions. /// /// # Pipeline Steps @@ -505,6 +520,10 @@ pub fn run_pipeline( prepare_output_dir(&config.output_dir)?; + let mut backend_options = backend_options_for(config); + let cluster_query_target = cluster_query_fallback_target(&backend_options); + let mut used_cluster_query_fallback = false; + let mut ctx = Context::new(); // Step 1: Register dialects @@ -541,7 +560,7 @@ pub fn run_pipeline( .body() .ok_or_else(|| PipelineError::NoBody(func.export_name.clone()))?; - let func_op_ptr = crate::translator::body::translate_body( + let (func_op_ptr, used_fallback) = crate::translator::body::translate_body( &mut ctx, &body, &func.instance, @@ -554,11 +573,13 @@ pub fn run_pipeline( config.debug_kind, func.debug_source_scopes.as_ref(), func.statement_debug_info.as_ref(), + cluster_query_target.is_some(), ) .map_err(|e| { // Use .disp(&ctx) for rich error formatting with location and backtrace PipelineError::Translation(format!("{}: {}", func.export_name, e.disp(&ctx))) })?; + used_cluster_query_fallback |= used_fallback; if config.debug_kind.variables_enabled() { attach_debug_global_variables(&mut ctx, func_op_ptr, &config.debug_global_variables); @@ -594,7 +615,15 @@ pub fn run_pipeline( .join(format!("{}.ptx", config.output_name)); let stale_artifacts = stale_compilation_artifact_paths(&config.output_dir, &config.output_name); - let backend_options = backend_options_for(config); + // Like IKET materialization, target-dependent query lowering must bind an + // advisory device hint before later feature-based target selection. A + // higher-floor operation must reject this target, not silently retarget + // a module whose cluster queries already use single-block semantics. + if used_cluster_query_fallback && backend_options.target_arch.is_none() { + backend_options.target_arch = cluster_query_target; + backend_options.target_arch_source = + "the device hint, pinned by pre-sm90 cluster query fallback"; + } let request = ModulePipelineRequest::for_rust_pipeline( device_externs, @@ -764,6 +793,28 @@ mod tests { use llvm_export::ops::DebugLocalTypeKind; use std::fs; + #[test] + fn cluster_query_fallback_requires_a_known_pre90_target() { + for (explicit, hint, expected) in [ + (None, None, None), + (None, Some("sm_80"), Some("sm_80")), + (None, Some("sm_120"), None), + (None, Some("invalid"), None), + (Some("sm_80"), Some("sm_120"), Some("sm_80")), + (Some("sm_120"), Some("sm_80"), None), + (Some("invalid"), Some("sm_80"), None), + ] { + let mut options = BackendOptions::default(); + options.target_arch = explicit.map(str::to_owned); + options.device_arch_hint = hint.map(|hint| DeviceArchHint::parse(hint.to_owned())); + assert_eq!( + cluster_query_fallback_target(&options).as_deref(), + expected, + "explicit={explicit:?}, hint={hint:?}" + ); + } + } + #[test] fn test_pipeline_config_default_values() { let config = PipelineConfig::default(); diff --git a/cuda-oxide/crates/mir-importer/src/translator/body.rs b/cuda-oxide/crates/mir-importer/src/translator/body.rs index cd49e86c91..a1babdcfc8 100644 --- a/cuda-oxide/crates/mir-importer/src/translator/body.rs +++ b/cuda-oxide/crates/mir-importer/src/translator/body.rs @@ -1883,6 +1883,10 @@ fn emit_entry_allocas( /// * `inline` - The `#[inline]` request, recorded as [`InlineIntentAttr`] on /// non-kernel functions (`set_inline_intent_attr`) /// * `override_name` - Custom export name (defaults to instance name) +/// * `pre90_cluster_queries` - Use single-block-cluster query semantics +/// +/// Returns the function operation and whether it used a pre-sm90 cluster +/// query fallback. The caller must preserve that target choice. pub fn translate_body( ctx: &mut Context, body: &mir::Body, @@ -1896,7 +1900,8 @@ pub fn translate_body( debug_kind: DebugKind, debug_source_scopes: Option<&DebugSourceScopeMap>, statement_debug_info: Option<&StatementDebugInfoMap>, -) -> TranslationResult> { + pre90_cluster_queries: bool, +) -> TranslationResult<(Ptr, bool)> { if let Some(statement_debug_info) = statement_debug_info { if statement_debug_info.blocks.len() != body.blocks.len() { return input_err_noloc!(TranslationErr::invalid_op(format!( @@ -1935,6 +1940,7 @@ pub fn translate_body( let num_locals = body.locals().len(); let mut value_map = ValueMap::new(num_locals); value_map.set_debug_variables(debug_kind.variables_enabled()); + value_map.pre90_cluster_queries = pre90_cluster_queries; // Resolve the per-body unchecked-indexing policy. Like the dynamic-shared // marker, the `#[kernel(unchecked_indexing)]` marker is scanned on any @@ -2497,7 +2503,7 @@ pub fn translate_body( } } - Ok(op_ptr) + Ok((op_ptr, value_map.used_cluster_query_fallback)) } /// Choose the debugger-visible spelling independently of a function's symbol. diff --git a/cuda-oxide/crates/mir-importer/src/translator/terminator/intrinsics/cluster_query.rs b/cuda-oxide/crates/mir-importer/src/translator/terminator/intrinsics/cluster_query.rs new file mode 100644 index 0000000000..1c605c982f --- /dev/null +++ b/cuda-oxide/crates/mir-importer/src/translator/terminator/intrinsics/cluster_query.rs @@ -0,0 +1,187 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +//! Single-block-cluster queries for known targets below sm90. +//! +//! Intercept compatibility paths before generated dispatch erases their +//! identity. Raw `cuda_intrinsics` calls keep their native target requirements. + +use crate::error::{TranslationErr, TranslationResult}; +use crate::translator::{terminator::helpers, values::ValueMap}; +use dialect_mir::ops::MirConstantOp; +use dialect_nvvm::ops::{ + ReadPtxSregCtaidXOp, ReadPtxSregCtaidYOp, ReadPtxSregCtaidZOp, ReadPtxSregNctaidXOp, + ReadPtxSregNctaidYOp, ReadPtxSregNctaidZOp, +}; +use pliron::basic_block::BasicBlock; +use pliron::builtin::attributes::IntegerAttr; +use pliron::builtin::types::{IntegerType, Signedness}; +use pliron::context::{Context, Ptr}; +use pliron::input_err; +use pliron::location::Location; +use pliron::op::Op; +use pliron::operation::Operation; +use pliron::utils::apint::APInt; +use rustc_public::mir; +use std::num::NonZeroUsize; + +#[derive(Debug, PartialEq, Eq)] +enum Fallback { + Constant(u32), + BlockX, + BlockY, + BlockZ, + GridX, + GridY, + GridZ, +} + +fn fallback(name: &str) -> Option { + Some(match name { + "cuda_device::cluster::cluster_ctaidX" + | "cuda_device::cluster::cluster_ctaidY" + | "cuda_device::cluster::cluster_ctaidZ" => Fallback::Constant(0), + "cuda_device::cluster::cluster_nctaidX" + | "cuda_device::cluster::cluster_nctaidY" + | "cuda_device::cluster::cluster_nctaidZ" => Fallback::Constant(1), + "cuda_device::cluster::__cluster_idxX" => Fallback::BlockX, + "cuda_device::cluster::__cluster_idxY" => Fallback::BlockY, + "cuda_device::cluster::__cluster_idxZ" => Fallback::BlockZ, + "cuda_device::cluster::__cluster_grid_dimX" => Fallback::GridX, + "cuda_device::cluster::__cluster_grid_dimY" => Fallback::GridY, + "cuda_device::cluster::__cluster_grid_dimZ" => Fallback::GridZ, + _ => return None, + }) +} + +#[allow(clippy::too_many_arguments)] +pub(crate) fn try_dispatch( + ctx: &mut Context, + body: &mir::Body, + name: &str, + args: &[mir::Operand], + destination: &mir::Place, + target: &Option, + block: Ptr, + prev_op: Option>, + values: &mut ValueMap, + blocks: &[Ptr], + loc: Location, +) -> TranslationResult>> { + if !values.pre90_cluster_queries { + return Ok(None); + } + let Some(fallback) = fallback(name) else { + return Ok(None); + }; + if !args.is_empty() { + return input_err!( + loc, + TranslationErr::unsupported(format!( + "cluster query `{name}` expects zero arguments, got {}", + args.len() + )) + ); + } + + macro_rules! register { + ($op:ty, $marker:literal) => { + helpers::emit_generated_nvvm_intrinsic( + ctx, + body, + <$op>::get_concrete_op_info(), + $marker, + destination, + target, + block, + prev_op, + values, + blocks, + loc, + )? + }; + } + let operation = match fallback { + Fallback::Constant(value) => { + let (destination, prev_op) = helpers::prepare_destination_write( + ctx, + body, + destination, + values, + block, + prev_op, + loc.clone(), + )?; + let ty = IntegerType::get(ctx, 32, Signedness::Unsigned); + let operation = Operation::new( + ctx, + MirConstantOp::get_concrete_op_info(), + vec![ty.into()], + vec![], + vec![], + 0, + ); + let value = APInt::from_u64(u64::from(value), NonZeroUsize::new(32).unwrap()); + MirConstantOp::new(operation).set_attr_value(ctx, IntegerAttr::new(ty, value)); + operation.deref_mut(ctx).set_loc(loc.clone()); + helpers::insert_op(ctx, operation, block, prev_op); + let result = operation.deref(ctx).get_result(0); + helpers::emit_prepared_result_and_goto( + ctx, + destination, + result, + target, + block, + operation, + values, + blocks, + loc, + "cluster query has no successor", + )? + } + Fallback::BlockX => register!(ReadPtxSregCtaidXOp, "v1:i0002"), + Fallback::BlockY => register!(ReadPtxSregCtaidYOp, "v1:i0006"), + Fallback::BlockZ => register!(ReadPtxSregCtaidZOp, "v1:i0010"), + Fallback::GridX => register!(ReadPtxSregNctaidXOp, "v1:i0004"), + Fallback::GridY => register!(ReadPtxSregNctaidYOp, "v1:i0008"), + Fallback::GridZ => register!(ReadPtxSregNctaidZOp, "v1:i0012"), + }; + values.used_cluster_query_fallback = true; + Ok(Some(operation)) +} + +#[cfg(test)] +mod tests { + use super::*; + + #[test] + fn only_compatibility_queries_receive_fallbacks() { + assert_eq!( + fallback("cuda_device::cluster::cluster_ctaidX"), + Some(Fallback::Constant(0)) + ); + assert_eq!( + fallback("cuda_device::cluster::cluster_nctaidZ"), + Some(Fallback::Constant(1)) + ); + assert_eq!( + fallback("cuda_device::cluster::__cluster_idxY"), + Some(Fallback::BlockY) + ); + assert_eq!( + fallback("cuda_device::cluster::__cluster_grid_dimZ"), + Some(Fallback::GridZ) + ); + for name in [ + "cuda_intrinsics::__cuda_oxide_intrinsic_abi_v1::i0266", + "cuda_intrinsics::sreg::cluster_dim_x", + "cuda_device::cluster::cluster_sync", + "cuda_device::cluster::map_shared_rank", + "other::cluster::cluster_ctaidX", + ] { + assert_eq!(fallback(name), None, "{name}"); + } + } +} diff --git a/cuda-oxide/crates/mir-importer/src/translator/terminator/intrinsics/mod.rs b/cuda-oxide/crates/mir-importer/src/translator/terminator/intrinsics/mod.rs index 5c50b07cde..175d25669d 100644 --- a/cuda-oxide/crates/mir-importer/src/translator/terminator/intrinsics/mod.rs +++ b/cuda-oxide/crates/mir-importer/src/translator/terminator/intrinsics/mod.rs @@ -46,6 +46,7 @@ pub mod asm; pub mod atomic; pub mod bigint; pub mod bitops; +pub(crate) mod cluster_query; pub mod debug; pub mod exact_div; pub mod float_math; 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 d88f37d4b0..ca9812195f 100644 --- a/cuda-oxide/crates/mir-importer/src/translator/terminator/mod.rs +++ b/cuda-oxide/crates/mir-importer/src/translator/terminator/mod.rs @@ -2701,6 +2701,22 @@ fn try_dispatch_intrinsic( return Ok(Some(operation)); } + if let Some(operation) = intrinsics::cluster_query::try_dispatch( + ctx, + body, + name, + args, + destination, + target, + block_ptr, + prev_op, + value_map, + block_map, + loc.clone(), + )? { + return Ok(Some(operation)); + } + if let Some(operation) = intrinsics::generated::try_dispatch_generated_intrinsic( ctx, body, diff --git a/cuda-oxide/crates/mir-importer/src/translator/values.rs b/cuda-oxide/crates/mir-importer/src/translator/values.rs index 450665aef0..9333a133be 100644 --- a/cuda-oxide/crates/mir-importer/src/translator/values.rs +++ b/cuda-oxide/crates/mir-importer/src/translator/values.rs @@ -77,6 +77,10 @@ pub struct ValueMap { /// When set, `translate_assert` elides `AssertMessage::BoundsCheck` /// terminators (out-of-bounds indexing becomes UB, like `get_unchecked`). unchecked_indexing: bool, + /// The pipeline selected single-block-cluster semantics for this body. + pub(crate) pre90_cluster_queries: bool, + /// A compatibility query actually used those semantics. + pub(crate) used_cluster_query_fallback: bool, } impl ValueMap { @@ -87,6 +91,8 @@ impl ValueMap { statement_debug_spills: vec![false; num_locals], debug_variables: false, unchecked_indexing: false, + pre90_cluster_queries: false, + used_cluster_query_fallback: false, } } From 038147676be4309e91e81654d295b6f38464d8c2 Mon Sep 17 00:00:00 2001 From: Trent Nelson Date: Mon, 19 Oct 2026 23:58:59 -0700 Subject: [PATCH 2/3] Test cluster query fallback through the MIR pipeline Signed-off-by: Trent Nelson --- cuda-oxide/crates/mir-importer/README.md | 10 + .../crates/mir-importer/src/pipeline.rs | 200 ++++++++++++++++++ 2 files changed, 210 insertions(+) diff --git a/cuda-oxide/crates/mir-importer/README.md b/cuda-oxide/crates/mir-importer/README.md index 6815539205..52c14a32aa 100644 --- a/cuda-oxide/crates/mir-importer/README.md +++ b/cuda-oxide/crates/mir-importer/README.md @@ -151,6 +151,16 @@ The pipeline inspects which intrinsics the code uses and selects a target: Override with `CUDA_OXIDE_TARGET=`. +From `cuda-oxide/`, run the source-to-NVVM-IR cluster-query regression with: + +```sh +cargo test -p mir-importer cluster_query_lowering_preserves_target_selection_and_raw_floors +``` + +It checks pre-sm90 compatibility lowering, conditional device-hint pinning, +and the raw intrinsic target floor using the pinned Rust toolchain. It requires +neither a GPU nor CUDA compiler tools. + ## Public API ### Types diff --git a/cuda-oxide/crates/mir-importer/src/pipeline.rs b/cuda-oxide/crates/mir-importer/src/pipeline.rs index 1d9af3ea3e..79b693960d 100644 --- a/cuda-oxide/crates/mir-importer/src/pipeline.rs +++ b/cuda-oxide/crates/mir-importer/src/pipeline.rs @@ -815,6 +815,206 @@ mod tests { } } + #[test] + fn cluster_query_lowering_preserves_target_selection_and_raw_floors() { + const CHILD: &str = "CUDA_OXIDE_TEST_CLUSTER_QUERY_PIPELINE"; + if std::env::var_os(CHILD).is_none() { + // Hint-only cases must not inherit an explicit target. Keep + // environment changes out of the parallel test process. + let output = std::process::Command::new(std::env::current_exe().unwrap()) + .args([ + "--exact", + "pipeline::tests::cluster_query_lowering_preserves_target_selection_and_raw_floors", + "--nocapture", + ]) + .env(CHILD, "1") + .env_remove("CUDA_OXIDE_TARGET") + .env_remove("CUDA_OXIDE_DEVICE_ARCH") + .output() + .unwrap(); + assert!( + output.status.success(), + "{}\n{}", + String::from_utf8_lossy(&output.stdout), + String::from_utf8_lossy(&output.stderr) + ); + return; + } + + let unique = std::time::SystemTime::now() + .duration_since(std::time::UNIX_EPOCH) + .unwrap() + .as_nanos(); + let root = std::env::temp_dir().join(format!( + "cuda_oxide_cluster_queries_{}_{}", + std::process::id(), + unique + )); + fs::create_dir_all(&root).unwrap(); + let fixture = root.join("queries.rs"); + // These non-inlined placeholders reproduce the defining paths at the + // importer boundary. Their bodies are not collected: actual dispatch + // must replace each call, never execute the sentinel value. + fs::write( + &fixture, + r#" +#![allow(dead_code, non_snake_case)] +macro_rules! queries { + ($($name:ident),*) => {$( + #[inline(never)] pub fn $name() -> u32 { 99 } + )*}; +} +mod cluster { + queries!(cluster_ctaidX, cluster_ctaidY, cluster_ctaidZ, + cluster_nctaidX, cluster_nctaidY, cluster_nctaidZ, + __cluster_idxX, __cluster_idxY, __cluster_idxZ, + __cluster_grid_dimX, __cluster_grid_dimY, __cluster_grid_dimZ); +} +mod warp { + #[inline(never)] pub fn elect_sync(_: u32) -> (u32, bool) { (99, false) } +} +mod __cuda_oxide_intrinsic_abi_v1 { queries!(i0266); } +fn compatibility() -> (u32, u32, u32, u32, u32, u32, u32, u32, u32, u32, u32, u32) { + use cluster::*; + (cluster_ctaidX(), cluster_ctaidY(), cluster_ctaidZ(), + cluster_nctaidX(), cluster_nctaidY(), cluster_nctaidZ(), + __cluster_idxX(), __cluster_idxY(), __cluster_idxZ(), + __cluster_grid_dimX(), __cluster_grid_dimY(), __cluster_grid_dimZ()) +} +fn compatibility_and_elect() -> (u32, u32) { + (cluster::cluster_nctaidX(), warp::elect_sync(u32::MAX).0) +} +fn elect_only() -> u32 { warp::elect_sync(u32::MAX).0 } +fn raw() -> u32 { __cuda_oxide_intrinsic_abi_v1::i0266() } +"#, + ) + .unwrap(); + let rustc = std::env::var_os("RUSTC").unwrap_or_else(|| "rustc".into()); + let sysroot = std::process::Command::new(rustc) + .args(["--print", "sysroot"]) + .output() + .unwrap(); + assert!(sysroot.status.success()); + let sysroot = String::from_utf8(sysroot.stdout).unwrap(); + + for crate_name in ["cuda_device", "cuda_intrinsics"] { + let args = vec![ + "rustc".to_string(), + "--edition=2024".to_string(), + "--crate-type=rlib".to_string(), + format!("--crate-name={crate_name}"), + "--emit=metadata".to_string(), + "-Zmir-opt-level=0".to_string(), + format!("--out-dir={}", root.display()), + format!("--sysroot={}", sysroot.trim()), + fixture.display().to_string(), + ]; + let output_dir = root.join(crate_name); + std::thread::Builder::new() + .stack_size(64 * 1024 * 1024) + .spawn(move || { + rustc_public::run!(&args, || { + use rustc_public::CrateDef as _; + + let compile = |name: &str, target: Option<&str>, hint: Option<&str>| { + let item = rustc_public::all_local_items() + .into_iter() + .find(|item| item.name() == format!("{crate_name}::{name}")) + .expect("fixture entry"); + let instance = Instance::try_from(item).unwrap(); + let body = instance.body().unwrap(); + let function = CollectedFunction { + instance, + rustc_mir_block_count: body.blocks.len(), + // These monomorphic, branch-free fixtures have + // no specialization-dependent dead successors. + rustc_mono_successors: body + .blocks + .iter() + .map(|block| block.terminator.successors()) + .collect(), + is_kernel: false, + export_name: name.to_string(), + debug_source_scopes: None, + statement_debug_info: None, + inline: None, + }; + let config = PipelineConfig { + output_dir: output_dir.clone(), + output_name: name.to_string(), + verbose: false, + emit_nvvm_ir: true, + target_arch: target.map(str::to_string), + device_arch_hint: hint + .map(|hint| DeviceArchHint::parse(hint.to_string())), + ..PipelineConfig::default() + }; + run_pipeline(&[function], &[], &config, Default::default()) + .map(|result| { + (result.target, fs::read_to_string(result.ll_path).unwrap()) + }) + .map_err(|error| error.to_string()) + }; + if crate_name == "cuda_device" { + for (target, hint) in + [(Some("sm_80"), Some("sm_120")), (None, Some("sm_80"))] + { + let (target, ir) = compile("compatibility", target, hint).unwrap(); + assert_eq!(target, "sm_80"); + for register in ["ctaid", "nctaid"] { + for axis in ["x", "y", "z"] { + assert!(ir.contains(&format!( + "call i32 @llvm.nvvm.read.ptx.sreg.{register}.{axis}()" + ))); + } + } + assert!(!ir.contains("sreg.cluster")); + assert!(!ir.contains("sreg.ncluster")); + for index in 0..6 { + let value = u32::from(index >= 3); + assert!( + ir.lines().any(|line| { + line.contains("insertvalue") + && line.contains(&format!("i32 {value}, {index}")) + }), + "{ir}" + ); + } + } + let error = compile("compatibility_and_elect", None, Some("sm_80")) + .expect_err("used fallback must pin the advisory hint"); + assert!(error.contains("sm_80") && error.contains("Sm90"), "{error}"); + let (target, _) = compile("elect_only", None, Some("sm_80")).unwrap(); + assert_eq!(target, "sm_90", "unused fallback must not pin the hint"); + let (target, ir) = compile("compatibility", None, None).unwrap(); + assert_eq!(target, "sm_90"); + assert!( + ir.contains("call i32 @llvm.nvvm.read.ptx.sreg.cluster.nctaid.x()") + ); + } else { + let error = compile("raw", Some("sm_80"), None) + .expect_err("raw intrinsics retain their native target floor"); + assert!( + error.contains("v1:i0266") && error.contains("sm_90"), + "{error}" + ); + let (target, ir) = compile("raw", Some("sm_120"), None).unwrap(); + assert_eq!(target, "sm_120"); + assert!( + ir.contains("call i32 @llvm.nvvm.read.ptx.sreg.cluster.nctaid.x()") + ); + } + std::ops::ControlFlow::<(), ()>::Continue(()) + }) + }) + .unwrap() + .join() + .unwrap() + .expect("in-process fixture compilation succeeds"); + } + fs::remove_dir_all(root).unwrap(); + } + #[test] fn test_pipeline_config_default_values() { let config = PipelineConfig::default(); From 333fc8d6473d03fbde812fa466cc5cdab4c6633c Mon Sep 17 00:00:00 2001 From: Trent Nelson Date: Tue, 20 Oct 2026 00:24:02 -0700 Subject: [PATCH 3/3] Verify cluster query axes and explain target pinning Signed-off-by: Trent Nelson --- cuda-oxide/crates/mir-importer/README.md | 5 +++- .../crates/mir-importer/src/pipeline.rs | 27 ++++++++++++++++--- 2 files changed, 27 insertions(+), 5 deletions(-) diff --git a/cuda-oxide/crates/mir-importer/README.md b/cuda-oxide/crates/mir-importer/README.md index 52c14a32aa..8c56d2cc6a 100644 --- a/cuda-oxide/crates/mir-importer/README.md +++ b/cuda-oxide/crates/mir-importer/README.md @@ -147,9 +147,12 @@ The pipeline inspects which intrinsics the code uses and selects a target: | tcgen05 / TMEM | sm_100a | Blackwell datacenter | | WGMMA | sm_90a | Hopper only | | TMA / mbarrier | sm_100 | Hopper+ compatible | +| Cluster compatibility queries | sm_90 by default | Known pre-sm90 targets use one-block clusters | | Basic CUDA | sm_80 | Ampere+ (max compat) | -Override with `CUDA_OXIDE_TARGET=`. +Override with `CUDA_OXIDE_TARGET=`. A device hint used for pre-sm90 +cluster-query fallback pins the artifact to that target; incompatible higher-floor +operations are rejected. Unused hints remain advisory. From `cuda-oxide/`, run the source-to-NVVM-IR cluster-query regression with: diff --git a/cuda-oxide/crates/mir-importer/src/pipeline.rs b/cuda-oxide/crates/mir-importer/src/pipeline.rs index 79b693960d..bbb3007ff2 100644 --- a/cuda-oxide/crates/mir-importer/src/pipeline.rs +++ b/cuda-oxide/crates/mir-importer/src/pipeline.rs @@ -961,11 +961,30 @@ fn raw() -> u32 { __cuda_oxide_intrinsic_abi_v1::i0266() } { let (target, ir) = compile("compatibility", target, hint).unwrap(); assert_eq!(target, "sm_80"); - for register in ["ctaid", "nctaid"] { - for axis in ["x", "y", "z"] { - assert!(ir.contains(&format!( + for (register_index, register) in + ["ctaid", "nctaid"].iter().enumerate() + { + for (axis_index, axis) in ["x", "y", "z"].iter().enumerate() { + let instruction = format!( "call i32 @llvm.nvvm.read.ptx.sreg.{register}.{axis}()" - ))); + ); + let value = ir + .lines() + .find_map(|line| { + line.contains(&instruction).then(|| { + line.trim().split_once(" = ").unwrap().0 + }) + }) + .expect("compatibility register read"); + let slot = 6 + register_index * 3 + axis_index; + assert!( + ir.lines().any(|line| { + line.contains("insertvalue") + && line + .contains(&format!("i32 {value}, {slot}")) + }), + "{ir}" + ); } } assert!(!ir.contains("sreg.cluster"));