2 Commits

Author	SHA1	Message	Date
jgrusewski	57a91567b4	fix(ci): PTX cache invalidation includes build.rs + all kernel dirs ... The script only hashed .cu/.cuh in crates/ml/. Missed: - build.rs changes (removing --use_fast_math → stale cubins cached) - crates/ml-dqn/src/*.cu (replay buffer, seg tree kernels) - crates/ml-core/src/cuda_autograd/*.cu - crates/ml-ppo/src/cuda_nn/*.cu Now hashes ALL kernel sources + ALL build.rs files. Any change to nvcc flags or kernel source invalidates the entire PTX cache. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>	2026-03-29 21:01:03 +02:00
jgrusewski	1fae917c22	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>	2026-03-15 11:58:40 +01:00