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>
cuMemcpyDtoHAsync with non-pinned host pointers blocks the CPU on
CUDA 13 / H100 driver 580+ until ALL pending GPU work completes.
With graph_forward + aux_ops + adam queued, this caused multi-minute
hangs between training steps.
Fix: allocate 12 bytes of pinned (page-locked) host memory via
cuMemHostAlloc(DEVICEMAP) for the 3 scalar readbacks (loss, mse_loss,
grad_norm). Pinned memory enables true async DtoH — CPU returns
immediately, GPU copies when it reaches that point in the stream.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
cuMemcpyDtoHAsync_v2 with non-pinned host pointer degrades to
synchronous copy on CUDA 13 / H100 driver 580+, blocking until
ALL pending GPU work completes. With graph_forward + aux_ops + adam
queued, this caused a multi-minute hang.
Fix: explicit cuStreamSynchronize + synchronous cuMemcpyDtoH.
Stream sync + log at timestep 0, 50%, and 100% to identify which
phase of experience collection hangs on H100 with 4096 episodes.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Both were unused (bf16 path replaced by f32 path in #30).
build_next_states_dtod had the same CPU loop anti-pattern (795K memcpy
calls at 4096 episodes) that caused the H100 hang.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Root cause of H100 training hang: build_next_states_f32() used a CPU
for-loop launching individual memcpy_dtod per episode per fill timestep.
At gpu_n_episodes=4096 × 2 counterfactual × 97 fill steps = 795K CUDA
API calls that saturated the driver command queue.
Replaced with single build_next_states_kernel launch — all episodes
processed in parallel, zero CPU loops. One kernel, done.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Moved init_from_fxcache BEFORE fold generation so GPU features buffer
exists when generate_folds_stratified_gpu is called. Features uploaded
as temporary f32 for regime classification (freed after fold gen).
Removes TODO — GPU prefix sum path is now the default when CUDA available.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
generate_folds_stratified_gpu uses GPU prefix sums for O(1) range
queries during boundary adjustment. 1.1M bars: ~3ms vs ~1000ms CPU.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
regime_classify_kernel: parallel classification of 1.1M bars
regime_prefix_sum_kernel: cumulative counts for O(1) range queries
classify_regimes_gpu(): Rust wrapper that runs both and returns host prefix sums
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Split graph_forward into graph_forward_main (Pass 1+2+loss+backward
on main stream) and graph_forward_ddqn (Pass 3 on double_dqn_stream).
Event synchronization happens in ungraphed replay_forward() code,
not inside captured functions. Each graph uses its own CublasForward
handle — no set_stream() during capture.
Fixes: H100 training hang with gpu_n_episodes=4096.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
submit_forward_ops_main: Pass 1 + 2 + loss + grad + backward (main stream)
submit_forward_ops_ddqn: Pass 3 Double DQN (double_dqn_stream)
No set_stream() or event ops inside either function.
Graph capture will record each on its own stream.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Each CUDA graph needs its own cuBLAS handle bound to its stream.
Eliminates set_stream() calls during graph capture that caused
the H100 training hang.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Was calling generate_folds() (non-stratified) instead of
generate_folds_stratified(). Walk-forward folds are now balanced
across trending/ranging/volatile regimes.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Three issues in one spec:
1. H100 hang: cuBLAS set_stream() inside CUDA graph capture → deadlock
2. Bug: train_walk_forward calls non-stratified generate_folds()
3. GPU regime: classification kernel + prefix sum for O(1) range queries
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
All training parameters now explicit in the TOML — no implicit defaults.
- batch_size=8192: maximizes H100 tensor core utilization
- max_bars=0: unlimited (train on full dataset)
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
- mbp10_data_dir: /mnt/training-data → /data/futures-baseline-mbp10
- Add missing trades_data_dir: /data/futures-baseline-trades
- Paths now match Argo workflow PVC mount at /data (CLI overrides still work)
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
- Preloaded fxcache truncated to max_bars from training profile
(1.1M → 5000 bars for smoketest, ~40x faster per trial)
- Consolidated 3 redundant DqnTrainingProfile::load() calls into 2
(VRAM gate reuses base_hp in train_with_params)
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
- Load TOML profile BEFORE VRAM budget gate so hidden_dim/batch_size
reflect the actual profile (smoketest=64 vs production=256)
- Fixes backtest_metrics=None: VRAM gate was pruning all trials using
conservative() defaults (hidden=256, batch=1024) on RTX 3050
- Add test_hyperopt_trial_produces_backtest_metrics regression test
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
- Add training_profile field to DQNTrainer (default: "dqn-production")
- Smoketest uses "dqn-smoketest" profile for RTX 3050 compatible sizes
- DQNTrainer fields (data_source, mbp10_data_dir, trades_data_dir) are
source of truth — TOML profile provides training params only
- Fix relative path resolution in preload_data for mbp10/trades dirs
- Add preload_data() call in hyperopt test (fxcache shared via Arc)
- Buffer size floor of 1024 in hyperopt param application
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
- Add mbp10_data_dir and trades_data_dir to smoketest TOML profile
- Fix data_loading path resolution: resolve relative paths against
workspace root via CARGO_MANIFEST_DIR (CWD is crates/ml during tests)
- Add buffer_size floor of 1024 in hyperopt adapter (GPU PER minimum)
- Regenerated fxcache with full OFI features (MBP-10 + trades)
All tests now consistently hit fxcache with OFI=true.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Removes run_nan_tracing (3 stream syncs per step on first 20 steps)
and unused grad_buf_ptr/params_buf_ptr accessors. Root causes are
fixed — runtime detection no longer needed. isfinite guards in CUDA
kernels remain as silent defense-in-depth.
17/17 smoke tests pass.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
- Reduce max_bars from 60000 to 5000 (smoketests ran 16min, now ~1min)
- Fix test_counterfactual_experiences_in_buffer: compute threshold from
actual config instead of hardcoded 4000, enlarge buffer to avoid wrap
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
test_data_dir() was returning the symbol-specific path
(test_data/futures-baseline/ES.FUT) causing fxcache key mismatch —
the cache was created with the base dir. Now returns the parent
directory so trainer.train(data_dir, "ES.FUT", ...) can discover
the fxcache via matching cache key.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Removes cql_grad_scratch_ptr, iqn_trunk_m_ptr, save_h_s1_ptr,
stream_cu — added during NaN investigation, no longer called.
Retains run_nan_tracing (lightweight GPU-side checks on first 20 steps)
and all isfinite guards as production safety nets.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Root cause #8: HER in-place relabel kernel read f32 PER states as
`const __nv_bfloat16*`, reinterpreting raw f32 bytes as garbage bf16
values (up to ~1e38). These were written into padded states_buf at HER
rows, corrupting the data read by the IQN trunk backward GEMM.
Root cause #9: HER dst_stride used unpadded state_dim (48) instead of
pad128(state_dim) (128), scattering goal columns into padding region.
Both fixes verified: 200/200 runs with zero NaN (was 4-15% before).
Removed diagnostic tracing code from investigation.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Added dqn_nan_check_f32/bf16 kernels for zero-overhead GPU-side NaN
detection between graph stages. Key findings:
- NaN appears in f32 params at step 0 (before first training forward)
- States (input data) are always clean
- The graph capture warmup launches corrupt params via:
C51 forward → NaN gradients → Adam warmup → NaN weights
- Removed Adam warmup launch, but NaN persists — the forward
graph warmup itself may corrupt buffers through bf16 overflow
in cuBLAS backward path
- Step 14/22 NaN in later runs is from accumulated corruption
Next: investigate whether cuBLAS backward with random Xavier weights
produces extreme gradients that overflow bf16 d_logits staging buffers.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Two fixes for training stability:
1. IQN kernel rewrite: all-bf16 → native f32 arithmetic.
BF16 exp() overflows at logit > 11.1, and bf16 accumulation
of 64 multiply-adds loses catastrophic precision. The IQN
head does NOT use cuBLAS tensor cores, so bf16 provided zero
performance benefit — only precision loss and NaN risk.
Pattern matches C51/MSE kernels: bf16 I/O at DQN trunk
boundary, f32 for all internal computation.
- Rust: all CudaSlice<half::bf16> → CudaSlice<f32>
- CUDA: all __nv_bfloat16 registers → float
- Adam, grad_norm, EMA: native f32
- IQN→PER boundary: f32→bf16 cast via cast_f32_to_td_errors()
2. Gradient budget alpha-blend: during MSE warmup (c51_alpha=0),
the budget was clipping MSE gradients to 60% of max_grad_norm
as if they were C51. Now blends: full budget at alpha=0,
C51 fraction at alpha=1.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Documented all findings from the training step performance session:
- Shared memory race was real but not the NaN source (racecheck: 0 hazards)
- NaN is non-deterministic numerical instability, NOT a concurrency bug
- Suspected math bug in gradient budget allocation during MSE warmup
- Investigation plan: GPU NaN detection kernels, gradient audit, git bisect
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Two optimizations:
1. Dual CUDA graph: graph_forward_mse skips C51 loss/grad/SAXPY blend
during MSE warmup (c51_alpha=0). Eliminates 7 kernel invocations
from the graph replay. replay_forward() routes by alpha automatically.
2. NaN-safe dqn_scale_f32_kernel: writes 0.0 directly when alpha==0
instead of multiplying (IEEE 754: 0*NaN=NaN). Defense-in-depth for
the C51 gradient buffer which can contain NaN from random logits.
Fold 3 epoch: 5.08s → 4.92s (under 5s target).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
compute-sanitizer --tool racecheck found 4000+ hazards per kernel invocation
in both c51_loss_batched and mse_loss_batched. The race is between consecutive
block_reduce calls sharing the same shmem_reduce buffer:
1. reduce_max returns shmem_reduce[0] → next reduce_sum writes shmem_reduce[warp]
2. reduce_sum returns shmem_reduce[0] → next reduce_max writes shmem_reduce[warp]
Fixed by adding __syncthreads() after reading the return value in both
block_reduce_max_f and block_reduce_sum_f (read-then-sync pattern). Also
added explicit barriers between spectral decoupling reduces and subsequent
softmax/log_softmax calls.
This race caused non-deterministic NaN in loss values. The effect was masked
by GPU printf statements (which acted as implicit memory barriers) and by
CUDA Graph replay timing differences.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Two remaining issues: (1) ~400ms/step instead of ~1-7ms on RTX 3050,
(2) Q-value explosion when C51 kicks in at epoch 2.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
The per-step GPU sync for 4-byte diversity_loss monitoring readback
was blocking the CPU on every training step, destroying the async
pipeline. Each step took ~350ms instead of ~1ms.
Fix: accumulate diversity_loss on GPU, defer readback to epoch boundary
via readback_diversity_loss(). The sync now happens ONCE per epoch
instead of per step.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
