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>
Replace 10 individual `spectral_norm_kernel` launches (each grid=(1,1,1))
with a single `spectral_norm_batched` launch (grid=(10,1,1)) that processes
all weight matrices in parallel across 10 blocks.
- Build descriptor buffer at construction (10 × 6 u64 entries with W/u/v
pointers, out_dim, in_dim per matrix) — pointers are stable
- Delete `spectral_norm_kernel` from dqn_utility_kernels.cu (replaced by
`spectral_norm_batched` which was already written)
- Remove spec_norm! macro and per-matrix launch loop
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Switch all 3 forward-pass cublasGemmEx calls from CUBLAS_COMPUTE_32F to
CUBLAS_COMPUTE_32F_FAST_TF32, enabling H100 TF32 tensor cores for 2-3x
throughput. Backward pass remains at full F32 precision for gradient stability.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Refactor spectral normalization to operate directly on params_bf16 via
raw u64 offset pointers (computed by bf16_weight_ptrs_from_base) instead
of separate DuelingWeightSet/BranchingWeightSet CudaSlice allocations.
This eliminates 10 memcpy_dtod_async calls per training step (~25MB/step,
25GB/epoch at 1000 steps) that previously synced spectrally-normalized
weights from per-layer slices back to the flat params_bf16 buffer.
Key changes:
- apply_spectral_norm now takes &DuelingWeightSet (immutable) instead of
&mut, only needed for first-call flatten into params_bf16
- Spectral norm kernels receive raw u64 pointers at GOFF_* offsets into
params_bf16, writing in-place — no sync copies needed
- Spectral norm u/v vectors remain separate allocations (unchanged)
- unflatten_online_weights promoted to pub(crate) for test access
- Tests updated to call unflatten after spectral norm for readback
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
seg_tree_update and seg_tree_insert used non-atomic
tree[node] = tree[2*node] + tree[2*node+1] with 8192 threads racing
on shared internal nodes. On H100 (132 SMs), all threads execute
simultaneously causing data races and hangs.
Replace with atomicAdd delta propagation: each thread computes
delta = new_leaf - old_leaf, then atomicAdd to every ancestor.
Commutative + associative = race-free. O(log n) per thread,
hardware-accelerated on SM 9.0.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Replace synchronous stream.synchronize() + memcpy_dtoh + CPU sum in
run_causal_intervention with a GPU reduction kernel (causal_mean_reduce)
and async DtoH to the pinned readback buffer at slot 8. The result is
double-buffered: the caller reads the previous iteration's value while
the current one transfers asynchronously.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Extend readback_pinned from 3 to 16 floats and write Q-stats DtoH
into slots [3..8] instead of the non-pinned q_stats_pending field.
On H100 CUDA 13 driver, non-pinned host memory causes cuMemcpyDtoHAsync
to degrade to a synchronous copy, creating a latent hang risk.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
GpuTrainResult was replaced by FusedStepResult (CPU-only f32 scalars).
F32ToU32Caster, get_cast_kernels_f32_to_u32, and f32_slice_to_gpu_tensor_gpu
were only used by the old pre-fused training path. Also removes
#[allow(dead_code)] from gather_f32 which is actively used.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Replace the orphaned per_update_kernel call site in fused_training.rs
with agent.update_priorities_from_td() which routes through
GpuReplayBuffer::update_priorities_gpu_raw -> seg_tree_update.
This kernel correctly writes priorities AND propagates the segment tree
from leaf to root, fixing stale PER sampling.
Add update_priorities_from_td() on ReplayBufferType and DqnAgentConfig.
Remove priorities_f32_ptr() which exposed raw pointers for the deleted
kernel. Remove batch_max epoch-boundary flush since seg_tree_update
propagates max_priority internally.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
The per_update_priorities_kernel scatter-wrote to the priorities[] buffer
but never updated the PER segment tree, leaving sampling priorities stale.
Remove the orphaned CUDA kernel, its cubin include, struct field,
compilation, and the update_priorities_cuda_raw / batch_max_readback_and_reset
methods from GpuDqnTrainer.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
GpuTrainResult::from_fused_scalars() allocated GPU memory and performed
two synchronous cuMemcpyHtoD copies after every training step. On H100
CUDA 13 driver, synchronous HtoD blocks until ALL pending async GPU work
completes — same class of bug as the DtoH pinned memory fix.
The uploaded scalars were never used: the training guard already reads
loss/grad_norm directly from GPU-resident buffers via loss_gpu_buf() /
grad_norm_gpu_buf(). Replace GpuTrainResult with lightweight
FusedStepResult (CPU-only f32 scalars, zero GPU ops).
Also adds H100_HANG4 diagnostic eprintln between PER update,
bookkeeping, varmap sync, and graph_aux capture to isolate if any
secondary hang exists.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
