diff --git a/crates/ringkernel-cuda/Cargo.toml b/crates/ringkernel-cuda/Cargo.toml index efda9d7..a2d5814 100644 --- a/crates/ringkernel-cuda/Cargo.toml +++ b/crates/ringkernel-cuda/Cargo.toml @@ -73,3 +73,28 @@ multi-gpu = ["cuda", "nvml-wrapper"] # When enabled, a build-time linker search path is added; users whose # NVSHMEM install lives outside /usr can set NVSHMEM_LIB_DIR. nvshmem = ["cuda", "multi-gpu"] + +# `cuda-1xxxx` feature names below are cudarc's own version features +# (propagated transitively via `cuda-version-from-build-system`), not +# declared here directly — silences rustc's `unexpected_cfgs` lint for +# `#[cfg(feature = "cuda-12040")]`-style checks against them in hopper/mod.rs. +[lints.rust] +unexpected_cfgs = { level = "allow", check-cfg = [ + 'cfg(feature, values("cuda-11040"))', + 'cfg(feature, values("cuda-11050"))', + 'cfg(feature, values("cuda-11060"))', + 'cfg(feature, values("cuda-11070"))', + 'cfg(feature, values("cuda-11080"))', + 'cfg(feature, values("cuda-12000"))', + 'cfg(feature, values("cuda-12010"))', + 'cfg(feature, values("cuda-12020"))', + 'cfg(feature, values("cuda-12030"))', + 'cfg(feature, values("cuda-12040"))', + 'cfg(feature, values("cuda-12050"))', + 'cfg(feature, values("cuda-12060"))', + 'cfg(feature, values("cuda-12080"))', + 'cfg(feature, values("cuda-12090"))', + 'cfg(feature, values("cuda-13000"))', + 'cfg(feature, values("cuda-13010"))', + 'cfg(feature, values("cuda-13020"))', +] } diff --git a/crates/ringkernel-cuda/README.md b/crates/ringkernel-cuda/README.md index b10799f..a6bc492 100644 --- a/crates/ringkernel-cuda/README.md +++ b/crates/ringkernel-cuda/README.md @@ -10,8 +10,11 @@ managing persistent GPU kernels. ## Requirements -- NVIDIA GPU with Compute Capability 7.0 or higher (Volta, Turing, Ampere, Ada, - Hopper, Blackwell) +- NVIDIA GPU with Compute Capability 6.0 or higher (Pascal and newer) for the + core persistent-actor/cooperative-groups path. Some features have higher + floors — Thread Block Clusters, DSMEM, TMA, and Green Contexts all require + Hopper (CC 9.0+). See the top-level README's "Feature to minimum compute + capability" table for the full per-feature breakdown. - CUDA Toolkit 12.x or later - cudarc 0.19.3 (pinned via workspace) - Linux (native) or Windows (WSL2, with cooperative-group limitations) diff --git a/crates/ringkernel-cuda/src/hopper/mod.rs b/crates/ringkernel-cuda/src/hopper/mod.rs index f140771..fd3c3ae 100644 --- a/crates/ringkernel-cuda/src/hopper/mod.rs +++ b/crates/ringkernel-cuda/src/hopper/mod.rs @@ -12,6 +12,22 @@ pub mod async_mem; pub mod cluster; pub mod dsmem; +// Green Contexts (CUDA 12.4+) — cudarc only exposes `CUgreenCtx` and its FFI +// functions behind these version features (matching NVIDIA's own +// introduction of the API in CUDA 12.4). Gating this at compile time, not +// just at runtime via `check_hopper_support()` below, so the crate still +// builds on any CUDA Toolkit older than 12.4 — which otherwise fails to +// compile at all, on any GPU, regardless of architecture. +#[cfg(any( + feature = "cuda-12040", + feature = "cuda-12050", + feature = "cuda-12060", + feature = "cuda-12080", + feature = "cuda-12090", + feature = "cuda-13000", + feature = "cuda-13010", + feature = "cuda-13020" +))] pub mod green_ctx; pub mod lifecycle; pub mod tma; @@ -43,3 +59,97 @@ pub const MAX_PORTABLE_CLUSTER_SIZE: u32 = 8; /// Maximum cluster size on Blackwell (B200). pub const MAX_BLACKWELL_CLUSTER_SIZE: u32 = 16; + +#[cfg(test)] +mod pascal_compat_tests { + //! Confirms Hopper-only feature checks degrade gracefully (return an + //! `Err`/`false`, not a panic or UB) on pre-Hopper hardware, now that the + //! compile-time gating fix above lets this module build at all on CUDA + //! Toolkits older than 12.4. Requires real GPU hardware — run with: + //! cargo test -p ringkernel-cuda --features cuda -- --ignored pascal_compat + use super::*; + use crate::device::CudaDevice; + + #[test] + #[ignore] // Requires CUDA hardware + fn pascal_compat_check_hopper_support_degrades_gracefully() { + let device = CudaDevice::new(0).expect("Failed to create device"); + let (major, minor) = device.compute_capability(); + println!("Device compute capability: {major}.{minor}"); + + let result = check_hopper_support(&device); + if major < 9 { + assert!( + result.is_err(), + "check_hopper_support returned Ok on pre-Hopper hardware" + ); + } else { + assert!(result.is_ok()); + } + } + + #[test] + #[ignore] // Requires CUDA hardware + fn pascal_compat_supports_cluster_launch_degrades_gracefully() { + let device = CudaDevice::new(0).expect("Failed to create device"); + let (major, _) = device.compute_capability(); + + let supported = supports_cluster_launch(&device); + if major < 9 { + assert!( + !supported, + "supports_cluster_launch returned true on pre-Hopper hardware" + ); + } + } + + #[test] + #[ignore] // Requires CUDA hardware + fn pascal_compat_dsmem_degrades_gracefully() { + let device = CudaDevice::new(0).expect("Failed to create device"); + let (major, _) = device.compute_capability(); + + let available = dsmem::is_dsmem_available(&device, 4); + if major < 9 { + assert!( + !available, + "is_dsmem_available returned true on pre-Hopper hardware" + ); + } + } + + #[test] + #[ignore] // Requires CUDA hardware + fn pascal_compat_cluster_size_degrades_to_one() { + let device = CudaDevice::new(0).expect("Failed to create device"); + let (major, _) = device.compute_capability(); + + let max_cluster = cluster::query_max_cluster_size(&device, std::ptr::null_mut()) + .expect("query_max_cluster_size should not error"); + if major < 9 { + assert_eq!( + max_cluster, 1, + "query_max_cluster_size did not degrade to 1 on pre-Hopper hardware" + ); + } + } + + #[test] + #[ignore] // Requires CUDA hardware + fn pascal_compat_cooperative_groups_still_available() { + // Cooperative groups (grid.sync()) have a lower floor (CC 6.0+, + // Pascal) than Hopper-specific features (CC 9.0+) — confirms the + // gating logic distinguishes per-feature floors correctly, rather + // than treating "not Hopper" as "nothing works". + let device = CudaDevice::new(0).expect("Failed to create device"); + let (major, _) = device.compute_capability(); + + let coop = device.supports_cooperative_groups(); + if major >= 6 { + assert!( + coop, + "supports_cooperative_groups returned false on CC 6.0+ hardware" + ); + } + } +} diff --git a/crates/ringkernel-cuda/src/lib.rs b/crates/ringkernel-cuda/src/lib.rs index b6a28bd..f941621 100644 --- a/crates/ringkernel-cuda/src/lib.rs +++ b/crates/ringkernel-cuda/src/lib.rs @@ -11,7 +11,11 @@ //! //! # Requirements //! -//! - NVIDIA GPU with Compute Capability 7.0+ +//! - NVIDIA GPU with Compute Capability 6.0+ (Pascal and newer) for the core +//! persistent-actor/cooperative-groups path. Some features have higher +//! floors (e.g. Thread Block Clusters/DSMEM/TMA/Green Contexts require +//! Hopper, CC 9.0+) — see the top-level README's "Feature to minimum +//! compute capability" table for the full breakdown. //! - CUDA Toolkit 11.0+ //! - Native Linux (persistent kernels) or WSL2 (event-driven fallback) //! @@ -227,12 +231,78 @@ pub fn compile_ptx(_cuda_source: &str) -> ringkernel_core::error::Result )) } -/// PTX kernel source template for persistent ring kernel. +/// PTX kernel source template for persistent ring kernel, generic over the +/// target compute capability (major, minor). /// /// This is a minimal kernel that immediately marks itself as terminated. -/// Uses PTX 8.0 / sm_75 as the lowest common denominator that supports -/// cooperative groups. PTX is forward-compatible, so sm_75 PTX runs on -/// sm_89/sm_90/sm_100 and newer GPUs. +/// +/// PTX `.target` directives are only forward-compatible (PTX built for +/// `sm_X` runs on `sm_X` and newer, never older) — so this must be +/// generated per-device rather than hardcoded to a single value. Maps the +/// device's actual compute capability to the corresponding PTX ISA +/// version, matching the minimum ISA version that introduced each `sm_XX` +/// target (per NVIDIA's PTX ISA documentation). +pub fn ring_kernel_ptx_template_for(major: u32, minor: u32) -> String { + let ptx_version = match (major, minor) { + // Pascal + (6, 0) | (6, 1) | (6, 2) => "5.0", + // Volta / Turing + (7, 0) | (7, 2) => "6.0", + (7, 5) => "6.3", + // Ampere + (8, 0) => "7.0", + (8, 6) | (8, 7) => "7.1", + (8, 9) => "8.0", + // Hopper + (9, 0) => "8.0", + // Blackwell and newer — fall through to the newest known ISA + (major, _) if major >= 10 => "8.5", + // Below Pascal (Maxwell/Kepler) or anything unrecognized: fall back + // to the oldest ISA version this template's instructions need. + _ => "5.0", + }; + format!( + r#" +.version {ptx_version} +.target sm_{major}{minor} +.address_size 64 + +.visible .entry ring_kernel_main( + .param .u64 control_block_ptr, + .param .u64 input_queue_ptr, + .param .u64 output_queue_ptr, + .param .u64 shared_state_ptr +) {{ + .reg .u64 %cb_ptr; + .reg .u32 %one; + + // Load control block pointer + ld.param.u64 %cb_ptr, [control_block_ptr]; + + // Mark as terminated immediately (offset 8) + mov.u32 %one, 1; + st.global.u32 [%cb_ptr + 8], %one; + + ret; +}} +"# + ) +} + +/// PTX kernel source template for persistent ring kernel — **deprecated** +/// fixed-`sm_75` fallback, kept only for API compatibility with existing +/// callers that don't have a device handle available. +/// +/// PTX built for `sm_75` will fail to load (`CUDA_ERROR_INVALID_PTX`) on any +/// GPU with a compute capability below 7.5 (e.g. Pascal, sm_61) — PTX +/// forward-compatibility only extends to equal-or-newer architectures. +/// Prefer [`ring_kernel_ptx_template_for`] with the actual target device's +/// compute capability (`CudaDevice::compute_capability`, not part of this +/// crate's public API surface). +#[deprecated( + since = "1.1.1", + note = "hardcodes sm_75, breaking on Pascal and older GPUs — use ring_kernel_ptx_template_for(major, minor) with the real device's compute capability instead" +)] pub const RING_KERNEL_PTX_TEMPLATE: &str = r#" .version 8.0 .target sm_75 diff --git a/crates/ringkernel-cuda/src/runtime.rs b/crates/ringkernel-cuda/src/runtime.rs index 4d6f937..946171a 100644 --- a/crates/ringkernel-cuda/src/runtime.rs +++ b/crates/ringkernel-cuda/src/runtime.rs @@ -20,7 +20,7 @@ use crate::device::CudaDevice; use crate::kernel::CudaKernel; use crate::memory::CudaMemoryPool; use crate::persistent::{PersistentSimulation, PersistentSimulationConfig}; -use crate::RING_KERNEL_PTX_TEMPLATE; +use crate::ring_kernel_ptx_template_for; /// CUDA runtime for RingKernel. pub struct CudaRuntime { @@ -263,6 +263,13 @@ impl RingKernelRuntime for CudaRuntime { let sim = PersistentSimulation::new(&self.device, sim_config)?; + // Generate PTX targeting the actual device's compute capability + // — PTX .target is forward-compatible only (never backward), so + // this must match (or be older than) the real device, not a + // hardcoded value. + let (cc_major, cc_minor) = self.device.compute_capability(); + let template_ptx = ring_kernel_ptx_template_for(cc_major, cc_minor); + // Determine PTX and function name for the persistent kernel. // Use cooperative kernel PTX from build.rs if available. #[cfg(feature = "cooperative")] @@ -274,10 +281,7 @@ impl RingKernelRuntime for CudaRuntime { "Cooperative PTX not available (nvcc not found at build time), \ falling back to template PTX" ); - ( - RING_KERNEL_PTX_TEMPLATE.to_string(), - "ring_kernel_main".to_string(), - ) + (template_ptx.clone(), "ring_kernel_main".to_string()) } else { (coop_ptx.to_string(), "coop_persistent_fdtd".to_string()) } @@ -288,15 +292,12 @@ impl RingKernelRuntime for CudaRuntime { tracing::warn!( "Cooperative feature not enabled, persistent simulation will use template PTX" ); - ( - RING_KERNEL_PTX_TEMPLATE.to_string(), - "ring_kernel_main".to_string(), - ) + (template_ptx.clone(), "ring_kernel_main".to_string()) }; // Still load PTX for the CudaKernel's module/function fields // (needed for state transition to Launched) - kernel.load_ptx(RING_KERNEL_PTX_TEMPLATE)?; + kernel.load_ptx(&template_ptx)?; let kernel = Arc::new(kernel); @@ -321,8 +322,9 @@ impl RingKernelRuntime for CudaRuntime { Ok(KernelHandle::new(id, kernel)) } else { - // Standard path: load template PTX - kernel.load_ptx(RING_KERNEL_PTX_TEMPLATE)?; + // Standard path: load template PTX, targeting the actual device. + let (cc_major, cc_minor) = self.device.compute_capability(); + kernel.load_ptx(&ring_kernel_ptx_template_for(cc_major, cc_minor))?; let kernel = Arc::new(kernel); self.kernels.write().insert(id.clone(), Arc::clone(&kernel));