diff --git a/cuda-oxide/crates/cuda-device/src/atomic.rs b/cuda-oxide/crates/cuda-device/src/atomic.rs index 84423472ed..7433fec228 100644 --- a/cuda-oxide/crates/cuda-device/src/atomic.rs +++ b/cuda-oxide/crates/cuda-device/src/atomic.rs @@ -201,8 +201,10 @@ macro_rules! define_integer_atomic { /// atomic and non-atomic access to the same location is UB. /// * For block-scoped types, only threads in the same thread block /// may access this memory. - /// * For system-scoped types, the memory must be reachable by both - /// the GPU and the CPU (unified memory). + /// * The memory and concurrently accessing threads must support + /// the chosen scope. System scope permits device memory accessed + /// only by threads on one GPU; CPU or peer-GPU access must satisfy + /// the [CUDA atomicity requirements](https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cuda-cpp-memory-model.html#atomicity). #[inline(always)] pub const unsafe fn from_ptr<'a>(ptr: *mut $ty) -> &'a Self { // SAFETY: `Self` is `#[repr(transparent)]` over `UnsafeCell<$ty>`, @@ -388,7 +390,9 @@ macro_rules! define_float_atomic { /// the entire lifetime `'a`. /// * No non-atomic accesses to `*ptr` may occur during `'a`. /// * Block-scoped: only same-block threads may access. - /// System-scoped: memory must be reachable by GPU and CPU. + /// * The memory and concurrently accessing threads must support + /// the chosen scope, as described on the integer atomic types. + /// System scope permits device memory accessed only by one GPU. #[inline(always)] pub const unsafe fn from_ptr<'a>(ptr: *mut $ty) -> &'a Self { // SAFETY: `Self` is `#[repr(transparent)]` over `UnsafeCell<$ty>`, diff --git a/cuda-oxide/crates/dialect-mir/src/attributes.rs b/cuda-oxide/crates/dialect-mir/src/attributes.rs index 1152f531e3..c5b82dce9a 100644 --- a/cuda-oxide/crates/dialect-mir/src/attributes.rs +++ b/cuda-oxide/crates/dialect-mir/src/attributes.rs @@ -79,19 +79,31 @@ pub struct MutabilityAttr(pub bool); /// Proven validity facts for one Rust-reference kernel parameter. /// /// Presence of this attribute means the source argument is a Rust reference -/// proven non-null by `rustc_public`. The payload is the pointee alignment in -/// bytes; it must be a non-zero power of two. LLVM `align 1` is redundant and +/// proven non-null by `rustc_public`. Pointee alignment is in bytes and must +/// be a non-zero power of two. LLVM `align 1` is redundant and /// is omitted during export, but retaining `1` here keeps the proof complete. /// /// The attribute is stored on `mir.func` under a source-argument-indexed key. -/// It intentionally carries no aliasing or dereferenceability promise. -#[pliron_attr( - name = "mir.reference_param_validity", - format = "$0", - verifier = "succ" -)] +/// Additional promises are independently proven by the importer. A zero +/// dereferenceable byte count records no dereferenceability promise. +#[pliron_attr(name = "mir.reference_param_validity", format, verifier = "succ")] #[derive(PartialEq, Eq, Clone, Copy, Debug, Hash)] -pub struct ReferenceParamValidityAttr(pub u64); +pub struct ReferenceParamValidityAttr { + pub pointee_alignment: u64, + pub no_alias: bool, + pub dereferenceable_bytes: u64, +} + +impl ReferenceParamValidityAttr { + /// Non-nullness and alignment, without additional reference promises. + pub const fn new(pointee_alignment: u64) -> Self { + Self { + pointee_alignment, + no_alias: false, + dereferenceable_bytes: 0, + } + } +} /// Structural field index for aggregate access ops /// (`mir.extract_field`, `mir.insert_field`, `mir.field_addr`, `mir.enum_payload`). diff --git a/cuda-oxide/crates/dialect-mir/src/ops/function.rs b/cuda-oxide/crates/dialect-mir/src/ops/function.rs index 8ed3037336..7716e39c78 100644 --- a/cuda-oxide/crates/dialect-mir/src/ops/function.rs +++ b/cuda-oxide/crates/dialect-mir/src/ops/function.rs @@ -66,7 +66,7 @@ fn reference_param_validity_key(index: usize) -> Identifier { /// |----------------|-----------|------------------------------------| /// | `sym_name` | StringAttr| Function name (from SymbolOpInterface) | /// | `mir_func_type`| TypeAttr | Function type (mir.func_type) | -/// | `reference_param_validity_N` | ReferenceParamValidityAttr | Proven nonnull/alignment for source argument `N` on a kernel entry | +/// | `reference_param_validity_N` | ReferenceParamValidityAttr | Proven reference facts for source argument `N` on a kernel entry | /// ``` /// /// # Verification @@ -116,8 +116,9 @@ impl MirFuncOp { /// Record rustc-proven validity for one source-level kernel argument. /// - /// Presence proves `nonnull`; the payload carries the rustc ABI alignment - /// of the pointee. The source argument index is deliberately retained until + /// Presence proves `nonnull`; the payload carries the pointee alignment + /// and independently proven aliasing and size facts. + /// The source argument index is deliberately retained until /// kernel ABI lowering decides which physical LLVM parameter represents it. pub fn set_reference_param_validity( &self, @@ -285,12 +286,12 @@ impl Verify for MirFuncOp { key_text ); }; - if validity.0 == 0 || !validity.0.is_power_of_two() { + if validity.pointee_alignment == 0 || !validity.pointee_alignment.is_power_of_two() { return verify_err!( op.loc(), "MirFuncOp reference validity alignment on argument {} must be a non-zero power of two, found {}", index, - validity.0 + validity.pointee_alignment ); } diff --git a/cuda-oxide/crates/dialect-mir/tests/ops_test/scalars_and_funcs.rs b/cuda-oxide/crates/dialect-mir/tests/ops_test/scalars_and_funcs.rs index d39db3e74e..e0ce9f0104 100644 --- a/cuda-oxide/crates/dialect-mir/tests/ops_test/scalars_and_funcs.rs +++ b/cuda-oxide/crates/dialect-mir/tests/ops_test/scalars_and_funcs.rs @@ -510,43 +510,50 @@ fn test_mir_func_reference_param_validity_verify() { }; let valid = make_func(&mut ctx, true); - valid.set_reference_param_validity(&mut ctx, 0, ReferenceParamValidityAttr(4)); - valid.set_reference_param_validity(&mut ctx, 1, ReferenceParamValidityAttr(4)); + valid.set_reference_param_validity( + &mut ctx, + 0, + ReferenceParamValidityAttr { + dereferenceable_bytes: 4, + ..ReferenceParamValidityAttr::new(4) + }, + ); + valid.set_reference_param_validity(&mut ctx, 1, ReferenceParamValidityAttr::new(4)); assert!( valid.verify(&ctx).is_ok(), "Rust reference facts on kernel source arguments are valid" ); let non_kernel = make_func(&mut ctx, false); - non_kernel.set_reference_param_validity(&mut ctx, 0, ReferenceParamValidityAttr(4)); + non_kernel.set_reference_param_validity(&mut ctx, 0, ReferenceParamValidityAttr::new(4)); assert!( non_kernel.verify(&ctx).is_err(), "reference validity is kernel-entry-only" ); let raw = make_func(&mut ctx, true); - raw.set_reference_param_validity(&mut ctx, 2, ReferenceParamValidityAttr(4)); + raw.set_reference_param_validity(&mut ctx, 2, ReferenceParamValidityAttr::new(4)); assert!( raw.verify(&ctx).is_err(), "raw pointers cannot carry Rust-reference validity" ); let scalar = make_func(&mut ctx, true); - scalar.set_reference_param_validity(&mut ctx, 3, ReferenceParamValidityAttr(4)); + scalar.set_reference_param_validity(&mut ctx, 3, ReferenceParamValidityAttr::new(4)); assert!( scalar.verify(&ctx).is_err(), "non-pointer parameters cannot carry Rust-reference validity" ); let bad_alignment = make_func(&mut ctx, true); - bad_alignment.set_reference_param_validity(&mut ctx, 0, ReferenceParamValidityAttr(3)); + bad_alignment.set_reference_param_validity(&mut ctx, 0, ReferenceParamValidityAttr::new(3)); assert!( bad_alignment.verify(&ctx).is_err(), "alignment must be a non-zero power of two" ); let out_of_range = make_func(&mut ctx, true); - out_of_range.set_reference_param_validity(&mut ctx, 9, ReferenceParamValidityAttr(4)); + out_of_range.set_reference_param_validity(&mut ctx, 9, ReferenceParamValidityAttr::new(4)); assert!( out_of_range.verify(&ctx).is_err(), "source argument index must be in range" diff --git a/cuda-oxide/crates/llvm-export/src/export/function.rs b/cuda-oxide/crates/llvm-export/src/export/function.rs index 324afd7996..3fa26fa0ae 100644 --- a/cuda-oxide/crates/llvm-export/src/export/function.rs +++ b/cuda-oxide/crates/llvm-export/src/export/function.rs @@ -61,6 +61,30 @@ fn sanitize_comment(text: &str) -> String { .collect() } +/// Spell only the facts already proven at the import boundary. Shared by +/// definitions and declarations, including the empty-region declaration path. +fn export_reference_parameter_attrs( + validity: ops::KernelReferenceParamValidityAttr, + include_alignment: bool, + output: &mut String, +) { + write!(output, " nonnull").unwrap(); + if validity.pointee_alignment > 1 && include_alignment { + write!(output, " align {}", validity.pointee_alignment).unwrap(); + } + if validity.no_alias { + write!(output, " noalias").unwrap(); + } + if validity.dereferenceable_bytes > 0 { + write!( + output, + " dereferenceable({})", + validity.dereferenceable_bytes + ) + .unwrap(); + } +} + /// Final LLVM spelling of a pliron function symbol. /// /// Shared-global owner attributes carry the raw symbol, so the module prepass, @@ -799,7 +823,7 @@ impl<'a> ModuleExportState<'a> { "kernel `@{fixed_func_name}` reference validity parameter {index} is not an LLVM pointer" )); } - reference_param_validities[index] = Some(validity.0); + reference_param_validities[index] = Some(validity); } let is_declaration = func.get_operation().deref(self.ctx).regions().count() == 0; @@ -975,11 +999,8 @@ impl<'a> ModuleExportState<'a> { let grid_constant = grid_constant_params .iter() .find(|parameter| parameter.index == i); - if let Some(alignment) = reference_param_validities[i] { - write!(output, " nonnull").unwrap(); - if alignment > 1 && grid_constant.is_none() { - write!(output, " align {alignment}").unwrap(); - } + if let Some(validity) = reference_param_validities[i] { + export_reference_parameter_attrs(validity, grid_constant.is_none(), output); } if let Some(parameter) = grid_constant { self.export_grid_constant_parameter_attrs( @@ -1046,15 +1067,9 @@ impl<'a> ModuleExportState<'a> { let block = entry_block.deref(self.ctx); let args = block.arguments(); let mut entry_prologue = String::new(); - // Kernel Rust references may carry importer-proven `nonnull` and - // pointee `align N`. Every other parameter remains bare. - // - // In particular this deliberately does NOT infer `noalias`, - // `readonly`, `dereferenceable`, or any validity fact from LLVM - // pointer shape. `DisjointSlice` and raw pointers therefore retain - // the conservative behavior required by their unsafe construction - // contracts. Future alias attributes need their own audited proof - // source rather than extending this mechanical export step. + // Kernel Rust references carry only importer-proven facts. + // LLVM pointer shape alone grants no reference promises, so raw + // pointers, slice metadata, and other parameters remain bare. for (i, arg) in args.enumerate() { if i > 0 { write!(output, ", ").unwrap(); @@ -1064,11 +1079,8 @@ impl<'a> ModuleExportState<'a> { let grid_constant = grid_constant_params .iter() .find(|parameter| parameter.index == i); - if let Some(alignment) = reference_param_validities[i] { - write!(output, " nonnull").unwrap(); - if alignment > 1 && grid_constant.is_none() { - write!(output, " align {alignment}").unwrap(); - } + if let Some(validity) = reference_param_validities[i] { + export_reference_parameter_attrs(validity, grid_constant.is_none(), output); } if let Some(parameter) = grid_constant { self.export_grid_constant_parameter_attrs( @@ -1409,11 +1421,8 @@ impl<'a> ModuleExportState<'a> { let grid_constant = grid_constant_params .iter() .find(|parameter| parameter.index == i); - if let Some(alignment) = reference_param_validities[i] { - write!(output, " nonnull").unwrap(); - if alignment > 1 && grid_constant.is_none() { - write!(output, " align {alignment}").unwrap(); - } + if let Some(validity) = reference_param_validities[i] { + export_reference_parameter_attrs(validity, grid_constant.is_none(), output); } if let Some(parameter) = grid_constant { self.export_grid_constant_parameter_attrs( diff --git a/cuda-oxide/crates/llvm-export/src/lib.rs b/cuda-oxide/crates/llvm-export/src/lib.rs index 6509d9277c..9f891f829e 100644 --- a/cuda-oxide/crates/llvm-export/src/lib.rs +++ b/cuda-oxide/crates/llvm-export/src/lib.rs @@ -210,16 +210,31 @@ pub mod ops { /// Kernel-entry parameter validity proven at the Rust MIR import boundary. /// - /// Presence proves `nonnull`; the payload is the rustc ABI alignment of - /// the pointee represented by this physical LLVM parameter. It deliberately - /// carries no aliasing, readonly, or dereferenceability promise. + /// Presence proves `nonnull`; each field transports an independently + /// proven promise for this physical LLVM parameter. A zero byte count + /// records no dereferenceability promise. #[pliron_attr( name = "llvm.kernel_reference_param_validity", - format = "$0", + format, verifier = "succ" )] #[derive(Clone, Copy, Debug, PartialEq, Eq, Hash)] - pub struct KernelReferenceParamValidityAttr(pub u64); + pub struct KernelReferenceParamValidityAttr { + pub pointee_alignment: u64, + pub no_alias: bool, + pub dereferenceable_bytes: u64, + } + + impl KernelReferenceParamValidityAttr { + /// Non-nullness and alignment, without additional reference promises. + pub const fn new(pointee_alignment: u64) -> Self { + Self { + pointee_alignment, + no_alias: false, + dereferenceable_bytes: 0, + } + } + } fn kernel_reference_param_validity_key(index: usize) -> Identifier { Identifier::try_new(reserved_oxide_symbols::kernel_reference_param_validity_key( @@ -271,10 +286,10 @@ pub mod ops { "kernel reference parameter validity attribute `{key_text}` has the wrong attribute type" )); }; - if validity.0 == 0 || !validity.0.is_power_of_two() { + if validity.pointee_alignment == 0 || !validity.pointee_alignment.is_power_of_two() { return Err(format!( "kernel reference parameter {index} alignment must be a non-zero power of two, found {}", - validity.0 + validity.pointee_alignment )); } result.push((index, validity)); diff --git a/cuda-oxide/crates/llvm-export/tests/export_test/reference_validity.rs b/cuda-oxide/crates/llvm-export/tests/export_test/reference_validity.rs index 58b3a36924..4767ff0c28 100644 --- a/cuda-oxide/crates/llvm-export/tests/export_test/reference_validity.rs +++ b/cuda-oxide/crates/llvm-export/tests/export_test/reference_validity.rs @@ -62,13 +62,13 @@ fn kernel_reference_validity_exports_nonnull_and_alignment() { &mut ctx, func.get_operation(), 0, - KernelReferenceParamValidityAttr(4), + KernelReferenceParamValidityAttr::new(4), ); set_kernel_reference_param_validity( &mut ctx, func.get_operation(), 1, - KernelReferenceParamValidityAttr(1), + KernelReferenceParamValidityAttr::new(1), ); func.get_operation().insert_at_back(module_block, &ctx); @@ -93,7 +93,7 @@ fn legacy_kernel_reference_validity_follows_typed_pointer_spelling() { &mut ctx, func.get_operation(), 0, - KernelReferenceParamValidityAttr(4), + KernelReferenceParamValidityAttr::new(4), ); func.get_operation().insert_at_back(module_block, &ctx); @@ -116,7 +116,7 @@ fn reference_validity_rejects_non_kernel_non_pointer_and_out_of_range() { &mut ctx, non_kernel.get_operation(), 0, - KernelReferenceParamValidityAttr(4), + KernelReferenceParamValidityAttr::new(4), ); non_kernel .get_operation() @@ -132,7 +132,7 @@ fn reference_validity_rejects_non_kernel_non_pointer_and_out_of_range() { &mut ctx, scalar.get_operation(), 0, - KernelReferenceParamValidityAttr(4), + KernelReferenceParamValidityAttr::new(4), ); scalar.get_operation().insert_at_back(scalar_block, &ctx); let error = export_module_to_string_with_config(&ctx, &scalar_module, &PtxExportConfig) @@ -146,7 +146,7 @@ fn reference_validity_rejects_non_kernel_non_pointer_and_out_of_range() { &mut ctx, range.get_operation(), 3, - KernelReferenceParamValidityAttr(4), + KernelReferenceParamValidityAttr::new(4), ); range.get_operation().insert_at_back(range_block, &ctx); let error = export_module_to_string_with_config(&ctx, &range_module, &PtxExportConfig) @@ -165,13 +165,98 @@ fn reference_validity_rejects_malformed_marker_indices() { let ptr = PointerType::get(&ctx, 0); let func = kernel_with_params(&mut ctx, "malformed_marker", vec![ptr.into()], true); let key = format!("{KERNEL_REFERENCE_PARAM_VALIDITY_KEY_PREFIX}{suffix}"); - func.get_operation() - .deref_mut(&ctx) - .attributes - .set(key.try_into().unwrap(), KernelReferenceParamValidityAttr(4)); + func.get_operation().deref_mut(&ctx).attributes.set( + key.try_into().unwrap(), + KernelReferenceParamValidityAttr::new(4), + ); func.get_operation().insert_at_back(module_block, &ctx); let error = export_module_to_string_with_config(&ctx, &module, &PtxExportConfig) .expect_err("malformed reference-validity marker must fail"); assert!(error.contains("invalid parameter index"), "{error}"); } } + +#[test] +fn reference_validity_exports_each_fact_independently_in_all_dialects() { + use llvm_export::export::ExportBackendConfig; + + let modern_nvvm = NvvmExportConfig::new(NvvmIrDialect::Modern); + let legacy_nvvm = NvvmExportConfig::new(NvvmIrDialect::LegacyLlvm7); + let configs: [&dyn ExportBackendConfig; 3] = [&PtxExportConfig, &modern_nvvm, &legacy_nvvm]; + for (config, pointer) in configs.into_iter().zip(["ptr", "ptr", "i8*"]) { + // Cover definitions, declarations, and declarations with an empty region. + for shape in 0..3 { + let definition = shape == 0; + let mut ctx = Context::new(); + let module = ModuleOp::new(&mut ctx, "independent_facts".try_into().unwrap()); + let module_block = module_top_block(&mut ctx, &module); + let ptr = PointerType::get(&ctx, 0); + let integer = IntegerType::get(&ctx, 64, Signedness::Signless); + let mut params = vec![ptr.into(); 6]; + params.push(integer.into()); + let func = if definition { + kernel_with_params(&mut ctx, "independent_facts", params, true) + } else { + let ty = FuncType::get(&ctx, VoidType::get(&ctx).into(), params, false); + let func = FuncOp::new(&mut ctx, "independent_facts".try_into().unwrap(), ty); + if shape == 2 { + pliron::operation::Operation::add_region(func.get_operation(), &mut ctx); + } + func.get_operation().deref_mut(&ctx).attributes.set( + "gpu_kernel".try_into().unwrap(), + StringAttr::new("true".to_string()), + ); + func + }; + for (index, (alignment, no_alias, bytes)) in [ + (4, true, 4), + (8, true, 8), + (1, false, 0), + (16, false, 128), + (4, true, 0), + ] + .into_iter() + .enumerate() + { + set_kernel_reference_param_validity( + &mut ctx, + func.get_operation(), + index, + KernelReferenceParamValidityAttr { + pointee_alignment: alignment, + no_alias, + dereferenceable_bytes: bytes, + }, + ); + } + func.get_operation().insert_at_back(module_block, &ctx); + let ir = export_module_to_string_with_config(&ctx, &module, config).unwrap(); + assert!(!ir.contains(" readonly"), "{ir}"); + let expected = [ + format!("{pointer} nonnull align 4 noalias dereferenceable(4)"), + format!("{pointer} nonnull align 8 noalias dereferenceable(8)"), + format!("{pointer} nonnull"), + format!("{pointer} nonnull align 16 dereferenceable(128)"), + format!("{pointer} nonnull align 4 noalias"), + pointer.to_string(), + "i64".to_string(), + ] + .into_iter() + .enumerate() + .map(|(index, parameter)| { + if definition { + format!("{parameter} %v{index}") + } else { + parameter + } + }) + .collect::>() + .join(", "); + assert!( + ir.contains(&format!("@independent_facts({expected})")), + "{ir}" + ); + assert!(!ir.contains("dereferenceable(0)"), "{ir}"); + } + } +} diff --git a/cuda-oxide/crates/mir-importer/src/translator/body.rs b/cuda-oxide/crates/mir-importer/src/translator/body.rs index 1a701eccac..93f0b22cf2 100644 --- a/cuda-oxide/crates/mir-importer/src/translator/body.rs +++ b/cuda-oxide/crates/mir-importer/src/translator/body.rs @@ -2141,16 +2141,33 @@ pub fn translate_body( matches!(signature.skip_binder().abi, rustc_public::ty::Abi::Rust) }); if is_kernel && has_rust_abi { + let fn_abi = instance.fn_abi().map_err(|error| { + input_error_noloc!(TranslationErr::unsupported(format!( + "could not query kernel reference argument ABI: {error:?}" + ))) + })?; + if fn_abi.fixed_count as usize != num_args { + return input_err_noloc!(TranslationErr::invalid_op(format!( + "kernel has {num_args} source arguments but rustc FnAbi declares {}", + fn_abi.fixed_count + ))); + } for arg_idx in 0..num_args { let local = mir::Local::from(arg_idx + 1); let source_ty = &body.locals()[local].ty; - let Some(validity) = facts::reference_param_validity(source_ty) else { + let Some(validity) = facts::reference_param_validity(source_ty, &fn_abi, arg_idx) + .map_err(|error| input_error_noloc!(TranslationErr::invalid_op(error)))? + else { continue; }; mir_func_op.set_reference_param_validity( ctx, arg_idx, - ReferenceParamValidityAttr(validity.pointee_alignment), + ReferenceParamValidityAttr { + pointee_alignment: validity.pointee_alignment, + no_alias: validity.no_alias, + dereferenceable_bytes: validity.dereferenceable_bytes, + }, ); } } diff --git a/cuda-oxide/crates/mir-importer/src/translator/facts.rs b/cuda-oxide/crates/mir-importer/src/translator/facts.rs index 2be0ceafe8..69302b9241 100644 --- a/cuda-oxide/crates/mir-importer/src/translator/facts.rs +++ b/cuda-oxide/crates/mir-importer/src/translator/facts.rs @@ -446,60 +446,93 @@ pub(crate) use pointer_origin::*; /// Validity facts that are safe to expose on a Rust-reference kernel parameter. /// -/// Presence of this fact proves non-nullness. `pointee_alignment` is rustc's -/// ABI alignment for the pointee represented by the physical data pointer. -/// No aliasing, lifetime, or dereferenceability guarantee is encoded here. +/// Presence proves non-nullness. Alignment comes from the pointee layout; +/// aliasing and dereferenceability come only from rustc's argument ABI. #[derive(Clone, Copy, Debug, PartialEq, Eq)] pub(crate) struct ReferenceParamValidity { pub(crate) pointee_alignment: u64, + pub(crate) no_alias: bool, + pub(crate) dereferenceable_bytes: u64, } -/// Return validity facts for the currently audited Rust-reference kernel scope. +/// Return typed rustc facts for a source Rust-ABI kernel reference argument. /// -/// The only accepted evidence is the typed `rustc_public` type/layout API: +/// Match the source index before physical ABI flattening: ignored ZST arguments +/// still occupy an `FnAbi` position. A sized reference uses its direct pointer +/// attributes; a slice uses the first member of its scalar pair. In particular, +/// an indirect by-value aggregate's carrier attributes never describe a Rust +/// reference or its backing allocation. /// -/// * `RigidTy::Ref` proves that the source value is a Rust reference and thus -/// that its data pointer is non-null. -/// * a sized pointee's `Ty::layout()` supplies its ABI alignment; -/// * for `&[T]` / `&mut [T]`, the physical data pointer points at `T`, so the -/// element layout supplies the alignment. -/// -/// Other DSTs, raw pointers, ADTs such as `DisjointSlice`, and any type whose -/// required typed layout is unavailable fail closed with no fact. +/// `RigidTy::Ref` proves non-nullness. Sized pointee layouts and slice element +/// layouts supply alignment. Other DSTs, raw pointers, and by-value ADTs such as +/// `DisjointSlice` produce no fact. Unsupported reference pass modes and source +/// argument mismatches are errors rather than guesses about the ABI mapping. pub(crate) fn reference_param_validity( ty: &rustc_public::ty::Ty, -) -> Option { + fn_abi: &rustc_public::abi::FnAbi, + source_index: usize, +) -> Result, String> { + use rustc_public::abi::PassMode; use rustc_public::ty::{RigidTy, TyKind}; - // Reuse #1186's typed pointer-origin oracle rather than independently - // recognizing reference provenance here. Raw pointers are therefore - // rejected by the same rustc_public-derived classification that mints - // concrete MIR pointer kinds. - let (pointee, origin) = pointer_origin_of_ty(ty)?; + let argument = fn_abi.args.get(source_index).ok_or_else(|| { + format!("kernel source argument index {source_index} is absent from rustc FnAbi") + })?; + if argument.ty != *ty { + return Err(format!( + "kernel source argument {source_index} type does not match rustc FnAbi" + )); + } + + // Reuse the typed pointer-origin oracle so reference provenance is never + // independently inferred from an ABI attribute or pointer-shaped carrier. + let Some((pointee, origin)) = pointer_origin_of_ty(ty) else { + return Ok(None); + }; if !origin.is_reference() { - return None; + return Ok(None); } - let alignment = match pointee.kind() { - TyKind::RigidTy(RigidTy::Slice(element)) => { - let shape = element.layout().ok()?.shape(); - if !shape.is_sized() { - return None; - } - shape.abi_align - } + let (layout_ty, is_slice) = match pointee.kind() { + TyKind::RigidTy(RigidTy::Slice(element)) => (element, true), + _ => (pointee, false), + }; + let Ok(layout) = layout_ty.layout() else { + return Ok(None); + }; + let shape = layout.shape(); + if !shape.is_sized() || shape.abi_align == 0 || !shape.abi_align.is_power_of_two() { + return Ok(None); + } + + let attrs = match (&argument.mode, is_slice) { + (PassMode::Direct(attrs), false) | (PassMode::Pair(attrs, _), true) => attrs, _ => { - let shape = pointee.layout().ok()?.shape(); - if !shape.is_sized() { - return None; - } - shape.abi_align + return Err(format!( + "kernel reference argument {source_index} has an unsupported rustc FnAbi pass mode" + )); } }; - - (alignment != 0 && alignment.is_power_of_two()).then_some(ReferenceParamValidity { - pointee_alignment: alignment, - }) + let flags = attrs.regular(); + // Do not consume rustc's `read_only`: NVPTX (including libNVVM) can lower + // volatile reads through a readonly + noalias kernel parameter to cached + // non-volatile loads. Let the optimizer infer read-only access from the + // body instead; volatile accesses prevent that inference. + // + // LLVM before 23 (including legacy libNVVM consumers) treats parameter + // dereferenceability as lasting for the whole call. rustc's NoFree is the + // proof needed by all supported consumers; producer LLVM version alone is + // insufficient. Zero covers both ZST pointees and unknown slice lengths. + let dereferenceable_bytes = if flags.no_free { + attrs.pointee_size().bytes() as u64 + } else { + 0 + }; + Ok(Some(ReferenceParamValidity { + pointee_alignment: shape.abi_align, + no_alias: flags.no_alias, + dereferenceable_bytes, + })) } // ============================================================================ @@ -880,8 +913,9 @@ pub fn kernel_caller(value: &u32) { another_kernel(value); } ); } - #[test] - fn reference_param_validity_uses_only_typed_rustc_public_facts() { + fn reference_validity_fixture( + opt_level: u8, + ) -> (Vec>, Vec) { let unique = std::time::SystemTime::now() .duration_since(std::time::UNIX_EPOCH) .expect("system time before unix epoch") @@ -896,6 +930,9 @@ pub fn kernel_caller(value: &u32) { another_kernel(value); } std::fs::write( &fixture, r#" +use core::cell::UnsafeCell; +use core::marker::PhantomPinned; + #[repr(align(16))] pub struct AlignedZst; @@ -903,7 +940,18 @@ pub struct ByValue { pub pointer: *const f32, } +pub struct LargeByValue { + pub values: [u64; 8], +} + +pub struct Pinned { + pub value: u32, + pub marker: PhantomPinned, +} + +#[track_caller] pub fn reference_validity( + _ignored: AlignedZst, _shared: &f32, _unique: &mut f32, _slice: &[f32], @@ -912,6 +960,13 @@ pub fn reference_validity( _zst: &AlignedZst, _raw: *const f32, _by_value: ByValue, + _cell: &UnsafeCell, + _mutable_cell: &mut UnsafeCell, + _pinned: &mut Pinned, + _pinned_shared: &Pinned, + _cell_slice: &[UnsafeCell], + _pinned_slice: &mut [Pinned], + _indirect: LargeByValue, ) {} "#, ) @@ -935,30 +990,57 @@ pub fn reference_validity( "--crate-name=reference_validity_fixture".to_string(), "--emit=metadata".to_string(), "-Zmir-opt-level=0".to_string(), + format!("-Copt-level={opt_level}"), format!("--out-dir={}", root.display()), format!("--sysroot={sysroot}"), fixture.display().to_string(), ]; - let facts = std::thread::Builder::new() + let result = std::thread::Builder::new() .stack_size(64 * 1024 * 1024) .spawn(move || { rustc_public::run!(&args, || { - use rustc_public::CrateDef; + use rustc_public::abi::PassMode; + use rustc_public::mir::mono; - let body = rustc_public::all_local_items() + let item = rustc_public::all_local_items() .into_iter() .find(|item| item.name().ends_with("::reference_validity")) - .and_then(|item| item.body()) - .expect("fixture function body"); - let facts = body - .locals() + .expect("fixture function"); + let instance = mono::Instance::try_from(item).expect("fixture instance"); + let body = instance.body().expect("fixture function body"); + let fn_abi = instance.fn_abi().expect("fixture function ABI"); + let source_types = body.locals()[1..17] + .iter() + .map(|decl| decl.ty) + .collect::>(); + let facts = source_types .iter() - .skip(1) - .take(8) - .map(|decl| reference_param_validity(&decl.ty)) + .enumerate() + .map(|(index, ty)| { + reference_param_validity(ty, &fn_abi, index) + .expect("source parameter matches ABI") + }) .collect::>(); - std::ops::ControlFlow::<(), _>::Continue(facts) + + // Reusing a neighboring argument after dropping an ignored + // ZST or flattening a slice must not mint reference facts. + let mut errors = vec![ + reference_param_validity(&source_types[1], &fn_abi, 0).unwrap_err(), + reference_param_validity(&source_types[1], &fn_abi, fn_abi.args.len()) + .unwrap_err(), + ]; + let mut wrong_mode = fn_abi.clone(); + wrong_mode.args[1].mode = fn_abi.args[15].mode.clone(); + assert!(matches!(wrong_mode.args[1].mode, PassMode::Indirect { .. })); + errors.push( + reference_param_validity(&source_types[1], &wrong_mode, 1).unwrap_err(), + ); + wrong_mode.args[3].mode = fn_abi.args[1].mode.clone(); + errors.push( + reference_param_validity(&source_types[3], &wrong_mode, 3).unwrap_err(), + ); + std::ops::ControlFlow::<(), _>::Continue((facts, errors)) }) }) .unwrap() @@ -967,31 +1049,59 @@ pub fn reference_validity( .expect("in-process fixture compilation succeeds"); std::fs::remove_dir_all(&root).ok(); + result + } + #[test] + fn reference_param_validity_consumes_typed_rustc_abi_facts() { + let (facts, errors) = reference_validity_fixture(3); + let observed = facts + .into_iter() + .map(|fact| { + fact.map(|fact| { + ( + fact.pointee_alignment, + fact.no_alias, + fact.dereferenceable_bytes, + ) + }) + }) + .collect::>(); assert_eq!( - facts, + observed, vec![ - Some(ReferenceParamValidity { - pointee_alignment: 4, - }), - Some(ReferenceParamValidity { - pointee_alignment: 4, - }), - Some(ReferenceParamValidity { - pointee_alignment: 4, - }), - Some(ReferenceParamValidity { - pointee_alignment: 4, - }), - Some(ReferenceParamValidity { - pointee_alignment: 1, - }), - Some(ReferenceParamValidity { - pointee_alignment: 16, - }), - None, - None, + None, // ignored by-value ZST + Some((4, true, 4)), // &f32 + Some((4, true, 4)), // &mut f32 + Some((4, true, 0)), // &[f32]: unknown length + Some((4, true, 0)), // &mut [f32] + Some((1, true, 1)), // &u8 + Some((16, true, 0)), // aligned reference to a ZST + None, // raw pointer + None, // direct by-value ADT + Some((4, false, 0)), // &UnsafeCell + Some((4, true, 4)), // &mut UnsafeCell + Some((4, false, 0)), // &mut !Unpin + Some((4, true, 4)), // shared frozen !Unpin + Some((4, false, 0)), // &[UnsafeCell] + Some((4, false, 0)), // &mut [!Unpin] + None, // indirect by-value carrier ] ); + assert!(errors[0].contains("type"), "{}", errors[0]); + assert!(errors[1].contains("index"), "{}", errors[1]); + assert!(errors[2].contains("pass mode"), "{}", errors[2]); + assert!(errors[3].contains("pass mode"), "{}", errors[3]); + } + + #[test] + fn reference_param_validity_retains_conservative_unoptimized_abi_facts() { + let (facts, _) = reference_validity_fixture(0); + assert_eq!(facts.iter().flatten().count(), 12); + for fact in facts.into_iter().flatten() { + assert!(!fact.no_alias); + assert_eq!(fact.dereferenceable_bytes, 0); + assert!(fact.pointee_alignment.is_power_of_two()); + } } } diff --git a/cuda-oxide/crates/mir-lower/src/lowering.rs b/cuda-oxide/crates/mir-lower/src/lowering.rs index a098ffab88..cd8310a489 100644 --- a/cuda-oxide/crates/mir-lower/src/lowering.rs +++ b/cuda-oxide/crates/mir-lower/src/lowering.rs @@ -419,7 +419,10 @@ fn kernel_reference_param_validities( mir_func: &MirFuncOp, parameter_layout: &[KernelParameterLayout], llvm_func_type: pliron::r#type::TypedHandle, -) -> std::result::Result, anyhow::Error> { +) -> std::result::Result< + Vec<(usize, dialect_mir::attributes::ReferenceParamValidityAttr)>, + anyhow::Error, +> { use pliron::builtin::type_interfaces::FunctionTypeInterface; let llvm_args = llvm_func_type.deref(ctx).arg_types().to_vec(); @@ -442,7 +445,7 @@ fn kernel_reference_param_validities( "kernel source argument {source_index} carries reference validity but its first ABI component is not an LLVM pointer" )); } - result.push((llvm_index, validity.0)); + result.push((llvm_index, validity)); } Ok(result) } @@ -450,14 +453,18 @@ fn kernel_reference_param_validities( fn propagate_kernel_reference_param_validities( ctx: &mut Context, llvm_func: &llvm::FuncOp, - validities: &[(usize, u64)], + validities: &[(usize, dialect_mir::attributes::ReferenceParamValidityAttr)], ) { - for &(index, alignment) in validities { + for &(index, validity) in validities { llvm::set_kernel_reference_param_validity( ctx, llvm_func.get_operation(), index, - llvm::KernelReferenceParamValidityAttr(alignment), + llvm::KernelReferenceParamValidityAttr { + pointee_alignment: validity.pointee_alignment, + no_alias: validity.no_alias, + dereferenceable_bytes: validity.dereferenceable_bytes, + }, ); } } @@ -1599,7 +1606,7 @@ mod reference_param_validity_tests { use super::*; use dialect_mir::{ attributes::ReferenceParamValidityAttr, - types::{MirPointerKind, MirPtrType, MirSliceType}, + types::{MirArrayType, MirPointerKind, MirPtrType, MirSliceType}, }; use pliron::{ builtin::{ @@ -1630,8 +1637,20 @@ mod reference_param_validity_tests { MirPtrType::get_generic_with_kind(&mut ctx, f32_ty, false, MirPointerKind::RawConst) .into(); - let mir_func_type = - FunctionType::get(&ctx, vec![shared_ref, shared_slice, raw_ptr], vec![]); + let zst: TypeHandle = MirArrayType::get(&mut ctx, f32_ty, 0).into(); + let mir_func_type = FunctionType::get( + &ctx, + vec![ + shared_ref, + zst, + shared_slice, + raw_ptr, + zst, + shared_slice, + shared_ref, + ], + vec![], + ); let op = Operation::new( &mut ctx, MirFuncOp::get_concrete_op_info(), @@ -1641,8 +1660,19 @@ mod reference_param_validity_tests { 1, ); let mir_func = MirFuncOp::new(&mut ctx, op, TypeAttr::new(mir_func_type.into())); - mir_func.set_reference_param_validity(&mut ctx, 0, ReferenceParamValidityAttr(4)); - mir_func.set_reference_param_validity(&mut ctx, 1, ReferenceParamValidityAttr(4)); + let sized = ReferenceParamValidityAttr { + pointee_alignment: 4, + no_alias: true, + dereferenceable_bytes: 4, + }; + let slice = ReferenceParamValidityAttr { + dereferenceable_bytes: 0, + ..sized + }; + let aliasable = ReferenceParamValidityAttr::new(4); + for (source_index, validity) in [(0, sized), (2, slice), (5, aliasable), (6, sized)] { + mir_func.set_reference_param_validity(&mut ctx, source_index, validity); + } let llvm_func_type = convert_function_type(&mut ctx, mir_func_type, true).expect("kernel type lowers"); @@ -1651,14 +1681,52 @@ mod reference_param_validity_tests { kernel_reference_param_validities(&mut ctx, &mir_func, &layout, llvm_func_type) .expect("reference validity mapping succeeds"); - assert_eq!(mapped, vec![(0, 4), (1, 4)]); + assert_eq!( + mapped, + vec![(0, sized), (1, slice), (4, aliasable), (6, sized)] + ); + let llvm_func = + llvm::FuncOp::new(&mut ctx, "transport".try_into().unwrap(), llvm_func_type); + propagate_kernel_reference_param_validities(&mut ctx, &llvm_func, &mapped); + let entries = + llvm::kernel_reference_param_validity_entries(&ctx, llvm_func.get_operation()).unwrap(); + assert_eq!( + entries, + vec![ + ( + 0, + llvm::KernelReferenceParamValidityAttr { + pointee_alignment: 4, + no_alias: true, + dereferenceable_bytes: 4, + } + ), + ( + 1, + llvm::KernelReferenceParamValidityAttr { + pointee_alignment: 4, + no_alias: true, + dereferenceable_bytes: 0, + } + ), + (4, llvm::KernelReferenceParamValidityAttr::new(4)), + ( + 6, + llvm::KernelReferenceParamValidityAttr { + pointee_alignment: 4, + no_alias: true, + dereferenceable_bytes: 4, + } + ), + ] + ); use pliron::builtin::type_interfaces::FunctionTypeInterface; let llvm_args = llvm_func_type.deref(&ctx).arg_types().to_vec(); assert_eq!( llvm_args.len(), - 4, - "slice lowers to data pointer + length while raw pointer stays one parameter" + 7, + "slices expand, ZSTs disappear, and raw pointers stay bare" ); assert!( llvm_args[0] @@ -1896,9 +1964,9 @@ mod grid_constant_abi_tests { 1, ); let function = MirFuncOp::new(&mut ctx, op, TypeAttr::new(function_type.into())); - function.set_reference_param_validity(&mut ctx, 0, ReferenceParamValidityAttr(1)); - function.set_reference_param_validity(&mut ctx, 2, ReferenceParamValidityAttr(64)); - function.set_reference_param_validity(&mut ctx, 4, ReferenceParamValidityAttr(64)); + function.set_reference_param_validity(&mut ctx, 0, ReferenceParamValidityAttr::new(1)); + function.set_reference_param_validity(&mut ctx, 2, ReferenceParamValidityAttr::new(64)); + function.set_reference_param_validity(&mut ctx, 4, ReferenceParamValidityAttr::new(64)); record(&mut ctx, op, 2, payload, 64); record(&mut ctx, op, 4, payload, 64); let llvm_type = convert_function_type(&mut ctx, function_type, true).unwrap(); @@ -1920,7 +1988,14 @@ mod grid_constant_abi_tests { assert_eq!(llvm_type_size_align(&ctx, grid[0].pointee).unwrap().0, 128); let validities = kernel_reference_param_validities(&mut ctx, &function, &layout, llvm_type).unwrap(); - assert_eq!(validities, vec![(0, 1), (2, 64), (4, 64)]); + assert_eq!( + validities, + vec![ + (0, ReferenceParamValidityAttr::new(1)), + (2, ReferenceParamValidityAttr::new(64)), + (4, ReferenceParamValidityAttr::new(64)), + ] + ); // A stale pointee annotation cannot silently change the launch packet size. record(&mut ctx, op, 2, byte, 64); diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/aligned_field_loads/verify-code-shape.sh b/cuda-oxide/crates/rustc-codegen-cuda/examples/aligned_field_loads/verify-code-shape.sh index a4724bdfa4..07569156a9 100755 --- a/cuda-oxide/crates/rustc-codegen-cuda/examples/aligned_field_loads/verify-code-shape.sh +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/aligned_field_loads/verify-code-shape.sh @@ -5,7 +5,7 @@ # The runtime SUCCESS gate only proves the kernels are bit-correct, which # holds with or without the optimization this example exists to demonstrate. # The perf property lives in the PTX shape: an #[repr(C, align(8))] pair must -# fuse its two adjacent f32 field reads into one ld.global.v2.b32, while the +# fuse its two adjacent f32 field reads into one vector load, while the # natural-align-4 control must keep two scalar loads. Assert both directions # so a silent regression of either the fusion or the control cannot pass. @@ -57,8 +57,10 @@ reject_entry_shape() { fi } -vector_load='ld\.global\.v2\.b32' -scalar_load='ld\.global\.b32' +# Read-only, non-aliasing inputs may use the non-coherent read-only cache +# (`ld.global.nc`). That changes the cache path, not the load width or alignment. +vector_load='ld\.global(\.nc)?\.v2\.b32' +scalar_load='ld\.global(\.nc)?\.b32' # The align(8) kernels must fuse the pair into a single wide load. The `lanes` # kernels reach the same two f32 through an array index rather than named diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/atomic_f16/src/bin/bench.rs b/cuda-oxide/crates/rustc-codegen-cuda/examples/atomic_f16/src/bin/bench.rs index ba4ca416b9..ba9f8d35c8 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/examples/atomic_f16/src/bin/bench.rs +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/atomic_f16/src/bin/bench.rs @@ -27,37 +27,41 @@ const BINS: &[u32] = &[1, 4, 16, 64, 256, 1024, 4096]; mod kernels { use super::*; + // Atomic storage arrives as a raw writable pointer: a shared &[f16] or + // &[f32] would promise immutable elements even after a cast to an atomic. + // The unsafe launch keeps storage alive and excludes non-atomic access. + #[kernel] - pub fn f32_add(hist: &[f32], n: u32, nbins: u32) { + pub fn f32_add(hist: *mut f32, n: u32, nbins: u32) { let i = thread::index_1d().get(); if i >= n as usize { return; } let bin = i % nbins as usize; - let cell = unsafe { &*(hist.as_ptr().add(bin) as *const DeviceAtomicF32) }; + let cell = unsafe { DeviceAtomicF32::from_ptr(hist.add(bin)) }; cell.fetch_add(1.0f32, AtomicOrdering::Relaxed); } #[kernel] - pub fn f16_add(hist: &[f16], n: u32, nbins: u32) { + pub fn f16_add(hist: *mut f16, n: u32, nbins: u32) { let i = thread::index_1d().get(); if i >= n as usize { return; } let bin = i % nbins as usize; - let cell = unsafe { &*(hist.as_ptr().add(bin) as *const DeviceAtomicF16) }; + let cell = unsafe { DeviceAtomicF16::from_ptr(hist.add(bin)) }; cell.fetch_add(1.0f16, AtomicOrdering::Relaxed); } #[kernel] - pub fn f32_add_return(hist: &[f32], mut old: DisjointSlice, n: u32, nbins: u32) { + pub fn f32_add_return(hist: *mut f32, mut old: DisjointSlice, n: u32, nbins: u32) { let gid = thread::index_1d(); let i = gid.get(); if i >= n as usize { return; } let bin = i % nbins as usize; - let cell = unsafe { &*(hist.as_ptr().add(bin) as *const DeviceAtomicF32) }; + let cell = unsafe { DeviceAtomicF32::from_ptr(hist.add(bin)) }; let prev = cell.fetch_add(1.0f32, AtomicOrdering::Relaxed); if let Some(out) = old.get_mut(gid) { *out = prev; @@ -65,14 +69,14 @@ mod kernels { } #[kernel] - pub fn f16_add_return(hist: &[f16], mut old: DisjointSlice, n: u32, nbins: u32) { + pub fn f16_add_return(hist: *mut f16, mut old: DisjointSlice, n: u32, nbins: u32) { let gid = thread::index_1d(); let i = gid.get(); if i >= n as usize { return; } let bin = i % nbins as usize; - let cell = unsafe { &*(hist.as_ptr().add(bin) as *const DeviceAtomicF16) }; + let cell = unsafe { DeviceAtomicF16::from_ptr(hist.add(bin)) }; let prev = cell.fetch_add(1.0f16, AtomicOrdering::Relaxed); if let Some(out) = old.get_mut(gid) { *out = prev; @@ -80,6 +84,12 @@ mod kernels { } } +// Request writable-buffer authority without creating a Rust slice over the +// device allocation. Callers keep it alive and synchronize before host reads. +fn atomic_storage(buffer: &mut DeviceBuffer) -> *mut T { + buffer.cu_deviceptr() as *mut T +} + fn main() -> Result<(), Box> { let ctx = CudaContext::new(0)?; let stream = ctx.default_stream(); @@ -101,20 +111,22 @@ fn bench_no_return( stream: &Arc, bins: u32, ) -> Result<(), Box> { - let hist_f32 = DeviceBuffer::::zeroed(stream, bins as usize)?; - let hist_f16 = DeviceBuffer::::zeroed(stream, bins as usize)?; + let mut hist_f32 = DeviceBuffer::::zeroed(stream, bins as usize)?; + let mut hist_f16 = DeviceBuffer::::zeroed(stream, bins as usize)?; let cfg = LaunchConfig::for_num_elems(N); let f32_ms = time_gpu_iters(stream, ITERS, || { reset(&hist_f32, stream)?; - // SAFETY: launch shape/resources match the kernel; buffers cover its accesses. - unsafe { module.f32_add(stream, cfg, &hist_f32, N, bins) }?; + // SAFETY: buffers cover all accesses and stay alive until synchronization. + // Atomic storage has no concurrent non-atomic accesses. + unsafe { module.f32_add(stream, cfg, atomic_storage(&mut hist_f32), N, bins) }?; Ok(()) })?; let f16_ms = time_gpu_iters(stream, ITERS, || { reset(&hist_f16, stream)?; - // SAFETY: launch shape/resources match the kernel; buffers cover its accesses. - unsafe { module.f16_add(stream, cfg, &hist_f16, N, bins) }?; + // SAFETY: buffers cover all accesses and stay alive until synchronization. + // Atomic storage has no concurrent non-atomic accesses. + unsafe { module.f16_add(stream, cfg, atomic_storage(&mut hist_f16), N, bins) }?; Ok(()) })?; @@ -128,22 +140,42 @@ fn bench_return( stream: &Arc, bins: u32, ) -> Result<(), Box> { - let hist_f32 = DeviceBuffer::::zeroed(stream, bins as usize)?; - let hist_f16 = DeviceBuffer::::zeroed(stream, bins as usize)?; + let mut hist_f32 = DeviceBuffer::::zeroed(stream, bins as usize)?; + let mut hist_f16 = DeviceBuffer::::zeroed(stream, bins as usize)?; let mut old_f32 = DeviceBuffer::::zeroed(stream, N as usize)?; let mut old_f16 = DeviceBuffer::::zeroed(stream, N as usize)?; let cfg = LaunchConfig::for_num_elems(N); let f32_ms = time_gpu_iters(stream, ITERS, || { reset(&hist_f32, stream)?; - // SAFETY: launch shape/resources match the kernel; buffers cover its accesses. - unsafe { module.f32_add_return(stream, cfg, &hist_f32, &mut old_f32, N, bins) }?; + // SAFETY: buffers cover all accesses and stay alive until synchronization. + // Atomic storage has no concurrent non-atomic accesses. + unsafe { + module.f32_add_return( + stream, + cfg, + atomic_storage(&mut hist_f32), + &mut old_f32, + N, + bins, + ) + }?; Ok(()) })?; let f16_ms = time_gpu_iters(stream, ITERS, || { reset(&hist_f16, stream)?; - // SAFETY: launch shape/resources match the kernel; buffers cover its accesses. - unsafe { module.f16_add_return(stream, cfg, &hist_f16, &mut old_f16, N, bins) }?; + // SAFETY: buffers cover all accesses and stay alive until synchronization. + // Atomic storage has no concurrent non-atomic accesses. + unsafe { + module.f16_add_return( + stream, + cfg, + atomic_storage(&mut hist_f16), + &mut old_f16, + N, + bins, + ) + }?; Ok(()) })?; diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/atomic_f16/src/main.rs b/cuda-oxide/crates/rustc-codegen-cuda/examples/atomic_f16/src/main.rs index 5e00e7c262..94a77623bb 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/examples/atomic_f16/src/main.rs +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/atomic_f16/src/main.rs @@ -19,8 +19,12 @@ use cuda_host::cuda_module; mod kernels { use super::*; + // Atomic storage arrives as a raw writable pointer: &[f16] would promise + // immutable elements even after a cast to an atomic. + // The unsafe launch keeps storage alive and excludes non-atomic access. + #[kernel] - pub fn device_hist_f16(hist: &[f16], mut old: DisjointSlice, n: u32, nbins: u32) { + pub fn device_hist_f16(hist: *mut f16, mut old: DisjointSlice, n: u32, nbins: u32) { let gid = thread::index_1d(); let i = gid.get(); if i >= n as usize { @@ -28,7 +32,7 @@ mod kernels { } let bin = i % nbins as usize; - let cell = unsafe { &*(hist.as_ptr().add(bin) as *const DeviceAtomicF16) }; + let cell = unsafe { DeviceAtomicF16::from_ptr(hist.add(bin)) }; let prev = cell.fetch_add(1.0f16, AtomicOrdering::Relaxed); if let Some(out_elem) = old.get_mut(gid) { *out_elem = prev; @@ -36,10 +40,10 @@ mod kernels { } #[kernel] - pub fn block_hist_f16(hist: &[f16], mut old: DisjointSlice) { + pub fn block_hist_f16(hist: *mut f16, mut old: DisjointSlice) { let gid = thread::index_1d(); let block = thread::blockIdx_x() as usize; - let cell = unsafe { &*(hist.as_ptr().add(block) as *const BlockAtomicF16) }; + let cell = unsafe { BlockAtomicF16::from_ptr(hist.add(block)) }; let prev = cell.fetch_add(1.0f16, AtomicOrdering::Relaxed); if let Some(out_elem) = old.get_mut(gid) { *out_elem = prev; @@ -47,9 +51,9 @@ mod kernels { } #[kernel] - pub fn device_sub_f16(counter: &[f16], mut old: DisjointSlice) { + pub fn device_sub_f16(counter: *mut f16, mut old: DisjointSlice) { let gid = thread::index_1d(); - let cell = unsafe { &*(counter.as_ptr() as *const DeviceAtomicF16) }; + let cell = unsafe { DeviceAtomicF16::from_ptr(counter) }; let prev = cell.fetch_sub(1.0f16, AtomicOrdering::Relaxed); if let Some(out_elem) = old.get_mut(gid) { *out_elem = prev; @@ -57,9 +61,11 @@ mod kernels { } #[kernel] - pub fn system_add_f16(counter: &[f16], mut old: DisjointSlice) { + pub fn system_add_f16(counter: *mut f16, mut old: DisjointSlice) { let gid = thread::index_1d(); - let cell = unsafe { &*(counter.as_ptr() as *const SystemAtomicF16) }; + // SAFETY: the live, aligned device allocation is accessed atomically + // only by this GPU; system scope does not require concurrent CPU access. + let cell = unsafe { SystemAtomicF16::from_ptr(counter) }; let prev = cell.fetch_add(1.0f16, AtomicOrdering::Relaxed); if let Some(out_elem) = old.get_mut(gid) { *out_elem = prev; @@ -67,13 +73,13 @@ mod kernels { } #[kernel] - pub fn device_swap_f16(cell: &[f16], vals: &[f16], mut old: DisjointSlice, n: u32) { + pub fn device_swap_f16(cell: *mut f16, vals: &[f16], mut old: DisjointSlice, n: u32) { let gid = thread::index_1d(); let i = gid.get(); if i >= n as usize { return; } - let cell = unsafe { &*(cell.as_ptr() as *const DeviceAtomicF16) }; + let cell = unsafe { DeviceAtomicF16::from_ptr(cell) }; let prev = cell.swap(vals[i], AtomicOrdering::Relaxed); if let Some(out_elem) = old.get_mut(gid) { *out_elem = prev; @@ -81,13 +87,13 @@ mod kernels { } #[kernel] - pub fn device_store_load_f16(cells: &[f16], mut out: DisjointSlice, n: u32) { + pub fn device_store_load_f16(cells: *mut f16, mut out: DisjointSlice, n: u32) { let gid = thread::index_1d(); let i = gid.get(); if i >= n as usize { return; } - let cell = unsafe { &*(cells.as_ptr().add(i) as *const DeviceAtomicF16) }; + let cell = unsafe { DeviceAtomicF16::from_ptr(cells.add(i)) }; let seen = cell.load(AtomicOrdering::Acquire); cell.store(seen + 1.0f16, AtomicOrdering::Release); if let Some(out_elem) = out.get_mut(gid) { @@ -96,6 +102,12 @@ mod kernels { } } +// Request writable-buffer authority without creating a Rust slice over the +// device allocation. Callers keep it alive and synchronize before host reads. +fn atomic_storage(buffer: &mut DeviceBuffer) -> *mut T { + buffer.cu_deviceptr() as *mut T +} + fn main() { println!("=== DeviceAtomicF16 / BlockAtomicF16 tests ===\n"); @@ -122,14 +134,15 @@ fn run_device_hist( ) { println!("--- DeviceAtomicF16 histogram: n={n}, bins={nbins} ---"); - let hist = DeviceBuffer::::zeroed(stream, nbins as usize).unwrap(); + let mut hist = DeviceBuffer::::zeroed(stream, nbins as usize).unwrap(); let mut old = DeviceBuffer::::zeroed(stream, n as usize).unwrap(); - // SAFETY: launch shape/resources match the kernel; buffers cover its accesses. + // SAFETY: buffers cover all accesses and stay alive until synchronization. + // Atomic storage has no concurrent non-atomic accesses. unsafe { module.device_hist_f16( stream, LaunchConfig::for_num_elems(n), - &hist, + atomic_storage(&mut hist), &mut old, n, nbins, @@ -158,12 +171,20 @@ fn run_device_sub(module: &kernels::LoadedModule, stream: &cuda_core::CudaStream println!("--- DeviceAtomicF16 fetch_sub: n={n} ---"); let initial = vec![n as f16]; - let counter = DeviceBuffer::from_host(stream, &initial).unwrap(); + let mut counter = DeviceBuffer::from_host(stream, &initial).unwrap(); let mut old = DeviceBuffer::::zeroed(stream, n as usize).unwrap(); - // SAFETY: launch shape/resources match the kernel; buffers cover its accesses. - unsafe { module.device_sub_f16(stream, LaunchConfig::for_num_elems(n), &counter, &mut old) } - .expect("Kernel launch failed"); + // SAFETY: buffers cover all accesses and stay alive until synchronization. + // Atomic storage has no concurrent non-atomic accesses. + unsafe { + module.device_sub_f16( + stream, + LaunchConfig::for_num_elems(n), + atomic_storage(&mut counter), + &mut old, + ) + } + .expect("Kernel launch failed"); stream.synchronize().unwrap(); @@ -180,12 +201,20 @@ fn run_device_sub(module: &kernels::LoadedModule, stream: &cuda_core::CudaStream fn run_system_add(module: &kernels::LoadedModule, stream: &cuda_core::CudaStream, n: u32) { println!("--- SystemAtomicF16 fetch_add: n={n} ---"); - let counter = DeviceBuffer::::zeroed(stream, 1).unwrap(); + let mut counter = DeviceBuffer::::zeroed(stream, 1).unwrap(); let mut old = DeviceBuffer::::zeroed(stream, n as usize).unwrap(); - // SAFETY: launch shape/resources match the kernel; buffers cover its accesses. - unsafe { module.system_add_f16(stream, LaunchConfig::for_num_elems(n), &counter, &mut old) } - .expect("Kernel launch failed"); + // SAFETY: buffers cover all accesses and stay alive until synchronization. + // Atomic storage has no concurrent non-atomic accesses. + unsafe { + module.system_add_f16( + stream, + LaunchConfig::for_num_elems(n), + atomic_storage(&mut counter), + &mut old, + ) + } + .expect("Kernel launch failed"); stream.synchronize().unwrap(); @@ -205,17 +234,18 @@ fn run_system_add(module: &kernels::LoadedModule, stream: &cuda_core::CudaStream fn run_device_swap(module: &kernels::LoadedModule, stream: &cuda_core::CudaStream, n: u32) { println!("--- DeviceAtomicF16 swap: n={n} ---"); - let cell = DeviceBuffer::::zeroed(stream, 1).unwrap(); + let mut cell = DeviceBuffer::::zeroed(stream, 1).unwrap(); let host_vals: Vec = (1..=n).map(|v| v as f16).collect(); let vals = DeviceBuffer::from_host(stream, &host_vals).unwrap(); let mut old = DeviceBuffer::::zeroed(stream, n as usize).unwrap(); - // SAFETY: launch shape/resources match the kernel; buffers cover its accesses. + // SAFETY: buffers cover all accesses and stay alive until synchronization. + // Atomic storage has no concurrent non-atomic accesses. unsafe { module.device_swap_f16( stream, LaunchConfig::for_num_elems(n), - &cell, + atomic_storage(&mut cell), &vals, &mut old, n, @@ -249,12 +279,19 @@ fn run_device_store_load(module: &kernels::LoadedModule, stream: &cuda_core::Cud println!("--- DeviceAtomicF16 load/store: n={n} ---"); let host_init: Vec = (0..n).map(|v| v as f16).collect(); - let cells = DeviceBuffer::from_host(stream, &host_init).unwrap(); + let mut cells = DeviceBuffer::from_host(stream, &host_init).unwrap(); let mut out = DeviceBuffer::::zeroed(stream, n as usize).unwrap(); - // SAFETY: launch shape/resources match the kernel; buffers cover its accesses. + // SAFETY: buffers cover all accesses and stay alive until synchronization. + // Atomic storage has no concurrent non-atomic accesses. unsafe { - module.device_store_load_f16(stream, LaunchConfig::for_num_elems(n), &cells, &mut out, n) + module.device_store_load_f16( + stream, + LaunchConfig::for_num_elems(n), + atomic_storage(&mut cells), + &mut out, + n, + ) } .expect("Kernel launch failed"); @@ -289,7 +326,7 @@ fn run_block_hist( "--- BlockAtomicF16 histogram: blocks={blocks}, threads/block={threads_per_block} ---" ); - let hist = DeviceBuffer::::zeroed(stream, blocks as usize).unwrap(); + let mut hist = DeviceBuffer::::zeroed(stream, blocks as usize).unwrap(); let mut old = DeviceBuffer::::zeroed(stream, (blocks * threads_per_block) as usize).unwrap(); let cfg = LaunchConfig { @@ -298,8 +335,10 @@ fn run_block_hist( shared_mem_bytes: 0, }; - // SAFETY: launch shape/resources match the kernel; buffers cover its accesses. - unsafe { module.block_hist_f16(stream, cfg, &hist, &mut old) }.expect("Kernel launch failed"); + // SAFETY: buffers cover all accesses and stay alive until synchronization. + // Atomic storage has no concurrent non-atomic accesses. + unsafe { module.block_hist_f16(stream, cfg, atomic_storage(&mut hist), &mut old) } + .expect("Kernel launch failed"); stream.synchronize().unwrap(); diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/kernel_reference_validity/src/main.rs b/cuda-oxide/crates/rustc-codegen-cuda/examples/kernel_reference_validity/src/main.rs index 32f62c9c19..eed9fcaa43 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/examples/kernel_reference_validity/src/main.rs +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/kernel_reference_validity/src/main.rs @@ -6,8 +6,8 @@ //! End-to-end LLVM-IR coverage for rustc-proven kernel reference validity. //! //! Build with `cargo oxide build kernel_reference_validity`, then run the host -//! binary (or `cargo oxide run kernel_reference_validity`) to inspect the -//! generated `.ll` parameter lines and execute valid slice-launch controls. +//! binary (or `cargo oxide run kernel_reference_validity`) for valid launches. +//! `verify-code-shape.sh` checks frontend attributes and optimized PTX loads. use cuda_device::{DisjointSlice, cuda_module, kernel, launch_contract, thread}; @@ -23,6 +23,19 @@ pub struct ByValue { pub pointer: *const f32, } +#[repr(C, align(16))] +pub struct GridConstants { + pub values: [u32; 4], +} + +/// # Safety +/// `output` must point to writable, aligned `u32` storage outside `constants`. +/// Concurrent invocations must not race on that storage. +#[kernel] +pub unsafe fn grid_constant(#[grid_constant] constants: &GridConstants, output: *mut u32) { + unsafe { *output = constants.values[0] }; +} + #[kernel] pub fn shared_ref(_value: &f32) {} @@ -41,6 +54,18 @@ pub fn align_one(_value: &u8) {} #[kernel] pub fn aligned_zst(_value: &AlignedZst) {} +// Interior mutability and pinning must follow rustc's flags, not `&` spelling. +#[kernel] +pub fn atomic_ref(_value: &core::sync::atomic::AtomicU32) {} + +pub struct PinnedValue { + pub value: u32, + pub pin: core::marker::PhantomPinned, +} + +#[kernel] +pub fn pinned_ref(_value: &mut PinnedValue) {} + #[kernel] pub fn raw_pointer(_value: *const f32) {} @@ -59,96 +84,62 @@ pub fn nvvm_slice(input: &[f32], mut output: DisjointSlice) { } } +/// With a valid shared input, the raw write cannot change the value read. +/// LLVM can reuse the first load once the reference guarantees reach it. +/// +/// # Safety +/// `scratch` and `output` must be writable, aligned `u32` pointers outside +/// `input`. Concurrent invocations must not race on their writes. #[kernel] -pub fn by_value(_value: ByValue) {} - -#[allow(improper_ctypes_definitions)] -#[kernel] -pub extern "C" fn c_reference(_value: &f32) {} - -fn kernel_header<'a>(llvm_ir: &'a str, name: &str) -> Result<&'a str, Box> { - let needle = format!("@{name}("); - llvm_ir - .lines() - .find(|line| line.trim_start().starts_with("define ") && line.contains(&needle)) - .ok_or_else(|| format!("missing LLVM kernel definition `{name}`").into()) -} - -fn require_reference_attrs( - llvm_ir: &str, - name: &str, - alignment: u64, -) -> Result<(), Box> { - let header = kernel_header(llvm_ir, name)?; - let expected = format!("nonnull align {alignment}"); - if !header.contains(&expected) { - return Err(format!( - "kernel `{name}` is missing `{expected}` on its reference parameter:\n{header}" - ) - .into()); +pub unsafe fn read_across_write(input: &u32, scratch: *mut u32, output: *mut u32) { + unsafe { + let before = read_input(input); + *scratch = 13; + *output = before.wrapping_add(read_input(input)); } - Ok(()) } -fn require_nonnull_without_alignment( - llvm_ir: &str, - name: &str, -) -> Result<(), Box> { - let header = kernel_header(llvm_ir, name)?; - if !header.contains("nonnull") || header.contains(" align ") { - return Err(format!( - "kernel `{name}` must carry nonnull without a redundant align-1 attribute:\n{header}" - ) - .into()); - } - Ok(()) +// Preserve two reads through MIR optimization; LLVM sees the reference facts +// on the entry parameter when it later optimizes these calls. +#[inline(never)] +unsafe fn read_input(input: *const u32) -> u32 { + unsafe { *input } } -fn require_single_slice_pointer_fact( - llvm_ir: &str, - name: &str, -) -> Result<(), Box> { - let header = kernel_header(llvm_ir, name)?; - if header.matches("nonnull").count() != 1 || header.matches("align 4").count() != 1 { - return Err(format!( - "kernel `{name}` must annotate only the slice data pointer:\n{header}" - ) - .into()); - } - Ok(()) +/// Even an immutable input must be read twice when both reads are volatile. +/// Explicit `readonly` plus `noalias` made NVPTX lose `volatile` and allowed +/// ptxas to merge the reads. Withholding `readonly` preserves both accesses. +/// +/// # Safety +/// `output` must be writable, aligned, outside `input`, and owned by this thread. +#[kernel] +pub unsafe fn volatile_direct(input: &u32, output: *mut u32) { + let before = unsafe { core::ptr::read_volatile(input) }; + let after = unsafe { core::ptr::read_volatile(input) }; + unsafe { *output = before.wrapping_add(after) }; } -fn require_bare(llvm_ir: &str, name: &str) -> Result<(), Box> { - let header = kernel_header(llvm_ir, name)?; - if header.contains("nonnull") || header.contains(" align ") { - return Err(format!( - "kernel `{name}` unexpectedly acquired Rust-reference validity:\n{header}" - ) - .into()); - } - Ok(()) +/// The same guarantee must survive calls that LLVM later inlines. +/// +/// # Safety +/// `output` must be writable, aligned, outside `input`, and owned by this thread. +#[kernel] +pub unsafe fn volatile_helper(input: &u32, output: *mut u32) { + unsafe { *output = read_volatile_input(input).wrapping_add(read_volatile_input(input)) }; } -fn verify_generated_llvm_ir() -> Result<(), Box> { - let path = - std::path::Path::new(env!("CARGO_MANIFEST_DIR")).join("kernel_reference_validity.ll"); - let llvm_ir = std::fs::read_to_string(&path)?; - - require_reference_attrs(&llvm_ir, "shared_ref", 4)?; - require_reference_attrs(&llvm_ir, "unique_ref", 4)?; - require_single_slice_pointer_fact(&llvm_ir, "shared_slice")?; - require_single_slice_pointer_fact(&llvm_ir, "unique_slice")?; - require_single_slice_pointer_fact(&llvm_ir, "nvvm_slice")?; - require_nonnull_without_alignment(&llvm_ir, "align_one")?; - require_reference_attrs(&llvm_ir, "aligned_zst", 16)?; +#[inline(never)] +fn read_volatile_input(input: &u32) -> u32 { + // SAFETY: the reference is aligned and initialized; reading does not mutate it. + unsafe { core::ptr::read_volatile(input) } +} - require_bare(&llvm_ir, "raw_pointer")?; - require_bare(&llvm_ir, "disjoint_slice")?; - require_bare(&llvm_ir, "by_value")?; - require_bare(&llvm_ir, "c_reference")?; +#[kernel] +pub fn by_value(_value: ByValue) {} - Ok(()) -} +#[allow(improper_ctypes_definitions)] +#[kernel] +pub extern "C" fn c_reference(_value: &f32) {} #[repr(C, align(32))] #[derive(Clone, Copy)] @@ -164,6 +155,16 @@ unsafe impl cuda_core::DeviceCopy for AlignedValue {} mod execution { use super::*; + #[kernel] + #[launch_contract(domain = 1, block = (32, 1, 1))] + pub fn add_inputs(a: &[f32], b: &[f32], mut output: DisjointSlice) { + let index = thread::index_1d(); + let i = index.get(); + if let Some(out) = output.get_mut(index) { + *out = a[i] + b[i]; + } + } + #[kernel] #[launch_contract(domain = 1, block = (1, 1, 1))] pub fn inspect_values(input: &[AlignedValue], mut output: DisjointSlice<[u64; 3]>) { @@ -197,6 +198,19 @@ fn verify_launch_packets() -> Result<(), Box> { let prepared = module.prepare_inspect_values(LaunchConfig1D::new(1, 1, 0))?; let prepared_zst = module.prepare_inspect_zsts(LaunchConfig1D::new(1, 1, 0))?; + // Two read-only parameters may share the same allocation across threads. + // Only the DisjointSlice output requires disjoint writable elements. + let values: Vec = (0..64).map(|i| i as f32).collect(); + let input = DeviceBuffer::from_host(&stream, &values)?; + let mut output = DeviceBuffer::::zeroed(&stream, values.len())?; + let add = module.prepare_add_inputs(LaunchConfig1D::new(2, 32, 0))?; + module.add_inputs(&stream, &add, &input, &input, &mut output)?; + assert_eq!( + output.to_host_vec(&stream)?, + values.iter().map(|value| value * 2.0).collect::>(), + "aliased read-only inputs with disjoint parallel output" + ); + for values in [vec![], vec![AlignedValue { values: [37; 8] }; 2]] { let input = DeviceBuffer::from_host(&stream, &values)?; let mut output = DeviceBuffer::<[u64; 3]>::zeroed(&stream, 1)?; @@ -261,10 +275,9 @@ fn verify_launch_packets() -> Result<(), Box> { } fn main() -> Result<(), Box> { - verify_generated_llvm_ir()?; verify_launch_packets()?; println!( - "SUCCESS: kernel reference validity checked in IR and sync/borrowed/owned async GPU launches" + "SUCCESS: kernel reference validity checked in sync/borrowed/owned async GPU launches" ); Ok(()) } diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/kernel_reference_validity/verify-code-shape.sh b/cuda-oxide/crates/rustc-codegen-cuda/examples/kernel_reference_validity/verify-code-shape.sh new file mode 100644 index 0000000000..9763ca878d --- /dev/null +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/kernel_reference_validity/verify-code-shape.sh @@ -0,0 +1,76 @@ +#!/usr/bin/env bash +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 +set -euo pipefail + +root="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +# The frontend .ll precedes LLVM optimization; .opt.ll may infer extra facts. +python3 - "${1:-${root}/kernel_reference_validity.ll}" \ + "${2:-${root}/kernel_reference_validity.ptx}" <<'PY' +from pathlib import Path +import re +import sys + +ir = Path(sys.argv[1]).read_text() +attribute = re.compile(r"\b(?:nonnull|noalias|readonly)\b|\balign \d+|\bdereferenceable(?:_or_null)?\(\d+\)") + +def check(name, expected): + header = next((line for line in ir.splitlines() + if line.startswith("define ") and f"@{name}(" in line), None) + assert header is not None, f"missing kernel {name}" + # Match balanced parentheses: dereferenceable(N) is inside the arguments. + start = header.index(f"@{name}(") + len(name) + 2 + depth = 1 + end = start + while depth: + depth += (header[end] == "(") - (header[end] == ")") + end += 1 + params = header[start:end - 1].split(",") + actual = [set(attribute.findall(param)) for param in params] + assert actual == expected, f"{name}: expected {expected}, got {actual}\n{header}" + +reference = {"nonnull", "align 4", "noalias"} +bare = set() +check("shared_ref", [reference | {"dereferenceable(4)"}]) +check("unique_ref", [reference | {"dereferenceable(4)"}]) +check("shared_slice", [reference, bare]) +check("unique_slice", [reference, bare]) +check("nvvm_slice", [reference, bare, bare, bare]) +check("align_one", [{"nonnull", "noalias", "dereferenceable(1)"}]) +check("aligned_zst", [{"nonnull", "align 16", "noalias"}]) +check("atomic_ref", [{"nonnull", "align 4"}]) +check("pinned_ref", [{"nonnull", "align 4"}]) +check("grid_constant", [{"nonnull", "align 16", "noalias", "dereferenceable(16)"}, bare]) +check("read_across_write", [reference | {"dereferenceable(4)"}, bare, bare]) +for name in ("volatile_direct", "volatile_helper"): + check(name, [reference | {"dereferenceable(4)"}, bare]) +check("raw_pointer", [bare]) +check("disjoint_slice", [bare, bare]) +check("by_value", [bare]) +check("c_reference", [bare]) +print("PASS: exact frontend reference attributes and bare ABI controls") + +ptx = Path(sys.argv[2]).read_text() + +def global_loads(name): + entry = re.search(r"\.entry\s+" + re.escape(name) + r"\([^;]*?\n\{(.*?)^\}", + ptx, re.MULTILINE | re.DOTALL) + assert entry is not None, f"missing PTX entry {name}" + loads = re.findall(r"^\s*(?:@!?%[\w.$]+\s+)?(ld(?:\.[\w:]+)+)\s", + entry[1], re.MULTILINE) + return [load for load in loads if "global" in load.split(".")] + +# Runtime values alone cannot detect two volatile reads being merged into one. +# Check both the direct route and the route through an inlined helper. +for name in ("volatile_direct", "volatile_helper"): + loads = global_loads(name) + # LLVM and libNVVM may spell the same 32-bit integer load b32, u32, or s32. + assert len(loads) == 2 and all(re.fullmatch(r"ld\.volatile\.global\.[bus]32", load) + for load in loads), f"{name}: {loads}" + +# Keeping noalias still lets LLVM reuse an ordinary read across a disjoint write. +loads = global_loads("read_across_write") +assert len(loads) == 1 and re.fullmatch(r"ld\.global(?:\.nc)?\.[bus]32", loads[0]), \ + f"read_across_write: {loads}" +print("PASS: volatile reads stay separate; ordinary reads still combine") +PY diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/scoped_atomic_load_store/src/main.rs b/cuda-oxide/crates/rustc-codegen-cuda/examples/scoped_atomic_load_store/src/main.rs index a82c664e4e..8674c63f1c 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/examples/scoped_atomic_load_store/src/main.rs +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/scoped_atomic_load_store/src/main.rs @@ -54,10 +54,12 @@ mod kernels { /// `fence.acq_rel.sys` / `llvm.nvvm.membar.sys`, and the fused-fence /// SeqCst forms `fence.sc.gpu; st.release.gpu.b32` and `fence.sc.gpu; /// ld.acquire.gpu.b32`. + // These buffers are atomically mutated. Raw writable pointers avoid the + // immutability promise that &[u32] and &[u64] would give the compiler. #[kernel] pub fn publish_and_observe( - flags: &[u32], - wide: &[u64], + flags: *mut u32, + wide: *mut u64, mut observed: DisjointSlice, mut observed_wide: DisjointSlice, mut fence_ok: DisjointSlice, @@ -70,12 +72,10 @@ mod kernels { // SAFETY: `tid` and `next` are both < N, the buffers have N elements, // and every access to them in this kernel is atomic. - let mine = unsafe { DeviceAtomicU32::from_ptr((flags.as_ptr() as *mut u32).add(tid)) }; - let mine_wide = unsafe { DeviceAtomicU64::from_ptr((wide.as_ptr() as *mut u64).add(tid)) }; - let neighbour = - unsafe { DeviceAtomicU32::from_ptr((flags.as_ptr() as *mut u32).add(next)) }; - let neighbour_wide = - unsafe { DeviceAtomicU64::from_ptr((wide.as_ptr() as *mut u64).add(next)) }; + let mine = unsafe { DeviceAtomicU32::from_ptr(flags.add(tid)) }; + let mine_wide = unsafe { DeviceAtomicU64::from_ptr(wide.add(tid)) }; + let neighbour = unsafe { DeviceAtomicU32::from_ptr(flags.add(next)) }; + let neighbour_wide = unsafe { DeviceAtomicU64::from_ptr(wide.add(next)) }; // Publish. The release store is what an acquire reader synchronises with. mine.store(tid as u32 + 1, AtomicOrdering::Release); @@ -122,20 +122,27 @@ mod kernels { } } +// The unsafe launch owns writable-buffer access until stream completion. +fn atomic_storage(buffer: &mut DeviceBuffer) -> *mut T { + buffer.cu_deviceptr() as *mut T +} + fn main() { println!("=== Scoped atomic load/store and fences ===\n"); let ctx = CudaContext::new(0).expect("Failed to create CUDA context"); let stream = ctx.default_stream(); - let flags = DeviceBuffer::::zeroed(&stream, N).unwrap(); - let wide = DeviceBuffer::::zeroed(&stream, N).unwrap(); + let mut flags = DeviceBuffer::::zeroed(&stream, N).unwrap(); + let mut wide = DeviceBuffer::::zeroed(&stream, N).unwrap(); let mut observed = DeviceBuffer::::zeroed(&stream, N).unwrap(); let mut observed_wide = DeviceBuffer::::zeroed(&stream, N).unwrap(); let mut fence_ok = DeviceBuffer::::zeroed(&stream, N).unwrap(); let module = kernels::load(&ctx).expect("Failed to load embedded CUDA module"); - // SAFETY: one block of N threads, and every buffer has N elements. + // SAFETY: one block of N threads, every buffer has N elements, and buffers + // stay alive through completion. Flags/wide have no non-atomic accesses + // while the kernel runs; output copies synchronize the same stream. unsafe { module.publish_and_observe( &stream, @@ -144,8 +151,8 @@ fn main() { block_dim: (N as u32, 1, 1), shared_mem_bytes: 0, }, - &flags, - &wide, + atomic_storage(&mut flags), + atomic_storage(&mut wide), &mut observed, &mut observed_wide, &mut fence_ok, diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/vectorization/src/main.rs b/cuda-oxide/crates/rustc-codegen-cuda/examples/vectorization/src/main.rs index 3fa9c5e123..6a5cf6115b 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/examples/vectorization/src/main.rs +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/vectorization/src/main.rs @@ -386,12 +386,22 @@ fn kernel_body<'source>(document: &ptx_parse::Document<'source>, name: &str) -> /// (`ld.global.v4.f32` on older llc, `ld.global.v2.b64` on llc-21/22); every /// accepted form is a single 128-bit access. fn find_128bit_global<'a>(body: &'a str, dir: &str) -> Option<&'a str> { - let prefix = format!("{dir}.global.v"); + let prefix = format!("{dir}.global."); body.lines().map(str::trim).find(|t| { - t.starts_with(&prefix) - && [".v4.f32", ".v4.b32", ".v2.f64", ".v2.b64"] - .iter() - .any(|w| t.contains(w)) + let Some(rest) = t.strip_prefix(&prefix) else { + return false; + }; + // Read-only inputs may use ld.global.nc; its cache choice preserves + // the global address space and transaction width checked here. + let rest = if dir == "ld" { + rest.strip_prefix("nc.").unwrap_or(rest) + } else { + rest + }; + matches!( + rest.split_whitespace().next(), + Some("v4.f32" | "v4.b32" | "v2.f64" | "v2.b64") + ) }) }