From ea0daf49c9ef768eaefcde65451331ef5e09045b Mon Sep 17 00:00:00 2001 From: Trent Nelson Date: Mon, 19 Oct 2026 21:39:54 -0700 Subject: [PATCH 1/2] feat(cargo-oxide): link external device LTO inputs Accept ordered external LTO files when materializing embedded cubins. Include their content in the codegen fingerprint and verify owned bytes before linking through the existing finalizer. Preserve inputs that overlap generated output paths and diagnose unresolved device symbols. Add a standalone CUDA FFI example and a native regression script. Signed-off-by: Trent Nelson --- cuda-oxide/crates/cargo-oxide/README.md | 30 +- .../cargo-oxide/src/commands/artifacts.rs | 17 +- .../cargo-oxide/src/commands/build_run.rs | 15 +- .../cargo-oxide/src/commands/context.rs | 6 + .../cargo-oxide/src/commands/fingerprint.rs | 5 + .../cargo-oxide/src/commands/interop.rs | 5 +- .../cargo-oxide/src/commands/materialize.rs | 16 +- .../src/commands/pipeline_debug.rs | 10 +- .../crates/cargo-oxide/src/commands/tests.rs | 105 ++- cuda-oxide/crates/cargo-oxide/src/main.rs | 114 ++- .../crates/cuda-artifact-finalizer/src/lib.rs | 16 + .../cuda-artifact-finalizer/src/ltoir_file.rs | 163 ++++ .../crates/mir-importer/src/pipeline.rs | 4 +- .../crates/reserved-oxide-symbols/src/lib.rs | 3 + .../examples/external_ltoir/Cargo.lock | 764 ++++++++++++++++++ .../examples/external_ltoir/Cargo.toml | 15 + .../examples/external_ltoir/README.md | 36 + .../examples/external_ltoir/add.cu | 12 + .../examples/external_ltoir/scale.cu | 10 + .../examples/external_ltoir/src/main.rs | 73 ++ .../crates/rustc-codegen-cuda/src/lib.rs | 12 +- .../rustc-codegen-cuda/src/materialize.rs | 150 +++- .../getting-started/installation.md | 16 + cuda-oxide/scripts/check-external-ltoir.sh | 106 +++ 24 files changed, 1657 insertions(+), 46 deletions(-) create mode 100644 cuda-oxide/crates/cuda-artifact-finalizer/src/ltoir_file.rs create mode 100644 cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/Cargo.lock create mode 100644 cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/Cargo.toml create mode 100644 cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/README.md create mode 100644 cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/add.cu create mode 100644 cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/scale.cu create mode 100644 cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/src/main.rs create mode 100755 cuda-oxide/scripts/check-external-ltoir.sh 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..b3f2b6042b --- /dev/null +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/Cargo.lock @@ -0,0 +1,764 @@ +# This file is automatically @generated by Cargo. +# It is not intended for manual editing. +version = 4 + +[[package]] +name = "aho-corasick" +version = "1.1.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c982642fa9e8606056828ee9a8505737230110bb1099153c79efe865c59d12ba" +dependencies = [ + "memchr", +] + +[[package]] +name = "anyhow" +version = "1.0.104" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "330a5ed07fa54e4702c9d6c4174f74427fc0ef6e214bbd677ae50a5099946470" + +[[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 2.0.119", + "which", +] + +[[package]] +name = "bitflags" +version = "2.13.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3ded4057c258ba199e2d26386d3af3780957ecaee6c4ef4041c6b4b8b97c0b06" + +[[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.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4e7648175b45a9a48536d676f68d918270699102aa8dab5496df06904c914600" + +[[package]] +name = "clang-sys" +version = "1.9.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "157a8ba7b480713b56f4c09fd13fc3e0a22a5dfab8097ba61cbc5feef950788a" +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.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "01a7799fd6b852db0e61728dde9a204c423b44d689dbd432522543614b490e78" +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 2.0.119", +] + +[[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 2.0.119", +] + +[[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 2.0.119", +] + +[[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.19.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0e9c71c2167ca323c882b99918929403426e2373ea17242ff5653e0d5e1058be" + +[[package]] +name = "equivalent" +version = "1.0.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "00d174d5400e5e8fd687ad1049e2f578285fa914201b1af7e8b112a4546bd826" + +[[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.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e4eba85ea1d0a966a983acd07deee566e67395d2d96b6fb39e62b5a833f1eb0b" + +[[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.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "cc4e190f5d26ca7051642629da2c52fc03bde85a03197c99408dcd291734c855" +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.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "20870f649af7073d53e38067b2a84312175d56ea15217e1b15bc83506ec50afb" + +[[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.191" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "65cd87c9a6d60536580b316ad985726104730a5151c99845f2490428ef91240a" + +[[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.34" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f9f8bd3e56ce4dfc153cf470fffbfa98c7620958b312ca5c3a4b8d5181fd13c6" + +[[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 2.0.119", +] + +[[package]] +name = "proc-macro2" +version = "1.0.107" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "985e7ec9bb745e6ce6535b544d84d6cd6f7ad8bd711c398938ae983b91a766d9" +dependencies = [ + "unicode-ident", +] + +[[package]] +name = "ptx-parse" +version = "0.1.0" + +[[package]] +name = "quote" +version = "1.0.47" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1fbf4db142a473a8d80c26bbf18454ed458bf8d26c8219c331daecfdbd079001" +dependencies = [ + "proc-macro2", +] + +[[package]] +name = "regex" +version = "1.13.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f020237b6c8eed93db2e2cb53c00c60a8e1bc73da7d073199a1180401450218d" +dependencies = [ + "aho-corasick", + "memchr", + "regex-automata", + "regex-syntax", +] + +[[package]] +name = "regex-automata" +version = "0.4.18" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ad8553b9b26413251cbf30e620595c7a41b3887f03da04579c0e6b0d6a06b4b2" +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.229" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4148590afebada386688f18773da617792bf2ef03ffc1e4cbd2b1d45b023e0ba" +dependencies = [ + "serde_core", + "serde_derive", +] + +[[package]] +name = "serde_core" +version = "1.0.229" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "67dca2c9c51e58a4791a4b1ed58308b39c64224d349a935ab5039aa360942a48" +dependencies = [ + "serde_derive", +] + +[[package]] +name = "serde_derive" +version = "1.0.229" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e7a5d71263a5a7d47b41f6b3f06ba276f10cc18b0931f1799f710578e2309348" +dependencies = [ + "proc-macro2", + "quote", + "syn 3.0.9", +] + +[[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.119" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "872831b642d1a07999a962a351ed35b955ea2cfc8f3862091e2a240a84f17297" +dependencies = [ + "proc-macro2", + "quote", + "unicode-ident", +] + +[[package]] +name = "syn" +version = "3.0.9" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d78c8dee4c7bf0e14673097256fed6142ce9d3b85a408189d07482442145823b" +dependencies = [ + "proc-macro2", + "quote", + "unicode-ident", +] + +[[package]] +name = "thiserror" +version = "2.0.21" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "09e52cb86a36cede5cb101bf8908837b3e4c6e5e59fe7fd85c23fb56200d189e" +dependencies = [ + "thiserror-impl", +] + +[[package]] +name = "thiserror-impl" +version = "2.0.21" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fe5197923287db20a58125f0bc85c062f7f2c892de97b18c356f9efb14b28524" +dependencies = [ + "proc-macro2", + "quote", + "syn 3.0.9", +] + +[[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.27" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a2c754d6c33795a1c324727428e5a7dedb5b06195f9890bdbcba760d3e246563" + +[[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.63" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e5fe1f8f1b06191a00962174c61aa5005e0bb391a6d80d07e24d115c01a92ed8" +dependencies = [ + "zerocopy-derive", +] + +[[package]] +name = "zerocopy-derive" +version = "0.8.63" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "863ad3ac83293fb4d740aedbfdc9240dd8d1a50c1099acd76ce80ce7c7230c7f" +dependencies = [ + "proc-macro2", + "quote", + "syn 2.0.119", +] + +[[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[@]}" From 79c0800425be486daad238eb2669d301c4cc3320 Mon Sep 17 00:00:00 2001 From: Trent Nelson Date: Mon, 19 Oct 2026 21:46:22 -0700 Subject: [PATCH 2/2] Align external LTO example with repository smoke checks Reuse existing example dependency pins and classify external LTO as an FFI fixture. Run its linking regression in native sweeps and verify relocatable debug output in the existing full-debug route. Signed-off-by: Trent Nelson --- .../examples/external_ltoir/Cargo.lock | 131 ++++++++---------- cuda-oxide/scripts/smoketest.sh | 19 ++- 2 files changed, 78 insertions(+), 72 deletions(-) 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 index b3f2b6042b..af40b6dd11 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/Cargo.lock +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/external_ltoir/Cargo.lock @@ -4,18 +4,18 @@ version = 4 [[package]] name = "aho-corasick" -version = "1.1.5" +version = "1.1.4" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "c982642fa9e8606056828ee9a8505737230110bb1099153c79efe865c59d12ba" +checksum = "ddd31a130427c27518df266943a5308ed92d4b226cc639f5a8f1002816174301" dependencies = [ "memchr", ] [[package]] name = "anyhow" -version = "1.0.104" +version = "1.0.102" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "330a5ed07fa54e4702c9d6c4174f74427fc0ef6e214bbd677ae50a5099946470" +checksum = "7f202df86484c868dbad7eaa557ef785d5c66295e41b460ef922eca0723b842c" [[package]] name = "bindgen" @@ -36,15 +36,15 @@ dependencies = [ "regex", "rustc-hash", "shlex", - "syn 2.0.119", + "syn", "which", ] [[package]] name = "bitflags" -version = "2.13.2" +version = "2.13.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "3ded4057c258ba199e2d26386d3af3780957ecaee6c4ef4041c6b4b8b97c0b06" +checksum = "b4388bee8683e3d04af747c73422af53102d2bd24d9eadb6cbc100baef4b43f8" [[package]] name = "block-buffer" @@ -66,15 +66,15 @@ dependencies = [ [[package]] name = "cfg-if" -version = "1.0.5" +version = "1.0.4" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "4e7648175b45a9a48536d676f68d918270699102aa8dab5496df06904c914600" +checksum = "9330f8b2ff13f34540b44e946ef35111825727b38d33286ef986142615121801" [[package]] name = "clang-sys" -version = "1.9.1" +version = "1.8.1" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "157a8ba7b480713b56f4c09fd13fc3e0a22a5dfab8097ba61cbc5feef950788a" +checksum = "0b023947811758c97c59bf9d1c188fd619ad4718dcaa767947df1cadb14f39f4" dependencies = [ "glob", "libc", @@ -92,9 +92,9 @@ dependencies = [ [[package]] name = "crc32fast" -version = "1.5.2" +version = "1.5.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "01a7799fd6b852db0e61728dde9a204c423b44d689dbd432522543614b490e78" +checksum = "9481c1c90cbf2ac953f07c8d4a58aa3945c425b7185c9154d67a65e4230da511" dependencies = [ "cfg-if", ] @@ -135,7 +135,7 @@ dependencies = [ "prettyplease", "proc-macro2", "quote", - "syn 2.0.119", + "syn", ] [[package]] @@ -155,7 +155,7 @@ version = "0.4.0" dependencies = [ "proc-macro2", "quote", - "syn 2.0.119", + "syn", ] [[package]] @@ -188,7 +188,7 @@ dependencies = [ "proc-macro2", "quote", "reserved-oxide-symbols", - "syn 2.0.119", + "syn", ] [[package]] @@ -207,15 +207,15 @@ dependencies = [ [[package]] name = "either" -version = "1.19.0" +version = "1.16.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "0e9c71c2167ca323c882b99918929403426e2373ea17242ff5653e0d5e1058be" +checksum = "91622ff5e7162018101f2fea40d6ebf4a78bbe5a49736a2020649edf9693679e" [[package]] name = "equivalent" -version = "1.0.3" +version = "1.0.2" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "00d174d5400e5e8fd687ad1049e2f578285fa914201b1af7e8b112a4546bd826" +checksum = "877a4ace8713b0bcf2a4e7eec82529c029f1d0619886d18145fea96c3ffe5c0f" [[package]] name = "errno" @@ -254,9 +254,9 @@ dependencies = [ [[package]] name = "glob" -version = "0.3.4" +version = "0.3.3" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "e4eba85ea1d0a966a983acd07deee566e67395d2d96b6fb39e62b5a833f1eb0b" +checksum = "0cc23270f6e1808e30a928bdc84dea0b9b4136a8bc82338574f23baf47bbd280" [[package]] name = "half" @@ -295,9 +295,9 @@ dependencies = [ [[package]] name = "indexmap" -version = "2.14.2" +version = "2.14.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "cc4e190f5d26ca7051642629da2c52fc03bde85a03197c99408dcd291734c855" +checksum = "d466e9454f08e4a911e14806c24e16fba1b4c121d1ea474396f396069cf949d9" dependencies = [ "equivalent", "hashbrown 0.17.1", @@ -320,9 +320,9 @@ checksum = "8f42a60cbdf9a97f5d2305f08a87dc4e09308d1276d28c869c684d7777685682" [[package]] name = "lazy_static" -version = "1.5.1" +version = "1.5.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "20870f649af7073d53e38067b2a84312175d56ea15217e1b15bc83506ec50afb" +checksum = "bbd2bcb4c963f2ddae06a2efc7e9f3591312473c50c6685e1f298068316e66fe" [[package]] name = "lazycell" @@ -332,9 +332,9 @@ checksum = "830d08ce1d1d941e6b30645f1a0eb5643013d835ce3779a5fc208261dbe10f55" [[package]] name = "libc" -version = "0.2.191" +version = "0.2.186" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "65cd87c9a6d60536580b316ad985726104730a5151c99845f2490428ef91240a" +checksum = "68ab91017fe16c622486840e4c83c9a37afeff978bd239b5293d61ece587de66" [[package]] name = "libloading" @@ -363,9 +363,9 @@ checksum = "d26c52dbd32dccf2d10cac7725f8eae5296885fb5703b261f7d0a0739ec807ab" [[package]] name = "log" -version = "0.4.34" +version = "0.4.29" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "f9f8bd3e56ce4dfc153cf470fffbfa98c7620958b312ca5c3a4b8d5181fd13c6" +checksum = "5e5032e24019045c762d3c0f28f5b6b8bbf38563a65908389bf7978758920897" [[package]] name = "memchr" @@ -440,14 +440,14 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "479ca8adacdd7ce8f1fb39ce9ecccbfe93a3f1344b3d0d97f20bc0196208f62b" dependencies = [ "proc-macro2", - "syn 2.0.119", + "syn", ] [[package]] name = "proc-macro2" -version = "1.0.107" +version = "1.0.106" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "985e7ec9bb745e6ce6535b544d84d6cd6f7ad8bd711c398938ae983b91a766d9" +checksum = "8fd00f0bb2e90d81d1044c2b32617f68fcb9fa3bb7640c23e9c748e53fb30934" dependencies = [ "unicode-ident", ] @@ -458,18 +458,18 @@ version = "0.1.0" [[package]] name = "quote" -version = "1.0.47" +version = "1.0.45" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "1fbf4db142a473a8d80c26bbf18454ed458bf8d26c8219c331daecfdbd079001" +checksum = "41f2619966050689382d2b44f664f4bc593e129785a36d6ee376ddf37259b924" dependencies = [ "proc-macro2", ] [[package]] name = "regex" -version = "1.13.1" +version = "1.12.3" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "f020237b6c8eed93db2e2cb53c00c60a8e1bc73da7d073199a1180401450218d" +checksum = "e10754a14b9137dd7b1e3e5b0493cc9171fdd105e0ab477f51b72e7f3ac0e276" dependencies = [ "aho-corasick", "memchr", @@ -479,9 +479,9 @@ dependencies = [ [[package]] name = "regex-automata" -version = "0.4.18" +version = "0.4.14" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "ad8553b9b26413251cbf30e620595c7a41b3887f03da04579c0e6b0d6a06b4b2" +checksum = "6e1dd4122fc1595e8162618945476892eefca7b88c52820e74af6262213cae8f" dependencies = [ "aho-corasick", "memchr", @@ -519,9 +519,9 @@ dependencies = [ [[package]] name = "serde" -version = "1.0.229" +version = "1.0.228" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "4148590afebada386688f18773da617792bf2ef03ffc1e4cbd2b1d45b023e0ba" +checksum = "9a8e94ea7f378bd32cbbd37198a4a91436180c5bb472411e48b5ec2e2124ae9e" dependencies = [ "serde_core", "serde_derive", @@ -529,22 +529,22 @@ dependencies = [ [[package]] name = "serde_core" -version = "1.0.229" +version = "1.0.228" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "67dca2c9c51e58a4791a4b1ed58308b39c64224d349a935ab5039aa360942a48" +checksum = "41d385c7d4ca58e59fc732af25c3983b67ac852c1a25000afe1175de458b67ad" dependencies = [ "serde_derive", ] [[package]] name = "serde_derive" -version = "1.0.229" +version = "1.0.228" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "e7a5d71263a5a7d47b41f6b3f06ba276f10cc18b0931f1799f710578e2309348" +checksum = "d540f220d3187173da220f885ab66608367b6574e925011a9353e4badda91d79" dependencies = [ "proc-macro2", "quote", - "syn 3.0.9", + "syn", ] [[package]] @@ -579,20 +579,9 @@ checksum = "0fda2ff0d084019ba4d7c6f371c95d8fd75ce3524c3cb8fb653a3023f6323e64" [[package]] name = "syn" -version = "2.0.119" +version = "2.0.117" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "872831b642d1a07999a962a351ed35b955ea2cfc8f3862091e2a240a84f17297" -dependencies = [ - "proc-macro2", - "quote", - "unicode-ident", -] - -[[package]] -name = "syn" -version = "3.0.9" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "d78c8dee4c7bf0e14673097256fed6142ce9d3b85a408189d07482442145823b" +checksum = "e665b8803e7b1d2a727f4023456bbbbe74da67099c585258af0ad9c5013b9b99" dependencies = [ "proc-macro2", "quote", @@ -601,22 +590,22 @@ dependencies = [ [[package]] name = "thiserror" -version = "2.0.21" +version = "2.0.18" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "09e52cb86a36cede5cb101bf8908837b3e4c6e5e59fe7fd85c23fb56200d189e" +checksum = "4288b5bcbc7920c07a1149a35cf9590a2aa808e0bc1eafaade0b80947865fbc4" dependencies = [ "thiserror-impl", ] [[package]] name = "thiserror-impl" -version = "2.0.21" +version = "2.0.18" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "fe5197923287db20a58125f0bc85c062f7f2c892de97b18c356f9efb14b28524" +checksum = "ebc4ee7f67670e9b64d05fa4253e753e016c6c95ff35b89b7941d6b856dec1d5" dependencies = [ "proc-macro2", "quote", - "syn 3.0.9", + "syn", ] [[package]] @@ -627,9 +616,9 @@ checksum = "b6f5e870be6c3b371b77fe0ee0bafb859fa4964b4404c27de1d380043c4dda20" [[package]] name = "unicode-ident" -version = "1.0.27" +version = "1.0.24" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "a2c754d6c33795a1c324727428e5a7dedb5b06195f9890bdbcba760d3e246563" +checksum = "e6e4313cd5fcd3dad5cafa179702e2b244f760991f45397d14d4ebf38247da75" [[package]] name = "version_check" @@ -739,22 +728,22 @@ checksum = "589f6da84c646204747d1270a2a5661ea66ed1cced2631d546fdfb155959f9ec" [[package]] name = "zerocopy" -version = "0.8.63" +version = "0.8.52" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "e5fe1f8f1b06191a00962174c61aa5005e0bb391a6d80d07e24d115c01a92ed8" +checksum = "ce1022995ff5ff5d841ad7d994facc23098cd40152f2c1d11cd607c6f530653f" dependencies = [ "zerocopy-derive", ] [[package]] name = "zerocopy-derive" -version = "0.8.63" +version = "0.8.52" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "863ad3ac83293fb4d740aedbfdc9240dd8d1a50c1099acd76ce80ce7c7230c7f" +checksum = "1ae7f38b72ec2a254e2b87ef277cf2cd4fb97cbebf944faa6f33354da0867930" dependencies = [ "proc-macro2", "quote", - "syn 2.0.119", + "syn", ] [[package]] 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