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

[EliasVahlberg]

Summary

Two real bugs blocked ringkernel-cuda (--features cuda) from building/running on any pre-Turing GPU or CUDA Toolkit below 12.4, found and fixed while getting the persistent-actor path running on a GTX 1080 (Pascal, sm_61). A third commit aligns crate-level docs with the fix.

Both bugs are toolkit-version / PTX-direction bugs, not GPU-architecture-specific in cause — they'd affect anyone on an older CUDA Toolkit or pre-Turing card, not just Pascal.

What's broken today (on main)

1. hopper::green_ctx is declared unconditionally (pub mod green_ctx; in hopper/mod.rs), but its FFI calls need cudarc's cuda-12040+ feature (matching NVIDIA's own introduction of CUgreenCtx in CUDA 12.4). On CUDA Toolkit 12.0.140, cudarc's own version-detection correctly does not enable that feature — so ringkernel-cuda fails to compile at all with --features cuda, on any GPU. The module's doc comment already claims Hopper features "gracefully fall back on older architectures" — true at runtime via check_hopper_support(), but there was no matching compile-time gate.

2. RING_KERNEL_PTX_TEMPLATE is hardcoded to .target sm_75, with a comment claiming this is a safe "lowest common denominator" because "PTX is forward-compatible, so sm_75 PTX runs on sm_89/sm_90/sm_100 and newer." The forward-compat direction claim is correct, but PTX built for sm_75 cannot run on anything older — so this produces CUDA_ERROR_INVALID_PTX on any pre-Turing card. This also directly contradicts the "Feature to minimum compute capability" table in the top-level README, which correctly states grid.sync()/persistent kernels float at CC 6.0 (Pascal).

What this PR does

• 86dd850: cfg-gates green_ctx behind the exact same cuda-12040..cuda-13020 feature set cudarc itself uses, so the crate builds cleanly on any CUDA Toolkit. Adds a [lints.rust] check-cfg allowance for the transitively-inherited feature names. Adds 5 #[ignore]-gated hardware tests (hopper::pascal_compat_tests, 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 (not panic) on real Pascal hardware, and that supports_cooperative_groups (its own, lower CC 6.0 floor) correctly stays available — i.e. the gating distinguishes per-feature floors rather than treating "not Hopper" as "nothing works."
• 7aa2e0e: replaces the single hardcoded RING_KERNEL_PTX_TEMPLATE const with ring_kernel_ptx_template_for(major, minor), generating the PTX 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 old const is kept, marked #[deprecated], for compatibility with any external callers.
• fa08c18: ringkernel-cuda's crate-level doc comment and README.md both claimed a blanket Compute Capability 7.0+ floor, which — after the above fixes — no longer matches the top-level README's own table for this exact feature. Corrected to CC 6.0+, with a pointer to the detailed table for higher per-feature floors (Hopper-only features are unaffected/unchanged).

Verification

All on a real GTX 1080 (Pascal, sm_61), CUDA Toolkit 12.0.140:

cargo run -p ringkernel --example basic_hello_kernel --features cuda

completes the full persistent-actor lifecycle (launch, activate, status check, suspend, resume, terminate) where it previously failed — first at compile time (14 unresolved-symbol errors from green_ctx), then at kernel launch (CUDA_ERROR_INVALID_PTX) once bug 1 was worked around.

cargo test -p ringkernel-cuda --features cuda --lib pascal_compat -- --ignored --nocapture

5 passed, 0 failed — confirms Hopper-only feature checks degrade gracefully rather than panicking, and that lower-floor features (cooperative groups) remain correctly available.

Also checked against the project's CI gates locally:

• cargo fmt --all -- --check — clean
• cargo clippy -p ringkernel-cuda -p ringkernel -p ringkernel-core --lib --bins -- -D warnings — clean
• cargo doc -p ringkernel-cuda --no-deps — clean, no broken intra-doc links
• Full-workspace clippy/MSRV-1.75 checks hit two pre-existing issues (ringkernel-accnet/ringkernel-procint GUI crates failing an f32: From<f64> lint under a newer rustc; a cpufeatures transitive dep needing edition2024 under Cargo 1.75) — confirmed both exist identically on unmodified main, unrelated to this PR.

Notes

• Not claiming Pascal as an officially supported/tested tier for the whole project — only that the specific feature these docs and code describe (persistent kernels via cooperative groups) demonstrably builds and runs correctly at its own stated CC 6.0 floor, per the project's existing compatibility table.
• Happy to split into smaller PRs, rename things, or adjust conventions if that's easier to review — this was originally a side investigation while comparing GPU-in-Rust toolchains on older hardware, not planned as a contribution from the start, so let me know if anything doesn't fit how you'd want it structured.