- Add n=0 early-return guard to reduce() to prevent CUDA divide-by-zero
- Add i32::MAX bounds check before casting n to prevent overflow
- Cache PTX compilation with OnceLock<Result<Ptx, String>> (compile once per process)
- Add Safety comment to unsafe block documenting slice/buffer invariants
- Fix doc comment: clarify 48-byte is GPU transfer size (12 × f32)
- Add shmem comment explaining kernel uses static __shared__ only
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
57fbf4d993