Files
foxhunt/crates/ml/tests
jgrusewski 3cbb0b761d fix: curiosity kernel — replace fused kernel with separate fwd_bwd + adam
The fused curiosity kernel (curiosity_fused_zero_fwd_bwd_adam) used
__threadfence() + __syncthreads() + atomicAdd block counter for
inter-block synchronization. This crashed on RTX 3050 (and potentially
other consumer GPUs) — the synchronize() after launch deadlocked.

Fix: use separate kernels instead:
1. curiosity_shift_states (shift next_states)
2. curiosity_forward_backward (accumulate gradients via atomicAdd)
3. curiosity_adam_step (4 launches, one per param group)

This adds 4 kernel launches (~5μs overhead) but eliminates the
fragile inter-block sync pattern. Gradients are zeroed via
memset_zeros before forward_backward (GPU-side, stream-ordered).

Curiosity re-enabled in smoke test. Full e2e DQN training passes
on RTX 3050 in 0.67s with branching + C51(11) + curiosity + NoisyNets.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
2026-03-20 11:15:03 +01:00
..