Merge pull request #12 from EliasVahlberg/pascal-compat-fixes

fix(cuda): compile-time gate for green_ctx, device-aware PTX target for pre-Turing GPUs

Author: mivertowski

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"))',
+] }

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)

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"
+            );
+        }
+    }
+}

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<String>
     ))
 }
 
-/// 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

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));