From 86dd850012b527af589ea58a67088966842a6ca9 Mon Sep 17 00:00:00 2001 From: Elias Vahlberg Date: Sun, 12 Jul 2026 20:16:05 +0200 Subject: [PATCH 1/3] fix(cuda): gate hopper::green_ctx behind cudarc's CUDA-12.4+ features green_ctx uses CUgreenCtx and its FFI functions, which cudarc only exposes behind the cuda-12040 through cuda-13020 feature flags (matching NVIDIA's own introduction of the Green Context API in CUDA 12.4). The hopper module declared 'pub mod green_ctx;' unconditionally, so ringkernel-cuda failed to compile at all with --features cuda on any CUDA Toolkit below 12.4 - regardless of GPU architecture. Confirmed on CUDA 12.0.140 (GTX 1080, Pascal): cargo build -p ringkernel --example basic_hello_kernel --features cuda failed with 14 unresolved-symbol errors before this fix. The module's own doc comment already claimed Hopper features 'gracefully fall back on older architectures' - true at runtime via check_hopper_support(), but there was no matching compile-time gate. This adds one, matching cudarc's own feature list exactly, plus a [lints.rust] check-cfg allowance so the transitively-inherited cudarc feature names don't trigger unexpected_cfgs warnings. Added 5 new #[ignore]-gated hardware tests (pascal_compat_tests module, run with --features cuda -- --ignored pascal_compat) confirming check_hopper_support, supports_cluster_launch, is_dsmem_available, and query_max_cluster_size all degrade gracefully (Err/false/1, not a panic) on real Pascal hardware now that the module compiles in. Also confirms supports_cooperative_groups (CC 6.0+ floor) correctly stays available, verifying the gating distinguishes per-feature floors rather than treating 'not Hopper' as 'nothing works'. All 5 passed on a GTX 1080. Co-Authored-By: Claude Sonnet 5 (1M context) --- crates/ringkernel-cuda/Cargo.toml | 25 ++++++ crates/ringkernel-cuda/src/hopper/mod.rs | 110 +++++++++++++++++++++++ 2 files changed, 135 insertions(+) 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/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" + ); + } + } +} From 7aa2e0e0184572117bcd019b889c570dbb0f62f5 Mon Sep 17 00:00:00 2001 From: Elias Vahlberg Date: Sun, 12 Jul 2026 20:16:16 +0200 Subject: [PATCH 2/3] fix(cuda): make RING_KERNEL_PTX_TEMPLATE target the actual device RING_KERNEL_PTX_TEMPLATE was hardcoded to '.target sm_75', with a comment claiming this was safe as a 'lowest common denominator' because 'PTX is forward-compatible, so sm_75 PTX runs on sm_89/sm_90/sm_100 and newer'. That forward-compatibility claim is correct in direction, but the lowest-common-denominator framing is backwards: PTX built for sm_75 can run on sm_75 and newer, never older. This produced CUDA_ERROR_INVALID_PTX on a GTX 1080 (Pascal, sm_61) and would affect any pre-Turing GPU (Pascal, Maxwell), despite the runtime's own supports_cooperative_groups() gating correctly reporting CC 6.0+ as the real floor for that feature. Replaces the single hardcoded const with ring_kernel_ptx_template_for(major, minor), which generates the PTX template with a .version/.target matched to the actual device's compute capability (queried via CudaDevice::compute_capability() at both call sites in runtime.rs - the cooperative/persistent-simulation launch path and the standard launch path). The old const is kept, marked #[deprecated], for API compatibility with any external callers that don't have a device handle available. Verified: cargo run -p ringkernel --example basic_hello_kernel --features cuda now completes the full persistent-actor lifecycle (launch, activate, status check, suspend, resume, terminate) on a GTX 1080, where it previously failed at kernel launch with CUDA_ERROR_INVALID_PTX. Co-Authored-By: Claude Sonnet 5 (1M context) --- crates/ringkernel-cuda/src/lib.rs | 74 +++++++++++++++++++++++++-- crates/ringkernel-cuda/src/runtime.rs | 26 +++++----- 2 files changed, 84 insertions(+), 16 deletions(-) diff --git a/crates/ringkernel-cuda/src/lib.rs b/crates/ringkernel-cuda/src/lib.rs index b6a28bd..e75565a 100644 --- a/crates/ringkernel-cuda/src/lib.rs +++ b/crates/ringkernel-cuda/src/lib.rs @@ -227,12 +227,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)); From fa08c18a2ebdf93cb02ed1d986fa50d73147996e Mon Sep 17 00:00:00 2001 From: Elias Vahlberg Date: Sun, 12 Jul 2026 20:47:31 +0200 Subject: [PATCH 3/3] docs(cuda): align crate-level Requirements with the actual CC floor ringkernel-cuda's own crate-level doc comment (rendered on docs.rs) and README.md (rendered on crates.io) both claimed 'Compute Capability 7.0+' as a blanket requirement. This contradicts the top-level project README's own 'Feature to minimum compute capability' table, which correctly states the core persistent-actor/cooperative-groups path floors at CC 6.0 (Pascal) - the same feature these two docs are describing. Before the two fixes earlier in this branch, CC 7.0+ was arguably an accurate practical statement (Pascal genuinely couldn't build or run this crate). After those fixes, it's simply incorrect - the core path now builds and runs correctly on Pascal (verified: this branch's basic_hello_kernel example completes its full lifecycle on a GTX 1080). Corrects both docs to CC 6.0+, with a pointer to the top-level README's detailed table for per-feature floors above that (Hopper-only features like clusters/DSMEM/TMA/Green Contexts still correctly require CC 9.0+ - unchanged, not part of this fix). Co-Authored-By: Claude Sonnet 5 (1M context) --- crates/ringkernel-cuda/README.md | 7 +++++-- crates/ringkernel-cuda/src/lib.rs | 6 +++++- 2 files changed, 10 insertions(+), 3 deletions(-) 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/lib.rs b/crates/ringkernel-cuda/src/lib.rs index e75565a..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) //!