diff --git a/cuda-oxide/crates/cargo-oxide/README.md b/cuda-oxide/crates/cargo-oxide/README.md index 2ed176b4ed..b4765d62c8 100644 --- a/cuda-oxide/crates/cargo-oxide/README.md +++ b/cuda-oxide/crates/cargo-oxide/README.md @@ -52,6 +52,7 @@ cargo oxide update --force # inside the workspace, run setup via update | Flag | Applies to | Description | |------------------------------|----------------------------------|-------------------------------------------------| +| `--link-ltoir ` | run, sanitize, build, test, pipeline, debug | Add an ordered external LTO-IR file to `--materialize-cubin` (repeatable) | | `--materialize-cubin` | run, sanitize, build, test, pipeline, debug | Finalize and embed target-specific native GPU code during the host build | | `--emit-nvvm-ir` | run, build, pipeline | Generate NVVM IR for libNVVM | | `--arch ` | run, sanitize, build, test, pipeline, emit-ltoir, tilelib, inspect, debug | Target architecture override | @@ -121,9 +122,32 @@ rechecks them inside the backend; use the wrapper flag instead of setting `CUDA_OXIDE_MATERIALIZE_CUBIN` around raw Cargo. That variable is an internal wrapper/backend signal, not a supported user interface: raw Cargo may reuse a previous artifact without running the backend, and a backend invocation without -the wrapper's fingerprint is rejected. Generic `#[cuda_module]` kernels and -`#[device] extern` declarations are rejected because their run-time -multi-artifact linking is not yet represented by this single-cubin path. +the wrapper's fingerprint is rejected. Generic `#[cuda_module]` kernels remain +unsupported by this single-cubin path because they merge PTX bundles at run time. + +For non-generic kernels calling `#[device] extern` functions, pass the definitions +as ordered, repeatable `--link-ltoir ` inputs alongside `--materialize-cubin`. +Inputs must be raw CUDA LTO-IR, not object files or fatbins. Each supplied file +is linked into every materialized device artifact in that Cargo invocation. +Keep these inputs separate from generated artifact paths; an overlap is rejected +before cleanup or device output emission. +Paths are relative to the invocation directory. The wrapper fingerprints their +canonical names and contents, and the backend verifies and owns those bytes +before linking. Changing an input rebuilds the device artifact; a change during +the wrapper/backend handoff is an error. The Rust module is linked first, followed +by the supplied files in command-line order. CUDA reports unresolved symbols, +invalid LTO-IR, and unsupported toolkit/target combinations during the build. +Use compatible CUDA compiler and linker versions when generating the inputs. + +```bash +cargo oxide run --materialize-cubin --arch sm_120 \ + --link-ltoir build/device-libs/scale.ltoir --link-ltoir build/device-libs/add.ltoir +``` + +The [`external_ltoir` example](../rustc-codegen-cuda/examples/external_ltoir/README.md) +includes C++ sources and commands to generate these files and run the ordinary +embedded-artifact loader. Generated binaries are not source dependencies. + Metadata-interop projects keep rejecting this global flag; they can instead request one independently finalized cubin per declared device crate with `artifact-kind = "cubin"`, as described below. Materialization is also a diff --git a/cuda-oxide/crates/cargo-oxide/src/commands/artifacts.rs b/cuda-oxide/crates/cargo-oxide/src/commands/artifacts.rs index 94e2caaa36..37b49a0072 100644 --- a/cuda-oxide/crates/cargo-oxide/src/commands/artifacts.rs +++ b/cuda-oxide/crates/cargo-oxide/src/commands/artifacts.rs @@ -134,12 +134,25 @@ pub(super) fn generated_artifact_paths(project_dir: &Path, package_name: &str) - /// Remove stale generated artifacts (`.ptx`, `.ll`, `.ltoir`, `.cubin`) from a /// previous run so we can verify the build produces fresh output. -pub(super) fn clean_generated_files(example_dir: &Path, example: &str) { - for file in generated_artifact_paths(example_dir, example) { +pub(super) fn clean_generated_files( + example_dir: &Path, + example: &str, + external_inputs: &[PathBuf], +) -> Result<(), String> { + let outputs = generated_artifact_paths(example_dir, example); + // Check the whole deletion set before removing even the first file. + for output in &outputs { + for input in external_inputs { + cuda_artifact_finalizer::check_ltoir_input_output(input, output) + .map_err(|error| error.to_string())?; + } + } + for file in outputs { if file.exists() { let _ = std::fs::remove_file(file); } } + Ok(()) } /// Human-readable label for the selected output format. diff --git a/cuda-oxide/crates/cargo-oxide/src/commands/build_run.rs b/cuda-oxide/crates/cargo-oxide/src/commands/build_run.rs index 27e7640d2b..0b74d36daf 100644 --- a/cuda-oxide/crates/cargo-oxide/src/commands/build_run.rs +++ b/cuda-oxide/crates/cargo-oxide/src/commands/build_run.rs @@ -87,7 +87,10 @@ pub fn codegen_run( std::process::exit(2); } - clean_generated_files(&example_dir, example); + clean_generated_files(&example_dir, example, &ctx.link_ltoir).unwrap_or_else(|error| { + eprintln!("Error: {error}"); + std::process::exit(2); + }); println!("========================================="); println!("RUSTC-CODEGEN-CUDA: {}", example); @@ -256,7 +259,10 @@ pub fn codegen_sanitize( return; } - clean_generated_files(&example_dir, example); + clean_generated_files(&example_dir, example, &ctx.link_ltoir).unwrap_or_else(|error| { + eprintln!("Error: {error}"); + std::process::exit(2); + }); println!("========================================="); println!("RUSTC-CODEGEN-CUDA SANITIZE: {}", example); @@ -381,7 +387,10 @@ pub fn codegen_build( std::process::exit(2); } - clean_generated_files(&example_dir, example); + clean_generated_files(&example_dir, example, &ctx.link_ltoir).unwrap_or_else(|error| { + eprintln!("Error: {error}"); + std::process::exit(2); + }); println!("========================================="); println!("RUSTC-CODEGEN-CUDA BUILD: {}", example); diff --git a/cuda-oxide/crates/cargo-oxide/src/commands/context.rs b/cuda-oxide/crates/cargo-oxide/src/commands/context.rs index 66fa592d43..120449f6d2 100644 --- a/cuda-oxide/crates/cargo-oxide/src/commands/context.rs +++ b/cuda-oxide/crates/cargo-oxide/src/commands/context.rs @@ -39,6 +39,8 @@ pub struct Context { pub is_workspace: bool, /// Project-local cuda-oxide defaults. pub config: OxideConfig, + /// Ordered external LTO-IR inputs from the command line. + pub link_ltoir: Vec, } /// Resolve the workspace root and backend, or exit with a helpful error. @@ -60,6 +62,7 @@ pub fn resolve_context() -> Context { let config = load_oxide_config(&workspace_root); let backend_so = backend::find_or_build_backend(&workspace_root, config.backend.as_deref()); return Context { + link_ltoir: Vec::new(), workspace_root, codegen_crate, examples_dir, @@ -78,6 +81,7 @@ pub fn resolve_context() -> Context { let config = load_oxide_config(&cwd); let backend_so = backend::find_or_build_backend(&cwd, config.backend.as_deref()); return Context { + link_ltoir: Vec::new(), workspace_root: cwd.clone(), codegen_crate: cwd.clone(), examples_dir: cwd.clone(), @@ -110,6 +114,7 @@ pub fn resolve_passive_context() -> Context { let config = load_oxide_config_lenient(&workspace_root); let backend_so = backend::backend_so_candidate(&workspace_root, config.backend.as_deref()); return Context { + link_ltoir: Vec::new(), workspace_root, codegen_crate, examples_dir, @@ -128,6 +133,7 @@ pub fn resolve_passive_context() -> Context { let config = load_oxide_config_lenient(&cwd); let backend_so = backend::backend_so_candidate(&cwd, config.backend.as_deref()); return Context { + link_ltoir: Vec::new(), workspace_root: cwd.clone(), codegen_crate: cwd.clone(), examples_dir: cwd.clone(), diff --git a/cuda-oxide/crates/cargo-oxide/src/commands/fingerprint.rs b/cuda-oxide/crates/cargo-oxide/src/commands/fingerprint.rs index 03a9e4ce24..2a84fc79d1 100644 --- a/cuda-oxide/crates/cargo-oxide/src/commands/fingerprint.rs +++ b/cuda-oxide/crates/cargo-oxide/src/commands/fingerprint.rs @@ -78,6 +78,7 @@ pub(super) fn passthrough_codegen_fingerprint_with_env( // Descriptor identity only accelerates verification; artifact identity is // already represented by the content-derived provenance above. effective_env.remove(MATERIALIZER_HANDSHAKE_ENV); + effective_env.remove(EXTERNAL_LTOIR_ENV); if opts.verbose { effective_env.insert("CUDA_OXIDE_VERBOSE".to_string(), b"1".to_vec()); @@ -95,6 +96,10 @@ pub(super) fn passthrough_codegen_fingerprint_with_env( effective_env.insert("CUDA_OXIDE_EMIT_NVVM_IR".to_string(), b"1".to_vec()); } if let Some(prepared) = &materialization.prepared { + effective_env.insert( + EXTERNAL_LTOIR_ENV.to_string(), + prepared.external_ltoir_json.as_bytes().to_vec(), + ); effective_env.insert(MATERIALIZE_ENV.to_string(), b"1".to_vec()); effective_env.insert( EXPECTED_PROVENANCE_ENV.to_string(), diff --git a/cuda-oxide/crates/cargo-oxide/src/commands/interop.rs b/cuda-oxide/crates/cargo-oxide/src/commands/interop.rs index 13d8c532ec..4ea4603dbd 100644 --- a/cuda-oxide/crates/cargo-oxide/src/commands/interop.rs +++ b/cuda-oxide/crates/cargo-oxide/src/commands/interop.rs @@ -573,7 +573,10 @@ fn build_interop_device_crate( } let artifact_name = interop_device_artifact_name(&manifest_path, device_crate); - clean_generated_files(&artifact_dir, &artifact_name); + clean_generated_files(&artifact_dir, &artifact_name, &ctx.link_ltoir).unwrap_or_else(|error| { + eprintln!("Error: {error}"); + std::process::exit(2); + }); if let Some(target) = &binary_target { touch_source_file(&target.source_path); } else { diff --git a/cuda-oxide/crates/cargo-oxide/src/commands/materialize.rs b/cuda-oxide/crates/cargo-oxide/src/commands/materialize.rs index 80dee3ee40..6b3b5beb56 100644 --- a/cuda-oxide/crates/cargo-oxide/src/commands/materialize.rs +++ b/cuda-oxide/crates/cargo-oxide/src/commands/materialize.rs @@ -9,6 +9,7 @@ use std::process::Command; use super::*; +pub(super) const EXTERNAL_LTOIR_ENV: &str = reserved_oxide_symbols::EXTERNAL_LTOIR_ENV; pub(super) const MATERIALIZE_ENV: &str = reserved_oxide_symbols::MATERIALIZE_CUBIN_ENV; pub(super) const EXPECTED_PROVENANCE_ENV: &str = reserved_oxide_symbols::MATERIALIZER_PROVENANCE_ENV; @@ -31,6 +32,7 @@ pub(super) struct MaterializationMode { pub(super) struct PreparedMaterialization { pub(super) provenance_sha256_hex: String, pub(super) tool_identity_handshake_json: String, + pub(super) external_ltoir_json: String, } impl MaterializationMode { @@ -39,10 +41,13 @@ impl MaterializationMode { } pub(super) fn apply(&self, cmd: &mut Command) { + // Never honor an inherited internal external-input list. + cmd.env_remove(EXTERNAL_LTOIR_ENV); if let Some(prepared) = &self.prepared { // These override inherited/project values: they are a single // wrapper-generated handshake tied to this Cargo invocation. - cmd.env(MATERIALIZE_ENV, "1") + cmd.env(EXTERNAL_LTOIR_ENV, &prepared.external_ltoir_json) + .env(MATERIALIZE_ENV, "1") .env(EXPECTED_PROVENANCE_ENV, &prepared.provenance_sha256_hex) .env( MATERIALIZER_HANDSHAKE_ENV, @@ -205,6 +210,14 @@ fn prepare_materialization_result_with_env( .parse() .map_err(|error| format!("invalid materialization target {arch:?}: {error}"))?; + let external_inputs = ctx + .link_ltoir + .iter() + .map(|path| cuda_artifact_finalizer::LtoirFile::capture(path)) + .collect::, _>>() + .map_err(|error| format!("could not capture external LTO-IR: {error}"))?; + let external_ltoir_json = serde_json::to_string(&external_inputs) + .map_err(|error| format!("could not encode external LTO-IR inputs: {error}"))?; let handshake = discover_materializer_handshake(ctx)?; let handshake_json = serde_json::to_string(&handshake) .map_err(|error| format!("could not encode materializer handshake: {error}"))?; @@ -212,6 +225,7 @@ fn prepare_materialization_result_with_env( prepared: Some(PreparedMaterialization { provenance_sha256_hex: digest_hex(&handshake.provenance_sha256), tool_identity_handshake_json: handshake_json, + external_ltoir_json, }), }) } diff --git a/cuda-oxide/crates/cargo-oxide/src/commands/pipeline_debug.rs b/cuda-oxide/crates/cargo-oxide/src/commands/pipeline_debug.rs index f84a54f918..43543c2436 100644 --- a/cuda-oxide/crates/cargo-oxide/src/commands/pipeline_debug.rs +++ b/cuda-oxide/crates/cargo-oxide/src/commands/pipeline_debug.rs @@ -41,7 +41,10 @@ pub fn codegen_show_pipeline( reject_interop_output_mode(emit_nvvm_ir, &materialization); } - clean_generated_files(&example_dir, example); + clean_generated_files(&example_dir, example, &ctx.link_ltoir).unwrap_or_else(|error| { + eprintln!("Error: {error}"); + std::process::exit(2); + }); println!("========================================="); println!("RUSTC-CODEGEN-CUDA PIPELINE: {}", example); @@ -196,7 +199,10 @@ pub fn codegen_debug( println!("Detected GPU arch: {dev} (via nvidia-smi)"); } - clean_generated_files(&example_dir, example); + clean_generated_files(&example_dir, example, &ctx.link_ltoir).unwrap_or_else(|error| { + eprintln!("Error: {error}"); + std::process::exit(2); + }); touch_main_rs(&example_dir); diff --git a/cuda-oxide/crates/cargo-oxide/src/commands/tests.rs b/cuda-oxide/crates/cargo-oxide/src/commands/tests.rs index bf0551deb2..b5e74e9f8c 100644 --- a/cuda-oxide/crates/cargo-oxide/src/commands/tests.rs +++ b/cuda-oxide/crates/cargo-oxide/src/commands/tests.rs @@ -142,6 +142,7 @@ fn cargo_artifact_freshness( fn test_context(config: OxideConfig) -> Context { Context { + link_ltoir: Vec::new(), workspace_root: PathBuf::from("/tmp/cargo-oxide-test-workspace"), codegen_crate: PathBuf::from("/tmp/cargo-oxide-test-codegen"), examples_dir: PathBuf::from("/tmp/cargo-oxide-test-examples"), @@ -447,7 +448,7 @@ fn generated_file_cleanup_preserves_ltoir_cubin_cache() { std::fs::create_dir_all(cached_cubin.parent().unwrap()).unwrap(); std::fs::write(&cached_cubin, b"persistent cache entry").unwrap(); - clean_generated_files(&root, "my-kernel"); + clean_generated_files(&root, "my-kernel", &[]).unwrap(); for extension in ["ptx", "ll", "ltoir", "cubin", "target"] { assert!(!root.join(format!("my_kernel.{extension}")).exists()); @@ -487,6 +488,7 @@ edition = "2024" std::fs::write(&cached_cubin, b"persistent cache").unwrap(); let ctx = Context { + link_ltoir: Vec::new(), workspace_root: root.clone(), codegen_crate: root.clone(), examples_dir: root.clone(), @@ -540,6 +542,7 @@ edition = "2024" std::fs::write(example_dir.join("demo.ptx"), b"generated").unwrap(); let ctx = Context { + link_ltoir: Vec::new(), workspace_root: root.clone(), codegen_crate: codegen_crate.clone(), examples_dir, @@ -586,6 +589,7 @@ edition = "2024" symlink(&external, root.join("target")).unwrap(); let ctx = Context { + link_ltoir: Vec::new(), workspace_root: root.clone(), codegen_crate: root.clone(), examples_dir: root.clone(), @@ -3097,6 +3101,7 @@ device-owner = { path = "../device-owner" } .unwrap(); let ctx = Context { + link_ltoir: Vec::new(), workspace_root: root.clone(), codegen_crate: root.join("unused-codegen-source"), examples_dir: root.join("unused-examples"), @@ -3260,6 +3265,7 @@ shared-dep = { path = "../shared-dep" } fn scratch_workspace_context(root: &Path) -> Context { Context { + link_ltoir: Vec::new(), workspace_root: root.to_path_buf(), codegen_crate: root.join("unused-codegen-source"), examples_dir: root.join("unused-examples"), @@ -3539,6 +3545,7 @@ fn passthrough_fingerprint_tracks_output_affecting_settings() { let materialized = MaterializationMode { prepared: Some(PreparedMaterialization { provenance_sha256_hex: "ab".repeat(32), + external_ltoir_json: "[]".to_string(), tool_identity_handshake_json: "{\"version\":1}".to_string(), }), }; @@ -3556,6 +3563,96 @@ fn passthrough_fingerprint_tracks_output_affecting_settings() { ); } +#[test] +fn cleanup_preserves_all_artifacts_when_a_link_input_overlaps() { + let root = unique_temp_dir("cargo_oxide_input_output_overlap"); + fs::create_dir_all(&root).unwrap(); + let first_output = root.join("my_kernel.ptx"); + let input = root.join("my_kernel.ltoir"); + fs::write(&first_output, b"previous output").unwrap(); + fs::write(&input, b"external input").unwrap(); + + let error = + clean_generated_files(&root, "my-kernel", std::slice::from_ref(&input)).unwrap_err(); + assert!(error.contains("move the external input"), "{error}"); + assert_eq!(fs::read(&input).unwrap(), b"external input"); + assert_eq!(fs::read(&first_output).unwrap(), b"previous output"); + + #[cfg(unix)] + for hard_link in [false, true] { + let alias = root.join("external.ltoir"); + if hard_link { + fs::hard_link(&input, &alias).unwrap(); + } else { + std::os::unix::fs::symlink(&input, &alias).unwrap(); + } + assert!(clean_generated_files(&root, "my-kernel", std::slice::from_ref(&alias)).is_err()); + assert_eq!(fs::read(&alias).unwrap(), b"external input"); + assert_eq!(fs::read(&first_output).unwrap(), b"previous output"); + fs::remove_file(alias).unwrap(); + } + + let separate = root.join("separate.ltoir"); + fs::rename(&input, &separate).unwrap(); + clean_generated_files(&root, "my-kernel", std::slice::from_ref(&separate)).unwrap(); + assert!(!first_output.exists()); + assert_eq!(fs::read(&separate).unwrap(), b"external input"); + fs::remove_dir_all(root).unwrap(); +} + +#[test] +fn external_ltoir_order_and_contents_change_the_codegen_fingerprint() { + let ctx = test_context(OxideConfig::default()); + let opts = CargoPassthroughOptions { + verbose: false, + emit_nvvm_ir: false, + arch: Some("sm_120"), + features: None, + cargo_target_dir: None, + device_codegen_crate: None, + device_cfgs: &[], + no_fmad: false, + unchecked_indexing: false, + materialize_cubin: true, + device_debug: DeviceDebug::Off, + debug_assertions: false, + }; + let root = unique_temp_dir("cargo_oxide_external_ltoir"); + fs::create_dir_all(&root).unwrap(); + let first = root.join("first.ltoir"); + let second = root.join("second.ltoir"); + fs::write(&first, b"first").unwrap(); + fs::write(&second, b"other").unwrap(); + let fingerprint = |paths: &[&Path]| { + let inputs: Vec<_> = paths + .iter() + .map(|path| cuda_artifact_finalizer::LtoirFile::capture(path).unwrap()) + .collect(); + let mode = MaterializationMode { + prepared: Some(PreparedMaterialization { + provenance_sha256_hex: "ab".repeat(32), + tool_identity_handshake_json: "{}".into(), + external_ltoir_json: serde_json::to_string(&inputs).unwrap(), + }), + }; + passthrough_codegen_fingerprint_with_env( + &ctx, + &opts, + None, + Some("sm_120"), + &mode, + &BTreeMap::new(), + ) + }; + let original = fingerprint(&[&first, &second]); + assert_eq!(original, fingerprint(&[&first, &second])); + assert_ne!(original, fingerprint(&[&second, &first])); + fs::write(&first, b"later").unwrap(); + assert_ne!(original, fingerprint(&[&first, &second])); + assert_ne!(original, fingerprint(&[&second])); + fs::remove_dir_all(root).unwrap(); +} + #[test] fn passthrough_fingerprint_tracks_non_unicode_presence_switch_bytes() { let ctx = test_context(OxideConfig::default()); @@ -3770,6 +3867,7 @@ fn materialization_forces_nvvm_ir_and_exact_provenance_handshake() { let materialization = MaterializationMode { prepared: Some(PreparedMaterialization { provenance_sha256_hex: "42".repeat(32), + external_ltoir_json: "[]".to_string(), tool_identity_handshake_json: "{\"version\":1}".to_string(), }), }; @@ -5213,6 +5311,7 @@ fn read_ptx_artifact_reports_missing_file() { #[test] fn nvvm_ir_requested_reads_project_configuration() { let ctx = Context { + link_ltoir: Vec::new(), workspace_root: PathBuf::from("/tmp/project"), codegen_crate: PathBuf::from("/tmp/project"), examples_dir: PathBuf::from("/tmp/project"), @@ -5230,6 +5329,7 @@ fn nvvm_ir_requested_reads_project_configuration() { #[test] fn nvvm_ir_requested_accepts_disabled_project_configuration() { let ctx = Context { + link_ltoir: Vec::new(), workspace_root: PathBuf::from("/tmp/project"), codegen_crate: PathBuf::from("/tmp/project"), examples_dir: PathBuf::from("/tmp/project"), @@ -5247,6 +5347,7 @@ fn nvvm_ir_requested_accepts_disabled_project_configuration() { #[test] fn nvvm_ir_requested_env_disable_overrides_enabled_project_configuration() { let ctx = Context { + link_ltoir: Vec::new(), workspace_root: PathBuf::from("/tmp/project"), codegen_crate: PathBuf::from("/tmp/project"), examples_dir: PathBuf::from("/tmp/project"), @@ -6382,7 +6483,7 @@ fn tilelib_artifacts_are_cleaned_between_builds() { for file in &files { std::fs::write(file, b"stale").unwrap(); } - clean_generated_files(&root, "my-lib"); + clean_generated_files(&root, "my-lib", &[]).unwrap(); for file in &files { assert!(!file.exists(), "{} survived the clean", file.display()); } diff --git a/cuda-oxide/crates/cargo-oxide/src/main.rs b/cuda-oxide/crates/cargo-oxide/src/main.rs index 9bdef87cfa..1d75410a15 100644 --- a/cuda-oxide/crates/cargo-oxide/src/main.rs +++ b/cuda-oxide/crates/cargo-oxide/src/main.rs @@ -75,6 +75,10 @@ enum Commands { DebugPolicy, /// Build and run an example or project Run { + /// Link an external LTO-IR file, in argument order (repeatable). + /// Requires --materialize-cubin; paths are relative to the invocation directory. + #[arg(long)] + link_ltoir: Vec, /// Example name (required in workspace, optional for standalone projects) example: Option, /// Generate NVVM IR (use with libNVVM -gen-lto) @@ -126,6 +130,10 @@ enum Commands { }, /// Build and run an example or project under NVIDIA Compute Sanitizer Sanitize { + /// Link an external LTO-IR file, in argument order (repeatable). + /// Requires --materialize-cubin; paths are relative to the invocation directory. + #[arg(long)] + link_ltoir: Vec, /// Example name (required in workspace, optional for standalone projects) example: Option, /// Compute Sanitizer tool to run @@ -171,6 +179,10 @@ enum Commands { }, /// Build an example or project (compile only, don't run) Build { + /// Link an external LTO-IR file, in argument order (repeatable). + /// Requires --materialize-cubin; paths are relative to the invocation directory. + #[arg(long)] + link_ltoir: Vec, /// Example name (required in workspace, optional for standalone projects) example: Option, /// Generate NVVM IR (use with libNVVM -gen-lto) @@ -265,6 +277,10 @@ enum Commands { }, /// Run Cargo tests through the cuda-oxide backend Test { + /// Link an external LTO-IR file, in argument order (repeatable). + /// Requires --materialize-cubin; paths are relative to the invocation directory. + #[arg(long)] + link_ltoir: Vec, /// Target architecture (e.g., sm_90, sm_100, sm_120) #[arg(long)] arch: Option, @@ -404,6 +420,10 @@ enum Commands { }, /// Show the full compilation pipeline (MIR -> PTX/NVVM IR) with verbose output Pipeline { + /// Link an external LTO-IR file, in argument order (repeatable). + /// Requires --materialize-cubin; paths are relative to the invocation directory. + #[arg(long)] + link_ltoir: Vec, /// Example name (required in workspace, optional for standalone projects) example: Option, /// Generate NVVM IR (use with libNVVM -gen-lto) @@ -434,6 +454,10 @@ enum Commands { }, /// Build with debug info and launch cuda-gdb Debug { + /// Link an external LTO-IR file, in argument order (repeatable). + /// Requires --materialize-cubin; paths are relative to the invocation directory. + #[arg(long)] + link_ltoir: Vec, /// Example name (required in workspace, optional for standalone projects) example: Option, /// Target architecture (e.g., sm_90, sm_100, sm_120). When omitted, @@ -604,7 +628,22 @@ fn build_passthrough_trigger( } } +fn external_ltoir_inputs(command: &Commands) -> &[PathBuf] { + match command { + Commands::Run { link_ltoir, .. } + | Commands::Sanitize { link_ltoir, .. } + | Commands::Build { link_ltoir, .. } + | Commands::Test { link_ltoir, .. } + | Commands::Pipeline { link_ltoir, .. } + | Commands::Debug { link_ltoir, .. } => link_ltoir, + _ => &[], + } +} + fn validate_materialization_cli(cli: &Cli) -> Result<(), String> { + if !external_ltoir_inputs(&cli.command).is_empty() && !cli.materialize_cubin { + return Err("--link-ltoir requires --materialize-cubin".to_string()); + } if !cli.materialize_cubin { return Ok(()); } @@ -731,6 +770,12 @@ fn main() { .exit(); } let materialize_cubin = cli.materialize_cubin; + let link_ltoir = external_ltoir_inputs(&cli.command).to_vec(); + let resolve_context = || { + let mut ctx = commands::resolve_context(); + ctx.link_ltoir = link_ltoir.clone(); + ctx + }; match cli.command { Commands::MaterializerHandshake => { @@ -740,6 +785,7 @@ fn main() { commands::print_debug_policy(); } Commands::Run { + link_ltoir: _, example, emit_nvvm_ir, arch, @@ -753,7 +799,7 @@ fn main() { device_debug, app_args, } => { - let ctx = commands::resolve_context(); + let ctx = resolve_context(); let example = resolve_example_name(example, &ctx, "run"); validate_output_arch( &ctx, @@ -779,6 +825,7 @@ fn main() { ); } Commands::Sanitize { + link_ltoir: _, example, tool, arch, @@ -791,7 +838,7 @@ fn main() { device_debug, sanitizer_args, } => { - let ctx = commands::resolve_context(); + let ctx = resolve_context(); let example = resolve_example_name(example, &ctx, "sanitize"); validate_output_arch(&ctx, &example, false, materialize_cubin, arch.as_deref()); let (sanitizer_args, application_args) = @@ -813,6 +860,7 @@ fn main() { ); } Commands::Build { + link_ltoir: _, example, emit_nvvm_ir, arch, @@ -829,7 +877,7 @@ fn main() { device_cfgs, cargo_args, } => { - let ctx = commands::resolve_context(); + let ctx = resolve_context(); let passthrough = use_build_passthrough( explicit_passthrough, cargo_target_dir.is_some(), @@ -919,7 +967,7 @@ fn main() { keep_going, fail_on_finding, } => { - let ctx = commands::resolve_context(); + let ctx = resolve_context(); let (seed_start, seed_end) = ptx_schedule::campaign::parse_seed_range(&seeds) .unwrap_or_else(|error| { eprintln!("Error: {error}"); @@ -983,6 +1031,7 @@ fn main() { } } Commands::Test { + link_ltoir: _, arch, cargo_target_dir, device_codegen_crate, @@ -994,7 +1043,7 @@ fn main() { device_debug, cargo_args, } => { - let ctx = commands::resolve_context(); + let ctx = resolve_context(); validate_output_arch( &ctx, "cargo test", @@ -1033,7 +1082,7 @@ fn main() { lineinfo, device_debug, } => { - let ctx = commands::resolve_context(); + let ctx = resolve_context(); let example = resolve_example_name(example, &ctx, "emit-ltoir"); commands::emit_ltoir( &ctx, @@ -1060,7 +1109,7 @@ fn main() { lineinfo, device_debug, } => { - let ctx = commands::resolve_context(); + let ctx = resolve_context(); let example = resolve_example_name(example, &ctx, "tilelib"); commands::tilelib( &ctx, @@ -1077,6 +1126,7 @@ fn main() { ); } Commands::Pipeline { + link_ltoir: _, example, emit_nvvm_ir, arch, @@ -1085,7 +1135,7 @@ fn main() { lineinfo, device_debug, } => { - let ctx = commands::resolve_context(); + let ctx = resolve_context(); let example = resolve_example_name(example, &ctx, "pipeline"); validate_output_arch( &ctx, @@ -1115,7 +1165,7 @@ fn main() { lineinfo, device_debug, } => { - let ctx = commands::resolve_context(); + let ctx = resolve_context(); let example = resolve_example_name(example, &ctx, "inspect"); commands::codegen_inspect_ptx( &ctx, @@ -1129,6 +1179,7 @@ fn main() { ); } Commands::Debug { + link_ltoir: _, example, arch, features, @@ -1136,7 +1187,7 @@ fn main() { cgdb, tui, } => { - let ctx = commands::resolve_context(); + let ctx = resolve_context(); let example = resolve_example_name(example, &ctx, "debug"); validate_output_arch(&ctx, &example, false, materialize_cubin, arch.as_deref()); commands::codegen_debug( @@ -1177,7 +1228,7 @@ fn main() { commands::doctor(&ctx); } Commands::Setup => { - let ctx = commands::resolve_context(); + let ctx = resolve_context(); commands::setup(&ctx); } Commands::Update { force } => { @@ -1374,6 +1425,47 @@ mod tests { assert_eq!(device_features.as_deref(), Some("tensor-cores")); } + #[test] + fn external_ltoir_requires_materialization_and_preserves_order() { + let unmaterialized = + Cli::try_parse_from(["cargo-oxide", "build", "--link-ltoir", "one"]).unwrap(); + assert!(validate_materialization_cli(&unmaterialized).is_err()); + let cli = Cli::try_parse_from([ + "cargo-oxide", + "build", + "--link-ltoir", + "one", + "--materialize-cubin", + "--link-ltoir", + "two", + ]) + .unwrap(); + assert_eq!( + external_ltoir_inputs(&cli.command), + [ + std::path::PathBuf::from("one"), + std::path::PathBuf::from("two") + ] + ); + validate_materialization_cli(&cli).unwrap(); + + let run = Cli::try_parse_from([ + "cargo-oxide", + "run", + "--link-ltoir", + "one", + "external_ltoir", + "--materialize-cubin", + ]) + .unwrap(); + validate_materialization_cli(&run).unwrap(); + assert_eq!(external_ltoir_inputs(&run.command), [PathBuf::from("one")]); + let Commands::Run { example, .. } = run.command else { + panic!("expected run command"); + }; + assert_eq!(example.as_deref(), Some("external_ltoir")); + } + #[test] fn materialize_cubin_is_a_global_codegen_option() { let after_subcommand = Cli::try_parse_from([ diff --git a/cuda-oxide/crates/cuda-artifact-finalizer/src/lib.rs b/cuda-oxide/crates/cuda-artifact-finalizer/src/lib.rs index 6a014422de..6e02f0c34b 100644 --- a/cuda-oxide/crates/cuda-artifact-finalizer/src/lib.rs +++ b/cuda-oxide/crates/cuda-artifact-finalizer/src/lib.rs @@ -12,6 +12,7 @@ mod diagnostics; mod link; +mod ltoir_file; mod nvvm; mod options; mod provenance; @@ -21,6 +22,7 @@ mod validation; pub use diagnostics::KernelResourceUsage; pub use libnvvm_sys::{CudaArch, CudaArchParseError, LibdeviceNotFound, NvvmError, find_libdevice}; pub use link::{LinkReport, LtoLinker}; +pub use ltoir_file::{LtoirFile, check_ltoir_input_output}; pub use nvjitlink_sys::NvJitLinkError; pub use nvvm::NvvmCompiler; pub use options::{DebugPolicy, FinalizationOptions, FinalizerOutput, NamedInput}; @@ -37,6 +39,20 @@ use thiserror::Error; /// Failures while compiling or linking CUDA artifacts. #[derive(Debug, Error)] pub enum FinalizerError { + /// A file path cannot be transported losslessly to the CUDA linker. + #[error("external LTO-IR path is not Unicode: {path:?}")] + InvalidLtoirPath { path: PathBuf }, + + /// A link input changed after the wrapper computed Cargo's fingerprint. + #[error("external LTO-IR input changed after Cargo fingerprinting: {path}; rerun cargo oxide")] + LtoirInputChanged { path: PathBuf }, + + /// An output would delete or overwrite an explicitly supplied link input. + #[error( + "external LTO-IR input {input} overlaps generated artifact {output}; move the external input to a separate directory" + )] + LtoirInputOutputOverlap { input: PathBuf, output: PathBuf }, + /// libNVVM failed to load, validate, or compile. #[error("libnvvm: {0}")] Nvvm(#[from] libnvvm_sys::NvvmError), diff --git a/cuda-oxide/crates/cuda-artifact-finalizer/src/ltoir_file.rs b/cuda-oxide/crates/cuda-artifact-finalizer/src/ltoir_file.rs new file mode 100644 index 0000000000..92aff18147 --- /dev/null +++ b/cuda-oxide/crates/cuda-artifact-finalizer/src/ltoir_file.rs @@ -0,0 +1,163 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +//! Content-bound external inputs shared by Cargo fingerprinting and finalization. + +use std::path::{Path, PathBuf}; + +use serde::{Deserialize, Serialize}; +use sha2::{Digest, Sha256}; + +use crate::FinalizerError; + +/// Identity of an external LTO-IR file captured before Cargo decides freshness. +/// +/// The absolute path is also the linker's diagnostic input name. Preserve the +/// order of these records when fingerprinting and linking a list of inputs. +/// [`Self::read_verified`] returns owned bytes so a later file change cannot +/// alter the data submitted to the linker. +#[derive(Clone, Debug, Eq, PartialEq, Serialize, Deserialize)] +#[serde(deny_unknown_fields)] +pub struct LtoirFile { + path: PathBuf, + sha256: [u8; 32], +} + +impl LtoirFile { + /// Capture the canonical path and SHA-256 of a nonempty input file. + /// + /// Input paths must be Unicode so the same name can reach nvJitLink and + /// the compiler process without lossy conversion. CUDA validates the + /// LTO-IR format and target compatibility during finalization. + pub fn capture(path: &Path) -> Result { + let path = path.canonicalize().map_err(|source| FinalizerError::Io { + path: path.to_owned(), + source, + })?; + if path.to_str().is_none() { + return Err(FinalizerError::InvalidLtoirPath { path }); + } + let bytes = read(&path)?; + Ok(Self { + path, + sha256: Sha256::digest(&bytes).into(), + }) + } + + /// Diagnostic name included in the ordered compilation identity. + pub fn name(&self) -> Result<&str, FinalizerError> { + self.path + .to_str() + .ok_or_else(|| FinalizerError::InvalidLtoirPath { + path: self.path.clone(), + }) + } + + /// Read the exact captured content or fail if it changed before codegen. + pub fn read_verified(&self) -> Result, FinalizerError> { + let bytes = read(&self.path)?; + let actual: [u8; 32] = Sha256::digest(&bytes).into(); + if actual != self.sha256 { + return Err(FinalizerError::LtoirInputChanged { + path: self.path.clone(), + }); + } + Ok(bytes) + } +} + +/// Reject a generated output that aliases an existing external link input. +/// +/// Call before deleting or writing outputs. Canonical paths cover symlinks; +/// Unix file identity also covers hard links with different canonical paths. +pub fn check_ltoir_input_output(input: &Path, output: &Path) -> Result<(), FinalizerError> { + let input_path = input.canonicalize().map_err(|source| FinalizerError::Io { + path: input.to_owned(), + source, + })?; + let output_path = match output.canonicalize() { + Ok(path) => path, + Err(error) if error.kind() == std::io::ErrorKind::NotFound => return Ok(()), + Err(source) => { + return Err(FinalizerError::Io { + path: output.to_owned(), + source, + }); + } + }; + if input_path == output_path || same_file(&input_path, &output_path) { + return Err(FinalizerError::LtoirInputOutputOverlap { + input: input.to_owned(), + output: output.to_owned(), + }); + } + Ok(()) +} + +#[cfg(unix)] +fn same_file(a: &Path, b: &Path) -> bool { + use std::os::unix::fs::MetadataExt; + match (std::fs::metadata(a), std::fs::metadata(b)) { + (Ok(a), Ok(b)) => a.dev() == b.dev() && a.ino() == b.ino(), + _ => false, + } +} + +#[cfg(not(unix))] +fn same_file(_a: &Path, _b: &Path) -> bool { + false +} + +fn read(path: &Path) -> Result, FinalizerError> { + let bytes = std::fs::read(path).map_err(|source| FinalizerError::Io { + path: path.to_owned(), + source, + })?; + if bytes.is_empty() { + return Err(FinalizerError::EmptyInput { + name: path.display().to_string(), + }); + } + Ok(bytes) +} + +#[cfg(test)] +mod tests { + use super::*; + + #[test] + fn captured_input_checks_content_and_owns_verified_bytes() { + let path = std::env::temp_dir().join(format!( + "oxide-ltoir-{}-{}.test", + std::process::id(), + std::time::SystemTime::now() + .duration_since(std::time::UNIX_EPOCH) + .unwrap() + .as_nanos() + )); + std::fs::write(&path, b"first").unwrap(); + let identity = LtoirFile::capture(&path).unwrap(); + let json = serde_json::to_string(&identity).unwrap(); + let restored: LtoirFile = serde_json::from_str(&json).unwrap(); + let owned = restored.read_verified().unwrap(); + std::fs::write(&path, b"other").unwrap(); + assert_eq!(owned, b"first"); + assert!(matches!( + restored.read_verified(), + Err(FinalizerError::LtoirInputChanged { .. }) + )); + assert_ne!(identity, LtoirFile::capture(&path).unwrap()); + std::fs::write(&path, []).unwrap(); + assert!(matches!( + LtoirFile::capture(&path), + Err(FinalizerError::EmptyInput { .. }) + )); + std::fs::remove_file(&path).unwrap(); + assert!(matches!( + restored.read_verified(), + Err(FinalizerError::Io { .. }) + )); + } +} diff --git a/cuda-oxide/crates/mir-importer/src/pipeline.rs b/cuda-oxide/crates/mir-importer/src/pipeline.rs index 12d74984ef..16f4534ef8 100644 --- a/cuda-oxide/crates/mir-importer/src/pipeline.rs +++ b/cuda-oxide/crates/mir-importer/src/pipeline.rs @@ -738,7 +738,9 @@ fn write_nvvm_compile_options_sidecar( }) } -fn stale_compilation_artifact_paths( +/// Paths the pipeline may remove or replace when compiling one module. +/// Callers can preflight these before compilation to protect external inputs. +pub fn stale_compilation_artifact_paths( output_dir: &Path, output_name: &str, ) -> Vec { diff --git a/cuda-oxide/crates/reserved-oxide-symbols/src/lib.rs b/cuda-oxide/crates/reserved-oxide-symbols/src/lib.rs index 010bc715d2..36ce23634c 100644 --- a/cuda-oxide/crates/reserved-oxide-symbols/src/lib.rs +++ b/cuda-oxide/crates/reserved-oxide-symbols/src/lib.rs @@ -72,6 +72,9 @@ pub const MATERIALIZER_PROVENANCE_ENV: &str = "CUDA_OXIDE_INTERNAL_MATERIALIZER_ /// code remains keyed by the content-derived provenance digest. pub const MATERIALIZER_HANDSHAKE_ENV: &str = "CUDA_OXIDE_INTERNAL_MATERIALIZER_HANDSHAKE"; +/// Ordered, content-bound external LTO-IR files supplied by cargo-oxide. +pub const EXTERNAL_LTOIR_ENV: &str = "CUDA_OXIDE_INTERNAL_EXTERNAL_LTOIR_V1"; + /// Optional comma-separated filter selecting crates that may own device code. pub const DEVICE_CODEGEN_CRATE_ENV: &str = "CUDA_OXIDE_DEVICE_CODEGEN_CRATE"; diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/Cargo.lock b/cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/Cargo.lock new file mode 100644 index 0000000000..af40b6dd11 --- /dev/null +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/Cargo.lock @@ -0,0 +1,753 @@ +# This file is automatically @generated by Cargo. +# It is not intended for manual editing. +version = 4 + +[[package]] +name = "aho-corasick" +version = "1.1.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ddd31a130427c27518df266943a5308ed92d4b226cc639f5a8f1002816174301" +dependencies = [ + "memchr", +] + +[[package]] +name = "anyhow" +version = "1.0.102" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7f202df86484c868dbad7eaa557ef785d5c66295e41b460ef922eca0723b842c" + +[[package]] +name = "bindgen" +version = "0.69.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "271383c67ccabffb7381723dea0672a673f292304fcb45c01cc648c7a8d58088" +dependencies = [ + "bitflags", + "cexpr", + "clang-sys", + "itertools", + "lazy_static", + "lazycell", + "log", + "prettyplease", + "proc-macro2", + "quote", + "regex", + "rustc-hash", + "shlex", + "syn", + "which", +] + +[[package]] +name = "bitflags" +version = "2.13.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b4388bee8683e3d04af747c73422af53102d2bd24d9eadb6cbc100baef4b43f8" + +[[package]] +name = "block-buffer" +version = "0.10.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3078c7629b62d3f0439517fa394996acacc5cbc91c5a20d8c658e77abd503a71" +dependencies = [ + "generic-array", +] + +[[package]] +name = "cexpr" +version = "0.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6fac387a98bb7c37292057cffc56d62ecb629900026402633ae9160df93a8766" +dependencies = [ + "nom", +] + +[[package]] +name = "cfg-if" +version = "1.0.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9330f8b2ff13f34540b44e946ef35111825727b38d33286ef986142615121801" + +[[package]] +name = "clang-sys" +version = "1.8.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0b023947811758c97c59bf9d1c188fd619ad4718dcaa767947df1cadb14f39f4" +dependencies = [ + "glob", + "libc", + "libloading", +] + +[[package]] +name = "cpufeatures" +version = "0.2.17" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "59ed5838eebb26a2bb2e58f6d5b5316989ae9d08bab10e0e6d103e656d1b0280" +dependencies = [ + "libc", +] + +[[package]] +name = "crc32fast" +version = "1.5.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9481c1c90cbf2ac953f07c8d4a58aa3945c425b7185c9154d67a65e4230da511" +dependencies = [ + "cfg-if", +] + +[[package]] +name = "crunchy" +version = "0.2.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "460fbee9c2c2f33933d720630a6a0bac33ba7053db5344fac858d4b8952d77d5" + +[[package]] +name = "crypto-common" +version = "0.1.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "78c8292055d1c1df0cce5d180393dc8cce0abec0a7102adb6c7b1eef6016d60a" +dependencies = [ + "generic-array", + "typenum", +] + +[[package]] +name = "cuda-artifact-finalizer" +version = "0.2.1" +dependencies = [ + "libnvvm-sys", + "nvjitlink-sys", + "serde", + "sha2", + "thiserror", +] + +[[package]] +name = "cuda-bindings" +version = "0.4.0" +dependencies = [ + "bindgen", + "libloading", + "prettyplease", + "proc-macro2", + "quote", + "syn", +] + +[[package]] +name = "cuda-core" +version = "0.4.0" +dependencies = [ + "anyhow", + "cuda-bindings", + "cuda-core-derive", + "half", + "oxide-artifacts", +] + +[[package]] +name = "cuda-core-derive" +version = "0.4.0" +dependencies = [ + "proc-macro2", + "quote", + "syn", +] + +[[package]] +name = "cuda-device" +version = "0.2.1" +dependencies = [ + "cuda-macros", +] + +[[package]] +name = "cuda-host" +version = "0.2.1" +dependencies = [ + "cuda-artifact-finalizer", + "cuda-core", + "cuda-macros", + "half", + "libc", + "oxide-artifacts", + "ptx-parse", + "sha2", + "thiserror", +] + +[[package]] +name = "cuda-macros" +version = "0.2.1" +dependencies = [ + "oxide-tile-export", + "proc-macro2", + "quote", + "reserved-oxide-symbols", + "syn", +] + +[[package]] +name = "cuda-target-spec" +version = "0.1.0" + +[[package]] +name = "digest" +version = "0.10.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9ed9a281f7bc9b7576e61468ba615a66a5c8cfdff42420a70aa82701a3b1e292" +dependencies = [ + "block-buffer", + "crypto-common", +] + +[[package]] +name = "either" +version = "1.16.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "91622ff5e7162018101f2fea40d6ebf4a78bbe5a49736a2020649edf9693679e" + +[[package]] +name = "equivalent" +version = "1.0.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "877a4ace8713b0bcf2a4e7eec82529c029f1d0619886d18145fea96c3ffe5c0f" + +[[package]] +name = "errno" +version = "0.3.14" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "39cab71617ae0d63f51a36d69f866391735b51691dbda63cf6f96d042b63efeb" +dependencies = [ + "libc", + "windows-sys 0.61.2", +] + +[[package]] +name = "external_ltoir" +version = "0.1.0" +dependencies = [ + "cuda-core", + "cuda-device", + "cuda-host", +] + +[[package]] +name = "foldhash" +version = "0.1.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d9c4f5dac5e15c24eb999c26181a6ca40b39fe946cbe4c263c7209467bc83af2" + +[[package]] +name = "generic-array" +version = "0.14.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "85649ca51fd72272d7821adaf274ad91c288277713d9c18820d8499a7ff69e9a" +dependencies = [ + "typenum", + "version_check", +] + +[[package]] +name = "glob" +version = "0.3.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0cc23270f6e1808e30a928bdc84dea0b9b4136a8bc82338574f23baf47bbd280" + +[[package]] +name = "half" +version = "2.7.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6ea2d84b969582b4b1864a92dc5d27cd2b77b622a8d79306834f1be5ba20d84b" +dependencies = [ + "cfg-if", + "crunchy", + "zerocopy", +] + +[[package]] +name = "hashbrown" +version = "0.15.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9229cfe53dfd69f0609a49f65461bd93001ea1ef889cd5529dd176593f5338a1" +dependencies = [ + "foldhash", +] + +[[package]] +name = "hashbrown" +version = "0.17.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ed5909b6e89a2db4456e54cd5f673791d7eca6732202bbf2a9cc504fe2f9b84a" + +[[package]] +name = "home" +version = "0.5.12" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "cc627f471c528ff0c4a49e1d5e60450c8f6461dd6d10ba9dcd3a61d3dff7728d" +dependencies = [ + "windows-sys 0.61.2", +] + +[[package]] +name = "indexmap" +version = "2.14.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d466e9454f08e4a911e14806c24e16fba1b4c121d1ea474396f396069cf949d9" +dependencies = [ + "equivalent", + "hashbrown 0.17.1", +] + +[[package]] +name = "itertools" +version = "0.12.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ba291022dbbd398a455acf126c1e341954079855bc60dfdda641363bd6922569" +dependencies = [ + "either", +] + +[[package]] +name = "itoa" +version = "1.0.18" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8f42a60cbdf9a97f5d2305f08a87dc4e09308d1276d28c869c684d7777685682" + +[[package]] +name = "lazy_static" +version = "1.5.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bbd2bcb4c963f2ddae06a2efc7e9f3591312473c50c6685e1f298068316e66fe" + +[[package]] +name = "lazycell" +version = "1.3.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "830d08ce1d1d941e6b30645f1a0eb5643013d835ce3779a5fc208261dbe10f55" + +[[package]] +name = "libc" +version = "0.2.186" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "68ab91017fe16c622486840e4c83c9a37afeff978bd239b5293d61ece587de66" + +[[package]] +name = "libloading" +version = "0.8.9" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d7c4b02199fee7c5d21a5ae7d8cfa79a6ef5bb2fc834d6e9058e89c825efdc55" +dependencies = [ + "cfg-if", + "windows-link", +] + +[[package]] +name = "libnvvm-sys" +version = "0.2.1" +dependencies = [ + "cuda-target-spec", + "libloading", + "thiserror", +] + +[[package]] +name = "linux-raw-sys" +version = "0.4.15" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d26c52dbd32dccf2d10cac7725f8eae5296885fb5703b261f7d0a0739ec807ab" + +[[package]] +name = "log" +version = "0.4.29" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5e5032e24019045c762d3c0f28f5b6b8bbf38563a65908389bf7978758920897" + +[[package]] +name = "memchr" +version = "2.8.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "cf8baf1c55e62ffcace7a9f06f4bd9cd3f0c4beb022d3b367256b91b87513d98" + +[[package]] +name = "minimal-lexical" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "68354c5c6bd36d73ff3feceb05efa59b6acb7626617f4962be322a825e61f79a" + +[[package]] +name = "nom" +version = "7.1.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d273983c5a657a70a3e8f2a01329822f3b8c8172b73826411a55751e404a0a4a" +dependencies = [ + "memchr", + "minimal-lexical", +] + +[[package]] +name = "nvjitlink-sys" +version = "0.2.1" +dependencies = [ + "libloading", + "thiserror", +] + +[[package]] +name = "object" +version = "0.36.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "62948e14d923ea95ea2c7c86c71013138b66525b86bdc08d2dcc262bdb497b87" +dependencies = [ + "crc32fast", + "hashbrown 0.15.5", + "indexmap", + "memchr", +] + +[[package]] +name = "once_cell" +version = "1.21.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9f7c3e4beb33f85d45ae3e3a1792185706c8e16d043238c593331cc7cd313b50" + +[[package]] +name = "oxide-artifacts" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "74d9a941c0c43d2dceaf25f5c9ca4c9f9cceb28dfa5e4e3cfd14e476f003e7c9" +dependencies = [ + "object", +] + +[[package]] +name = "oxide-tile-export" +version = "0.2.1" +dependencies = [ + "serde", + "serde_json", + "thiserror", +] + +[[package]] +name = "prettyplease" +version = "0.2.37" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "479ca8adacdd7ce8f1fb39ce9ecccbfe93a3f1344b3d0d97f20bc0196208f62b" +dependencies = [ + "proc-macro2", + "syn", +] + +[[package]] +name = "proc-macro2" +version = "1.0.106" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8fd00f0bb2e90d81d1044c2b32617f68fcb9fa3bb7640c23e9c748e53fb30934" +dependencies = [ + "unicode-ident", +] + +[[package]] +name = "ptx-parse" +version = "0.1.0" + +[[package]] +name = "quote" +version = "1.0.45" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "41f2619966050689382d2b44f664f4bc593e129785a36d6ee376ddf37259b924" +dependencies = [ + "proc-macro2", +] + +[[package]] +name = "regex" +version = "1.12.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e10754a14b9137dd7b1e3e5b0493cc9171fdd105e0ab477f51b72e7f3ac0e276" +dependencies = [ + "aho-corasick", + "memchr", + "regex-automata", + "regex-syntax", +] + +[[package]] +name = "regex-automata" +version = "0.4.14" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6e1dd4122fc1595e8162618945476892eefca7b88c52820e74af6262213cae8f" +dependencies = [ + "aho-corasick", + "memchr", + "regex-syntax", +] + +[[package]] +name = "regex-syntax" +version = "0.8.11" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d6f6ff9a378485b298a5286656da665ba74413d36db0979633275d2e708145d4" + +[[package]] +name = "reserved-oxide-symbols" +version = "0.2.1" + +[[package]] +name = "rustc-hash" +version = "1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "08d43f7aa6b08d49f382cde6a7982047c3426db949b1424bc4b7ec9ae12c6ce2" + +[[package]] +name = "rustix" +version = "0.38.44" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fdb5bc1ae2baa591800df16c9ca78619bf65c0488b41b96ccec5d11220d8c154" +dependencies = [ + "bitflags", + "errno", + "libc", + "linux-raw-sys", + "windows-sys 0.59.0", +] + +[[package]] +name = "serde" +version = "1.0.228" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9a8e94ea7f378bd32cbbd37198a4a91436180c5bb472411e48b5ec2e2124ae9e" +dependencies = [ + "serde_core", + "serde_derive", +] + +[[package]] +name = "serde_core" +version = "1.0.228" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "41d385c7d4ca58e59fc732af25c3983b67ac852c1a25000afe1175de458b67ad" +dependencies = [ + "serde_derive", +] + +[[package]] +name = "serde_derive" +version = "1.0.228" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d540f220d3187173da220f885ab66608367b6574e925011a9353e4badda91d79" +dependencies = [ + "proc-macro2", + "quote", + "syn", +] + +[[package]] +name = "serde_json" +version = "1.0.154" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e7e9cc8b1b85264074fbcc02a88680c4096b1e47df8f739dceb03bf482f04bd6" +dependencies = [ + "itoa", + "memchr", + "serde", + "serde_core", + "zmij", +] + +[[package]] +name = "sha2" +version = "0.10.9" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a7507d819769d01a365ab707794a4084392c824f54a7a6a7862f8c3d0892b283" +dependencies = [ + "cfg-if", + "cpufeatures", + "digest", +] + +[[package]] +name = "shlex" +version = "1.3.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0fda2ff0d084019ba4d7c6f371c95d8fd75ce3524c3cb8fb653a3023f6323e64" + +[[package]] +name = "syn" +version = "2.0.117" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e665b8803e7b1d2a727f4023456bbbbe74da67099c585258af0ad9c5013b9b99" +dependencies = [ + "proc-macro2", + "quote", + "unicode-ident", +] + +[[package]] +name = "thiserror" +version = "2.0.18" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4288b5bcbc7920c07a1149a35cf9590a2aa808e0bc1eafaade0b80947865fbc4" +dependencies = [ + "thiserror-impl", +] + +[[package]] +name = "thiserror-impl" +version = "2.0.18" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ebc4ee7f67670e9b64d05fa4253e753e016c6c95ff35b89b7941d6b856dec1d5" +dependencies = [ + "proc-macro2", + "quote", + "syn", +] + +[[package]] +name = "typenum" +version = "1.20.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b6f5e870be6c3b371b77fe0ee0bafb859fa4964b4404c27de1d380043c4dda20" + +[[package]] +name = "unicode-ident" +version = "1.0.24" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e6e4313cd5fcd3dad5cafa179702e2b244f760991f45397d14d4ebf38247da75" + +[[package]] +name = "version_check" +version = "0.9.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a" + +[[package]] +name = "which" +version = "4.4.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "87ba24419a2078cd2b0f2ede2691b6c66d8e47836da3b6db8265ebad47afbfc7" +dependencies = [ + "either", + "home", + "once_cell", + "rustix", +] + +[[package]] +name = "windows-link" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f0805222e57f7521d6a62e36fa9163bc891acd422f971defe97d64e70d0a4fe5" + +[[package]] +name = "windows-sys" +version = "0.59.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1e38bc4d79ed67fd075bcc251a1c39b32a1776bbe92e5bef1f0bf1f8c531853b" +dependencies = [ + "windows-targets", +] + +[[package]] +name = "windows-sys" +version = "0.61.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ae137229bcbd6cdf0f7b80a31df61766145077ddf49416a728b02cb3921ff3fc" +dependencies = [ + "windows-link", +] + +[[package]] +name = "windows-targets" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9b724f72796e036ab90c1021d4780d4d3d648aca59e491e6b98e725b84e99973" +dependencies = [ + "windows_aarch64_gnullvm", + "windows_aarch64_msvc", + "windows_i686_gnu", + "windows_i686_gnullvm", + "windows_i686_msvc", + "windows_x86_64_gnu", + "windows_x86_64_gnullvm", + "windows_x86_64_msvc", +] + +[[package]] +name = "windows_aarch64_gnullvm" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "32a4622180e7a0ec044bb555404c800bc9fd9ec262ec147edd5989ccd0c02cd3" + +[[package]] +name = "windows_aarch64_msvc" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "09ec2a7bb152e2252b53fa7803150007879548bc709c039df7627cabbd05d469" + +[[package]] +name = "windows_i686_gnu" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8e9b5ad5ab802e97eb8e295ac6720e509ee4c243f69d781394014ebfe8bbfa0b" + +[[package]] +name = "windows_i686_gnullvm" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0eee52d38c090b3caa76c563b86c3a4bd71ef1a819287c19d586d7334ae8ed66" + +[[package]] +name = "windows_i686_msvc" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "240948bc05c5e7c6dabba28bf89d89ffce3e303022809e73deaefe4f6ec56c66" + +[[package]] +name = "windows_x86_64_gnu" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "147a5c80aabfbf0c7d901cb5895d1de30ef2907eb21fbbab29ca94c5b08b1a78" + +[[package]] +name = "windows_x86_64_gnullvm" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "24d5b23dc417412679681396f2b49f3de8c1473deb516bd34410872eff51ed0d" + +[[package]] +name = "windows_x86_64_msvc" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "589f6da84c646204747d1270a2a5661ea66ed1cced2631d546fdfb155959f9ec" + +[[package]] +name = "zerocopy" +version = "0.8.52" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ce1022995ff5ff5d841ad7d994facc23098cd40152f2c1d11cd607c6f530653f" +dependencies = [ + "zerocopy-derive", +] + +[[package]] +name = "zerocopy-derive" +version = "0.8.52" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1ae7f38b72ec2a254e2b87ef277cf2cd4fb97cbebf944faa6f33354da0867930" +dependencies = [ + "proc-macro2", + "quote", + "syn", +] + +[[package]] +name = "zmij" +version = "1.0.23" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "29666d0abbfad1e3dc4dcf6144730dd3a3ab225bbbdac83319345b1b44ccfc1b" diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/Cargo.toml b/cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/Cargo.toml new file mode 100644 index 0000000000..0daa15796f --- /dev/null +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/Cargo.toml @@ -0,0 +1,15 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 +[package] +name = "external_ltoir" +version = "0.1.0" +edition = "2024" +license = "Apache-2.0" +publish = false + +[workspace] + +[dependencies] +cuda-device = { path = "../../../cuda-device" } +cuda-host = { path = "../../../cuda-host" } +cuda-core = { path = "../../../../../cuda-core" } diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/README.md b/cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/README.md new file mode 100644 index 0000000000..a81f51b014 --- /dev/null +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/README.md @@ -0,0 +1,36 @@ +# External device LTO-IR + +This example calls a CUDA C++ device function from a Rust kernel. The C++ +function calls another function in a second LTO-IR file. `cargo oxide` links +both files into the embedded cubin during the build, and the program uses the +ordinary `kernels::load` API. It checks `output[i] = 2 * input[i] + 3` on the GPU. + +From `cuda-oxide/`, with the CUDA Toolkit and a matching GPU available: + +```bash +example=crates/rustc-codegen-cuda/examples/external_ltoir +mkdir -p build/device-libs +nvcc --ltoir -dlto -arch=sm_120 \ + "$example/scale.cu" -o build/device-libs/scale.ltoir +nvcc --ltoir -dlto -arch=sm_120 \ + "$example/add.cu" -o build/device-libs/add.ltoir +cargo oxide run external_ltoir --arch sm_120 --materialize-cubin \ + --link-ltoir build/device-libs/scale.ltoir \ + --link-ltoir build/device-libs/add.ltoir +``` + +Use the same architecture for both compilers. Each `--link-ltoir` adds one +file in argument order; paths are relative to the command's working directory. +The inputs must be raw LTO-IR, not object files or fatbins. `--ltoir -dlto` +writes the explicit output path, including with architecture names such as +`sm_90a`. All supplied files are linked into each materialized device artifact +in the Cargo invocation. The executable contains the finished cubin and needs no external `.ltoir` files +or runtime link step. + +`scripts/check-external-ltoir.sh sm_120` exercises run, build and test, repeated +build freshness, content changes at the same path, a missing transitive symbol, +execution with the link inputs removed, corrupt LTO-IR, and rejection of an +input/output collision without changing either file. It prints the +temporary directory containing its logs and generated inputs; it uses the +same nvcc commands with those temporary output paths. The check requires a GPU +matching the requested architecture. diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/add.cu b/cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/add.cu new file mode 100644 index 0000000000..d74d9bd981 --- /dev/null +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/add.cu @@ -0,0 +1,12 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#ifndef EXTERNAL_ADDEND +#define EXTERNAL_ADDEND 3 +#endif + +extern "C" __device__ int external_add(int value) { + return value + EXTERNAL_ADDEND; +} diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/scale.cu b/cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/scale.cu new file mode 100644 index 0000000000..ea5be973c4 --- /dev/null +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/scale.cu @@ -0,0 +1,10 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +extern "C" __device__ int external_add(int value); + +extern "C" __device__ int external_scale_add(int value) { + return external_add(value * 2); +} diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/src/main.rs b/cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/src/main.rs new file mode 100644 index 0000000000..04efbfab78 --- /dev/null +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/src/main.rs @@ -0,0 +1,73 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +//! Link two CUDA C++ device libraries during compilation and load the embedded cubin. + +use cuda_core::simt::LaunchConfig; +use cuda_core::{CudaContext, DeviceBuffer}; +use cuda_device::{DisjointSlice, device, kernel, thread}; +use cuda_host::cuda_module; + +#[device] +unsafe extern "C" { + fn external_scale_add(value: i32) -> i32; +} + +#[cuda_module] +mod kernels { + use super::*; + + #[kernel] + pub fn scale_add(input: &[i32], mut output: DisjointSlice) { + let index = thread::index_1d(); + let offset = index.get(); + if let Some(value) = output.get_mut(index) { + // SAFETY: the C++ definition takes and returns an int32 scalar; + // it accesses no pointers and has no participation requirement. + *value = unsafe { external_scale_add(input[offset]) }; + } + } +} + +fn verify(expected_addend: i32) -> Result<(), Box> { + let context = CudaContext::new(0)?; + let stream = context.default_stream(); + let input: Vec = (-128..129).collect(); + let device_input = DeviceBuffer::from_host(&stream, &input)?; + let mut device_output = DeviceBuffer::::zeroed(&stream, input.len())?; + let module = kernels::load(&context)?; + // SAFETY: the equal-length buffers cover every access, the bounds check + // handles the partial last block, and the kernel needs no shared memory. + unsafe { + module.scale_add( + &stream, + LaunchConfig::for_num_elems(input.len() as u32), + &device_input, + &mut device_output, + )?; + } + let actual = device_output.to_host_vec(&stream)?; + let expected: Vec = input + .iter() + .map(|value| value * 2 + expected_addend) + .collect(); + assert_eq!(actual, expected); + println!("PASS external LTO-IR: 257 values, addend={expected_addend}"); + Ok(()) +} + +fn main() -> Result<(), Box> { + let expected_addend = std::env::args() + .nth(1) + .map(|value| value.parse()) + .transpose()? + .unwrap_or(3); + verify(expected_addend) +} + +#[test] +fn linked_device_libraries_execute() { + verify(3).unwrap(); +} diff --git a/cuda-oxide/crates/rustc-codegen-cuda/src/lib.rs b/cuda-oxide/crates/rustc-codegen-cuda/src/lib.rs index 600b298617..1fe8466077 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/src/lib.rs +++ b/cuda-oxide/crates/rustc-codegen-cuda/src/lib.rs @@ -628,7 +628,7 @@ impl CodegenBackend for CudaCodegenBackend { ); materialize::validate_collection( - materialization_request, + materialization_request.as_ref(), !collection_result.device_externs.is_empty(), collection_result.requires_ptx_bundle_merge, ) @@ -664,6 +664,16 @@ impl CodegenBackend for CudaCodegenBackend { dump_mir_dialect: self.config.dump_mir_dialect, dump_llvm_dialect: self.config.dump_llvm_dialect, }; + materialize::validate_output_paths( + materialization_request.as_ref(), + &device_config.output_dir, + &device_config.output_name, + ) + .unwrap_or_else(|error| { + tcx.dcx().fatal(format!( + "[rustc_codegen_cuda] Cannot write device artifacts: {error}" + )) + }); // Run the cuda-oxide pipeline, catching backend panics and // re-emitting them as a cuda-oxide diagnostic. A panic diff --git a/cuda-oxide/crates/rustc-codegen-cuda/src/materialize.rs b/cuda-oxide/crates/rustc-codegen-cuda/src/materialize.rs index 387c4ce4f6..c2f813e0b0 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/src/materialize.rs +++ b/cuda-oxide/crates/rustc-codegen-cuda/src/materialize.rs @@ -27,10 +27,11 @@ pub(crate) const MATERIALIZER_HANDSHAKE_ENV: &str = reserved_oxide_symbols::MATERIALIZER_HANDSHAKE_ENV; pub(crate) const CODEGEN_FINGERPRINT_ENV: &str = reserved_oxide_symbols::CODEGEN_FINGERPRINT_ENV; -#[derive(Clone, Copy, Debug, Eq, PartialEq)] +#[derive(Debug)] pub(crate) struct MaterializationRequest { expected_provenance: [u8; 32], tool_identity_handshake: cuda_artifact_finalizer::MaterializerHandshakeV1, + external_inputs: Vec<(String, Vec)>, } /// Cubin bytes plus ptxas resource diagnostics from the final link. @@ -97,10 +98,13 @@ pub(crate) enum MaterializeError { RequiresPtxBundleMerge, #[error( - "build-time cubin materialization does not yet support #[device] extern declarations because their ordered external link inputs are not available to the backend" + "build-time cubin materialization of #[device] extern declarations requires their definitions; supply ordered --link-ltoir inputs" )] HasDeviceExterns, + #[error("invalid cargo-oxide external LTO-IR input list: {reason}")] + InvalidExternalInputs { reason: String }, + #[error( "build-time cubin materialization requires an NVVM IR or LTOIR artifact, but codegen produced PTX; use cargo-oxide so materialization can force NVVM IR emission" )] @@ -147,12 +151,37 @@ pub(crate) fn request_from_env() -> Result, Mater })?; validate_tool_identity_handshake(expected_provenance, &handshake)?; validate_codegen_fingerprint()?; + let external_inputs = external_inputs_from_env()?; Ok(Some(MaterializationRequest { expected_provenance, tool_identity_handshake: handshake, + external_inputs, })) } +fn external_inputs_from_env() -> Result)>, MaterializeError> { + let value = match std::env::var(reserved_oxide_symbols::EXTERNAL_LTOIR_ENV) { + Ok(value) => value, + Err(std::env::VarError::NotPresent) => return Ok(Vec::new()), + Err(error) => { + return Err(MaterializeError::InvalidExternalInputs { + reason: error.to_string(), + }); + } + }; + let files: Vec = + serde_json::from_str(&value).map_err(|error| MaterializeError::InvalidExternalInputs { + reason: error.to_string(), + })?; + files + .into_iter() + .map(|file| { + let bytes = file.read_verified()?; + Ok((file.name()?.to_owned(), bytes)) + }) + .collect() +} + fn validate_tool_identity_handshake( expected_provenance: [u8; 32], handshake: &cuda_artifact_finalizer::MaterializerHandshakeV1, @@ -174,22 +203,46 @@ fn validate_tool_identity_handshake( /// Reject artifact-loading models the finalizer cannot reproduce, before any /// CUDA compiler library is discovered or loaded. pub(crate) fn validate_collection( - request: Option, + request: Option<&MaterializationRequest>, has_device_externs: bool, requires_ptx_bundle_merge: bool, ) -> Result<(), MaterializeError> { - if request.is_none() { + let Some(request) = request else { return Ok(()); - } + }; if requires_ptx_bundle_merge { return Err(MaterializeError::RequiresPtxBundleMerge); } - if has_device_externs { + if has_device_externs && request.external_inputs.is_empty() { return Err(MaterializeError::HasDeviceExterns); } Ok(()) } +/// Use the actual rustc crate name and configured output directory, including +/// for Cargo passthrough builds whose working directory differs from the CLI's. +pub(crate) fn validate_output_paths( + request: Option<&MaterializationRequest>, + output_dir: &std::path::Path, + output_name: &str, +) -> Result<(), MaterializeError> { + let Some(request) = request else { + return Ok(()); + }; + let mut outputs = + mir_importer::pipeline::stale_compilation_artifact_paths(output_dir, output_name); + outputs.push(output_dir.join(format!( + "{output_name}.{}", + oxide_tile_export::DESCRIPTIONS_SUFFIX + ))); + for (input, _) in &request.external_inputs { + for output in &outputs { + cuda_artifact_finalizer::check_ltoir_input_output(std::path::Path::new(input), output)?; + } + } + Ok(()) +} + pub(crate) fn nvvm_ir_to_cubin( request: MaterializationRequest, nvvm_ir: &[u8], @@ -199,8 +252,13 @@ pub(crate) fn nvvm_ir_to_cubin( debug_policy: DebugPolicy, ) -> Result { let options = options(target, allow_fma_contraction, debug_policy)?; - let finalizer = checked_finalizer(request)?; - let report = finalizer.materialize_nvvm_ir_with_report(module_name, nvvm_ir, &options)?; + let finalizer = checked_finalizer(&request)?; + let ltoir = finalizer + .compiler() + .compile_nvvm_ir_to_ltoir(module_name, nvvm_ir, &options)?; + let name = format!("{module_name}.ltoir"); + let inputs = link_inputs(&request, &name, <oir); + let report = finalizer.link_ltoir_with_report(&inputs, &options, FinalizerOutput::Cubin)?; Ok(MaterializedCubin { bytes: report.image, resource_usage: report.resource_usage, @@ -216,18 +274,30 @@ pub(crate) fn ltoir_to_cubin( debug_policy: DebugPolicy, ) -> Result { let options = options(target, allow_fma_contraction, debug_policy)?; - let finalizer = checked_finalizer(request)?; - let report = finalizer.link_ltoir_with_report( - &[NamedInput::new(module_name, ltoir)], - &options, - FinalizerOutput::Cubin, - )?; + let finalizer = checked_finalizer(&request)?; + let inputs = link_inputs(&request, module_name, ltoir); + let report = finalizer.link_ltoir_with_report(&inputs, &options, FinalizerOutput::Cubin)?; Ok(MaterializedCubin { bytes: report.image, resource_usage: report.resource_usage, }) } +fn link_inputs<'a>( + request: &'a MaterializationRequest, + name: &'a str, + bytes: &'a [u8], +) -> Vec> { + std::iter::once(NamedInput::new(name, bytes)) + .chain( + request + .external_inputs + .iter() + .map(|(name, bytes)| NamedInput::new(name, bytes)), + ) + .collect() +} + fn options( target: &str, allow_fma_contraction: bool, @@ -239,7 +309,7 @@ fn options( .with_debug_policy(debug_policy)) } -fn checked_finalizer(request: MaterializationRequest) -> Result { +fn checked_finalizer(request: &MaterializationRequest) -> Result { let finalizer = Finalizer::discover_with_handshake(&request.tool_identity_handshake)?; let actual = finalizer .provenance_digest() @@ -384,22 +454,64 @@ mod tests { )); } + #[test] + fn output_overlap_uses_the_actual_crate_and_output_directory() { + let root = std::env::temp_dir().join(format!( + "oxide-materialize-overlap-{}-{}", + std::process::id(), + std::time::SystemTime::now() + .duration_since(std::time::UNIX_EPOCH) + .unwrap() + .as_nanos() + )); + std::fs::create_dir_all(&root).unwrap(); + let input = root.join("custom_target.ll"); + std::fs::write(&input, b"external input").unwrap(); + let handshake = test_handshake(); + let request = MaterializationRequest { + expected_provenance: handshake.provenance_sha256, + tool_identity_handshake: handshake, + external_inputs: vec![( + input.to_str().unwrap().to_owned(), + b"external input".to_vec(), + )], + }; + assert!(validate_output_paths(Some(&request), &root, "package_name").is_ok()); + assert!(matches!( + validate_output_paths(Some(&request), &root, "custom_target"), + Err(MaterializeError::Finalizer( + FinalizerError::LtoirInputOutputOverlap { .. } + )) + )); + assert_eq!(std::fs::read(&input).unwrap(), b"external input"); + std::fs::remove_dir_all(root).unwrap(); + } + #[test] fn unsupported_collection_models_fail_without_tools() { let handshake = test_handshake(); - let request = Some(MaterializationRequest { + let mut request = MaterializationRequest { expected_provenance: handshake.provenance_sha256, tool_identity_handshake: handshake, - }); + external_inputs: Vec::new(), + }; assert!(matches!( - validate_collection(request, false, true), + validate_collection(Some(&request), false, true), Err(MaterializeError::RequiresPtxBundleMerge) )); assert!(matches!( - validate_collection(request, true, false), + validate_collection(Some(&request), true, false), Err(MaterializeError::HasDeviceExterns) )); assert!(validate_collection(None, true, true).is_ok()); + request + .external_inputs + .push(("definition.ltoir".into(), vec![1])); + assert!(validate_collection(Some(&request), true, false).is_ok()); + assert!(matches!( + validate_collection(Some(&request), true, true), + Err(MaterializeError::RequiresPtxBundleMerge) + )); } #[test] diff --git a/cuda-oxide/cuda-oxide-book/getting-started/installation.md b/cuda-oxide/cuda-oxide-book/getting-started/installation.md index 4a39d2ea3a..6d2971f6b5 100644 --- a/cuda-oxide/cuda-oxide-book/getting-started/installation.md +++ b/cuda-oxide/cuda-oxide-book/getting-started/installation.md @@ -373,6 +373,22 @@ cargo oxide emit-ltoir my_kernels --arch sm_90 Every command takes `--help`, which lists the flags each one accepts. +For non-generic kernels that call external device functions, `--materialize-cubin` +links their CUDA LTO-IR definitions during the ordinary Cargo build. Pass each +input with `--link-ltoir` after the subcommand, in link order: + +```bash +cargo oxide run --materialize-cubin --arch sm_120 \ + --link-ltoir build/device-libs/scale.ltoir --link-ltoir build/device-libs/add.ltoir +``` + +Inputs must be raw CUDA LTO-IR, not object files or fatbins. Every supplied +file is linked into each materialized device artifact in the Cargo invocation. +This requires a build-host CUDA Toolkit compatible with the external inputs. +The resulting executable embeds the final cubin. Input names and contents are +part of Cargo's device-code fingerprint, so a changed definition rebuilds the +artifact. See the `external_ltoir` example for C++ sources and generation commands. + --- ## Verifying your installation diff --git a/cuda-oxide/scripts/check-external-ltoir.sh b/cuda-oxide/scripts/check-external-ltoir.sh new file mode 100755 index 0000000000..ff2f87dbc5 --- /dev/null +++ b/cuda-oxide/scripts/check-external-ltoir.sh @@ -0,0 +1,106 @@ +#!/usr/bin/env bash +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 +# Run with a CUDA Toolkit and GPU matching the requested architecture. +set -euo pipefail + +cd "$(dirname "$0")/.." +arch="${1:-sm_120}" +example="$PWD/crates/rustc-codegen-cuda/examples/external_ltoir" +artifacts="$(mktemp -d "${TMPDIR:-/tmp}/oxide-external-ltoir.XXXXXX")" +echo "External LTO-IR regression artifacts: $artifacts" +# Keep the example's executable location independent of the caller's Cargo cache. +unset CARGO_TARGET_DIR +oxide=(cargo oxide) +if [[ -n "${CARGO_OXIDE_BIN:-}" ]]; then + oxide=("$CARGO_OXIDE_BIN") +fi + +nvcc_flags=(--ltoir -dlto -arch="$arch") +if [[ -n "${NVCC_CCBIN:-${CUDAHOSTCXX:-}}" ]]; then + nvcc_flags+=(-ccbin="${NVCC_CCBIN:-$CUDAHOSTCXX}") +fi +nvcc "${nvcc_flags[@]}" "$example/scale.cu" -o "$artifacts/scale.ltoir" +nvcc "${nvcc_flags[@]}" -DEXTERNAL_ADDEND=3 "$example/add.cu" -o "$artifacts/add.ltoir" + +flags=(--arch "$arch" --materialize-cubin + --link-ltoir "$artifacts/scale.ltoir" --link-ltoir "$artifacts/add.ltoir") +cargo_args=(--manifest-path "$example/Cargo.toml" --release) + +run() { + local label="$1" + shift + if ! "$@" >"$artifacts/$label.log" 2>&1; then + cat "$artifacts/$label.log" >&2 + return 1 + fi + echo "PASS $label" +} + +run run "${oxide[@]}" run external_ltoir "${flags[@]}" +run test "${oxide[@]}" test "${flags[@]}" -- "${cargo_args[@]}" -- --nocapture +run build "${oxide[@]}" build "${flags[@]}" -- "${cargo_args[@]}" -v +run fresh "${oxide[@]}" build "${flags[@]}" -- "${cargo_args[@]}" -v +grep -q 'Fresh external_ltoir ' "$artifacts/fresh.log" + +# Change just the library content, retaining its path, size and modification +# time. The next passthrough build must invalidate Cargo's cached device code. +cp -p "$artifacts/add.ltoir" "$artifacts/add-original.ltoir" +nvcc "${nvcc_flags[@]}" -DEXTERNAL_ADDEND=1 "$example/add.cu" -o "$artifacts/add.ltoir" +if [[ "$(wc -c < "$artifacts/add-original.ltoir")" != "$(wc -c < "$artifacts/add.ltoir")" ]]; then + echo "error: this nvcc did not produce equal-sized addend variants" >&2 + exit 1 +fi +if cmp -s "$artifacts/add-original.ltoir" "$artifacts/add.ltoir"; then + echo "error: the addend variants have identical bytes" >&2 + exit 1 +fi +touch -r "$artifacts/add-original.ltoir" "$artifacts/add.ltoir" +run content-change "${oxide[@]}" build "${flags[@]}" -- "${cargo_args[@]}" -v +mv "$artifacts/scale.ltoir" "$artifacts/scale.saved" +mv "$artifacts/add.ltoir" "$artifacts/add.saved" +run changed-result "$example/target/release/external_ltoir" 1 +mv "$artifacts/scale.saved" "$artifacts/scale.ltoir" +mv "$artifacts/add.saved" "$artifacts/add.ltoir" + +if "${oxide[@]}" build --arch "$arch" --materialize-cubin \ + --link-ltoir "$artifacts/scale.ltoir" -- "${cargo_args[@]}" \ + >"$artifacts/missing-symbol.log" 2>&1; then + echo "error: a missing transitive device definition unexpectedly linked" >&2 + exit 1 +fi +grep -q 'external_add' "$artifacts/missing-symbol.log" +grep -qi 'nvJitLink' "$artifacts/missing-symbol.log" +echo "PASS missing-symbol" + +printf 'not an LTO-IR module\n' >"$artifacts/corrupt.ltoir" +if "${oxide[@]}" build --arch "$arch" --materialize-cubin \ + --link-ltoir "$artifacts/scale.ltoir" --link-ltoir "$artifacts/corrupt.ltoir" \ + -- "${cargo_args[@]}" >"$artifacts/corrupt-input.log" 2>&1; then + echo "error: corrupt LTO-IR unexpectedly linked" >&2 + exit 1 +fi +grep -qi 'nvJitLink' "$artifacts/corrupt-input.log" +grep -q 'corrupt.ltoir' "$artifacts/corrupt-input.log" +echo "PASS corrupt-input" + +# Exercise the backend guard using a disposable output directory. The supplied +# input and a previous output must both survive rejection before codegen writes. +collision="$artifacts/collision" +mkdir -p "$collision" +cp "$artifacts/scale.ltoir" "$collision/external_ltoir.ll" +printf 'previous output sentinel\n' >"$artifacts/previous-output" +cp "$artifacts/previous-output" "$collision/external_ltoir.ptx" +if CUDA_OXIDE_PTX_DIR="$collision" "${oxide[@]}" build --arch "$arch" --materialize-cubin \ + --link-ltoir "$collision/external_ltoir.ll" --link-ltoir "$artifacts/add.ltoir" \ + -- "${cargo_args[@]}" >"$artifacts/input-output-overlap.log" 2>&1; then + echo "error: an external input overlapping generated output was accepted" >&2 + exit 1 +fi +grep -q 'overlaps generated artifact' "$artifacts/input-output-overlap.log" +cmp "$artifacts/scale.ltoir" "$collision/external_ltoir.ll" +cmp "$artifacts/previous-output" "$collision/external_ltoir.ptx" +echo "PASS input-output-overlap" + +cp -p "$artifacts/add-original.ltoir" "$artifacts/add.ltoir" +run restored "${oxide[@]}" run external_ltoir "${flags[@]}" diff --git a/cuda-oxide/scripts/smoketest.sh b/cuda-oxide/scripts/smoketest.sh index 35d8eb57a2..18a82fe8cb 100755 --- a/cuda-oxide/scripts/smoketest.sh +++ b/cuda-oxide/scripts/smoketest.sh @@ -90,7 +90,7 @@ set -uo pipefail TCGEN05_EXAMPLES=(gemm_sol gemm_sol_final tcgen05 tcgen05_matmul) WGMMA_EXAMPLES=(wgmma) BLACKWELL_MMA_EXAMPLES=(mma_mxf8f6f4) -LTOIR_EXAMPLES=(addressof_sharedarray constant_memory cpp_consumes_rust_device device_ffi_test legacy_atomic_fadd legacy_atomic_rmw_cas legacy_nvvm_pointer_shapes manual_launch_libdevice mathdx_ffi_test primitive_stress) +LTOIR_EXAMPLES=(addressof_sharedarray constant_memory cpp_consumes_rust_device device_ffi_test external_ltoir legacy_atomic_fadd legacy_atomic_rmw_cas legacy_nvvm_pointer_shapes manual_launch_libdevice mathdx_ffi_test primitive_stress) LTOIR_MODERN_EXAMPLES=(small_type_ffi_test) TILELIB_EXAMPLES=(tilelib_relu) AUTO_NVVM_EXAMPLES=(libdevice_math) @@ -1038,6 +1038,11 @@ verdict_compile() { "${ex}" "${ex_dir}/${artifact}.ptx" return ;; + external_ltoir) + full_debug_relocatable_ffi_verdict \ + "${ex}" "${ex_dir}/${artifact}.ptx" + return + ;; cutile_inter_kernel) full_debug_ptx_verdict \ "${ex_dir}/simt/cutile_inter_kernel_simt.ptx" \ @@ -1180,6 +1185,9 @@ full_debug_relocatable_ffi_verdict() { mathdx_ffi_test) expected_undefined=$'cublasdx_gemm_32x32x32_block_dim_x\ncublasdx_gemm_32x32x32_block_dim_y\ncublasdx_gemm_32x32x32_block_dim_z\ncublasdx_gemm_32x32x32_f32\ncublasdx_gemm_32x32x32_f32_alphabeta\ncublasdx_gemm_32x32x32_smem_size\ncublasdx_gemm_32x32x32_smem_size_ab\ncufftdx_fft_16_c2c_f32_forward\ncufftdx_fft_16_c2c_f32_inverse\ncufftdx_fft_16_elements_per_thread\ncufftdx_fft_16_storage_size\ncufftdx_fft_32_storage_size\ncufftdx_fft_8_c2c_f32_forward\ncufftdx_fft_8_c2c_f32_inverse\ncufftdx_fft_8_elements_per_thread\ncufftdx_fft_8_storage_size\ndebug_extern_double_array' ;; + external_ltoir) + expected_undefined='external_scale_add' + ;; small_type_ffi_test) expected_undefined=$'small_add_u16\nsmall_add_u8\nsmall_half_add\nsmall_not_bool\nsmall_scale_i16\nsmall_scale_i8\nsmall_widen_i8\nsmall_widen_u16' ;; @@ -1405,6 +1413,15 @@ run_cargo() { run_full_debug_build "${ex}" "${log}" "${cat}" return fi + # This fixture exercises compile-time linking of external device libraries. + # Its script supplies those inputs and verifies the embedded executable. + # Compile-only and full-debug retain the ordinary artifact-only FFI routes. + if [[ ${COMPILE_ONLY} -eq 0 && "${ex}" == "external_ltoir" ]]; then + bash scripts/check-external-ltoir.sh "${LTOIR_ARCH}" >"${log}" 2>&1 + CARGO_EC=$? + if [[ ${VERBOSE} -eq 1 ]]; then cat "${log}"; fi + return + fi # Rust `char` must remain plain i32 in both NVVM dialects. A runtime # round-trip cannot distinguish it from an incorrect zeroext i32 ABI. if [[ ${COMPILE_ONLY} -eq 1 && "${ex}" == "device_ffi_test" ]]; then