perf(cuda): H100 kernel optimizations — nvcc pipeline, kernel fusion, GPU-only training

- Migrate from NVRTC JIT to cached nvcc -O3 for all CUDA kernels
- Fuse guard kernels, increase prefetch chunk, eliminate per-step GPU alloc
- H100-specific: fused Adam, warp reductions, shmem tiling, PPO occupancy
- Vectorize gather_states with __ldg() and 4x unroll
- sincosf() Box-Muller + paired Gaussian generation in noisy nets
- Shared-memory tiled branching DQN forward pass for sm_<90
- GPU-resident training guard kernel replacing Candle tensor ops
- Eliminate all to_vec1/to_vec2 CPU roundtrips, DtoD weight copy
- GPU PER mandatory everywhere — kill CPU replay path on CUDA
- Full GPU action masking — eliminate CPU fallback path
- Fix cuBLAS handle sharing via OnceLock (root cause of 49 cascade failures)
- Fix ILLEGAL_ADDRESS: scratch1_dist buffer overflow, stack sizing, curand determinism
- Fix CudaStream lifetime: bind before .context() to extend lifetime
- Keep raw cudarc buffers alive across epochs
- Add gpu-hotpath-guard.sh (37 patterns) and ptx-cache-invalidate.sh

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>

scripts/gpu-hotpath-guard.sh

@@ -93,13 +93,9 @@ LEAK_PATTERNS=(
     # Audit those separately with: scripts/cuda-perf-audit.sh
 )
 
-# ── Excluded paths (test-only infra that deliberately uses CPU) ────────
+# ── Excluded paths ────────────────────────────────────────────────────
+# NO test exclusions — all DQN/CUDA code must be GPU-only, including tests.
 EXCLUDE_PATHS=(
-    "smoke_tests/"
-    "_test.rs"
-    "_tests.rs"
-    "performance_tests.rs"
-    "performance_validation.rs"
     "demo_dqn.rs"
 )
 
@@ -127,10 +123,6 @@ check_file() {
         return 0
     fi
 
-    # Find the line number where #[cfg(test)] starts — everything after is test code
-    local test_start
-    test_start=$(grep -n '#\[cfg(test)\]' "$file" 2>/dev/null | head -1 | cut -d: -f1 || echo "999999")
-
     for pattern in "${LEAK_PATTERNS[@]}"; do
         # Find matches, exclude: doc comments, Storage match arms, regular comments
         local hits
@@ -142,27 +134,14 @@ check_file() {
             | grep -v 'Storage::Metal' \
             || true)
 
-        # Filter out lines in test code (line number >= test_start)
         if [ -n "$hits" ]; then
-            local filtered=""
-            while IFS= read -r line; do
-                local lineno
-                lineno=$(echo "$line" | cut -d: -f1)
-                if [ "$lineno" -lt "$test_start" ] 2>/dev/null; then
-                    filtered="${filtered}${line}"$'\n'
-                fi
-            done <<< "$hits"
-            filtered="${filtered%$'\n'}"  # trim trailing newline
-
-            if [ -n "$filtered" ]; then
-                if [ "$found" -eq 0 ]; then
-                    echo "GPU HOT-PATH LEAK: $file"
-                    found=1
-                fi
-                echo "$filtered" | while IFS= read -r line; do
-                    echo "  $line"
-                done
+            if [ "$found" -eq 0 ]; then
+                echo "GPU HOT-PATH LEAK: $file"
+                found=1
             fi
+            echo "$hits" | while IFS= read -r line; do
+                echo "  $line"
+            done
         fi
     done
 

scripts/ptx-cache-invalidate.sh

@@ -0,0 +1,52 @@
+#!/usr/bin/env bash
+# ptx-cache-invalidate.sh — Invalidate PTX cache when CUDA kernels change.
+#
+# The Rust runtime caches compiled PTX in $CARGO_TARGET_DIR/.ptx_cache/ (or
+# /tmp/.ptx_cache/).  The cache key includes a SHA-256 of the kernel source,
+# so changed .cu files naturally miss.  However, if the CI PVC preserves the
+# cache across runs and the build binary is stale, old PTX can survive.
+#
+# This script computes a content hash of all .cu/.cuh kernel source files and
+# compares it against a stamp file in the cache directory.  If the hash differs
+# (any kernel source changed), the entire PTX cache is purged, forcing nvcc
+# recompilation on the next run.
+#
+# Usage:  scripts/ptx-cache-invalidate.sh [cache_dir]
+#   cache_dir  defaults to $CARGO_TARGET_DIR/.ptx_cache, then /tmp/.ptx_cache
+
+set -euo pipefail
+
+KERNEL_DIR="crates/ml/src/cuda_pipeline"
+CACHE_DIR="${1:-${CARGO_TARGET_DIR:+${CARGO_TARGET_DIR}/.ptx_cache}}"
+CACHE_DIR="${CACHE_DIR:-/tmp/.ptx_cache}"
+
+# Compute content hash of all CUDA source files
+HASH=$(find "$KERNEL_DIR" -type f \( -name '*.cu' -o -name '*.cuh' \) \
+       -exec sha256sum {} + 2>/dev/null | sort | sha256sum | cut -d' ' -f1)
+
+STAMP_FILE="${CACHE_DIR}/.kernel_hash"
+
+if [ -f "$STAMP_FILE" ]; then
+    OLD_HASH=$(cat "$STAMP_FILE")
+    if [ "$HASH" = "$OLD_HASH" ]; then
+        echo "PTX cache valid — kernel hash unchanged ($HASH)"
+        exit 0
+    fi
+    echo "PTX cache STALE — kernel hash changed"
+    echo "  old: $OLD_HASH"
+    echo "  new: $HASH"
+else
+    echo "PTX cache stamp missing — first run or cache cleared"
+fi
+
+# Purge stale PTX cache
+if [ -d "$CACHE_DIR" ]; then
+    PTX_COUNT=$(find "$CACHE_DIR" -name '*.ptx' -type f 2>/dev/null | wc -l)
+    echo "Purging ${PTX_COUNT} cached PTX files from ${CACHE_DIR}"
+    rm -f "${CACHE_DIR}"/*.ptx
+fi
+
+# Write new stamp
+mkdir -p "$CACHE_DIR"
+echo "$HASH" > "$STAMP_FILE"
+echo "PTX cache stamp updated: $HASH"