The data preload created a dummy DQNTrainer with buffer_size=1, which
fails the GPU PER capacity check (capacity must be >= batch_size).
This caused the preload to fail silently, falling back to per-trial
data loading from disk — 50 × 5s = 250s wasted.
Fix: set buffer_size = max(batch_size, 1024) so the PER allocation
succeeds. The preload trainer doesn't train — it just loads data into
a shared Arc for all trials.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Three issues fixed:
1. Best checkpoint was loaded via safetensors → Candle VarMap, but the
evaluator now reads from fused_ctx GPU buffers. The VarMap weights
were never used — evaluator always saw last-epoch weights, not best.
2. Added save_best_params/restore_best_params to FusedTrainingCtx:
- save: async DtoD copy params_bf16 → best_params_snapshot
- restore: async DtoD copy best_params_snapshot → params_bf16 + unflatten
Zero sync, zero Candle, zero safetensors roundtrip.
3. Hyperopt evaluation now calls restore_best_gpu_params() instead of
loading safetensors checkpoint. Pure GPU path end-to-end.
Added params_bf16()/params_bf16_mut() accessors to GpuDqnTrainer.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Same bug as the training loop validation: hyperopt's backtest evaluator
read weights from agent.get_q_network_vars() (Candle VarMap) which is
NOT synced during fused CUDA training. All previous hyperopt runs were
evaluating with epoch-0 weights — Sharpe scores were meaningless.
Fix: read directly from fused_ctx.online_dueling_ref() /
online_branching_ref() via new DQNTrainer::fused_online_weights() method.
Falls back to error if fused_ctx not available (required in CUDA builds).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Verified: grad_norm=1.03 stable across all 10 epochs on H100. The
previous grad_norm=0.0 at epoch 5 was a transient issue, not a kernel
bug. The two-phase reduction is correct.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Temporary sync+readback in compute_grad_norm_outside_graph to catch
when grad_norm_f32_buf is exactly 0. Also reads grad_buf[0..4] to
determine if gradients themselves are zero vs finalize kernel bug.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Changed best_epoch<=1 check from tracing::warn to assert!. With 3+
epochs, best_epoch must be >1 — if it's always 1, the evaluator has
stale weights or state. This would have caught both the VarMap bug
and the done_flags persistence bug.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
The GPU backtest evaluator's done_flags persisted across epochs —
after epoch 1's evaluation, done_flags[0]=1, causing subsequent
evaluations to skip the entire simulation (constant val_loss).
Added reset_evaluation_state() called before each evaluate_dqn_graphed:
- Re-init portfolio to initial capital
- Zero done_flags, step_returns, step_rewards, actions_history
This was the root cause of val_loss=0.828125 every epoch.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Warns when best_epoch==1 with 3+ epochs trained, which indicates the
evaluator may be reading stale weights (constant val_loss). This would
have caught the Candle VarMap bug where val_loss=0.804688 every epoch.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
The GPU evaluator was reading weights from agent.get_q_network_vars()
(Candle VarMap) which is NOT synced during fused CUDA training. Weights
stayed at epoch-0 values → val_loss=0.804688 every epoch (constant).
Fix: read weights directly from fused_ctx.online_dueling_ref() and
fused_ctx.online_branching_ref() — the GPU-resident buffers updated
by Adam each training step.
Added online_dueling_ref() and online_branching_ref() accessor methods
to FusedTrainingCtx.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Three optimizations to IQN forward+loss and backward kernels:
1. Block size 32→256 (8 warps): occupancy 3.1%→25% on H100.
Inner hidden_dim loops now stride by 256 instead of 32.
Block-level reduction via shared memory replaces warp-only shuffle.
2. Precomputed cosine features [N, embed_dim]: eliminates 16.7M
cosf() calls per step. cos(π·(d+1)·τ_i) computed once at
construction (τ are fixed QR-DQN midpoints).
3. Quantile Huber loss distributed across 256 threads: was single-lane
serial (32×32=1024 iterations on lane 0). Now each thread handles
~4 pairs.
Expected impact: IQN from ~80ms to ~10-15ms per step (occupancy +
cosine elimination + parallel loss).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
New backtest_env_step_batch kernel processes all chunk_len steps in a
single launch. Each thread handles one window, loops over steps
sequentially reading from the chunked_actions buffer. Portfolio state
stays in registers across the step loop — zero global memory round-trips
between steps.
Eliminates 512 individual kernel launches per chunk (was: DtoD copy +
env_step per step = 1024 launches per chunk). With 301 chunks for 154K
bars: ~154K launches → 301 launches. Kernel launch overhead drops from
~1.15s to ~2.3ms.
Combined with chunk 64→512 + sync reduction: validation expected to drop
from 2.5s to <0.5s.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
8× fewer cuBLAS forward passes for validation (2406→301 chunks for 154K
bars). Removed 3 chunk-0 debug syncs that stalled the pipeline. Reduced
per-chunk sync from every chunk to every 10th (301→31 pipeline stalls).
Portfolio features stale for ~8.5h within each chunk — acceptable since
market features (42/48 dims) dominate action selection. Causal
sensitivity analysis confirms portfolio features have <2% importance.
Expected validation: 2.5s → ~0.5-1.0s per epoch.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Replaced the CPU validation path (compute_validation_loss) with the
existing GpuBacktestEvaluator. The CPU path downloaded Q-values to host,
did argmax + PnL + Sharpe on CPU — 2s/epoch from the GPU→CPU sync.
The GPU evaluator does everything on-device: cuBLAS forward, argmax,
PnL simulation, Sharpe/drawdown/PF computation. Lazy-initialized once
per fold from val_data, called each epoch with current weights.
Deleted:
- val_features_gpu, val_closes_gpu, val_ofi_gpu fields (CPU caches)
- CPU state tensor assembly (80-line loop)
- clone_htod_f32_to_bf16 upload
- dtoh_bf16_to_f32 download + GPU→CPU sync
- CPU argmax loop
- CPU Sharpe computation
- Epsilon save/restore
Net -61 lines of dead CPU code.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Changed 2 backward cuBLAS GEMMs from CUBLAS_COMPUTE_32F to
CUBLAS_COMPUTE_32F_FAST_TF32. Both activation gradient (BF16→BF16)
and weight gradient (BF16→F32) paths now use H100 tensor cores.
TF32's 10-bit mantissa is sufficient for gradients — Adam's momentum
exponential moving average smooths any additional noise. Forward
GEMMs already use TF32 (unchanged).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Removed 19 diagnostic eprintln + 7 stream.synchronize() calls that were
added during H100 hang debugging. All hang root causes are now fixed:
- grad_norm atomicAdd → two-phase reduction
- vaccine dot+norm atomicAdd → two-phase reduction
- C51 mixup spin-wait barrier → split into two kernels
Zero diagnostic overhead in the training hot loop.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
The C51 manifold mixup used a spin-wait inter-block barrier that
deadlocks on H100 when batch_size > co-resident blocks. No grid cap
or grid-striding can reliably fix this — shared memory per block limits
actual occupancy below theoretical max.
Proper fix: split into two separate kernel launches:
1. c51_loss_batched: projection + CE loss (grid=batch_size, no barrier)
2. c51_mixup_ce: reads save_projected, mixes with partner, overwrites
loss (grid=batch_size, no barrier)
CUDA stream ordering guarantees kernel 1 completes before kernel 2
reads the projections. Zero spin-waits, zero atomicAdd barriers,
zero deadlock risk on any GPU architecture.
Removed: mixup_barrier_buf memset, all barrier code from C51 kernel.
Added: c51_mixup_ce_kernel field, compile, launch method.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Root cause: C51 loss kernel launched with grid=batch_size (8192) blocks
and a spin-wait inter-block barrier for manifold mixup. On H100 only
~500 blocks can be co-resident → deadlock (blocks 500-8191 wait for
scheduling while blocks 0-499 spin-wait for block 8191).
Fix (NVIDIA two-phase approach):
- Grid capped to min(batch_size, 512) co-resident blocks
- Grid-strided sample loop: each block processes multiple samples
- Phase 1: all blocks write save_projected for ALL samples (no barrier)
- Barrier: waits for gridDim.x (all co-resident), not batch_size
- Phase 2: new grid-strided loop reads partner projections, mixes, computes CE
Non-mixup path (mixup_alpha=0): CE loss computed inline, no barrier.
Mixup path: two separate sample loops with single barrier between them.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
gradient_dot_and_norm kernel used atomicAdd with 1264 blocks to compute
dot(g_train, g_val) and |g_val|² — same H100 L2 contention pattern as
the grad_norm hang. Replaced with two-phase: per-block partials + single-
block finalize (gradient_dot_norm_finalize). Zero atomicAdd.
Also added step-0-only per-phase GPU syncs to catch any remaining hangs
immediately (forward, aux, conditional, Adam, PER — only on step 0).
TODO: Causal intervention runs 14 full cuBLAS forward passes per step.
Should be optimized (batched or reduced frequency).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Syncs after forward, HER, aux, conditional, Adam, PER — only on step 0.
Steps 1+ run fully async. This will pinpoint the exact hanging phase
without impacting overall training speed.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Previous commit missed 2 IQN/ensemble trunk grad_norm sites that still
passed a single-float norm buffer to the new block_sums kernel — buffer
overflow on GPU. Now all 4 call sites (main, CQL, IQN trunk, ensemble
trunk) use the pre-allocated grad_norm_partials buffer with two-phase
reduction. Zero atomicAdd, zero memset across the entire training step.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Root cause: dqn_grad_norm_kernel used atomicAdd to a single float with
1264 blocks. On H100 (132 SMs, all blocks co-resident), this created
extreme L2 cache contention causing the kernel to take seconds/hang.
On RTX 3050 (20 SMs), only ~200 blocks were resident, masking the issue.
Fix: two-phase reduction with zero atomicAdd:
- Phase 1: each block writes its partial sum to block_sums[blockIdx.x]
- Phase 2: single block reduces partials → f32 sum-of-squares + bf16 norm
Also removed:
- 3x memset_zeros on grad_norm_f32_buf (no longer needed — no atomicAdd)
- All diagnostic eprintln and stream.synchronize() from fused_training.rs
- All FUSED_DIAG/TRAIN_DIAG/ADAM_DIAG logging
- TF32 confirmed NOT the cause — kept CUBLAS_COMPUTE_32F_FAST_TF32
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
TF32 confirmed NOT the cause — graph_forward GPU sync passes in 3.5ms
with TF32 enabled. The hang is in the Adam phase (after aux ops sync).
Added granular sync inside replay_adam_and_readback to isolate whether
grad_norm reduction or graph_adam replay hangs.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
TF32 (CUBLAS_COMPUTE_32F_FAST_TF32) causes kernel hang inside CUDA Graph
replay on H100 SM 9.0. Reverted forward GEMMs to CUBLAS_COMPUTE_32F.
Backward GEMMs already use F32 for gradient precision.
Added per-phase stream.synchronize() with timing after graph_forward,
aux ops (IQL/IQN/attention), and Adam to isolate any remaining hang.
If TF32 was the sole cause, all syncs will pass and we'll see epoch
timing data. If another op hangs, the logs pinpoint which phase.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Temporary diagnostics between each fused training step to identify
which operation hangs after replay_forward on H100.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Eliminated 5 per-step GPU allocations:
- insert_prio_buf, insert_idx_buf, insert_ep_buf: 3x cuMemAlloc per
insert_batch/insert_batch_bf16 call (up to 4M elements each)
- scratch_f32: reusable 1-element temp for flush_max_priority and
apply_max_priority_scalar (was 2x a32f per call)
- update_batch_spare: pre-allocated spare for None→Some priority swap,
replenished by flush_max_priority (was 1x a32f per epoch start)
cuMemAlloc is a synchronizing driver call that stalls the GPU pipeline.
With 8192 batch_size and multiple inserts per epoch, this was measurable
overhead on H100.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
The per_prefix_scan kernel was processing capacity/256 tiles (96K on H100
with 24.7M auto-sized capacity) instead of size/256 tiles (~3200 for 819K
experiences). This caused the scan to take ~50ms per step, appearing as a
hang (6.5 min/epoch). Now computes actual_tiles from self.size at scan
time — all tiles fit in one round of 512 resident blocks (<1ms).
Removed dead num_scan_tiles field. Added regression test with 1M capacity
and 256 experiences to catch the high capacity/size ratio failure mode.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
With 96K tiles on H100 (132 SMs, ~500 resident blocks), the original
per_prefix_scan launched one block per tile. Blocks waiting on
predecessors that weren't scheduled created circular deadlocks.
Fix: blocks grab tiles dynamically via atomicAdd on a global counter.
Grid size capped at 512 (max resident). Each block processes multiple
tiles in a while loop, ensuring all tiles complete without preemption.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Remove diagnostic eprintln calls from the H100 hang investigation:
- H100_STEP, H100_LOOP, H100_HANG4 prefixes in fused_training.rs
- adam_readback prefix + ADAM_CTR static counter in gpu_dqn_trainer.rs
- H100_LOOP + FUSED STEP ERROR in training_loop.rs
- Debug cuStreamSynchronize block after PER update in fused_training.rs
Integration tests and smoke tests already match current signatures
(ext_stream: Option<&Arc<CudaStream>> with None). seg_tree_kernel.cu
already deleted in prior commit.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Replace segment tree sampling with prefix scan + binary search:
- sample_proportional: per_prefix_scan builds prefix sum, per_sample does
Philox RNG threshold search with binary search over prefix sum
- insert_batch/insert_batch_bf16: per_insert_pa writes priority^alpha
directly to priorities_pa (no rebuild needed, scan runs at sample time)
- Delete seg_tree_kernel.cu (no longer referenced)
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
The full tree rebuild (25 levels × 33M nodes) took too long on H100.
Only 8192 of 33M leaves change per step — rebuilding all nodes is
4000× wasted work.
Restore atomicAdd delta propagation in seg_tree_update_leaves:
O(B × log N) = 8192 × 25 = 200K atomics vs O(N) = 33M reads.
atomicExch for leaf writes (duplicate-index safe).
Keep seg_tree_rebuild_level for insert_batch (new slots need full rebuild).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
cuModuleLoadData hangs on H100 CUDA 13 when called after CUDA graph
capture — the context state appears incompatible with module loading.
Pre-initialize the OnceLock during GpuReplayBuffer::new() when the
context is clean and no graphs have been captured.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
get_cast_kernels uses OnceLock + cuModuleLoadData which requires the
CUDA context to be bound to the current thread. When ext_stream (the
trainer's stream) was passed, its context wasn't bound — causing a hang
in cuModuleLoadData on H100.
Fix: always use self.stream (replay buffer's stream, with bound context)
for kernel compilation. Only use ext_stream for the actual kernel launches.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
The atomicAdd delta propagation caused severe L2 cache contention on
H100: 8192 threads all atomicAdd to tree[1] (root) serialized through
a single cacheline, taking 10+ seconds per PER update.
Split into two phases:
1. seg_tree_update_leaves: write all leaves in parallel (zero contention)
2. seg_tree_rebuild_level: rebuild one tree level per kernel launch,
bottom-up. 20 levels × ~2µs = ~40µs total. Each level is fully
parallel — nodes at the same level are independent.
Also applies to seg_tree_insert (insert_batch path).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
The replay buffer used its own CUDA stream for seg_tree_update, but
the td_errors buffer was computed on the trainer's stream. Without
synchronization, the replay stream could read incomplete data — causing
a hang on H100 where streams execute truly concurrently.
Fix: update_priorities_gpu now accepts an optional ext_stream parameter.
The fused training path passes the trainer's main stream, ensuring all
GPU work (td_errors computation + seg_tree_update) runs on the same
stream with implicit ordering.
Also pre-allocate update_td_f32, update_batch_max, update_max_merge
scratch buffers at construction — eliminates 3 cuMemAlloc calls per
training step that could trigger implicit device synchronization.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
PER samples with replacement, so a batch can contain duplicate indices.
With non-atomic leaf writes, two threads processing the same index both
read the same old_pa, but only one write survives — internal nodes get
double the delta while the leaf changed once, permanently corrupting
the tree sum.
atomicExch returns the exact replaced value so each thread propagates
the correct delta regardless of concurrent duplicate writes.
Also fix readback_pinned layout doc: slot 9 is used for diversity loss.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Capture the batched spectral norm kernel (10 matrices) + bf16→f32 sync
into a CUDA graph on step 1, replayed on steps 2+ with zero launch
overhead. Spectral norm uses fixed device pointers (params_bf16,
spec_u/v descriptors) that never reallocate, making it a perfect graph
capture candidate. Graph is invalidated on fold reset.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Remove 3 memset_zeros calls on buffers that are fully overwritten by
their consumer kernels (no atomicAdd accumulation):
- q_stats_buf (2 sites): q_stats_reduce is a single-thread kernel that
directly assigns all 5 output floats — no accumulation pattern
- causal_mean_scratch (1 site): causal_mean_reduce writes out[0] via
direct assignment in thread 0 — no atomicAdd
All remaining memsets are documented as REQUIRED with the specific
accumulation pattern (atomicAdd / beta=1.0 GEMM) that necessitates them.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Audit all grid=(1,1,1) kernels called per training step:
- grad_norm_finalize: cannot fuse into grad_norm_kernel because the
reduction uses multi-block atomicAdd — no block knows it holds the
final sum without cooperative launch. 1-thread finalize is negligible.
- stochastic_depth_rng: standalone RNG writing 3 floats, not a finalizer
after a reduction — nothing to fuse into.
- causal_reduce, causal_mean_reduce, q_stats: only called at
intervention-interval or epoch-boundary, not per training step.
Document rationale inline so future reviewers don't re-investigate.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
