- gradient_utils: Add TensorId-based fallback when Var identity mismatch
causes 0/N vars to match GradStore (Candle Adam clones Arcs). Fallback
computes norm AND clips via insert_id. Throttled warning (1st + every
1000th). 7 unit tests including mismatch-actually-clips.
- monitoring: Track full 45-action factored space (5 exposure × 3 order
× 3 urgency). Fix validate_rewards false alarm on GPU path where single
aggregated mean_reward per epoch gives N=1 → std=0.
- trainer: GPU experience collection routes exposure actions through
route_action() for factored tracking instead of exposure-only counts.
Applied in both per-step and epoch-summary paths.
- train_baseline_rl: Auto-detect VRAM <8GB → disable GPU replay buffer
to prevent OOM on RTX 3050 Ti class GPUs.
- smoke_test_real_data: E2E DQN training test with 6 assertions (epoch
completion, loss decrease, finite losses, Q-value divergence, 45-action
space, finite gradient norms).
Validated: 1642 tests pass (ml=915, ml-core=311, ml-dqn=416), 0 clippy
warnings, baseline RL trains 10 epochs on CUDA with Sharpe +5.45.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Move // gpu-ok: annotations to same line as violation patterns so the
guard script's grep -v filter actually suppresses them. Fixes 6 false
positives in ensemble adapters (dqn, ppo, liquid, kan, tggn, diffusion).
Replace hardcoded 48KB shmem limit in compile_forward_kernel() with
GPU-aware query (max_shared_memory_kb) — matches gpu_experience_collector
pattern. H100 now gets 128-row tiles (was 64), eliminating tile loops
for ≤128-dim layers in the backtest forward kernel.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Capture the full DQN backtest step loop (gather → forward → DtoD →
env_step × max_len) as a replayable CUDA Graph:
- evaluate_dqn_graphed(): captures on first call via
CudaStream::begin_capture/end_capture, replays cached graph on
subsequent calls. Falls back to evaluate_dqn() on any failure.
- invalidate_dqn_graph(): discards cached graph when weights change.
- SendSyncGraph: newtype wrapper for CudaGraph (single-threaded use).
- Full unrolled capture: each step's step_i32 argument is baked in at
capture time, avoiding GPU-resident step counter complexity.
Eliminates ~5-10μs per kernel launch × 4 kernels × max_len steps
of CUDA driver overhead per evaluation.
evaluate_baseline.rs: added --cuda-graphs CLI flag to opt in.
14 backtest evaluator tests pass, 0 clippy errors.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Multi-stream pipeline (evaluate() closure path):
- Secondary env_stream with CudaEvent sync allows env_step to overlap
with gather/forward of the next iteration on H100's 132 SMs
Pure-CUDA DQN forward (evaluate_dqn):
- backtest_forward_kernel.cu: warp-cooperative dueling Q forward via NVRTC
- Eliminates candle per-op dispatch overhead, enables future CUDA Graph capture
- evaluate_baseline.rs: auto-detects dueling network and uses pure-CUDA path
Pure-CUDA PPO forward (evaluate_ppo):
- backtest_forward_ppo_kernel.cu: actor MLP → softmax → 45→5 exposure collapse → argmax
- One thread per window, bypasses candle entirely
CUDA supervised signal→action (evaluate_supervised):
- backtest_forward_supervised_kernel.cu: threshold bucketing kernel
- Candle still used for model forward (TFT/Mamba/etc), but argmax is GPU-native
Shared helpers: launch_gather, launch_env_step_on, launch_metrics_and_download
deduplicate code across all four evaluation paths.
914 tests pass, 0 clippy errors.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
- Lower warp kernel SM threshold from sm_90 to sm_70 (covers all modern GPUs)
- Add q_forward_branching_warp_shmem: warp-cooperative branching DQN forward
with 3 independent advantage heads + shared memory tiling
- Fix online/target dispatch to use warp branching variant when available
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Restore 45-action factored space via Branching DQN (Tavakoli 2018),
outputting 11 Q-values (5+3+3) instead of 45. This was reduced to 5
exposure-only actions during debugging and was never intended as permanent.
- Enable use_branching: true by default in DQNConfig and DQNHyperparameters
- Add branching paths to select_action_with_confidence and select_action_inference
- Update agent.rs select_action_factored for branching-aware selection
- Expand CountBonus to per-branch tracking with bonuses_branched()
- Add order_type + urgency distribution tracking in monitoring
- Add DQN_ORDER_ACTIONS=3, DQN_URGENCY_ACTIONS=3, DQN_TOTAL_ACTIONS=45 to CUDA header
- Fix 7 pre-existing clippy doc_markdown errors in regime_conditional.rs
- Fix pre-existing cognitive_complexity in replay_buffer_type.rs (extract helpers)
- Fix flaky GPU test OOM under parallel execution (CPU fallback + test VRAM safety)
- Delete unused flash_attention submodules (block_sparse, causal_masking, etc.)
- Add GPU hot-path guard scripts and ensemble/hyperopt adapter improvements
Tests: ml-dqn 416/0, ml 905/0, clippy 0 errors on both crates
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Continuous vectors and expected counts updated from 12 to 14 params
after adding signal_high_bps and signal_low_bps to Mamba2Params.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
All supervised hyperopt adapters (TFT, Mamba2, Liquid, KAN, xLSTM, TGGN,
TLOB, Diffusion) now include:
- signal_high_bps and signal_low_bps in ParameterSpace (tunable thresholds)
- backtest_sharpe and backtest_trades fields in Metrics structs
- extract_objective prefers GPU backtest fitness over val_loss when available
- All tests updated (101 passing)
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Wire GpuBacktestEvaluator into PPO hyperopt trials so the optimizer
ranks candidates by walk-forward Sharpe ratio instead of raw episode
reward. The PPO actor's 45-action softmax probabilities are collapsed
to 5 exposure scores via ppo_to_exposure_scores before the backtest
loop. When CUDA is unavailable or the backtest fails, the adapter
falls back to the original -avg_episode_reward objective.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
- Flip --gpu-eval default to true (--no-gpu-eval to opt out)
- Move actions_history scatter-write into env kernel (zero CPU accumulation)
- Remove done-flag periodic download (env kernel handles per-thread)
- Only GPU→CPU transfer is final metrics readback (n_windows × 40 bytes)
- Add spread cost discrepancy warning when GPU path is active
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Document known benign race on reset_flags (cross-block __threadfence
not being a grid-wide barrier) and writeback picking an arbitrary
episode's DSR/EMA as the representative seed. Remove dead reads for
epoch_port_value/pos/cash in both standard and warp-cooperative kernel
variants.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Add --gpu-eval flag that routes DQN fold evaluation through
GpuBacktestEvaluator instead of the per-bar CPU path. The GPU path
uploads the full test fold as a single window, runs the env kernel
loop on-device, and downloads only the final WindowMetrics (10 floats).
Key design choices:
- GPU path uses greedy argmax (not hierarchical softmax); results differ
slightly from CPU path — this is documented and intentional
- Falls back to CPU path on any GPU error (no silent failures)
- Also adds --initial-capital flag needed by GpuBacktestConfig
- Fix pre-existing clippy::wildcard_enum_match_arm in
GpuBacktestEvaluator::new (Device::_ → explicit variants)
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Implements backtest_gather_kernel.cu (gather_states) which reads directly
from the pre-uploaded features buffer and live portfolio state on GPU,
eliminating the large GPU→CPU download of the full features buffer that
the old gather_states() path performed each step (n_windows×max_len×feat_dim
floats, e.g. 134 MB for 8 windows × 100k steps × 42 features).
The new path: kernel writes [n_windows, state_dim] into states_buf, then
only that tiny buffer (~1.5 KB for 8×48) is downloaded to create the
Candle tensor — a ~100,000x reduction in per-step data transfer.
Wiring changes in GpuBacktestEvaluator:
- Added GATHER_PTX OnceLock + compile_gather_ptx()
- Added gather_kernel (CudaFunction) and states_buf (CudaSlice<f32>) fields
- Added portfolio_dim field (always 3, validated in gather_states())
- Allocates states_buf = n_windows * (feature_dim + 3) in new()
- gather_states() now launches kernel then downloads small output buffer
- metrics download updated to 10 floats/window (Task 13 extended metrics)
- Added 3 new tests: gather PTX compilation, portfolio_dim validation,
state_dim calculation; all 10 gpu_backtest tests pass
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Add GPU-accelerated backtest path that runs walk-forward evaluation
entirely on GPU (env step kernel + metrics reduction), falling back
to the existing CPU path if CUDA is unavailable or the GPU path fails.
Changes:
- Make DQN::q_values_for_batch() public for external Q-value access
- Add RegimeConditionalDQN::batch_q_values() for regime-routed Q-values
- Add DQNAgentType::batch_q_values() dispatch method
- Add DQNTrainer::evaluate_gpu() method (#[cfg(feature = "cuda")])
- Wire GPU-first backtest at the decision point with CPU fallback
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Implements the GPU backtest evaluator (Task 9): uploads walk-forward
window data once, runs the step loop with Candle forward pass + env
kernel, then a single metrics reduction kernel with one final download.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
compute_epoch_q_diagnostics now uses compute_q_diagnostics_gpu() free
function with Candle tensor ops. Batches gap stats + per-action means
into single 8-float readback instead of full N×5 to_vec2 download.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
One block per window. Parallel reduction for Sharpe, Sortino,
total PnL, max drawdown, win rate, trade count. Single kernel
launch reduces all windows simultaneously.
One thread per walk-forward window, parallel across all windows.
Handles: action→exposure mapping, trade execution with tx costs,
mark-to-market, step return calculation, drawdown tracking.
Portfolio state persists across steps in GPU global memory.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
DSR portfolio reset and normalizer reset now happen via kernel flags
instead of CPU state mutation. Eliminates cudaStreamSynchronize at
epoch boundaries.
Also fix brace mismatch in gpu_training_guard.rs that left the
accumulate_q_value/read_q_accumulator/reset_q_accumulator methods
outside the impl block (introduced by Task 4 in-progress work).
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Eliminates N*8 bytes of memcpy_dtoh per experience kernel launch.
MonitoringReducer accumulates stats on GPU, single 48-byte download
at epoch boundary. Zero cudaStreamSynchronize during experience collection.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
- 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>
Replaces per-launch rewards/actions download with single epoch-end
reduction. monitoring_reduce kernel computes mean, std, min, max,
Sharpe estimate, and per-action counts via parallel reduction.
Single 48-byte download instead of N*8 bytes per kernel launch.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Episode-boundary resets previously used cold hardcoded constants (0.0f,
1e-8f) for DSR accumulators and EMA normalizer. Now warm-starts from
the persistent epoch state so all episodes within an epoch benefit from
accumulated statistics, not just the first.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
The epoch_vol_ema, epoch_dsr_mean, and epoch_dsr_var variables were loaded
from global memory at kernel start but never wired into the actual computation.
The EMA normalizer (ema_mean/ema_var) and DSR accumulators (dsr_A/dsr_B) were
initialized with hardcoded constants, so epoch state round-tripped unchanged.
Changes:
- Seed ema_mean/ema_var from epoch_dsr_mean/epoch_dsr_var at kernel start
- Seed dsr_A/dsr_B from epoch_dsr_mean/epoch_dsr_var at kernel start
- Seed ema_init/dsr_initialized from epoch_step_count > 0 (skip cold start
on subsequent epochs)
- Add local_vol_ema/local_median_vol seeded from epoch_vol_ema/epoch_median_vol
- Update vol EMA each timestep from market feature index 3 (log close return),
mirroring CPU DQNTrainer::vol_ema / median_vol logic
- At writeback, write actual computed dsr_A/dsr_B or ema_mean/ema_var (conditioned
on use_dsr), and computed local_vol_ema/local_median_vol, instead of unmodified
loaded epoch values
- Applied identically to both dqn_full_experience_kernel (standard) and
dqn_full_experience_kernel_warp (warp-cooperative) variants
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
The GPU PER replay buffer had a hardcoded 4 GB MAX_BYTES limit that
rejected the auto-sizer's 10M-entry proposal on H100 (80 GB VRAM).
Now per_max_buffer_bytes() computes 20% of total VRAM (min 1 GB) and
flows through OptimalReplayConfig → DQNConfig → GpuReplayBufferConfig
so both subsystems agree on the budget.
Also fixes misleading regime detection log (indices 211/219 → 40/41)
and renames dqn_config_2025 → dqn_default_config.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Eliminates cudaStreamSynchronize between epochs by keeping vol EMA,
portfolio state, and DSR normalizer in persistent CudaSlice buffers.
Kernel reads initial state at launch, writes final state at exit.
- Add epoch_state CudaSlice<f32>[8] field and reset_flags u32 bitfield
to GpuExperienceCollector struct
- Allocate epoch_state in new() with sensible defaults (vol_ema=0.01,
initial_capital for portfolio, dsr_var=1.0 to avoid div-by-zero)
- Pass epoch_state and reset_flags as final args to both kernel variants
(standard per-thread and warp-cooperative)
- Kernel: thread/lane 0 of block 0 applies reset flags atomically with
__threadfence, all threads read 8-float epoch state from L1-cached
global memory, last block writes back updated values at exit
- Auto-clear reset_flags after each launch (one-shot semantics)
- Add set_reset_flags(), clear_reset_flags(), epoch_state_gpu(),
rewards_gpu(), actions_gpu() public methods
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Three fixes for GPU-accelerated Branching DQN training:
1. **GPU experience collector**: NoisyLinear creates standalone Vars via
Var::from_tensor(), bypassing VarMap registration. The GPU collector
looks up weights by name ("value_fc.weight") from VarMap and falls
back to CPU (~5x slower) when missing. Fix: register mu vars in
VarMap at construction, keep sigma vars standalone.
2. **Optimizer device mismatch**: Using only vars().all_vars() left
NoisyLinear head params frozen. backward() produces gradients the
optimizer doesn't know about → device mismatch in clip_grad_norm.
Fix: all_trainable_vars() = VarMap (shared+mu) + sigma.
3. **Single-threaded CPU bottleneck**: Runtime::new() creates a
current-thread scheduler → 1 OS thread → all async work serialized.
Fix: multi-thread runtime (4 workers) created once in DQNTrainer::new(),
shared across preload/training/backtest phases. Eliminates 3 fallback
Runtime::new() callsites.
Also: polyak_update_var_pairs with debug_assert_eq, two-phase target
network sync (VarMap Polyak + sigma var_pairs Polyak), copy_weights_from
handles NoisyLinear heads.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Three optimizations targeting GPU allocation churn and lock contention:
1. In-place polyak update: Var::set() reuses existing GPU buffer instead
of Var::from_tensor() which allocates a new one per call. Eliminates
~10,000 cudaMalloc/cudaFree per epoch (20 params × 500 steps).
2. Fused affine ops: Replace 13 Tensor::full()/Tensor::ones() constant
tensor allocations per step with tensor.affine(mul, add) — a single
fused kernel. Patterns: 1-x → x.affine(-1,1), γ*x → x.affine(γ,0),
0.5*x² → (x*x).affine(0.5,0). Applied across all 4 loss paths
(branching Bellman/Huber, standard Bellman/Huber, IQN). Eliminates
~6,500 GPU allocs/epoch.
3. Batch pre-sampling (K=8): Sample 8 batches under one READ lock, train
all 8 under one WRITE lock. Reduces async lock acquisitions from
2×N to 2×ceil(N/8). Priority staleness across 8 steps is negligible.
Combined estimated impact: 20-35% H100 throughput improvement.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Candle's Tensor::cumsum(0) internally allocates an [n,n] upper-triangular
matrix (9.3 GB for n=50K) causing OOM on GPUs ≤48 GB. Replace with a
block-parallel Hillis-Steele scan kernel that is O(n) in time and memory.
Additional fixes in this commit:
- Break autograd chain leak in loss/grad accumulation via .detach()
(was leaking ~32 MB/step across entire training run)
- Release features_raw_cuda/targets_raw_cuda after GPU experience
collection (~164 MB VRAM reclaimed)
- Use softmax eval (temp=0.3) in walk-forward backtest to prevent
action collapse causing trades=0 on early-stage models
Validated: 1286 tests pass (408 ml-dqn + 878 ml), 0 failures.
VRAM stable at 754 MB across 45K+ training steps on RTX 3050 Ti.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
- Staging/GPU replay buffer: truncate batch when it exceeds ring buffer
capacity (CPU fallback can flush 100K+ experiences at once)
- GpuExperienceCollector: compute shmem tile rows dynamically to stay
under 48 KB default limit instead of hardcoded 64-row constant
- Auto batch size: subtract concurrent VRAM consumers (~530 MB model +
optimizer + PER + data) before budgeting experience collector at 40%
- Hyperopt DQN adapter: GPU PER always on, include replay buffer in
VRAM estimate for small GPUs (was excluded assuming CPU-only PER)
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
- DQN hyperopt adapter: remove static VRAM gate for GPU experience
collector and GPU PER — dynamic scaling handles constraints at
runtime with graceful CPU fallback on init failure
- Remove is_parquet_file branching from hyperopt eval path (all data
loads via DBN pipeline now)
- Update training example CLI args for consistency
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Eliminate all per-step GPU→CPU synchronization barriers from the
training guard. Replace device+host buffer pairs and memcpy_dtoh
with cuMemHostAlloc(DEVICEMAP) mapped pinned memory.
Key changes:
- MappedBuffer struct: cuMemHostAlloc + cuMemHostGetDevicePointer_v2
allocates memory visible to both CPU and GPU simultaneously
- Double-buffering: kernel writes to buffer[N%2], CPU reads buffer
[(N-1)%2] — one-step delayed halt detection, zero sync
- __threadfence_system() in CUDA kernels ensures writes visible to
CPU across PCIe without explicit memcpy
- read_volatile on host pointer prevents CPU-side caching
Eliminated:
- check_and_accumulate: 28-byte memcpy_dtoh (every training step)
- qvalue_stats: 16-byte memcpy_dtoh (every 50 steps)
- qvalue_divergence: 20-byte memcpy_dtoh (every 50 steps)
Kept: read_accumulators memcpy_dtoh (12 bytes, epoch boundary only —
accumulator buffer stays in device memory for kernel read-modify-write).
1286 tests pass (878 ml + 408 ml-dqn), 0 failures.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Two remaining GPU→CPU synchronization barriers in the CUDA-active
training loop:
1. detect_dead_neurons() → to_vec0() called every training step from
log_diagnostics() in the guard path. Added check_gradient_collapse()
method that performs the same collapse detection logic but skips the
expensive per-parameter weight scan. Dead neuron detection now only
runs at epoch boundary via log_diagnostics().
2. forward() Q-value clipping monitoring → to_vec2() called every 1000
steps during compute_loss_internal() and Q-value estimation forward
passes. Added training_forward_active flag that gates the monitoring
block; set to true during all training-path forward() calls
(compute_loss_internal + Q-value estimation), false during
inference/evaluation.
All 1577 tests pass (878 ml + 408 ml-dqn + 291 ml-core), 0 failures.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Wire route_exposure_to_factored() into both select_actions_batch_gpu
and select_actions_batch GPU paths, eliminating per-item CPU routing.
Previously, exposure indices (0-4) were downloaded from GPU and routed
to factored indices (0-44) one-by-one on CPU via route_action(). Now
the exposure→factored mapping runs entirely on GPU via the routing
kernel, with a single batch readback of the final factored indices.
Branching DQN path unchanged (already produces factored indices 0-44).
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Replace CPU-bound Q-value monitoring with GPU-resident qvalue_stats and
qvalue_divergence kernels from GpuTrainingGuard. The two readback sites
(estimate_avg_q_value_with_early_stopping's mean_all().to_scalar() and
log_q_values' to_vec2()) are now bypassed when the GPU training guard is
active. CPU fallback path preserved for non-CUDA builds and guard-absent
scenarios.
Changes:
- Add DQN::log_q_values_from_stats() accepting pre-computed GPU stats
- Add DQNAgentType::log_q_values_from_stats() delegate
- Replace Q-estimation in train_step_single_batch with GPU kernel path
- Replace Q-estimation in train_step_with_accumulation with GPU kernel path
- Import IndexOp trait for Tensor::i() in trainer.rs
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Adds a standalone CUDA kernel that converts DQN exposure indices (0-4)
to factored action indices (0-44) entirely on GPU, reusing the existing
route_order() device function from common_device_functions.cuh. Wired
into GpuActionSelector as route_exposure_to_factored() method, following
the same DtoD copy pattern as the existing select_actions methods. This
eliminates a GPU->CPU->GPU roundtrip when post-hoc routing is needed
after epsilon_greedy_select.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Replace the GPU->CPU sync barrier (to_vec1 readback) in the DQN training
hot path with the GpuTrainingGuard CUDA kernel that performs NaN detection,
loss clipping, and gradient collapse checks entirely on-device. The guard
is lazy-initialized on first training step and falls back to the original
CPU readback path if CUDA kernel compilation fails.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Wraps 4 CUDA kernels (training_guard_check, training_guard_accumulate,
qvalue_stats_reduce, qvalue_divergence_check) with a Rust struct that
uses OnceLock PTX caching, pre-allocated device buffers, and host-side
Vec mirrors for zero-allocation readbacks per training step.
Accumulator (3-float acc_buf) stays on-device for epoch-boundary
averaging without CPU roundtrips.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
