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
..
2026-03-12 01:44:03 +01:00
2026-03-13 11:56:56 +01:00
2026-03-12 01:44:03 +01:00
2026-02-25 23:20:44 +01:00
2026-03-04 15:29:15 +01:00
2026-03-01 23:41:13 +01:00
2026-02-25 21:19:32 +01:00
2025-10-30 01:26:02 +01:00
2026-03-04 23:46:36 +01:00
2026-02-25 23:24:53 +01:00
2026-03-04 23:46:36 +01:00
2026-02-26 09:04:58 +01:00
2026-03-15 11:58:40 +01:00
2026-03-11 22:00:13 +01:00
2026-02-24 10:32:21 +01:00
2026-02-26 00:50:25 +01:00
2026-02-26 00:50:25 +01:00
2026-03-11 22:00:13 +01:00
2026-02-26 00:50:25 +01:00
2026-03-15 11:58:40 +01:00
2026-03-13 11:54:22 +01:00
2026-03-04 22:08:31 +01:00
2026-02-24 10:32:21 +01:00
2026-03-04 22:08:31 +01:00
2026-03-04 22:08:31 +01:00
2026-03-04 22:08:31 +01:00