From d485dde5a28ecfae9324342fa4e0c2d96db3a51f Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 27 Mar 2026 00:33:05 +0100 Subject: [PATCH] fix: eliminate NVRTC, fix evaluator SIGSEGV, 22D search space, trade physics MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Major changes: - Eliminate ALL runtime NVRTC: c51_loss + mse_loss kernels converted to precompiled cubins with runtime num_atoms/v_min/v_max/branch params - Fix evaluator SIGSEGV: rng_states/q_gaps sized for chunked batch (cn) not n_windows; NULL pointer guard in action_select kernel - Add trade_physics.cuh shared header for train/eval consistency - Add equity circuit breaker (25% DD from peak) + margin-aware position cap - Consolidate 46D→22D hyperopt search space, enable ensemble by default - Fix trade counting: use exposure index not factored action - Fix Calmar overflow: clamp to ±100 in kernel - Softer CVaR penalty (cap 3.0 not 10.0) for undertrained models - Fix win_rate display (ratio→percentage) - Remove dead code (normalize_reward, calculate_completion_penalty) - Add tracing subscriber to hyperopt test for visible metrics - Per-chunk sync in evaluator for reliable error reporting Co-Authored-By: Claude Opus 4.6 (1M context) --- crates/ml-hyperopt/src/optimizer.rs | 1 - crates/ml/build.rs | 8 +- .../src/cuda_pipeline/backtest_env_kernel.cu | 169 +- .../cuda_pipeline/backtest_metrics_kernel.cu | 22 +- .../ml/src/cuda_pipeline/c51_loss_kernel.cu | 319 ++-- .../src/cuda_pipeline/experience_kernels.cu | 33 +- .../cuda_pipeline/gpu_backtest_evaluator.rs | 187 +-- .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 87 +- crates/ml/src/cuda_pipeline/mod.rs | 4 +- .../ml/src/cuda_pipeline/mse_loss_kernel.cu | 196 ++- crates/ml/src/cuda_pipeline/trade_physics.cuh | 159 ++ crates/ml/src/hyperopt/adapters/dqn.rs | 1447 ++++++----------- crates/ml/src/hyperopt/campaign.rs | 6 + crates/ml/src/trainers/dqn/fused_training.rs | 2 +- crates/ml/src/trainers/dqn/trainer/tests.rs | 27 +- .../src/trainers/dqn/trainer/training_loop.rs | 9 +- 16 files changed, 1187 insertions(+), 1489 deletions(-) create mode 100644 crates/ml/src/cuda_pipeline/trade_physics.cuh diff --git a/crates/ml-hyperopt/src/optimizer.rs b/crates/ml-hyperopt/src/optimizer.rs index ae4b2d13c..d2f3c396e 100644 --- a/crates/ml-hyperopt/src/optimizer.rs +++ b/crates/ml-hyperopt/src/optimizer.rs @@ -498,7 +498,6 @@ impl ArgminOptimizer { } Err(e) => { tracing::error!("Trial {} training FAILED: {:#}", trial_num, e); - eprintln!("!!! TRIAL {} TRAINING FAILED: {:#}", trial_num, e); eprintln!("!!! FULL ERROR CHAIN: {:?}", e); (1e6, None) } diff --git a/crates/ml/build.rs b/crates/ml/build.rs index ff1d6d64c..a37a494d8 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -12,6 +12,10 @@ fn main() { let out_dir = PathBuf::from(std::env::var("OUT_DIR").unwrap()); let kernel_dir = Path::new("src/cuda_pipeline"); let common_header = kernel_dir.join("common_device_functions.cuh"); + let trade_physics_header = kernel_dir.join("trade_physics.cuh"); + + // Rebuild cubins when shared headers change + println!("cargo:rerun-if-changed={}", trade_physics_header.display()); // Detect GPU architecture from env or default to sm_80 let cuda_compute_cap = std::env::var("CUDA_COMPUTE_CAP").unwrap_or_else(|_| "80".to_string()); @@ -144,7 +148,8 @@ fn compile_kernel( let tmp_src = out_dir.join(format!("_{kernel_name}")); std::fs::write(&tmp_src, &full_source).unwrap(); - // Compile with nvcc + // Compile with nvcc (-I kernel_dir for shared .cuh includes) + let include_flag = format!("-I{}", kernel_dir.display()); let status = Command::new(nvcc) .args([ "-cubin", @@ -153,6 +158,7 @@ fn compile_kernel( "--use_fast_math", "--ftz=true", "--fmad=true", + &include_flag, "-o", cubin_path.to_str().unwrap(), tmp_src.to_str().unwrap(), ]) diff --git a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu index 8536f1196..d184e016a 100644 --- a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu +++ b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu @@ -10,8 +10,12 @@ // [5] hold_time - bars held in current position (0 = flat) // [6] cum_return - cumulative log return // [7] step_count - number of completed steps +// +// Trade physics (action decoding, hold enforcement, tx costs, capital floor) +// are shared with the training kernel via trade_physics.cuh. #define PORTFOLIO_STATE_SIZE 8 +#include "trade_physics.cuh" extern "C" __global__ void backtest_env_step( // Market data (read-only, uploaded once) @@ -37,20 +41,17 @@ extern "C" __global__ void backtest_env_step( float tx_cost_bps, float spread_cost, int current_step, - // Branching DQN action space sizes - int b0_size, // exposure branch size (default 9) - int b1_size, // order type branch size (default 3) - int b2_size, // urgency branch size (default 3) - int min_hold_bars // minimum bars to hold before exit/reversal + int b0_size, + int b1_size, + int b2_size, + int min_hold_bars ) { - // Shared memory tile for coalesced portfolio reads/writes __shared__ float shmem_pf[256 * PORTFOLIO_STATE_SIZE]; int w = blockIdx.x * blockDim.x + threadIdx.x; int local_tid = threadIdx.x; - // Cooperative tile load: all active threads load their portfolio slice - // in one coalesced warp transaction before divergent early-exit checks. + // Cooperative tile load int ps = w * PORTFOLIO_STATE_SIZE; if (w < n_windows) { for (int i = 0; i < PORTFOLIO_STATE_SIZE; i++) { @@ -84,64 +85,55 @@ extern "C" __global__ void backtest_env_step( float hold_time = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 5]; float cum_return = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 6]; - // Decode factored action to exposure index (matches experience_env_step). - // Factored action = exposure_idx * b1_size * b2_size + order_idx * b2_size + urgency_idx - int action_val = actions[w]; - int exposure_idx; - if (b1_size > 0 && b2_size > 0) { - int denom = b1_size * b2_size; - exposure_idx = action_val / denom; - } else { - exposure_idx = action_val; + // ── Capital floor circuit breaker (shared: trade_physics.cuh) ──────── + if (check_capital_floor(value, max_equity)) { + // Force flat: close any open position at current price + if (fabsf(position) > 0.001f && close > 0.0f && entry_price > 0.0f) { + float pnl = position * (close - entry_price); + cash += pnl; + } + float liq_value = cash; + float liq_ret = (value > 0.01f) ? (liq_value - value) / value : 0.0f; + step_rewards[w] = liq_ret; + step_returns[w * max_len + current_step] = liq_ret; + actions_history[w * max_len + current_step] = b0_size / 2; // Flat + portfolio_state[ps + 0] = liq_value; + portfolio_state[ps + 1] = 0.0f; + portfolio_state[ps + 2] = liq_value; + portfolio_state[ps + 3] = 0.0f; + portfolio_state[ps + 4] = max_equity; + portfolio_state[ps + 5] = 0.0f; + portfolio_state[ps + 6] = cum_return + liq_ret; + portfolio_state[ps + 7] += 1.0f; + done_flags[w] = 1; + return; } - // Clamp to valid range - if (exposure_idx < 0) exposure_idx = 0; - if (exposure_idx >= b0_size) exposure_idx = b0_size - 1; - // Map exposure index to fraction: -1.0 + idx * (2.0 / (b0_size - 1)) - // With b0_size=9: step=0.25, range [-1.0, +1.0] - float step_size = (b0_size > 1) ? 2.0f / (float)(b0_size - 1) : 0.0f; - float target_exposure = -1.0f + (float)exposure_idx * step_size; - target_exposure *= max_position; + // ── Decode action (shared: trade_physics.cuh) ──────────────────────── + int action_val = actions[w]; + int exposure_idx = decode_exposure_index(action_val, b0_size, b1_size, b2_size); + float target_exposure = compute_target_position(exposure_idx, b0_size, max_position); + int order_type_idx = decode_order_type(action_val, b1_size, b2_size); + + // ── Margin-aware position cap (shared: trade_physics.cuh) ──────────── + // ES futures: ~$15K initial margin per contract. Prevents overleveraging + // when equity is depleted — a depleted account can't hold the same + // position as a full account. close * 50 ≈ contract notional for ES. + float margin_per_contract = close * 50.0f * 0.06f; // ~6% of notional ≈ $15K for ES at $5000 + target_exposure = apply_margin_cap(target_exposure, value, margin_per_contract); // Suppress unused variable warnings (void)open; (void)high; (void)low; - // Decode order_type and urgency from factored action (for tx cost differentiation) - int order_type_idx = (b1_size > 0 && b2_size > 0) - ? (action_val / b2_size) % b1_size : 0; + // ── Hold enforcement (shared: trade_physics.cuh) ───────────────────── + int is_last_bar = (current_step >= wlen - 1) ? 1 : 0; + target_exposure = enforce_hold(position, target_exposure, hold_time, min_hold_bars, is_last_bar); - // Hold enforcement: block exits and reversals during minimum hold period. - // Matches the training kernel's hold enforcement to eliminate train/eval mismatch. - { - float prev_position = position; // position before any action is applied - int prev_sign = (prev_position > 0.001f) ? 1 : ((prev_position < -0.001f) ? -1 : 0); - int curr_sign = (target_exposure > 0.001f) ? 1 : ((target_exposure < -0.001f) ? -1 : 0); - int wants_exit = (prev_sign != 0 && curr_sign == 0); - int wants_reversal = (prev_sign != 0 && curr_sign != 0 && prev_sign != curr_sign); - if (hold_time > 0.0f && hold_time < (float)min_hold_bars - && (wants_exit || wants_reversal) - && current_step < wlen - 1) { - target_exposure = prev_position; // Override: keep position - } - } - - // Execute trade if position changes + // ── Execute trade (shared: trade_physics.cuh) ──────────────────────── float delta = target_exposure - position; - float trade_cost = 0.0f; + float prev_position = position; if (fabsf(delta) > 0.001f && close > 0.0f) { - /* EXACT match of training kernel (experience_env_step) tx cost formula: - * tx_cost = |delta| * close * (multiplier * 0.0001 * spread_scale * impact + premium) - * - spread_scale: sqrt(|delta|/max_pos) (Almgren-Chriss square-root impact) - * - order_premium: Market=0, IoC=+2bps, LimitMaker=-5bps rebate */ - float spread_scale = sqrtf(fabsf(delta) / fmaxf(max_position, 0.01f)); - if (spread_scale < 1.0f) spread_scale = 1.0f; - float impact_scale = 1.0f; /* simplified: no dynamic spread widening in backtest */ - float order_premium = (order_type_idx == 0) ? 0.0f - : (order_type_idx == 1) ? 0.0002f /* IoC: +2 bps */ - : -0.0005f; /* LimitMaker: -5 bps rebate */ - trade_cost = fabsf(delta) * close * (tx_cost_bps * 0.0001f * spread_scale * impact_scale + order_premium) - + fabsf(delta) * spread_cost * 0.5f; + float trade_cost = compute_tx_cost(delta, close, tx_cost_bps, spread_cost, max_position, order_type_idx); cash -= trade_cost; // Mark-to-market old position @@ -154,27 +146,8 @@ extern "C" __global__ void backtest_env_step( entry_price = close; } - // Update hold_time based on resulting position. - // Read prev_position from local var (before trade block may have updated `position`). - // Note: `position` after the trade block reflects the new position. - { - float prev_pos_for_hold = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 1]; - if (fabsf(position) > 0.001f) { - if (fabsf(prev_pos_for_hold) < 0.001f) { - hold_time = 1.0f; // Entry from flat - } else { - int prev_sign_h = (prev_pos_for_hold > 0.001f) ? 1 : -1; - int new_sign_h = (position > 0.001f) ? 1 : -1; - if (new_sign_h == prev_sign_h) { - hold_time += 1.0f; // Same direction — increment - } else { - hold_time = 1.0f; // Reversal allowed — new trade - } - } - } else { - hold_time = 0.0f; // Flat - } - } + // ── Hold time tracking (shared: trade_physics.cuh) ─────────────────── + hold_time = update_hold_time(prev_position, position, hold_time); // Mark-to-market current position float unrealized = 0.0f; @@ -183,23 +156,51 @@ extern "C" __global__ void backtest_env_step( } float new_value = cash + unrealized; + // ── Post-trade floor check: catch intra-step breaches ──────────── + // The pre-trade check (top of kernel) catches breaches from the + // PREVIOUS step. This catches breaches from THIS step's price move + // or transaction costs. Without it, max_dd overshoots the 75% floor + // because the metrics kernel records the breached value before the + // next step's pre-trade check fires. + if (check_capital_floor(new_value, max_equity)) { + // Emergency liquidation: close position at current price + if (fabsf(position) > 0.001f) { + // Position already marked-to-market above; just go flat + new_value = cash + unrealized; // already computed + position = 0.0f; + cash = new_value; + entry_price = 0.0f; + } + float step_ret = (value > 0.01f) ? (new_value - value) / value : 0.0f; + step_rewards[w] = step_ret; + step_returns[w * max_len + current_step] = step_ret; + actions_history[w * max_len + current_step] = b0_size / 2; + portfolio_state[ps + 0] = new_value; + portfolio_state[ps + 1] = 0.0f; + portfolio_state[ps + 2] = new_value; + portfolio_state[ps + 3] = 0.0f; + portfolio_state[ps + 4] = max_equity; + portfolio_state[ps + 5] = 0.0f; + portfolio_state[ps + 6] = cum_return + step_ret; + portfolio_state[ps + 7] += 1.0f; + done_flags[w] = 1; + return; + } + // Step return float step_ret = (value > 0.0f) ? (new_value - value) / value : 0.0f; float new_cum_return = cum_return + step_ret; - - // Update max equity for drawdown float new_max = fmaxf(max_equity, new_value); - // Write portfolio state directly to global memory (no shmem write-back needed — - // threads with done_flags[w]==1 exited early and cannot participate in __syncthreads) + // Write portfolio state portfolio_state[ps + 0] = new_value; portfolio_state[ps + 1] = position; portfolio_state[ps + 2] = cash; portfolio_state[ps + 3] = entry_price; portfolio_state[ps + 4] = new_max; - portfolio_state[ps + 5] = hold_time; // hold_time (bars in current position) + portfolio_state[ps + 5] = hold_time; portfolio_state[ps + 6] = new_cum_return; - portfolio_state[ps + 7] += 1.0f; // step count + portfolio_state[ps + 7] += 1.0f; // Outputs step_rewards[w] = step_ret; diff --git a/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu index 9f9a1b94e..fdc6fdaaf 100644 --- a/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu +++ b/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu @@ -108,13 +108,16 @@ extern "C" __global__ void compute_backtest_metrics( if (exp_idx >= 0 && exp_idx < num_actions) local_action_mask |= (1 << exp_idx); - // Boundary-aware trade tracking + // Boundary-aware trade tracking: count EXPOSURE changes only. + // The factored action encodes (exposure, order_type, urgency) but only + // exposure changes trigger actual position changes in env_step. + // Using exp_idx prevents order/urgency flips from inflating trade count. if (i == chunk_start) { - bnd_first_action = act; - bnd_cur_action = act; + bnd_first_action = exp_idx; + bnd_cur_action = exp_idx; } - if (act != bnd_cur_action) { + if (exp_idx != bnd_cur_action) { bnd_num_changes++; if (bnd_num_changes == 1) { // First change: everything accumulated so far is the prefix @@ -125,11 +128,11 @@ extern "C" __global__ void compute_backtest_metrics( if (bnd_cur_return > 0.0f) bnd_complete_wins++; } bnd_cur_return = 0.0f; - bnd_cur_action = act; + bnd_cur_action = exp_idx; } bnd_cur_return += r; - bnd_last_action = act; + bnd_last_action = exp_idx; } // Post-loop: assign suffix and handle no-change edge case @@ -393,9 +396,11 @@ extern "C" __global__ void compute_backtest_metrics( } float daily_mean = total_return / (float)sort_len; float max_dd = metrics_out[out_base + 2]; - float calmar = (max_dd > 1e-8f) + float calmar = (max_dd > 0.001f) ? (daily_mean * annualization_factor * annualization_factor) / max_dd : 0.0f; + // Clamp Calmar to prevent inf/NaN from dominating the objective + calmar = fmaxf(-100.0f, fminf(100.0f, calmar)); // Omega ratio: sum of positive returns / sum of |negative returns| float gain_sum = 0.0f, loss_sum = 0.0f; @@ -403,7 +408,8 @@ extern "C" __global__ void compute_backtest_metrics( if (s_sorted[i] > 0.0f) gain_sum += s_sorted[i]; else loss_sum -= s_sorted[i]; } - float omega = (loss_sum > 1e-10f) ? gain_sum / loss_sum : 0.0f; + float omega = (loss_sum > 0.001f) ? gain_sum / loss_sum : 0.0f; + omega = fmaxf(0.0f, fminf(100.0f, omega)); metrics_out[out_base + 6] = var_95; metrics_out[out_base + 7] = cvar_95; diff --git a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu index 68c9748a3..ca7b70af5 100644 --- a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu +++ b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu @@ -10,89 +10,38 @@ * - Input: value_logits + advantage_logits per branch from both online, * target, and online_next networks (pre-computed by cuBLAS forward pass). * - Grid: (batch_size, 1, 1) — one block per sample. - * - Block: (256, 1, 1) — 256 threads cooperate on NUM_ATOMS=51 work. + * - Block: (256, 1, 1) — 256 threads cooperate on num_atoms work. * - Shared memory: ~2.8 KB per block (vs 46 KB/thread in fused kernel). * * Does NOT modify dqn_forward_loss_kernel — that kernel remains as fallback. * - * Requires common_device_functions.cuh prepended via NVRTC source concatenation - * (same mechanism as dqn_training_kernel.cu). The following constants must be - * injected via NVRTC #define before this source is compiled: - * NUM_ATOMS, V_MIN, V_MAX, BRANCH_0_SIZE, BRANCH_1_SIZE, BRANCH_2_SIZE + * Requires common_device_functions.cuh prepended via build.rs source + * concatenation. All variable constants (NUM_ATOMS, V_MIN, V_MAX, + * BRANCH_*_SIZE) are passed as runtime kernel parameters. */ -/* ── Compile-time defaults (overridden by NVRTC injection) ───────────── */ +/* ── Fixed compile-time constants ─────────────────────────────────────── */ #ifndef MAX_PER_SAMPLE_CE #define MAX_PER_SAMPLE_CE 50.0f /* Layer 1: 10x converged loss — breaks PER feedback loop */ #endif #ifndef LABEL_SMOOTHING_EPS #define LABEL_SMOOTHING_EPS 0.01f /* Layer 2: 1% uniform mix — prevents target collapse */ #endif -#ifndef NUM_ATOMS -#define NUM_ATOMS 51 -#endif -#ifndef BRANCH_0_SIZE -#define BRANCH_0_SIZE 9 -#endif -#ifndef BRANCH_1_SIZE -#define BRANCH_1_SIZE 3 -#endif -#ifndef BRANCH_2_SIZE -#define BRANCH_2_SIZE 3 -#endif -#ifndef V_MIN -#define V_MIN (-25.0f) -#endif -#ifndef V_MAX -#define V_MAX (25.0f) -#endif -#define DELTA_Z ((V_MAX - V_MIN) / (float)(NUM_ATOMS - 1)) -/* Total atoms per branch output */ -#define BRANCH_0_ATOMS (BRANCH_0_SIZE * NUM_ATOMS) -#define BRANCH_1_ATOMS (BRANCH_1_SIZE * NUM_ATOMS) -#define BRANCH_2_ATOMS (BRANCH_2_SIZE * NUM_ATOMS) -/* Largest branch output size (floats) — used to size shmem regions */ -#define MAX_BRANCH_ATOMS (MAX_BRANCH_SIZE * NUM_ATOMS) +/* Fixed upper bound for register-allocated arrays — covers num_atoms up to 128 */ +#define MAX_ATOMS 128 + +/* Largest branch output size (floats) — used to size __shared__ arrays */ #define MAX_BRANCH_SIZE 9 /* max(BRANCH_0_SIZE, BRANCH_1_SIZE, BRANCH_2_SIZE) */ #define NUM_BRANCHES 3 /* Maximum warp count for BLOCK_THREADS = 256 */ #define NUM_WARPS 8 /* 256 / 32 */ /* ── Block-level thread count (must match launch config) ─────────────── */ -/* 256 threads: 8 warps, enough to saturate the issue queue for small NUM_ATOMS. - * At NUM_ATOMS=51, each thread handles at most 1 atom in most passes. */ +/* 256 threads: 8 warps, enough to saturate the issue queue for small num_atoms. + * At num_atoms=51, each thread handles at most 1 atom in most passes. */ #define BLOCK_THREADS 256 -/* ── Shared memory layout (all regions in floats) ─────────────────────── - * - * Region Size (floats) Purpose - * ───────────────────────────────────────────────────────────────── - * shmem_support NUM_ATOMS=51 Precomputed z_j = V_MIN + j*DELTA_Z - * shmem_val NUM_ATOMS=51 Value logits for current network/state - * shmem_adv MAX_BRANCH_ATOMS=255 Adv logits for all actions in branch - * shmem_lp NUM_ATOMS=51 log_softmax result (current/target action) - * shmem_proj NUM_ATOMS=51 Bellman-projected target distribution - * shmem_current_lp NUM_ATOMS=51 Saved current action log-probs (for CE) - * shmem_blk_reduce NUM_WARPS=8 Block-level reduction scratch - * ───────────────────────────────────────────────────────────────── - * Total floats: 51+51+255+51+51+51+8 = 518 - * Total bytes: 518 * 4 = 2072 bytes ≈ 2.0 KB per block - * H100 SM: 228 KB / 2.0 KB = 114 concurrent blocks per SM - * - * Note: shmem_blk_reduce is NUM_WARPS floats and reused for all block - * reductions (max-finding, sum-finding, cross-entropy). One reduction - * at a time — never overlapping. - */ -#define SHMEM_SUPPORT_OFF 0 -#define SHMEM_VAL_OFF (SHMEM_SUPPORT_OFF + NUM_ATOMS) -#define SHMEM_ADV_OFF (SHMEM_VAL_OFF + NUM_ATOMS) -#define SHMEM_LP_OFF (SHMEM_ADV_OFF + MAX_BRANCH_SIZE * NUM_ATOMS) -#define SHMEM_PROJ_OFF (SHMEM_LP_OFF + NUM_ATOMS) -#define SHMEM_CURRENT_LP_OFF (SHMEM_PROJ_OFF + NUM_ATOMS) -#define SHMEM_REDUCE_OFF (SHMEM_CURRENT_LP_OFF + NUM_ATOMS) -#define SHMEM_TOTAL_FLOATS (SHMEM_REDUCE_OFF + NUM_WARPS) - /* ── Device helpers ───────────────────────────────────────────────────── */ /** @@ -157,35 +106,36 @@ __device__ __forceinline__ float block_reduce_sum( } /** - * Block-cooperative log_softmax over NUM_ATOMS elements. + * Block-cooperative log_softmax over num_atoms elements. * - * Reads from `logits[0..NUM_ATOMS]` in shared memory. - * Writes result to `out[0..NUM_ATOMS]` in shared memory. + * Reads from `logits[0..num_atoms]` in shared memory. + * Writes result to `out[0..num_atoms]` in shared memory. * Both `logits` and `out` may alias the same region. * `shmem_reduce[NUM_WARPS]` is used as scratch. * Caller must __syncthreads() before reading `out`. */ __device__ void block_log_softmax( - const float* __restrict__ logits, /* [NUM_ATOMS] shmem */ - float* __restrict__ out, /* [NUM_ATOMS] shmem output */ + const float* __restrict__ logits, /* [num_atoms] shmem */ + float* __restrict__ out, /* [num_atoms] shmem output */ float* __restrict__ shmem_reduce, /* [NUM_WARPS] scratch */ - int tid + int tid, + int num_atoms ) { /* Pass 1: block-max */ float local_max = -1e30f; - for (int i = tid; i < NUM_ATOMS; i += BLOCK_THREADS) + for (int i = tid; i < num_atoms; i += BLOCK_THREADS) local_max = fmaxf(local_max, logits[i]); float block_max = block_reduce_max(local_max, shmem_reduce, tid); /* Pass 2: sum(exp(x - max)) */ float local_sum = 0.0f; - for (int i = tid; i < NUM_ATOMS; i += BLOCK_THREADS) + for (int i = tid; i < num_atoms; i += BLOCK_THREADS) local_sum += expf(logits[i] - block_max); float block_sum = block_reduce_sum(local_sum, shmem_reduce, tid); float log_sum_exp = logf(block_sum + 1e-10f); /* Pass 3: write log(softmax) = x - max - log_sum_exp */ - for (int i = tid; i < NUM_ATOMS; i += BLOCK_THREADS) + for (int i = tid; i < num_atoms; i += BLOCK_THREADS) out[i] = logits[i] - block_max - log_sum_exp; __syncthreads(); } @@ -197,13 +147,14 @@ __device__ void block_log_softmax( * Returns the scalar expected Q (identical on all threads after call). */ __device__ float block_expected_q( - const float* __restrict__ log_probs, /* [NUM_ATOMS] shmem */ - const float* __restrict__ support, /* [NUM_ATOMS] shmem */ + const float* __restrict__ log_probs, /* [num_atoms] shmem */ + const float* __restrict__ support, /* [num_atoms] shmem */ float* __restrict__ shmem_reduce, /* [NUM_WARPS] scratch */ - int tid + int tid, + int num_atoms ) { float local_sum = 0.0f; - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) local_sum += expf(log_probs[j]) * support[j]; return block_reduce_sum(local_sum, shmem_reduce, tid); } @@ -211,44 +162,48 @@ __device__ float block_expected_q( /** * Block-cooperative Bellman projection. * - * Projects `target_probs[NUM_ATOMS]` (probabilities, not log-probs) onto the - * C51 support using T_z = clamp(r + gamma * z_j * (1 - done), V_MIN, V_MAX). - * Result accumulated into `projected[NUM_ATOMS]` (must be zero-initialized). + * Projects `target_probs[num_atoms]` (probabilities, not log-probs) onto the + * C51 support using T_z = clamp(r + gamma * z_j * (1 - done), v_min, v_max). + * Result accumulated into `projected[num_atoms]` (must be zero-initialized). * - * With BLOCK_THREADS=256 and NUM_ATOMS=51, each thread handles at most 1 source - * atom. The per-thread `local_proj[NUM_ATOMS]` register array therefore has at + * With BLOCK_THREADS=256 and num_atoms=51, each thread handles at most 1 source + * atom. The per-thread `local_proj[MAX_ATOMS]` register array therefore has at * most 2 non-zero entries (lower and upper neighbors for that one atom). * The warp butterfly + block accumulate then merges them. */ __device__ void block_bellman_project( - const float* __restrict__ target_probs, /* [NUM_ATOMS] shmem probabilities */ + const float* __restrict__ target_probs, /* [num_atoms] shmem probabilities */ float reward, float done, float gamma, - float* __restrict__ projected, /* [NUM_ATOMS] shmem, zero-initialized */ - const float* __restrict__ support, /* [NUM_ATOMS] shmem */ + float* __restrict__ projected, /* [num_atoms] shmem, zero-initialized */ + const float* __restrict__ support, /* [num_atoms] shmem */ float* __restrict__ shmem_reduce, /* [NUM_WARPS] scratch */ - int tid + int tid, + int num_atoms, + float v_min, + float v_max, + float delta_z ) { /* Phase 1: each thread accumulates its atoms' contributions in registers. - * With BLOCK_THREADS=256 and NUM_ATOMS=51, each thread processes at most + * With BLOCK_THREADS=256 and num_atoms=51, each thread processes at most * one source atom (tid 0..50) or zero atoms (tid 51..255). */ - float local_proj[NUM_ATOMS]; - for (int k = 0; k < NUM_ATOMS; k++) local_proj[k] = 0.0f; + float local_proj[MAX_ATOMS]; + for (int k = 0; k < num_atoms; k++) local_proj[k] = 0.0f; - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) { + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) { float z_j = support[j]; float t_z = reward + gamma * z_j * (1.0f - done); - t_z = fminf(fmaxf(t_z, V_MIN), V_MAX); + t_z = fminf(fmaxf(t_z, v_min), v_max); - float b_pos = (t_z - V_MIN) / DELTA_Z; + float b_pos = (t_z - v_min) / delta_z; /* Keep b_pos away from exact upper boundary to prevent phantom * upper==lower when ceilf(N-1) == floorf(N-1) = N-1. * Without this, probability mass piles on boundary atoms. */ - b_pos = fminf(b_pos, (float)(NUM_ATOMS - 1) - 1e-6f); + b_pos = fminf(b_pos, (float)(num_atoms - 1) - 1e-6f); b_pos = fmaxf(b_pos, 1e-6f); int lower = (int)floorf(b_pos); int upper = (int)ceilf(b_pos); - lower = max(min(lower, NUM_ATOMS - 1), 0); - upper = max(min(upper, NUM_ATOMS - 1), 0); + lower = max(min(lower, num_atoms - 1), 0); + upper = max(min(upper, num_atoms - 1), 0); float frac = b_pos - floorf(b_pos); float p_j = target_probs[j]; @@ -264,7 +219,7 @@ __device__ void block_bellman_project( int lane = tid & 31; int warp = tid >> 5; - for (int k = 0; k < NUM_ATOMS; k++) { + for (int k = 0; k < num_atoms; k++) { /* Warp butterfly sum for atom k */ float val = local_proj[k]; for (int offset = 16; offset > 0; offset >>= 1) @@ -310,22 +265,22 @@ __device__ void block_bellman_project( extern "C" __global__ void c51_loss_batched( /* ── Online network outputs on current STATES ─────────────────── */ - const float* __restrict__ on_value_logits, /* [B, NUM_ATOMS] */ - const float* __restrict__ on_adv_logits_b0, /* [B, BRANCH_0_SIZE * NUM_ATOMS] */ - const float* __restrict__ on_adv_logits_b1, /* [B, BRANCH_1_SIZE * NUM_ATOMS] */ - const float* __restrict__ on_adv_logits_b2, /* [B, BRANCH_2_SIZE * NUM_ATOMS] */ + const float* __restrict__ on_value_logits, /* [B, num_atoms] */ + const float* __restrict__ on_adv_logits_b0, /* [B, b0_size * num_atoms] */ + const float* __restrict__ on_adv_logits_b1, /* [B, b1_size * num_atoms] */ + const float* __restrict__ on_adv_logits_b2, /* [B, b2_size * num_atoms] */ /* ── Target network outputs on NEXT_STATES ────────────────────── */ - const float* __restrict__ tg_value_logits, /* [B, NUM_ATOMS] */ - const float* __restrict__ tg_adv_logits_b0, /* [B, BRANCH_0_SIZE * NUM_ATOMS] */ - const float* __restrict__ tg_adv_logits_b1, /* [B, BRANCH_1_SIZE * NUM_ATOMS] */ - const float* __restrict__ tg_adv_logits_b2, /* [B, BRANCH_2_SIZE * NUM_ATOMS] */ + const float* __restrict__ tg_value_logits, /* [B, num_atoms] */ + const float* __restrict__ tg_adv_logits_b0, /* [B, b0_size * num_atoms] */ + const float* __restrict__ tg_adv_logits_b1, /* [B, b1_size * num_atoms] */ + const float* __restrict__ tg_adv_logits_b2, /* [B, b2_size * num_atoms] */ /* ── Online network outputs on NEXT_STATES (Double DQN selector) */ - const float* __restrict__ on_next_value_logits, /* [B, NUM_ATOMS] */ - const float* __restrict__ on_next_adv_logits_b0, /* [B, BRANCH_0_SIZE * NUM_ATOMS] */ - const float* __restrict__ on_next_adv_logits_b1, /* [B, BRANCH_1_SIZE * NUM_ATOMS] */ - const float* __restrict__ on_next_adv_logits_b2, /* [B, BRANCH_2_SIZE * NUM_ATOMS] */ + const float* __restrict__ on_next_value_logits, /* [B, num_atoms] */ + const float* __restrict__ on_next_adv_logits_b0, /* [B, b0_size * num_atoms] */ + const float* __restrict__ on_next_adv_logits_b1, /* [B, b1_size * num_atoms] */ + const float* __restrict__ on_next_adv_logits_b2, /* [B, b2_size * num_atoms] */ /* ── Batch data ───────────────────────────────────────────────── */ const int* __restrict__ actions, /* [B] factored action indices 0-44 */ @@ -339,69 +294,92 @@ extern "C" __global__ void c51_loss_batched( float* __restrict__ total_loss, /* [1] batch mean loss (atomicAdd) */ /* ── Saved tensors for backward pass ─────────────────────────── */ - float* __restrict__ save_current_lp, /* [B, NUM_BRANCHES, NUM_ATOMS] */ - float* __restrict__ save_projected, /* [B, NUM_BRANCHES, NUM_ATOMS] */ + float* __restrict__ save_current_lp, /* [B, NUM_BRANCHES, num_atoms] */ + float* __restrict__ save_projected, /* [B, NUM_BRANCHES, num_atoms] */ /* ── Config ───────────────────────────────────────────────────── */ float gamma, - int batch_size + int batch_size, + int num_atoms, + float v_min, + float v_max, + int b0_size, + int b1_size, + int b2_size ) { extern __shared__ float shmem[]; + /* ── Compute runtime shared memory offsets ────────────────────── */ + const int max_branch = max(max(b0_size, b1_size), b2_size); + int off_support = 0; + int off_val = off_support + num_atoms; + int off_adv = off_val + num_atoms; + int off_lp = off_adv + max_branch * num_atoms; + int off_proj = off_lp + num_atoms; + int off_current_lp = off_proj + num_atoms; + int off_reduce = off_current_lp + num_atoms; + /* Named pointers into the flat shared memory region */ - float* shmem_support = shmem + SHMEM_SUPPORT_OFF; /* [NUM_ATOMS] z_j */ - float* shmem_val = shmem + SHMEM_VAL_OFF; /* [NUM_ATOMS] value logits */ - float* shmem_adv = shmem + SHMEM_ADV_OFF; /* [MAX_BRANCH_SIZE*NUM_ATOMS] adv logits */ - float* shmem_lp = shmem + SHMEM_LP_OFF; /* [NUM_ATOMS] log-prob result */ - float* shmem_proj = shmem + SHMEM_PROJ_OFF; /* [NUM_ATOMS] Bellman projection */ - float* shmem_current_lp = shmem + SHMEM_CURRENT_LP_OFF; /* [NUM_ATOMS] taken action lp */ - float* shmem_reduce = shmem + SHMEM_REDUCE_OFF; /* [NUM_WARPS] reduction scratch */ + float* shmem_support = shmem + off_support; /* [num_atoms] z_j */ + float* shmem_val = shmem + off_val; /* [num_atoms] value logits */ + float* shmem_adv = shmem + off_adv; /* [max_branch*num_atoms] adv logits */ + float* shmem_lp = shmem + off_lp; /* [num_atoms] log-prob result */ + float* shmem_proj = shmem + off_proj; /* [num_atoms] Bellman projection */ + float* shmem_current_lp = shmem + off_current_lp; /* [num_atoms] taken action lp */ + float* shmem_reduce = shmem + off_reduce; /* [NUM_WARPS] reduction scratch */ int sample_id = blockIdx.x; int tid = threadIdx.x; if (sample_id >= batch_size) return; + /* ── Compute runtime derived values ──────────────────────────── */ + float delta_z = (v_max - v_min) / (float)(num_atoms - 1); + + int b0_atoms = b0_size * num_atoms; + int b1_atoms = b1_size * num_atoms; + int b2_atoms = b2_size * num_atoms; + /* ── Precompute C51 support atoms once per block ─────────────── */ - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) - shmem_support[j] = V_MIN + (float)j * DELTA_Z; + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) + shmem_support[j] = v_min + (float)j * delta_z; __syncthreads(); /* ── Decode factored action into per-branch indices ──────────── */ int factored_action = actions[sample_id]; /* Clamp to valid range — guards against uninitialized replay buffer entries */ - int max_action = BRANCH_0_SIZE * BRANCH_1_SIZE * BRANCH_2_SIZE; + int max_action = b0_size * b1_size * b2_size; if (factored_action < 0 || factored_action >= max_action) factored_action = 0; - int a0 = factored_action / (BRANCH_1_SIZE * BRANCH_2_SIZE); - int a1 = (factored_action / BRANCH_2_SIZE) % BRANCH_1_SIZE; - int a2 = factored_action % BRANCH_2_SIZE; + int a0 = factored_action / (b1_size * b2_size); + int a1 = (factored_action / b2_size) % b1_size; + int a2 = factored_action % b2_size; - /* Per-branch metadata (compile-time sizes, runtime-selected values) */ + /* Per-branch metadata (runtime sizes, runtime-selected values) */ const int branch_action[NUM_BRANCHES] = { a0, a1, a2 }; - const int branch_sizes[NUM_BRANCHES] = { BRANCH_0_SIZE, BRANCH_1_SIZE, BRANCH_2_SIZE }; + const int branch_sizes[NUM_BRANCHES] = { b0_size, b1_size, b2_size }; /* Per-sample row pointers into each branch's advantage logit tensors */ const float* on_adv_row[NUM_BRANCHES] = { - on_adv_logits_b0 + (long long)sample_id * BRANCH_0_ATOMS, - on_adv_logits_b1 + (long long)sample_id * BRANCH_1_ATOMS, - on_adv_logits_b2 + (long long)sample_id * BRANCH_2_ATOMS + on_adv_logits_b0 + (long long)sample_id * b0_atoms, + on_adv_logits_b1 + (long long)sample_id * b1_atoms, + on_adv_logits_b2 + (long long)sample_id * b2_atoms }; const float* tg_adv_row[NUM_BRANCHES] = { - tg_adv_logits_b0 + (long long)sample_id * BRANCH_0_ATOMS, - tg_adv_logits_b1 + (long long)sample_id * BRANCH_1_ATOMS, - tg_adv_logits_b2 + (long long)sample_id * BRANCH_2_ATOMS + tg_adv_logits_b0 + (long long)sample_id * b0_atoms, + tg_adv_logits_b1 + (long long)sample_id * b1_atoms, + tg_adv_logits_b2 + (long long)sample_id * b2_atoms }; const float* on_next_adv_row[NUM_BRANCHES] = { - on_next_adv_logits_b0 + (long long)sample_id * BRANCH_0_ATOMS, - on_next_adv_logits_b1 + (long long)sample_id * BRANCH_1_ATOMS, - on_next_adv_logits_b2 + (long long)sample_id * BRANCH_2_ATOMS + on_next_adv_logits_b0 + (long long)sample_id * b0_atoms, + on_next_adv_logits_b1 + (long long)sample_id * b1_atoms, + on_next_adv_logits_b2 + (long long)sample_id * b2_atoms }; - /* Value logit row pointers (same NUM_ATOMS for all branches) */ - const float* on_val_row = on_value_logits + (long long)sample_id * NUM_ATOMS; - const float* tg_val_row = tg_value_logits + (long long)sample_id * NUM_ATOMS; - const float* on_next_val_row = on_next_value_logits + (long long)sample_id * NUM_ATOMS; + /* Value logit row pointers (same num_atoms for all branches) */ + const float* on_val_row = on_value_logits + (long long)sample_id * num_atoms; + const float* tg_val_row = tg_value_logits + (long long)sample_id * num_atoms; + const float* on_next_val_row = on_next_value_logits + (long long)sample_id * num_atoms; float reward = rewards[sample_id]; float done = dones[sample_id]; @@ -416,14 +394,14 @@ extern "C" __global__ void c51_loss_batched( for (int d = 0; d < NUM_BRANCHES; d++) { int n_d = branch_sizes[d]; int a_d = branch_action[d]; - int n_atoms = n_d * NUM_ATOMS; /* total adv logit elements for this branch */ + int n_atoms = n_d * num_atoms; /* total adv logit elements for this branch */ /* Global memory offset for save tensors: [sample, branch, atom] */ - long long save_off = ((long long)sample_id * NUM_BRANCHES + d) * NUM_ATOMS; + long long save_off = ((long long)sample_id * NUM_BRANCHES + d) * num_atoms; /* ── Helper: compute dueling logits for a single action into shmem_lp - * shmem_val must already hold value logits [NUM_ATOMS] - * shmem_adv must already hold all adv logits [n_d * NUM_ATOMS] + * shmem_val must already hold value logits [num_atoms] + * shmem_adv must already hold all adv logits [n_d * num_atoms] * Result: dueling(a_sel)[j] = V[j] + A[a_sel,j] - mean_a(A[*,j]) * written into shmem_lp[j] */ /* (Inlined below to avoid device-function overhead for this hot path) */ @@ -434,7 +412,7 @@ extern "C" __global__ void c51_loss_batched( * ═══════════════════════════════════════════════════════════ */ /* Load value logits */ - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) shmem_val[j] = on_val_row[j]; __syncthreads(); @@ -444,25 +422,25 @@ extern "C" __global__ void c51_loss_batched( __syncthreads(); /* Dueling for action a_d: Q[j] = V[j] + A[a_d,j] - mean_a(A[*,j]) */ - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) { + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) { float a_mean = 0.0f; for (int a = 0; a < n_d; a++) - a_mean += shmem_adv[a * NUM_ATOMS + j]; + a_mean += shmem_adv[a * num_atoms + j]; a_mean /= (float)n_d; - shmem_lp[j] = shmem_val[j] + shmem_adv[a_d * NUM_ATOMS + j] - a_mean; + shmem_lp[j] = shmem_val[j] + shmem_adv[a_d * num_atoms + j] - a_mean; } __syncthreads(); /* log_softmax → result back into shmem_lp */ - block_log_softmax(shmem_lp, shmem_lp, shmem_reduce, tid); - /* shmem_lp[0..NUM_ATOMS] = log_probs for taken action a_d */ + block_log_softmax(shmem_lp, shmem_lp, shmem_reduce, tid, num_atoms); + /* shmem_lp[0..num_atoms] = log_probs for taken action a_d */ /* Copy into shmem_current_lp for use in cross-entropy (step e) */ - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) shmem_current_lp[j] = shmem_lp[j]; /* Save for backward kernel */ - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) save_current_lp[save_off + j] = shmem_lp[j]; __syncthreads(); @@ -474,7 +452,7 @@ extern "C" __global__ void c51_loss_batched( * ═══════════════════════════════════════════════════════════ */ /* Load online_next value and advantage logits */ - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) shmem_val[j] = on_next_val_row[j]; __syncthreads(); @@ -484,10 +462,10 @@ extern "C" __global__ void c51_loss_batched( /* Compute average advantage across actions for dueling (same for all a) */ /* shmem_proj temporarily holds per-atom mean advantage — cleared next step */ - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) { + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) { float a_mean = 0.0f; for (int a = 0; a < n_d; a++) - a_mean += shmem_adv[a * NUM_ATOMS + j]; + a_mean += shmem_adv[a * num_atoms + j]; a_mean /= (float)n_d; shmem_proj[j] = a_mean; /* reuse proj region for mean advantage */ } @@ -501,15 +479,15 @@ extern "C" __global__ void c51_loss_batched( for (int a = 0; a < n_d; a++) { /* Dueling logits for action a: V[j] + A[a,j] - mean_j(A[*,j]) */ - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) - shmem_lp[j] = shmem_val[j] + shmem_adv[a * NUM_ATOMS + j] - shmem_proj[j]; + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) + shmem_lp[j] = shmem_val[j] + shmem_adv[a * num_atoms + j] - shmem_proj[j]; __syncthreads(); /* log_softmax in-place on shmem_lp */ - block_log_softmax(shmem_lp, shmem_lp, shmem_reduce, tid); + block_log_softmax(shmem_lp, shmem_lp, shmem_reduce, tid, num_atoms); /* Expected Q for action a */ - float eq = block_expected_q(shmem_lp, shmem_support, shmem_reduce, tid); + float eq = block_expected_q(shmem_lp, shmem_support, shmem_reduce, tid, num_atoms); if (tid == 0) eq_per_action[a] = eq; __syncthreads(); } @@ -530,7 +508,7 @@ extern "C" __global__ void c51_loss_batched( * ═══════════════════════════════════════════════════════════ */ /* Load target value and advantage logits */ - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) shmem_val[j] = tg_val_row[j]; __syncthreads(); @@ -539,26 +517,26 @@ extern "C" __global__ void c51_loss_batched( __syncthreads(); /* Per-atom mean advantage for target network */ - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) { + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) { float a_mean = 0.0f; for (int a = 0; a < n_d; a++) - a_mean += shmem_adv[a * NUM_ATOMS + j]; + a_mean += shmem_adv[a * num_atoms + j]; a_mean /= (float)n_d; - shmem_lp[j] = shmem_val[j] + shmem_adv[best_next_a * NUM_ATOMS + j] - a_mean; + shmem_lp[j] = shmem_val[j] + shmem_adv[best_next_a * num_atoms + j] - a_mean; } __syncthreads(); /* log_softmax → log-probs for target's best_next action */ - block_log_softmax(shmem_lp, shmem_lp, shmem_reduce, tid); + block_log_softmax(shmem_lp, shmem_lp, shmem_reduce, tid, num_atoms); /* Convert to probabilities for Bellman projection */ - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) shmem_proj[j] = expf(shmem_lp[j]); __syncthreads(); /* ═══════════════════════════════════════════════════════════ * STEP d: Bellman projection - * T_z_j = clamp(r + gamma * z_j * (1-done), V_MIN, V_MAX) + * T_z_j = clamp(r + gamma * z_j * (1-done), v_min, v_max) * Project shmem_proj (target probs) onto shmem_proj (output). * * The projection writes back to shmem_proj in-place (the input @@ -567,28 +545,29 @@ extern "C" __global__ void c51_loss_batched( /* Zero-initialize projected output — block_bellman_project accumulates */ /* We use shmem_lp as the output projected array (shmem_proj still holds probs) */ - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) shmem_lp[j] = 0.0f; __syncthreads(); block_bellman_project(shmem_proj, reward, done, gamma, - shmem_lp, shmem_support, shmem_reduce, tid); + shmem_lp, shmem_support, shmem_reduce, tid, + num_atoms, v_min, v_max, delta_z); __syncthreads(); /* Layer 2: Label smoothing — mix ε uniform into projected target. * Prevents any atom from having zero probability → no log(0) in CE. */ { - const float ls_uniform = LABEL_SMOOTHING_EPS / (float)NUM_ATOMS; + const float ls_uniform = LABEL_SMOOTHING_EPS / (float)num_atoms; const float ls_scale = 1.0f - LABEL_SMOOTHING_EPS; - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) shmem_lp[j] = ls_scale * shmem_lp[j] + ls_uniform; __syncthreads(); } - /* shmem_lp[0..NUM_ATOMS] = smoothed projected target distribution */ + /* shmem_lp[0..num_atoms] = smoothed projected target distribution */ /* Save projected for backward kernel */ - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) save_projected[save_off + j] = shmem_lp[j]; __syncthreads(); @@ -598,7 +577,7 @@ extern "C" __global__ void c51_loss_batched( * ═══════════════════════════════════════════════════════════ */ float local_ce = 0.0f; - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) local_ce -= shmem_lp[j] * shmem_current_lp[j]; float branch_ce = block_reduce_sum(local_ce, shmem_reduce, tid); diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index ede205d38..3d6112e80 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -54,6 +54,10 @@ /* ------------------------------------------------------------------ */ #define PORTFOLIO_STRIDE 20 +/* Shared trade physics (action decoding, hold enforcement, tx costs, + * capital floor). Single source of truth — also used by backtest_env_kernel. */ +#include "trade_physics.cuh" + /* ------------------------------------------------------------------ */ /* Compile-time constants — overridable via NVRTC #define injection. */ /* ------------------------------------------------------------------ */ @@ -378,12 +382,16 @@ extern "C" __global__ void experience_action_select( /* Hold enforcement: during minimum hold period, force current exposure. * Prevents both greedy and random exploration from changing exposure. - * Portfolio state: ps[0] = position, ps[10] = hold_time. */ - int ps_base = i * 20; /* PORTFOLIO_STRIDE = 20 */ - float hold_time_val = portfolio_states[ps_base + 10]; - float cur_position = portfolio_states[ps_base + 0]; - int in_hold = (hold_time_val > 0.0f && hold_time_val < (float)min_hold_bars - && min_hold_bars > 0); + * Portfolio state: ps[0] = position, ps[10] = hold_time. + * When portfolio_states is NULL (backtest evaluator), skip hold enforcement. */ + int in_hold = 0; + float cur_position = 0.0f; + if (min_hold_bars > 0 && portfolio_states != NULL) { + int ps_base = i * 20; /* PORTFOLIO_STRIDE = 20 */ + float hold_time_val = portfolio_states[ps_base + 10]; + cur_position = portfolio_states[ps_base + 0]; + in_hold = (hold_time_val > 0.0f && hold_time_val < (float)min_hold_bars); + } if (in_hold) { /* Compute current exposure index from position */ @@ -667,8 +675,17 @@ extern "C" __global__ void experience_env_step( } /* end Kelly block */ } /* end Kelly scope */ - /* Final NaN/Inf guard on target_position. - * If exposure_idx_to_fraction produced NaN, zero the position. */ + /* Margin-aware position cap (shared: trade_physics.cuh). + * Prevents overleveraging when equity is depleted — a depleted account + * can't hold the same position size as a full account. + * margin ≈ 6% of notional (CME ES initial margin ~$15K per contract). */ + { + float margin_per_contract = raw_close * 50.0f * 0.06f; + float portfolio_val = ps[2]; + target_position = apply_margin_cap(target_position, portfolio_val, margin_per_contract); + } + + /* Final NaN/Inf guard on target_position. */ if (isnan(target_position) || isinf(target_position)) { target_position = 0.0f; } diff --git a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs index 99d4d87c4..18a2953ee 100644 --- a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs +++ b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs @@ -33,7 +33,7 @@ use std::mem::ManuallyDrop; use std::sync::Arc; use cudarc::driver::{CudaFunction, CudaGraph, CudaSlice, CudaStream, DevicePtr, LaunchConfig, PushKernelArg}; use cudarc::nvrtc::Ptx; -use tracing::{info, warn}; +use tracing::info; use ml_core::nvtx::NvtxRange; use crate::MLError; @@ -854,6 +854,15 @@ impl GpuBacktestEvaluator { ) -> Result, MLError> { let _nvtx = NvtxRange::new("backtest_evaluate_dqn_graphed"); + // Sync stream and drain any CUDA errors before evaluator starts. + // If training left deferred errors, log them explicitly. + if let Err(sync_err) = self.stream.synchronize() { + eprintln!("EVALUATOR PRE-SYNC: training left stale CUDA error: {sync_err}"); + } + if let Err(ctx_err) = self.stream.context().check_err() { + eprintln!("EVALUATOR PRE-SYNC: context error drained: {ctx_err}"); + } + // Lazy-initialise cuBLAS context + buffers self.ensure_cublas_ready(dqn_cfg)?; @@ -869,97 +878,19 @@ impl GpuBacktestEvaluator { super::batched_forward::f32_weight_ptrs(flat, ¶m_sizes, &self.stream) }; - // ── Replay cached graph if available ──────────────────────────────── - if let Some(ref graph) = self.dqn_graph { - graph.0.launch().map_err(|e| { - MLError::ModelError(format!("CUDA graph replay failed: {e}")) - })?; - return self.launch_metrics_and_download(); - } - - // ── First call: capture the step loop into a CUDA graph ───────────── - info!( - "GpuBacktestEvaluator: capturing CUDA graph for {} steps × {} windows", - self.max_len, self.n_windows, - ); - - // Synchronize the stream before capture to ensure all pending work - // is complete. CUDA Graph capture requires no in-flight work. - self.stream.synchronize().map_err(|e| { - MLError::ModelError(format!("stream sync before capture: {e}")) - })?; - - // Disable cudarc's automatic event tracking during graph capture. - // After fork(), cudarc injects cuStreamWaitEvent/cuEventRecord into - // every buffer access, creating cross-stream event references that - // invalidate CUDA graph capture. Safe: single-stream, all work synced. - // SAFETY: single-stream capture, all work synchronized, re-enabled below. - unsafe { self.stream.context().disable_event_tracking(); } - - // Begin stream capture in THREAD_LOCAL mode (only captures work - // submitted from this thread on this stream — safe for single-stream use). - let capture_result = self.stream.begin_capture( - cudarc::driver::sys::CUstreamCaptureMode::CU_STREAM_CAPTURE_MODE_THREAD_LOCAL, - ); - if let Err(e) = capture_result { - unsafe { self.stream.context().enable_event_tracking(); } - warn!( - "CUDA graph capture begin failed ({e}), falling back to non-graphed evaluate_dqn" - ); - return self.submit_dqn_step_loop_cublas(&w_ptrs, dqn_cfg) - .and_then(|()| self.launch_metrics_and_download()); - } - - // Submit the step loop to the stream (captured, NOT executed yet). - let capture_err = self.submit_dqn_step_loop_cublas(&w_ptrs, dqn_cfg); - - // End capture — get the graph regardless of whether submit succeeded, - // because the stream is in capture mode and must be ended. - let graph_result = self.stream.end_capture( - cudarc::driver::sys::CUgraphInstantiate_flags::CUDA_GRAPH_INSTANTIATE_FLAG_AUTO_FREE_ON_LAUNCH, - ); - - // Re-enable event tracking now that capture is complete. - // SAFETY: restores normal multi-stream synchronization. - unsafe { self.stream.context().enable_event_tracking(); } - - // If the step loop submission had an error, propagate it - if let Err(e) = capture_err { - warn!("CUDA graph capture step loop failed ({e}), falling back to non-graphed path"); - return self.submit_dqn_step_loop_cublas(&w_ptrs, dqn_cfg) - .and_then(|()| self.launch_metrics_and_download()); - } - - // Process graph instantiation result - let graph = match graph_result { - Ok(Some(g)) => g, - Ok(None) => { - warn!("CUDA graph capture returned empty graph, falling back to non-graphed path"); - return self.submit_dqn_step_loop_cublas(&w_ptrs, dqn_cfg) - .and_then(|()| self.launch_metrics_and_download()); - } - Err(e) => { - warn!( - "CUDA graph instantiation failed ({e}), falling back to non-graphed path" - ); - return self.submit_dqn_step_loop_cublas(&w_ptrs, dqn_cfg) - .and_then(|()| self.launch_metrics_and_download()); - } - }; - - // Launch the graph (this actually executes the captured step loop) - if let Err(e) = graph.launch() { - warn!("CUDA graph launch failed ({e}), falling back to non-graphed path"); - return self.submit_dqn_step_loop_cublas(&w_ptrs, dqn_cfg) - .and_then(|()| self.launch_metrics_and_download()); - } - - info!("GpuBacktestEvaluator: CUDA graph captured and launched successfully"); - - // Cache the graph for subsequent calls - self.dqn_graph = Some(SendSyncGraph(graph)); - - // Metrics are NOT in the graph — run them normally + // Direct kernel launches (no CUDA Graph for the backtest step loop). + // + // The step loop captures 500K+ kernel launches (10K steps × 10 windows + // × ~5 ops/step). cuGraphLaunch SIGSEGVs on graphs this large due to + // CUDA driver-internal memory limits. Direct launches are fine because: + // - Eval runs ONCE per hyperopt trial (not per training step) + // - Training is 99%+ of hyperopt wall time; eval is negligible + // - cuBLAS forward passes dominate eval throughput, not launch overhead + // + // The training pipeline (graph_forward + graph_adam in GpuDqnTrainer) + // still uses CUDA Graphs — those capture ~20 operations per replay, + // well within driver limits. + self.submit_dqn_step_loop_cublas(&w_ptrs, dqn_cfg)?; self.launch_metrics_and_download() } @@ -1057,10 +988,12 @@ impl GpuBacktestEvaluator { let state_row_bytes = n * state_dim * std::mem::size_of::(); let action_row_bytes = n * std::mem::size_of::(); + let total_chunks = (self.max_len + DQN_BACKTEST_CHUNK_SIZE - 1) / DQN_BACKTEST_CHUNK_SIZE; for chunk_start in (0..self.max_len).step_by(DQN_BACKTEST_CHUNK_SIZE) { let chunk_end = (chunk_start + DQN_BACKTEST_CHUNK_SIZE).min(self.max_len); let chunk_len = chunk_end - chunk_start; let batch = n * chunk_len; // total rows for this chunk's forward pass + let chunk_idx = chunk_start / DQN_BACKTEST_CHUNK_SIZE; // ── Phase 1: Gather states for all steps in the chunk ──────────── // @@ -1088,6 +1021,13 @@ impl GpuBacktestEvaluator { } } + // Sync after Phase 1 gather — catch any gather/DtoD errors + if chunk_idx == 0 { + self.stream.synchronize().map_err(|e| MLError::ModelError(format!( + "backtest chunk {chunk_idx}: sync after gather phase: {e}" + )))?; + } + // ── Phase 2: Single cuBLAS forward on [batch, state_dim] ───────── // // For the last chunk, batch may be < n * CHUNK_SIZE. The cuBLAS forward @@ -1107,6 +1047,12 @@ impl GpuBacktestEvaluator { ch_v_logits, ch_b_logits, )?; + if chunk_idx == 0 { + self.stream.synchronize().map_err(|e| MLError::ModelError(format!( + "backtest chunk 0: sync after cuBLAS forward: {e}" + )))?; + } + // ── Phase 3: C51 expected Q-values on chunked output ───────────── let batch_i32 = batch as i32; let blocks_256 = ((batch + 255) / 256) as u32; @@ -1135,15 +1081,15 @@ impl GpuBacktestEvaluator { )))?; } + if chunk_idx == 0 { + self.stream.synchronize().map_err(|e| MLError::ModelError(format!( + "backtest chunk 0: sync after expected_q: {e}" + )))?; + } + // ── Phase 4: Greedy action selection on chunked Q-values ───────── - // Kernel signature: (q_values, out_actions, rng_states, out_q_gaps, epsilon, N, - // b0, b1, b2, q_gap_threshold, portfolio_states, min_hold_bars, max_position) - // Hold enforcement handled by backtest_env_step (Layer 2), not action masking. - // The greedy argmax (epsilon=0) doesn't need masking — any hold-violating - // action is overridden by the env_step kernel. This avoids the stride-8/stride-20 - // portfolio layout incompatibility between backtest and training buffers. - let null_portfolio: u64 = 0; // NULL — Layer 2 in env_step handles holds - let eval_min_hold: i32 = 0; // 0 = disabled here (enforced in env_step instead) + let null_portfolio: u64 = 0; + let eval_min_hold: i32 = 0; let eval_max_pos: f32 = 0.0; unsafe { self.stream @@ -1158,15 +1104,21 @@ impl GpuBacktestEvaluator { .arg(&b1) .arg(&b2) .arg(&self.q_gap_threshold) - .arg(&null_portfolio) // portfolio_states (NULL — no hold masking in eval) - .arg(&eval_min_hold) // min_hold_bars (0 = disabled) - .arg(&eval_max_pos) // max_position (unused when disabled) + .arg(&null_portfolio) + .arg(&eval_min_hold) + .arg(&eval_max_pos) .launch(launch_cfg) .map_err(|e| MLError::ModelError(format!( "chunked action_select chunk_start={chunk_start}: {e}" )))?; } + if chunk_idx == 0 { + self.stream.synchronize().map_err(|e| MLError::ModelError(format!( + "backtest chunk 0: sync after action_select: {e}" + )))?; + } + // ── Phase 5: Per-step env_step (sequential, stateful) ──────────── // // For each step in the chunk, extract the n_windows actions from the @@ -1194,6 +1146,13 @@ impl GpuBacktestEvaluator { self.launch_env_step(step)?; } + + // Per-chunk sync: catch deferred CUDA errors from this chunk's + // kernel launches. Without this, an error in chunk N surfaces as + // a confusing failure in chunk N+1 or later. + self.stream.synchronize().map_err(|e| MLError::ModelError(format!( + "backtest chunk {chunk_idx}/{total_chunks} (steps {chunk_start}..{chunk_end}) sync: {e}" + )))?; } Ok(()) @@ -1428,17 +1387,6 @@ impl GpuBacktestEvaluator { .load_function("experience_action_select") .map_err(|e| MLError::ModelError(format!("experience_action_select load: {e}")))?; - // RNG states (initialised with random seeds for greedy mode — epsilon=0 means - // the RNG values are never actually used, but the kernel reads them). - let rng_seeds: Vec = (0..n).map(|_| fastrand::u32(..)).collect(); - let rng_states = self.stream.clone_htod(&rng_seeds) - .map_err(|e| MLError::ModelError(format!("alloc rng_states: {e}")))?; - - // Q-gap output buffer — kernel writes per-episode Q-gaps (unused by backtest - // but required by the experience_action_select kernel signature). - let q_gaps_buf = self.stream.alloc_zeros::(n) - .map_err(|e| MLError::ModelError(format!("alloc q_gaps_buf: {e}")))?; - // ── Chunked cuBLAS forward (batch_size = n_windows * CHUNK_SIZE) ───── // // Separate CublasForward + scratch buffers to amortise kernel-launch @@ -1446,6 +1394,17 @@ impl GpuBacktestEvaluator { let chunk = DQN_BACKTEST_CHUNK_SIZE; let cn = n * chunk; // chunked batch dimension + // RNG states: must be large enough for chunked batch (n * CHUNK_SIZE), + // not just n_windows, because chunked action_select launches with batch = n * chunk. + let rng_seeds: Vec = (0..cn).map(|_| fastrand::u32(..)).collect(); + let rng_states = self.stream.clone_htod(&rng_seeds) + .map_err(|e| MLError::ModelError(format!("alloc rng_states: {e}")))?; + + // Q-gap output buffer: same sizing as rng_states — chunked action_select + // writes batch = n * chunk_len entries per call. + let q_gaps_buf = self.stream.alloc_zeros::(cn) + .map_err(|e| MLError::ModelError(format!("alloc q_gaps_buf: {e}")))?; + let chunked_cublas = CublasForward::new( &self.stream, cn, diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 1bcf15077..50f8cfd67 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -3025,9 +3025,14 @@ impl GpuDqnTrainer { // The experience collector pre-computes R_n = sum(gamma^i * r_i). let gamma = self.config.gamma.powi(self.config.n_steps as i32); let batch_i32 = b as i32; + let na_i32 = na as i32; + let v_min_f = self.config.v_min; + let v_max_f = self.config.v_max; + let b0_i32 = b0 as i32; + let b1_i32 = b1 as i32; + let b2_i32 = b2 as i32; - // Shared memory: SHMEM_TOTAL_FLOATS * 4 bytes - // 51+51+255+51+51+51+8 = 518 floats = 2072 bytes + // Shared memory: support + val + adv + lp + proj + current_lp + reduce let max_branch = b0.max(b1).max(b2); let shmem_floats = na + na + max_branch * na + na + na + na + 8; let shmem_bytes = (shmem_floats * std::mem::size_of::()) as u32; @@ -3062,9 +3067,15 @@ impl GpuDqnTrainer { // ── Saved for backward (2) ── .arg(&self.save_current_lp) .arg(&self.save_projected) - // ── Config (2) ── + // ── Config (8) ── .arg(&gamma) .arg(&batch_i32) + .arg(&na_i32) + .arg(&v_min_f) + .arg(&v_max_f) + .arg(&b0_i32) + .arg(&b1_i32) + .arg(&b2_i32) .launch(LaunchConfig { grid_dim: (b as u32, 1, 1), block_dim: (256, 1, 1), @@ -3159,6 +3170,12 @@ impl GpuDqnTrainer { // N-step returns: use gamma^n for the Bellman projection. let gamma = self.config.gamma.powi(self.config.n_steps as i32); let batch_i32 = b as i32; + let na_i32 = na as i32; + let v_min_f = self.config.v_min; + let v_max_f = self.config.v_max; + let b0_i32 = b0 as i32; + let b1_i32 = b1 as i32; + let b2_i32 = b2 as i32; // Shared memory: same layout as C51 kernel (support + val + adv + lp + proj + current_lp + reduce) let max_branch = b0.max(b1).max(b2); @@ -3196,9 +3213,15 @@ impl GpuDqnTrainer { // save_projected = per-branch E[Q] values (3 floats per sample per branch) .arg(&self.save_current_lp) .arg(&self.save_projected) - // ── Config (2) ── + // ── Config (8) ── .arg(&gamma) .arg(&batch_i32) + .arg(&na_i32) + .arg(&v_min_f) + .arg(&v_max_f) + .arg(&b0_i32) + .arg(&b1_i32) + .arg(&b2_i32) .launch(LaunchConfig { grid_dim: (b as u32, 1, 1), block_dim: (256, 1, 1), @@ -3775,35 +3798,16 @@ fn compile_per_update_kernel(stream: &Arc) -> Result, - config: &GpuDqnTrainConfig, + _config: &GpuDqnTrainConfig, ) -> Result { - // c51_loss uses NUM_ATOMS for BOTH loop bounds AND array sizing — - // MUST be runtime-compiled with the correct #define injected. - let defines = format!( - "#define NUM_ATOMS {na}\n\ - #define V_MIN ({v_min:.1}f)\n\ - #define V_MAX ({v_max:.1}f)\n\ - #define BRANCH_0_SIZE {b0}\n\ - #define BRANCH_1_SIZE {b1}\n\ - #define BRANCH_2_SIZE {b2}\n", - na = config.num_atoms, - v_min = config.v_min, - v_max = config.v_max, - b0 = config.branch_0_size, - b1 = config.branch_1_size, - b2 = config.branch_2_size, - ); - let kernel_src = include_str!("c51_loss_kernel.cu"); - let full_source = format!("{defines}\n{kernel_src}"); let context = stream.context(); - let ptx = crate::cuda_pipeline::compile_ptx_for_device(&full_source, &context) - .map_err(|e| MLError::ModelError(format!("c51_loss_kernel compilation: {e}")))?; + let ptx = Ptx::from_binary(C51_LOSS_CUBIN.to_vec()); let module = context.load_module(ptx) - .map_err(|e| MLError::ModelError(format!("c51_loss module load: {e}")))?; + .map_err(|e| MLError::ModelError(format!("c51_loss cubin load: {e}")))?; module.load_function("c51_loss_batched") .map_err(|e| MLError::ModelError(format!("c51_loss_batched load: {e}"))) } @@ -3823,31 +3827,16 @@ fn compile_c51_grad_kernel( .map_err(|e| MLError::ModelError(format!("c51_grad_kernel load: {e}"))) } -/// Compile MSE loss kernel at runtime with correct NUM_ATOMS injected. -/// MSE uses NUM_ATOMS for loop bounds — MUST match runtime config. +/// Load the MSE loss kernel from precompiled cubin. +/// +/// All variable constants (num_atoms, v_min, v_max, branch sizes) are passed +/// as runtime kernel parameters — no NVRTC compilation needed. fn compile_mse_loss_kernel( stream: &Arc, - config: &GpuDqnTrainConfig, + _config: &GpuDqnTrainConfig, ) -> Result { - let defines = format!( - "#define NUM_ATOMS {na}\n\ - #define V_MIN ({v_min:.1}f)\n\ - #define V_MAX ({v_max:.1}f)\n\ - #define BRANCH_0_SIZE {b0}\n\ - #define BRANCH_1_SIZE {b1}\n\ - #define BRANCH_2_SIZE {b2}\n", - na = config.num_atoms, - v_min = config.v_min, - v_max = config.v_max, - b0 = config.branch_0_size, - b1 = config.branch_1_size, - b2 = config.branch_2_size, - ); - let kernel_src = include_str!("mse_loss_kernel.cu"); - let full_source = format!("{defines}\n{kernel_src}"); let context = stream.context(); - let ptx = crate::cuda_pipeline::compile_ptx_for_device(&full_source, &context) - .map_err(|e| MLError::ModelError(format!("mse_loss_kernel compilation: {e}")))?; + let ptx = Ptx::from_binary(MSE_LOSS_CUBIN.to_vec()); let module = context.load_module(ptx) .map_err(|e| MLError::ModelError(format!("mse_loss cubin load: {e}")))?; module.load_function("mse_loss_batched") diff --git a/crates/ml/src/cuda_pipeline/mod.rs b/crates/ml/src/cuda_pipeline/mod.rs index 228ac9328..0e1365448 100644 --- a/crates/ml/src/cuda_pipeline/mod.rs +++ b/crates/ml/src/cuda_pipeline/mod.rs @@ -100,9 +100,7 @@ pub fn optimal_launch_dims(n_items: u32, max_threads_per_block: u32) -> (u32, u3 /// so any change to the kernel source or network dimensions invalidates the cache. /// /// The fused experience collector kernel (4490 lines, branching+C51+NoisyNets) -// Re-export precompilation from ml-core so existing callers -// (`crate::cuda_pipeline::compile_ptx_for_device`) keep working. -pub use ml_core::cuda_compile::compile_ptx_for_device; +// compile_ptx_for_device re-export removed: all kernels now use precompiled cubins. /// Clone a `CudaSlice` via DtoD memcpy. /// diff --git a/crates/ml/src/cuda_pipeline/mse_loss_kernel.cu b/crates/ml/src/cuda_pipeline/mse_loss_kernel.cu index 37bd58ded..ebd0e35b0 100644 --- a/crates/ml/src/cuda_pipeline/mse_loss_kernel.cu +++ b/crates/ml/src/cuda_pipeline/mse_loss_kernel.cu @@ -12,47 +12,15 @@ * Grid = (batch_size, 1, 1) * Block = (BLOCK_THREADS=256, 1, 1) * - * Requires #define: NUM_ATOMS, V_MIN, V_MAX, BRANCH_0_SIZE, BRANCH_1_SIZE, BRANCH_2_SIZE + * All variable constants (NUM_ATOMS, V_MIN, V_MAX, BRANCH_*_SIZE) are + * passed as runtime kernel parameters. */ -#ifndef NUM_ATOMS -#define NUM_ATOMS 51 -#endif -#ifndef BRANCH_0_SIZE -#define BRANCH_0_SIZE 9 -#endif -#ifndef BRANCH_1_SIZE -#define BRANCH_1_SIZE 3 -#endif -#ifndef BRANCH_2_SIZE -#define BRANCH_2_SIZE 3 -#endif -#ifndef V_MIN -#define V_MIN (-25.0f) -#endif -#ifndef V_MAX -#define V_MAX (25.0f) -#endif -#define DELTA_Z ((V_MAX - V_MIN) / (float)(NUM_ATOMS - 1)) - -#define BRANCH_0_ATOMS (BRANCH_0_SIZE * NUM_ATOMS) -#define BRANCH_1_ATOMS (BRANCH_1_SIZE * NUM_ATOMS) -#define BRANCH_2_ATOMS (BRANCH_2_SIZE * NUM_ATOMS) #define MAX_BRANCH_SIZE 9 #define NUM_BRANCHES 3 #define NUM_WARPS 8 #define BLOCK_THREADS 256 -/* ── Shared memory layout (same as c51_loss_kernel.cu) ─────────────── */ -#define SHMEM_SUPPORT_OFF 0 -#define SHMEM_VAL_OFF (SHMEM_SUPPORT_OFF + NUM_ATOMS) -#define SHMEM_ADV_OFF (SHMEM_VAL_OFF + NUM_ATOMS) -#define SHMEM_LP_OFF (SHMEM_ADV_OFF + MAX_BRANCH_SIZE * NUM_ATOMS) -#define SHMEM_PROJ_OFF (SHMEM_LP_OFF + NUM_ATOMS) -#define SHMEM_CURRENT_LP_OFF (SHMEM_PROJ_OFF + NUM_ATOMS) -#define SHMEM_REDUCE_OFF (SHMEM_CURRENT_LP_OFF + NUM_ATOMS) -#define SHMEM_TOTAL_FLOATS (SHMEM_REDUCE_OFF + NUM_WARPS) - /* ── Device helpers (same as c51_loss_kernel.cu) ───────────────────── */ __device__ __forceinline__ float block_reduce_max( @@ -102,28 +70,29 @@ __device__ __forceinline__ float block_reduce_sum( * Also writes softmax probs to save_probs if save_probs != NULL. */ __device__ float block_softmax_expected_q( - float* __restrict__ shmem_logits, /* [NUM_ATOMS] in shmem (modified in-place) */ - const float* __restrict__ shmem_support, /* [NUM_ATOMS] z_j */ + float* __restrict__ shmem_logits, /* [num_atoms] in shmem (modified in-place) */ + const float* __restrict__ shmem_support, /* [num_atoms] z_j */ float* __restrict__ shmem_reduce, /* [NUM_WARPS] scratch */ - float* __restrict__ save_probs_out, /* [NUM_ATOMS] global mem or NULL */ - int tid + float* __restrict__ save_probs_out, /* [num_atoms] global mem or NULL */ + int tid, + int num_atoms ) { /* Pass 1: block-max */ float local_max = -1e30f; - for (int i = tid; i < NUM_ATOMS; i += BLOCK_THREADS) + for (int i = tid; i < num_atoms; i += BLOCK_THREADS) local_max = fmaxf(local_max, shmem_logits[i]); float block_max = block_reduce_max(local_max, shmem_reduce, tid); /* Pass 2: sum(exp(x - max)) */ float local_sum = 0.0f; - for (int i = tid; i < NUM_ATOMS; i += BLOCK_THREADS) + for (int i = tid; i < num_atoms; i += BLOCK_THREADS) local_sum += expf(shmem_logits[i] - block_max); float block_sum = block_reduce_sum(local_sum, shmem_reduce, tid); float log_sum_exp = logf(block_sum + 1e-10f); /* Pass 3: softmax probs + expected Q */ /* Write probs to shmem_logits in-place */ - for (int i = tid; i < NUM_ATOMS; i += BLOCK_THREADS) { + for (int i = tid; i < num_atoms; i += BLOCK_THREADS) { float lp = shmem_logits[i] - block_max - log_sum_exp; float prob = expf(lp); shmem_logits[i] = prob; /* overwrite logits with probs */ @@ -133,7 +102,7 @@ __device__ float block_softmax_expected_q( /* E[Q] = sum_j(prob_j * z_j) */ float local_eq = 0.0f; - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) local_eq += shmem_logits[j] * shmem_support[j]; return block_reduce_sum(local_eq, shmem_reduce, tid); } @@ -155,22 +124,22 @@ __device__ float block_softmax_expected_q( extern "C" __global__ void mse_loss_batched( /* ── Online network outputs on current STATES ─────────────────── */ - const float* __restrict__ on_value_logits, /* [B, NUM_ATOMS] */ - const float* __restrict__ on_adv_logits_b0, /* [B, BRANCH_0_SIZE * NUM_ATOMS] */ - const float* __restrict__ on_adv_logits_b1, /* [B, BRANCH_1_SIZE * NUM_ATOMS] */ - const float* __restrict__ on_adv_logits_b2, /* [B, BRANCH_2_SIZE * NUM_ATOMS] */ + const float* __restrict__ on_value_logits, /* [B, num_atoms] */ + const float* __restrict__ on_adv_logits_b0, /* [B, b0_size * num_atoms] */ + const float* __restrict__ on_adv_logits_b1, /* [B, b1_size * num_atoms] */ + const float* __restrict__ on_adv_logits_b2, /* [B, b2_size * num_atoms] */ /* ── Target network outputs on NEXT_STATES ────────────────────── */ - const float* __restrict__ tg_value_logits, /* [B, NUM_ATOMS] */ - const float* __restrict__ tg_adv_logits_b0, /* [B, BRANCH_0_SIZE * NUM_ATOMS] */ - const float* __restrict__ tg_adv_logits_b1, /* [B, BRANCH_1_SIZE * NUM_ATOMS] */ - const float* __restrict__ tg_adv_logits_b2, /* [B, BRANCH_2_SIZE * NUM_ATOMS] */ + const float* __restrict__ tg_value_logits, /* [B, num_atoms] */ + const float* __restrict__ tg_adv_logits_b0, /* [B, b0_size * num_atoms] */ + const float* __restrict__ tg_adv_logits_b1, /* [B, b1_size * num_atoms] */ + const float* __restrict__ tg_adv_logits_b2, /* [B, b2_size * num_atoms] */ /* ── Online network outputs on NEXT_STATES (Double DQN selector) */ - const float* __restrict__ on_next_value_logits, /* [B, NUM_ATOMS] */ - const float* __restrict__ on_next_adv_logits_b0, /* [B, BRANCH_0_SIZE * NUM_ATOMS] */ - const float* __restrict__ on_next_adv_logits_b1, /* [B, BRANCH_1_SIZE * NUM_ATOMS] */ - const float* __restrict__ on_next_adv_logits_b2, /* [B, BRANCH_2_SIZE * NUM_ATOMS] */ + const float* __restrict__ on_next_value_logits, /* [B, num_atoms] */ + const float* __restrict__ on_next_adv_logits_b0, /* [B, b0_size * num_atoms] */ + const float* __restrict__ on_next_adv_logits_b1, /* [B, b1_size * num_atoms] */ + const float* __restrict__ on_next_adv_logits_b2, /* [B, b2_size * num_atoms] */ /* ── Batch data ───────────────────────────────────────────────── */ const int* __restrict__ actions, /* [B] factored action indices 0-44 */ @@ -184,64 +153,87 @@ extern "C" __global__ void mse_loss_batched( float* __restrict__ total_loss, /* [1] batch mean loss (atomicAdd) */ /* ── Saved tensors for backward pass ─────────────────────────── */ - float* __restrict__ save_current_lp, /* [B, NUM_BRANCHES, NUM_ATOMS] online probs */ - float* __restrict__ save_projected, /* [B, NUM_BRANCHES, NUM_ATOMS] [td, E_Q, 0..] */ + float* __restrict__ save_current_lp, /* [B, NUM_BRANCHES, num_atoms] online probs */ + float* __restrict__ save_projected, /* [B, NUM_BRANCHES, num_atoms] [td, E_Q, 0..] */ /* ── Config ───────────────────────────────────────────────────── */ float gamma, - int batch_size + int batch_size, + int num_atoms, + float v_min, + float v_max, + int b0_size, + int b1_size, + int b2_size ) { extern __shared__ float shmem[]; - float* shmem_support = shmem + SHMEM_SUPPORT_OFF; - float* shmem_val = shmem + SHMEM_VAL_OFF; - float* shmem_adv = shmem + SHMEM_ADV_OFF; - float* shmem_lp = shmem + SHMEM_LP_OFF; - float* shmem_proj = shmem + SHMEM_PROJ_OFF; - float* shmem_current_lp = shmem + SHMEM_CURRENT_LP_OFF; - float* shmem_reduce = shmem + SHMEM_REDUCE_OFF; + /* ── Compute runtime shared memory offsets ────────────────────── */ + const int max_branch = max(max(b0_size, b1_size), b2_size); + int off_support = 0; + int off_val = off_support + num_atoms; + int off_adv = off_val + num_atoms; + int off_lp = off_adv + max_branch * num_atoms; + int off_proj = off_lp + num_atoms; + int off_current_lp = off_proj + num_atoms; + int off_reduce = off_current_lp + num_atoms; + + float* shmem_support = shmem + off_support; + float* shmem_val = shmem + off_val; + float* shmem_adv = shmem + off_adv; + float* shmem_lp = shmem + off_lp; + float* shmem_proj = shmem + off_proj; + float* shmem_current_lp = shmem + off_current_lp; + float* shmem_reduce = shmem + off_reduce; int sample_id = blockIdx.x; int tid = threadIdx.x; if (sample_id >= batch_size) return; + /* ── Compute runtime derived values ──────────────────────────── */ + float delta_z = (v_max - v_min) / (float)(num_atoms - 1); + + int b0_atoms = b0_size * num_atoms; + int b1_atoms = b1_size * num_atoms; + int b2_atoms = b2_size * num_atoms; + /* ── Precompute C51 support atoms once per block ─────────────── */ - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) - shmem_support[j] = V_MIN + (float)j * DELTA_Z; + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) + shmem_support[j] = v_min + (float)j * delta_z; __syncthreads(); /* ── Decode factored action into per-branch indices ──────────── */ int factored_action = actions[sample_id]; - int max_action = BRANCH_0_SIZE * BRANCH_1_SIZE * BRANCH_2_SIZE; + int max_action = b0_size * b1_size * b2_size; if (factored_action < 0 || factored_action >= max_action) factored_action = 0; - int a0 = factored_action / (BRANCH_1_SIZE * BRANCH_2_SIZE); - int a1 = (factored_action / BRANCH_2_SIZE) % BRANCH_1_SIZE; - int a2 = factored_action % BRANCH_2_SIZE; + int a0 = factored_action / (b1_size * b2_size); + int a1 = (factored_action / b2_size) % b1_size; + int a2 = factored_action % b2_size; const int branch_action[NUM_BRANCHES] = { a0, a1, a2 }; - const int branch_sizes[NUM_BRANCHES] = { BRANCH_0_SIZE, BRANCH_1_SIZE, BRANCH_2_SIZE }; + const int branch_sizes[NUM_BRANCHES] = { b0_size, b1_size, b2_size }; const float* on_adv_row[NUM_BRANCHES] = { - on_adv_logits_b0 + (long long)sample_id * BRANCH_0_ATOMS, - on_adv_logits_b1 + (long long)sample_id * BRANCH_1_ATOMS, - on_adv_logits_b2 + (long long)sample_id * BRANCH_2_ATOMS + on_adv_logits_b0 + (long long)sample_id * b0_atoms, + on_adv_logits_b1 + (long long)sample_id * b1_atoms, + on_adv_logits_b2 + (long long)sample_id * b2_atoms }; const float* tg_adv_row[NUM_BRANCHES] = { - tg_adv_logits_b0 + (long long)sample_id * BRANCH_0_ATOMS, - tg_adv_logits_b1 + (long long)sample_id * BRANCH_1_ATOMS, - tg_adv_logits_b2 + (long long)sample_id * BRANCH_2_ATOMS + tg_adv_logits_b0 + (long long)sample_id * b0_atoms, + tg_adv_logits_b1 + (long long)sample_id * b1_atoms, + tg_adv_logits_b2 + (long long)sample_id * b2_atoms }; const float* on_next_adv_row[NUM_BRANCHES] = { - on_next_adv_logits_b0 + (long long)sample_id * BRANCH_0_ATOMS, - on_next_adv_logits_b1 + (long long)sample_id * BRANCH_1_ATOMS, - on_next_adv_logits_b2 + (long long)sample_id * BRANCH_2_ATOMS + on_next_adv_logits_b0 + (long long)sample_id * b0_atoms, + on_next_adv_logits_b1 + (long long)sample_id * b1_atoms, + on_next_adv_logits_b2 + (long long)sample_id * b2_atoms }; - const float* on_val_row = on_value_logits + (long long)sample_id * NUM_ATOMS; - const float* tg_val_row = tg_value_logits + (long long)sample_id * NUM_ATOMS; - const float* on_next_val_row = on_next_value_logits + (long long)sample_id * NUM_ATOMS; + const float* on_val_row = on_value_logits + (long long)sample_id * num_atoms; + const float* tg_val_row = tg_value_logits + (long long)sample_id * num_atoms; + const float* on_next_val_row = on_next_value_logits + (long long)sample_id * num_atoms; float reward = rewards[sample_id]; float done = dones[sample_id]; @@ -253,16 +245,16 @@ extern "C" __global__ void mse_loss_batched( for (int d = 0; d < NUM_BRANCHES; d++) { int n_d = branch_sizes[d]; int a_d = branch_action[d]; - int n_atoms = n_d * NUM_ATOMS; + int n_atoms = n_d * num_atoms; - long long save_off = ((long long)sample_id * NUM_BRANCHES + d) * NUM_ATOMS; + long long save_off = ((long long)sample_id * NUM_BRANCHES + d) * num_atoms; /* ═══════════════════════════════════════════════════════════ * STEP a: Online E[Q] for the taken action a_d * ═══════════════════════════════════════════════════════════ */ /* Load online value logits */ - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) shmem_val[j] = on_val_row[j]; __syncthreads(); @@ -272,26 +264,26 @@ extern "C" __global__ void mse_loss_batched( __syncthreads(); /* Dueling for action a_d: Q[j] = V[j] + A[a_d,j] - mean_a(A[*,j]) */ - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) { + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) { float a_mean = 0.0f; for (int a = 0; a < n_d; a++) - a_mean += shmem_adv[a * NUM_ATOMS + j]; + a_mean += shmem_adv[a * num_atoms + j]; a_mean /= (float)n_d; - shmem_lp[j] = shmem_val[j] + shmem_adv[a_d * NUM_ATOMS + j] - a_mean; + shmem_lp[j] = shmem_val[j] + shmem_adv[a_d * num_atoms + j] - a_mean; } __syncthreads(); /* softmax + E[Q], saving probs to save_current_lp */ float online_eq = block_softmax_expected_q( shmem_lp, shmem_support, shmem_reduce, - save_current_lp + save_off, tid + save_current_lp + save_off, tid, num_atoms ); /* ═══════════════════════════════════════════════════════════ * STEP b: Double DQN — argmax action on online_next network * ═══════════════════════════════════════════════════════════ */ - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) shmem_val[j] = on_next_val_row[j]; __syncthreads(); @@ -300,10 +292,10 @@ extern "C" __global__ void mse_loss_batched( __syncthreads(); /* Per-atom mean advantage for dueling */ - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) { + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) { float a_mean = 0.0f; for (int a = 0; a < n_d; a++) - a_mean += shmem_adv[a * NUM_ATOMS + j]; + a_mean += shmem_adv[a * num_atoms + j]; a_mean /= (float)n_d; shmem_proj[j] = a_mean; } @@ -313,13 +305,13 @@ extern "C" __global__ void mse_loss_batched( __shared__ int best_next_a_shared; for (int a = 0; a < n_d; a++) { - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) - shmem_lp[j] = shmem_val[j] + shmem_adv[a * NUM_ATOMS + j] - shmem_proj[j]; + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) + shmem_lp[j] = shmem_val[j] + shmem_adv[a * num_atoms + j] - shmem_proj[j]; __syncthreads(); float eq = block_softmax_expected_q( shmem_lp, shmem_support, shmem_reduce, - (float*)0, tid /* no save needed for action selection */ + (float*)0, tid, num_atoms /* no save needed for action selection */ ); if (tid == 0) eq_per_action[a] = eq; __syncthreads(); @@ -338,7 +330,7 @@ extern "C" __global__ void mse_loss_batched( * STEP c: Target E[Q] for best_next action * ═══════════════════════════════════════════════════════════ */ - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) shmem_val[j] = tg_val_row[j]; __syncthreads(); @@ -346,18 +338,18 @@ extern "C" __global__ void mse_loss_batched( shmem_adv[i] = tg_adv_row[d][i]; __syncthreads(); - for (int j = tid; j < NUM_ATOMS; j += BLOCK_THREADS) { + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) { float a_mean = 0.0f; for (int a = 0; a < n_d; a++) - a_mean += shmem_adv[a * NUM_ATOMS + j]; + a_mean += shmem_adv[a * num_atoms + j]; a_mean /= (float)n_d; - shmem_lp[j] = shmem_val[j] + shmem_adv[best_next_a * NUM_ATOMS + j] - a_mean; + shmem_lp[j] = shmem_val[j] + shmem_adv[best_next_a * num_atoms + j] - a_mean; } __syncthreads(); float target_eq = block_softmax_expected_q( shmem_lp, shmem_support, shmem_reduce, - (float*)0, tid + (float*)0, tid, num_atoms ); /* ═══════════════════════════════════════════════════════════ @@ -379,7 +371,7 @@ extern "C" __global__ void mse_loss_batched( save_projected[save_off + 1] = online_eq; } /* Zero padding */ - for (int j = tid + 2; j < NUM_ATOMS; j += BLOCK_THREADS) + for (int j = tid + 2; j < num_atoms; j += BLOCK_THREADS) save_projected[save_off + j] = 0.0f; __syncthreads(); diff --git a/crates/ml/src/cuda_pipeline/trade_physics.cuh b/crates/ml/src/cuda_pipeline/trade_physics.cuh new file mode 100644 index 000000000..afb47e863 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/trade_physics.cuh @@ -0,0 +1,159 @@ +/** + * Shared trade physics for DQN environment step kernels. + * + * This header defines the SINGLE SOURCE OF TRUTH for: + * - Factored action decoding (exposure index extraction) + * - Exposure-to-position mapping (index → fraction → contracts) + * - Hold enforcement (block exits/reversals during min_hold_bars) + * - Transaction cost computation (Almgren-Chriss impact model) + * - Capital floor circuit breaker (75% of peak equity) + * - Hold time tracking (bars in current position) + * + * Both backtest_env_kernel.cu and experience_kernels.cu MUST include + * this header to ensure train/eval consistency. Any change here + * automatically propagates to both. + * + * All functions are __device__ __forceinline__ for zero call overhead. + */ + +#ifndef TRADE_PHYSICS_CUH +#define TRADE_PHYSICS_CUH + +/* ── Factored action → exposure index ─────────────────────────────────── */ + +__device__ __forceinline__ int decode_exposure_index( + int action_val, int b0_size, int b1_size, int b2_size +) { + int exposure_idx; + if (b1_size > 0 && b2_size > 0) { + int denom = b1_size * b2_size; + exposure_idx = action_val / denom; + } else { + exposure_idx = action_val; + } + if (exposure_idx < 0) exposure_idx = 0; + if (exposure_idx >= b0_size) exposure_idx = b0_size - 1; + return exposure_idx; +} + +/* ── Exposure index → target position ─────────────────────────────────── */ + +__device__ __forceinline__ float compute_target_position( + int exposure_idx, int b0_size, float max_position +) { + float step_size = (b0_size > 1) ? 2.0f / (float)(b0_size - 1) : 0.0f; + float fraction = -1.0f + (float)exposure_idx * step_size; + return fraction * max_position; +} + +/* ── Margin-aware position cap ───────────────────────────────────────── */ +/* Real brokers enforce margin requirements: you can't hold more contracts + * than your equity supports. ES futures require ~$15K initial margin per + * contract ($12K maintenance). Without this, a depleted account holds the + * same position as a full account → catastrophic overleveraging. + * + * This caps the position at equity / margin_per_contract, ensuring the + * portfolio can't take outsized risk with depleted capital. + * + * margin_per_contract: CME initial margin per contract (~$15,000 for ES) + * Passed as 0 to disable (backward compat with training kernel). */ + +__device__ __forceinline__ float apply_margin_cap( + float target_position, + float equity, + float margin_per_contract +) { + if (margin_per_contract <= 0.0f || equity <= 0.0f) return target_position; + float max_contracts = equity / margin_per_contract; + if (max_contracts < 0.01f) max_contracts = 0.0f; /* can't afford any position */ + float abs_target = fabsf(target_position); + if (abs_target > max_contracts) { + float sign = (target_position > 0.0f) ? 1.0f : -1.0f; + return sign * max_contracts; + } + return target_position; +} + +/* ── Order type index from factored action ────────────────────────────── */ + +__device__ __forceinline__ int decode_order_type( + int action_val, int b1_size, int b2_size +) { + return (b1_size > 0 && b2_size > 0) ? (action_val / b2_size) % b1_size : 0; +} + +/* ── Hold enforcement: block exits and reversals ─────────────────────── */ + +__device__ __forceinline__ float enforce_hold( + float prev_position, + float target_exposure, + float hold_time, + int min_hold_bars, + int is_last_bar /* bypass hold on last bar of episode/window */ +) { + int prev_sign = (prev_position > 0.001f) ? 1 : ((prev_position < -0.001f) ? -1 : 0); + int curr_sign = (target_exposure > 0.001f) ? 1 : ((target_exposure < -0.001f) ? -1 : 0); + int wants_exit = (prev_sign != 0 && curr_sign == 0); + int wants_reversal = (prev_sign != 0 && curr_sign != 0 && prev_sign != curr_sign); + + if (hold_time > 0.0f && hold_time < (float)min_hold_bars + && (wants_exit || wants_reversal) + && !is_last_bar) { + return prev_position; /* Override: keep position */ + } + return target_exposure; +} + +/* ── Transaction cost (Almgren-Chriss impact model) ──────────────────── */ + +__device__ __forceinline__ float compute_tx_cost( + float delta, /* position change (signed) */ + float close, /* current price */ + float tx_cost_bps, /* base cost multiplier */ + float spread_cost, /* bid-ask spread cost per unit */ + float max_position, /* for impact scaling */ + int order_type_idx /* 0=Market, 1=IoC, 2=LimitMaker */ +) { + float abs_delta = fabsf(delta); + float spread_scale = sqrtf(abs_delta / fmaxf(max_position, 0.01f)); + if (spread_scale < 1.0f) spread_scale = 1.0f; + float impact_scale = 1.0f; + float order_premium = (order_type_idx == 0) ? 0.0f + : (order_type_idx == 1) ? 0.0002f /* IoC: +2 bps */ + : -0.0005f; /* LimitMaker: -5 bps rebate */ + return abs_delta * close * (tx_cost_bps * 0.0001f * spread_scale * impact_scale + order_premium) + + abs_delta * spread_cost * 0.5f; +} + +/* ── Capital floor circuit breaker (PDT $25K rule) ───────────────────── */ +/* Triggers when equity drops 25% from peak (floor = 75% of peak). */ +/* With $35K initial: floor = $26.25K — above $25K PDT minimum. */ +/* Matches training kernel (experience_env_step) capital_floor logic. */ + +__device__ __forceinline__ int check_capital_floor( + float value, float max_equity +) { + float capital_floor = max_equity * 0.75f; + return (value < capital_floor && max_equity > 1.0f) ? 1 : 0; +} + +/* ── Hold time tracking ──────────────────────────────────────────────── */ + +__device__ __forceinline__ float update_hold_time( + float prev_position, float new_position, float hold_time +) { + if (fabsf(new_position) > 0.001f) { + if (fabsf(prev_position) < 0.001f) { + return 1.0f; /* Entry from flat */ + } + int prev_sign = (prev_position > 0.001f) ? 1 : -1; + int new_sign = (new_position > 0.001f) ? 1 : -1; + if (new_sign == prev_sign) { + return hold_time + 1.0f; /* Same direction — increment */ + } + return 1.0f; /* Reversal — new trade */ + } + return 0.0f; /* Flat */ +} + +#endif /* TRADE_PHYSICS_CUH */ diff --git a/crates/ml/src/hyperopt/adapters/dqn.rs b/crates/ml/src/hyperopt/adapters/dqn.rs index e5dcb4d03..1220ae262 100644 --- a/crates/ml/src/hyperopt/adapters/dqn.rs +++ b/crates/ml/src/hyperopt/adapters/dqn.rs @@ -3,7 +3,7 @@ //! This module provides a production-ready adapter for optimizing DQN //! hyperparameters using the generic optimization framework. It implements: //! -//! - 46D continuous parameter space with log-scale handling +//! - 22D continuous parameter space with log-scale handling (consolidated from 46D) //! - Training wrapper that integrates with existing DQN pipeline //! - Backtest-based objective using Sharpe ratio (production metric) //! - Prioritized Experience Replay (PER) tuning for Rainbow DQN performance (+25-40% convergence speed) @@ -94,7 +94,7 @@ pub struct BacktestMetrics { pub information_ratio: f64, /// Omega Ratio (upside vs. downside potential) pub omega_ratio: f64, - /// Number of unique actions used during eval (out of 5 exposure actions) + /// Number of unique exposure levels used during eval (out of 9: S100..L100) pub unique_actions: usize, /// Buy action percentage in backtest [0.0, 1.0] (Long50 + Long100) pub buy_action_pct: f64, @@ -144,30 +144,23 @@ const TRIAL_VRAM_MB: f64 = 7000.0; /// Corrected: was 0.02 (20KB) — 100x too high const MB_PER_SAMPLE: f64 = 0.0005; -/// DQN hyperparameter space (46D continuous) +/// DQN hyperparameter space (22D continuous, consolidated from 46D) /// -/// Fixed params (not in search space): -/// - noisy_epsilon_floor: 0.10 (minimum random exploration to prevent action collapse) +/// Only 22 parameters are in the search space; all others are fixed at defaults. /// -/// Tuned params restored in C1-C4: -/// - entropy_coefficient: (0.05, 0.5) — 10x wider for Q-values ~7.0 -/// - weight_decay: log(1e-4, 1e-2) — 10x stronger to prevent overfitting -/// - count_bonus_coefficient: (0.0, 0.3) — UCB exploration bonus, complementary to NoisyNet -/// - sharpe_weight: (0.0, 0.5) — tunable Sharpe ratio blending in composite reward -/// -/// Composite reward weights (indices 31-38): -/// - w_dsr: DSR (Sharpe) weight -/// - w_pnl: Normalized PnL weight -/// - w_dd: Drawdown penalty weight -/// - w_idle: Idle penalty weight (replaces hold_penalty_weight) -/// - dd_threshold: Drawdown tolerance before penalty -/// - loss_aversion: Asymmetric loss scaling (prospect theory) -/// - time_decay_rate: Position staleness rent per step +/// ## Parameters in search space (22D): +/// - learning_rate, batch_size, gamma, buffer_size, max_position_absolute +/// - huber_delta, entropy_coefficient, transaction_cost_multiplier +/// - per_alpha, per_beta_start, dueling_hidden_dim, n_steps, tau +/// - num_atoms, v_max, noisy_sigma_init, minimum_profit_factor +/// - weight_decay, hidden_dim_base, min_hold_bars, c51_warmup_epochs, iqn_lambda /// /// ## Parameter Scaling /// -/// - **Log-scale**: Learning rate, buffer_size, huber_delta, noisy_sigma_init, tau, weight_decay -/// - **Linear scale**: Batch size, gamma, max_position, entropy, tx_cost, per_alpha, per_beta, v_min, v_max, dueling_hidden_dim, n_steps, num_atoms, reward weights +/// - **Log-scale**: learning_rate, buffer_size, huber_delta, tau, weight_decay +/// - **Linear scale**: batch_size, gamma, max_position, entropy, tx_cost, per_alpha, +/// per_beta, dueling_hidden_dim, n_steps, num_atoms, v_max, noisy_sigma_init, +/// minimum_profit_factor, hidden_dim_base, min_hold_bars, c51_warmup_epochs, iqn_lambda /// /// This scaling ensures efficient exploration by egobox's Bayesian optimization. #[derive(Debug, Clone, PartialEq, Serialize, Deserialize)] @@ -509,148 +502,35 @@ impl Default for DQNParams { impl ParameterSpace for DQNParams { fn continuous_bounds() -> Vec<(f64, f64)> { - // 46D search space (Task 11: +4 dims at indices 41-44, Task 4: +1 dim at index 45) - // All bounds loaded from config/training/dqn-hyperopt.toml [search_space]. - // To change search ranges, edit the TOML — no code changes needed. - // Log-scale transforms are applied here; TOML stores linear values. + // 22D search space (consolidated from 46D — marginal/dead params fixed at defaults) + // Log-scale transforms are applied here for parameters that span orders of magnitude. let hp = crate::training_profile::HyperoptProfile::load("dqn-hyperopt"); let b = |name: &str, default: (f64, f64)| hp.bound(name, default); - let lr = b("learning_rate", (1e-5, 3e-4)); + let lr = b("learning_rate", (1e-5, 1e-3)); let bs = b("batch_size", (64.0, 512.0)); - let gm = b("gamma", (0.88, 0.99)); - let bf = b("buffer_size", (50_000.0, 100_000.0)); + let gm = b("gamma", (0.95, 0.999)); + let bf = b("buffer_size", (10_000.0, 100_000.0)); let mp = b("max_position_absolute", (1.0, 4.0)); - let hd = b("huber_delta", (10.0, 40.0)); + let hd = b("huber_delta", (0.1, 20.0)); let ec = b("entropy_coefficient", (0.05, 0.5)); let tc = b("transaction_cost_multiplier", (0.5, 2.0)); let pa = b("per_alpha", (0.4, 0.8)); let pb = b("per_beta_start", (0.2, 0.6)); - let mut vr = b("v_range", (10.0, 50.0)); + let dh = b("dueling_hidden_dim", (64.0, 256.0)); + let nst = b("n_steps", (1.0, 10.0)); + let tau = b("tau", (0.0001, 0.01)); + let na = b("num_atoms", (21.0, 101.0)); + let vm = b("v_max", (5.0, 300.0)); let ns = b("noisy_sigma_init", (0.1, 1.0)); - let mut dh = b("dueling_hidden_dim", (128.0, 512.0)); - let nst = b("n_steps", (3.0, 5.0)); - let mut na = b("num_atoms", (11.0, 101.0)); - let wd = b("weight_decay", (1e-4, 1e-2)); - let kf = b("kelly_fractional", (0.25, 0.75)); - let km = b("kelly_max_fraction", (0.1, 0.5)); - let vw = b("volatility_window", (10.0, 30.0)); - let tau = b("tau", (0.005, 0.01)); - let mut hdim = b("hidden_dim_base", (128.0, 512.0)); - let cql = b("cql_alpha", (0.0, 1.0)); - let lrd = b("lr_decay_type", (0.0, 2.0)); - let dsr = b("dsr_eta", (0.001, 0.05)); let mpf = b("minimum_profit_factor", (1.1, 2.0)); - let cb = b("count_bonus_coefficient", (0.0, 0.3)); - let sw = b("sharpe_weight", (0.0, 0.5)); - let mut bh = b("branch_hidden_dim", (64.0, 256.0)); - let ga = b("gradient_accumulation_steps", (1.0, 1.0)); - let iq = b("iqn_lambda", (0.0, 2.0)); - - // Composite reward weights (indices 31-38) - let mut rw_dsr = b("w_dsr", (0.1, 2.0)); - let mut rw_pnl = b("w_pnl", (0.0, 1.0)); - let mut rw_dd = b("w_dd", (0.0, 5.0)); - let mut rw_idle = b("w_idle", (0.0, 0.1)); - let mut rw_ddt = b("dd_threshold", (0.01, 0.10)); - let mut rw_la = b("loss_aversion", (1.0, 3.0)); - let mut rw_tdr = b("time_decay_rate", (0.0001, 0.005)); - let mut rw_qgt = b("q_gap_threshold", (0.0, 0.5)); - let sn_sigma = b("spectral_norm_sigma_max", (1.0, 10.0)); - let cw = b("c51_warmup_epochs", (0.0, 20.0)); - // Task 11: New hyperopt-tunable parameters - let hr = b("her_ratio", (0.0, 0.5)); - let cwt = b("curiosity_weight_tunable", (0.0, 0.2)); - let ucvar = b("use_cvar_action_selection", (0.0, 1.0)); - let cva = b("cvar_alpha", (0.01, 0.2)); + let wd = b("weight_decay", (1e-5, 1e-2)); + let hdim = b("hidden_dim_base", (64.0, 256.0)); let mhb = b("min_hold_bars", (2.0, 20.0)); + let cw = b("c51_warmup_epochs", (3.0, 15.0)); + let iq = b("iqn_lambda", (0.0, 1.0)); - // Two-phase hyperopt: fix architecture OR dynamics params to single-point bounds. - // Single-point bounds (v, v) make the optimizer trivially converge on those - // dimensions, focusing search power on the free dimensions. - use crate::training_profile::{get_hyperopt_phase, HyperoptPhase}; - let phase = get_hyperopt_phase(); - - match phase { - HyperoptPhase::Fast => { - // Phase 1: fix architecture to small network from [phase_fast] TOML. - // Also fix reward weights to defaults (search learning dynamics only). - // All values MUST exist in [phase_fast] — no silent defaults. - let pf = hp.phase_fast.as_ref() - .expect("[phase_fast] section missing from dqn-hyperopt.toml — required for --phase fast"); - let fix_hdim = pf.hidden_dim_base - .expect("phase_fast.hidden_dim_base missing from dqn-hyperopt.toml"); - let fix_na = pf.num_atoms - .expect("phase_fast.num_atoms missing from dqn-hyperopt.toml"); - let fix_bh = pf.branch_hidden_dim - .expect("phase_fast.branch_hidden_dim missing from dqn-hyperopt.toml"); - let fix_dh = pf.dueling_hidden_dim - .expect("phase_fast.dueling_hidden_dim missing from dqn-hyperopt.toml"); - let fix_vr = pf.v_range - .expect("phase_fast.v_range missing from dqn-hyperopt.toml"); - - hdim = (fix_hdim, fix_hdim); - na = (fix_na, fix_na); - bh = (fix_bh, fix_bh); - dh = (fix_dh, fix_dh); - vr = (fix_vr, fix_vr); - - // Fix reward weights to defaults in Phase Fast (not searched) - let fix_w_dsr = pf.w_dsr - .expect("phase_fast.w_dsr missing from dqn-hyperopt.toml"); - let fix_w_pnl = pf.w_pnl - .expect("phase_fast.w_pnl missing from dqn-hyperopt.toml"); - let fix_w_dd = pf.w_dd - .expect("phase_fast.w_dd missing from dqn-hyperopt.toml"); - let fix_w_idle = pf.w_idle - .expect("phase_fast.w_idle missing from dqn-hyperopt.toml"); - let fix_ddt = pf.dd_threshold - .expect("phase_fast.dd_threshold missing from dqn-hyperopt.toml"); - let fix_la = pf.loss_aversion - .expect("phase_fast.loss_aversion missing from dqn-hyperopt.toml"); - let fix_tdr = pf.time_decay_rate - .expect("phase_fast.time_decay_rate missing from dqn-hyperopt.toml"); - let fix_qgt = pf.q_gap_threshold.unwrap_or(0.0); - - rw_dsr = (fix_w_dsr, fix_w_dsr); - rw_pnl = (fix_w_pnl, fix_w_pnl); - rw_dd = (fix_w_dd, fix_w_dd); - rw_idle = (fix_w_idle, fix_w_idle); - rw_ddt = (fix_ddt, fix_ddt); - rw_la = (fix_la, fix_la); - rw_tdr = (fix_tdr, fix_tdr); - rw_qgt = (fix_qgt, fix_qgt); - - tracing::info!( - "Phase FAST: fixed architecture (hidden_dim={}, num_atoms={}, branch_hidden={}, dueling_hidden={}, v_range={}) + reward weights", - fix_hdim, fix_na, fix_bh, fix_dh, fix_vr - ); - } - HyperoptPhase::Full(ref best_vec) => { - // Phase 2: fix learning dynamics from Phase 1 best params. - // Search only architecture dims (~5D free, ~33D fixed). - // Architecture dims (indices 10, 13, 15, 21, 28) stay free. - tracing::info!("Phase FULL: fixing {} dynamics params from Phase 1 best", best_vec.len()); - // Bounds will be applied after vec construction below. - } - HyperoptPhase::Reward(ref best_vec) => { - // Phase 3: fix dynamics + architecture from Phase 2 best params. - // Search ONLY reward weights (indices 31-38, 7D). - // All other dims fixed to Phase 2's best values. - tracing::info!("Phase REWARD: fixing {} dynamics+architecture params, searching 7 reward dims", best_vec.len()); - // Bounds will be applied after vec construction below. - } - HyperoptPhase::Risk(ref best_vec) => { - // Phase 4: fix dynamics + architecture from Phase 2 best params. - // Search ONLY risk params (8D): - // kelly_fractional (17), dd_threshold (35), loss_aversion (36), - // time_decay_rate (37), q_gap_threshold (38), w_dsr (31), w_pnl (32), w_dd (33). - tracing::info!("Phase RISK: fixing {} dynamics+architecture params, searching 8 risk dims", best_vec.len()); - // Bounds will be applied after vec construction below. - } - } - - let mut bounds = vec![ + let bounds = vec![ (lr.0.ln(), lr.1.ln()), // 0: learning_rate (log scale) bs, // 1: batch_size gm, // 2: gamma @@ -661,204 +541,55 @@ impl ParameterSpace for DQNParams { tc, // 7: transaction_cost_multiplier pa, // 8: per_alpha pb, // 9: per_beta_start - vr, // 10: v_range - (1.0, 1.0), // 11: use_branching (FIXED: always true) - (ns.0.ln(), ns.1.ln()), // 12: noisy_sigma_init (log scale) - dh, // 13: dueling_hidden_dim - nst, // 14: n_steps - na, // 15: num_atoms - (wd.0.ln(), wd.1.ln()), // 16: weight_decay (log scale) - kf, // 17: kelly_fractional - km, // 18: kelly_max_fraction - vw, // 19: volatility_window - (tau.0.ln(), tau.1.ln()), // 20: tau (log scale) - hdim, // 21: hidden_dim_base - cql, // 22: cql_alpha - lrd, // 23: lr_decay_type - (dsr.0.ln(), dsr.1.ln()), // 24: dsr_eta (log scale) - mpf, // 25: minimum_profit_factor - cb, // 26: count_bonus_coefficient - sw, // 27: sharpe_weight - bh, // 28: branch_hidden_dim - ga, // 29: gradient_accumulation_steps - iq, // 30: iqn_lambda - // Composite reward weights (indices 31-38) - rw_dsr, // 31: w_dsr - rw_pnl, // 32: w_pnl - rw_dd, // 33: w_dd - rw_idle, // 34: w_idle - rw_ddt, // 35: dd_threshold - rw_la, // 36: loss_aversion - rw_tdr, // 37: time_decay_rate - rw_qgt, // 38: q_gap_threshold - sn_sigma, // 39: spectral_norm_sigma_max - cw, // 40: c51_warmup_epochs - // Task 11: New hyperopt-tunable parameters - hr, // 41: her_ratio - cwt, // 42: curiosity_weight_tunable - ucvar, // 43: use_cvar_action_selection - cva, // 44: cvar_alpha - mhb, // 45: min_hold_bars [2, 20] + dh, // 10: dueling_hidden_dim + nst, // 11: n_steps + (tau.0.ln(), tau.1.ln()), // 12: tau (log scale) + na, // 13: num_atoms + vm, // 14: v_max (v_min = -v_max) + ns, // 15: noisy_sigma_init + mpf, // 16: minimum_profit_factor + (wd.0.ln(), wd.1.ln()), // 17: weight_decay (log scale) + hdim, // 18: hidden_dim_base + mhb, // 19: min_hold_bars + cw, // 20: c51_warmup_epochs + iq, // 21: iqn_lambda ]; - // Phase FULL: fix all non-architecture dims to Phase 1 best values. - // Architecture dims that stay FREE: 10 (v_range), 13 (dueling_hidden), - // 15 (num_atoms), 21 (hidden_dim_base), 28 (branch_hidden_dim). - if let HyperoptPhase::Full(ref best_vec) = phase { - let arch_dims: &[usize] = &[10, 13, 15, 21, 28]; - for (i, bv) in best_vec.iter().enumerate() { - if i < bounds.len() && !arch_dims.contains(&i) { - bounds[i] = (*bv, *bv); - } - } - } - - // Phase REWARD: fix ALL dims EXCEPT reward weights (indices 31-38). - // This is the fastest phase — only 7D search, ~10s/trial. - if let HyperoptPhase::Reward(ref best_vec) = phase { - let reward_dims: &[usize] = &[31, 32, 33, 34, 35, 36, 37, 38]; - for (i, bv) in best_vec.iter().enumerate() { - if i < bounds.len() && !reward_dims.contains(&i) { - bounds[i] = (*bv, *bv); - } - } - } - - // Phase RISK: fix ALL dims EXCEPT risk params (8D). - // Risk dims: kelly_fractional (17), w_dsr (31), w_pnl (32), w_dd (33), - // dd_threshold (35), loss_aversion (36), time_decay_rate (37), - // q_gap_threshold (38). - if let HyperoptPhase::Risk(ref best_vec) = phase { - let risk_dims: &[usize] = &[17, 31, 32, 33, 35, 36, 37, 38]; - for (i, bv) in best_vec.iter().enumerate() { - if i < bounds.len() && !risk_dims.contains(&i) { - bounds[i] = (*bv, *bv); - } - } - } - bounds } fn from_continuous(x: &[f64]) -> Result { - if x.len() < 39 { - return Err(MLError::ConfigError(format!("Expected at least 39 continuous parameters, got {}", x.len()))); + if x.len() < 22 { + return Err(MLError::ConfigError(format!("Expected at least 22 continuous parameters, got {}", x.len()))); } - // Task 11: New parameters (indices 41-44) — backwards compatible via .get() - let her_ratio = x.get(41).copied().unwrap_or(0.2).clamp(0.0, 0.5); - let curiosity_weight_tunable = x.get(42).copied().unwrap_or(0.1).clamp(0.0, 0.2); - let use_cvar_action_selection = x.get(43).copied().unwrap_or(1.0).clamp(0.0, 1.0); - let cvar_alpha = x.get(44).copied().unwrap_or(0.05).clamp(0.01, 0.2); - // min_hold_bars (index 45): integer, rounded from continuous [2, 20] - let min_hold_bars = (x.get(45).copied().unwrap_or(5.0).round() as usize).clamp(2, 20); - // === 26 TUNED parameters (from search space) === - // C2: Removed curiosity_weight and noisy_epsilon_floor (exploration stacking) + // === 22 TUNED parameters (from search space) === let learning_rate = x[0].exp(); let mut batch_size = x[1].round().clamp(64.0, 1024.0) as usize; - let buffer_size = x[3].exp().round().max(50_000.0) as usize; + let gamma = x[2].clamp(0.95, 0.999); + let buffer_size = x[3].exp().round().max(10_000.0) as usize; let max_position_absolute = x[4].clamp(1.0, 4.0); let huber_delta = x[5].exp(); let entropy_coefficient = x[6]; let transaction_cost_multiplier = x[7]; let per_alpha = x[8].clamp(0.4, 0.8); let per_beta_start = x[9].clamp(0.2, 0.6); + let dueling_hidden_dim = (x[10].round() / 64.0).round() * 64.0; + let dueling_hidden_dim = dueling_hidden_dim.clamp(64.0, 256.0) as usize; + let n_steps = x[11].round().clamp(1.0, 10.0) as usize; + let tau = x[12].exp().clamp(0.0001, 0.01); + let num_atoms = (x[13].round() as usize).max(21); + let v_max = x[14].clamp(5.0, 300.0); + let v_min = -v_max; + let noisy_sigma_init = x[15].clamp(0.1, 1.0); + let minimum_profit_factor = x[16].clamp(1.1, 2.0); + let weight_decay = x[17].exp().clamp(1e-5, 1e-2); + let hidden_dim_base = ((x[18].round() / 64.0).round() * 64.0).clamp(64.0, 256.0) as usize; + let min_hold_bars = (x[19].round() as usize).clamp(2, 20); + let c51_warmup_epochs = x[20].round().clamp(3.0, 15.0) as usize; + let iqn_lambda = x[21].clamp(0.0, 1.0); - // Rainbow DQN extensions — DYNAMIC v_range from gamma + reward_scale. - // V_max = reward_scale / (1 - gamma) with 20% headroom. - // Q* ~ reward_scale / (1-gamma). gamma=0.95, reward_scale=10 → Q*=200. Headroom 1.2x. - let gamma_val = x[2].clamp(0.88, 0.99); - // Reward v6: REWARD_SCALE=10 (single source of truth via DQNHyperparameters::reward_scale). - let reward_scale = 10.0_f64; - let v_range = (reward_scale / (1.0 - gamma_val) * 1.2).clamp(20.0, 300.0); - let v_min = -v_range; - let v_max = v_range; - // x[11] was use_branching — always true, slot kept for index compatibility - let noisy_sigma_init = x[12].exp().clamp(0.1, 1.0); - let dueling_hidden_dim = (x[13].round() / 128.0).round() * 128.0; - // Upper clamp matches large_gpu expansion in continuous_bounds_for (768); - // VRAM estimator already caps the search bound for smaller GPUs. - let dueling_hidden_dim = dueling_hidden_dim.clamp(128.0, 768.0) as usize; - let n_steps = x[14].round().clamp(3.0, 5.0) as usize; - // Phase Fast fixes num_atoms=11. The old (round/50, clamp 51-301) silently - // rounded 11→0→51, making Phase Fast's small-network config dead code. - let num_atoms = (x[15].round() as usize).max(11); - - // Weight decay - let weight_decay = x[16].exp().clamp(1e-5, 1e-3); - - // Kelly risk parameters (kelly_min_trades fixed to 20) - let kelly_fractional = x[17].clamp(0.25, 0.75); - let kelly_max_fraction = x[18].clamp(0.1, 0.5); - - // Volatility window - let volatility_window = x[19].round().clamp(10.0, 30.0) as usize; - - // Soft update (C2: curiosity_weight removed from index 20, tau shifted from 21 to 20) - let tau = x[20].exp().clamp(0.005, 0.01); - - // GPU-dynamic hidden dims (C2: shifted from 22 to 21) - // Upper clamp matches large_gpu expansion (2048); VRAM estimator caps for smaller GPUs. - let hidden_dim_base = ((x[21].round() / 64.0).round() * 64.0).clamp(64.0, 2048.0) as usize; - - // noisy_epsilon_floor FIXED to 0.10 below (C2: was index 23, not in search space) - - // CQL regularization (C2: shifted from 24 to 22) - let cql_alpha = x[22].clamp(0.0, 1.0); - - // Training dynamics (3D) (C2: shifted from 25-27 to 23-25) - let lr_decay_type = x[23].round().clamp(0.0, 2.0); - let dsr_eta = x[24].exp().clamp(0.001, 0.05); - let minimum_profit_factor = x[25].clamp(1.1, 2.0); - let count_bonus_coefficient = x[26].clamp(0.0, 0.3); - - // === FIXED parameters (validated defaults, removed from search) === - let kelly_min_trades: usize = 20; - let ensemble_size = 5.0; - let beta_variance = 0.5; - let beta_disagreement = 0.5; - let beta_entropy = 0.2; - let td_error_clamp_max = 10.0; - let batch_diversity_cooldown = 50.0; - let sharpe_weight = x[27].clamp(0.0, 0.5); // C4 FIX: tunable Sharpe blending - // Upper clamp matches large_gpu expansion (512); VRAM estimator caps for smaller GPUs. - let branch_hidden_dim = ((x[28].round() / 64.0).round() * 64.0).clamp(64.0, 512.0) as usize; - // C6: Gradient accumulation — round to nearest power of 2 in {1,2,4,8} - let gradient_accumulation_steps = { - let raw = x[29].round().clamp(1.0, 8.0) as usize; - // Snap to nearest power of 2: 1,2,3→2, 4,5,6→4, 7,8→8 - match raw { - 0..=1 => 1, - 2..=3 => 2, - 4..=5 => 4, - _ => 8, - } - }; - let eval_softmax_temp = 1.0; // Fixed: backtest uses greedy argmax (Fix A) - let gae_lambda = 0.95; - let noisy_sigma_initial = 0.5; - let noisy_sigma_final = 0.3; - let norm_type = 1.0; // RMSNorm - let activation_type = 1.0; // LeakyReLU - // M1: Enable QR-DQN when num_atoms > 100 (high atom count signals distributional preference) - let num_quantiles: usize = 32; // IQN default quantiles - let qr_kappa = 1.0; - // C7: IQN dual-head lambda (idx 30) - let iqn_lambda = x[30].clamp(0.0, 2.0); - - // C8: Composite reward weights (indices 31-38) - let w_dsr = x[31].clamp(0.1, 2.0); - let w_pnl = x[32].clamp(0.0, 1.0); - let w_dd = x[33].clamp(0.0, 5.0); - let w_idle = x[34].clamp(0.0, 0.1); - let dd_threshold = x[35].clamp(0.01, 0.10); - let loss_aversion = x[36].clamp(1.0, 3.0); - let time_decay_rate = x[37].clamp(0.0001, 0.005); - let q_gap_threshold = x[38].clamp(0.0, 0.5); - let spectral_norm_sigma_max = if x.len() > 39 { x[39].clamp(1.0, 10.0) } else { 3.0 }; - let c51_warmup_epochs = x.get(40).copied().unwrap_or(5.0).round().clamp(0.0, 20.0) as usize; - - // WAVE 6 FIX #2: Batch size floor for high learning rates + // Batch size floor for high learning rates if learning_rate > 2e-4 && batch_size < 120 { tracing::info!( "Adjusting batch_size from {} to 120 (LR={:.2e} requires larger batches)", @@ -868,10 +599,11 @@ impl ParameterSpace for DQNParams { batch_size = 120; } + // === ALL other fields FIXED at defaults (removed from search space) === let params = Self { learning_rate, batch_size, - gamma: x[2].clamp(0.88, 0.99), + gamma, buffer_size, max_position_absolute, huber_delta, @@ -887,170 +619,120 @@ impl ParameterSpace for DQNParams { noisy_sigma_init, minimum_profit_factor, weight_decay, - kelly_fractional, - kelly_max_fraction, - kelly_min_trades, - volatility_window, - use_ensemble_uncertainty: false, - ensemble_size, - beta_variance, - beta_disagreement, - beta_entropy, - warmup_ratio: 0.0, // Fixed: batch training never increments total_steps - curiosity_weight: 0.0, // C2: fixed to 0.0 (conflicts with NoisyNet) + hidden_dim_base, + min_hold_bars, + c51_warmup_epochs, + iqn_lambda, tau, - td_error_clamp_max, - batch_diversity_cooldown, - lr_decay_type, - sharpe_weight, - gae_lambda, - noisy_sigma_initial, - noisy_sigma_final, - // Architecture — always on. Built, tested, use it. + // --- Fixed reward weights --- + w_dsr: 1.0, + w_pnl: 0.3, + w_dd: 1.0, + w_idle: 0.01, + dd_threshold: 0.01, + loss_aversion: 1.5, + time_decay_rate: 0.0005, + q_gap_threshold: 0.1, + // --- Fixed Kelly params --- + kelly_fractional: 0.5, + kelly_max_fraction: 0.12, + kelly_min_trades: 20, + volatility_window: 23, + // --- Fixed ensemble params (ENABLED by default) --- + use_ensemble_uncertainty: true, + ensemble_size: 5.0, + beta_variance: 0.5, + beta_disagreement: 0.5, + beta_entropy: 0.2, + // --- Fixed dead/secondary params --- + warmup_ratio: 0.0, + curiosity_weight: 0.0, + td_error_clamp_max: 10.0, + batch_diversity_cooldown: 50.0, + lr_decay_type: 0.0, + sharpe_weight: 0.1, + gae_lambda: 0.95, + // --- Fixed network arch flags --- + noisy_sigma_initial: 0.5, + noisy_sigma_final: 0.3, use_spectral_norm: true, use_attention: true, use_residual: true, - norm_type, // 1.0 = RMSNorm - activation_type, // 1.0 = LeakyReLU - num_quantiles, - qr_kappa, - iqn_lambda, // C7: IQN dual-head loss weight - spectral_norm_sigma_max, // dim 39: Lipschitz constraint [1.0, 10.0] - hidden_dim_base, - noisy_epsilon_floor: 0.10, // Minimum 10% random exploration to prevent action collapse - count_bonus_coefficient, // C3 FIX: re-enabled UCB exploration bonus (complementary to NoisyNet) - cql_alpha, - eval_softmax_temp, - dsr_eta, - branch_hidden_dim, - gradient_accumulation_steps, - // C8: Composite reward weights (indices 31-38) - w_dsr, - w_pnl, - w_dd, - w_idle, - dd_threshold, - loss_aversion, - time_decay_rate, - q_gap_threshold, - c51_warmup_epochs, // Index 40: C51 warmup epochs [0, 20] - dt_pretrain_epochs: 0, // Fixed: DT pre-training not in hyperopt search space - // Task 11: New hyperopt-tunable parameters (indices 41-44) - her_ratio, - curiosity_weight_tunable, - use_cvar_action_selection, - cvar_alpha, - min_hold_bars, // Index 45: min_hold_bars [2, 20] + norm_type: 1.0, + activation_type: 1.0, + noisy_epsilon_floor: 0.1, + count_bonus_coefficient: 0.2, + cql_alpha: 0.5, + eval_softmax_temp: 1.0, + dsr_eta: 0.004, + branch_hidden_dim: 64, + gradient_accumulation_steps: 1, + dt_pretrain_epochs: 0, + spectral_norm_sigma_max: 5.0, + num_quantiles: 32, + qr_kappa: 1.0, + her_ratio: 0.3, + curiosity_weight_tunable: 0.1, + use_cvar_action_selection: 0.0, + cvar_alpha: 0.05, }; Ok(params) } fn to_continuous(&self) -> Vec { - // 46D search space (Task 11: +4 dims at indices 41-44, Task 4: +1 dim at index 45) + // 22D search space (consolidated from 46D) vec![ - self.learning_rate.ln(), // 0 - self.batch_size as f64, // 1 - self.gamma, // 2 - (self.buffer_size as f64).ln(), // 3 - self.max_position_absolute, // 4 - self.huber_delta.ln(), // 5 - self.entropy_coefficient, // 6 - self.transaction_cost_multiplier, // 7 - self.per_alpha, // 8 - self.per_beta_start, // 9 - self.v_max, // 10: v_range (symmetric: v_min=-v_range) - 1.0, // 11: use_branching (always on — fixed at 1.0) - self.noisy_sigma_init.ln(), // 12 - self.dueling_hidden_dim as f64, // 13 - self.n_steps as f64, // 14 - self.num_atoms as f64, // 15 - self.weight_decay.ln(), // 16 - self.kelly_fractional, // 17 - self.kelly_max_fraction, // 18 - self.volatility_window as f64, // 19 - self.tau.ln(), // 20 - self.hidden_dim_base as f64, // 21 - self.cql_alpha, // 22 - self.lr_decay_type, // 23 - self.dsr_eta.ln(), // 24 - self.minimum_profit_factor, // 25 - self.count_bonus_coefficient, // 26 - self.sharpe_weight, // 27: C4 FIX - self.branch_hidden_dim as f64, // 28: C5 branch_hidden_dim - self.gradient_accumulation_steps as f64, // 29: C6 gradient accumulation - self.iqn_lambda, // 30: C7 IQN dual-head loss weight - // C8: Composite reward weights (indices 31-38) - self.w_dsr, // 31 - self.w_pnl, // 32 - self.w_dd, // 33 - self.w_idle, // 34 - self.dd_threshold, // 35 - self.loss_aversion, // 36 - self.time_decay_rate, // 37 - self.q_gap_threshold, // 38 - self.spectral_norm_sigma_max, // 39 - self.c51_warmup_epochs as f64, // 40 - // Task 11: New hyperopt-tunable parameters - self.her_ratio, // 41 - self.curiosity_weight_tunable, // 42 - self.use_cvar_action_selection, // 43 - self.cvar_alpha, // 44 - self.min_hold_bars as f64, // 45 + self.learning_rate.ln(), // 0: learning_rate (log) + self.batch_size as f64, // 1: batch_size + self.gamma, // 2: gamma + (self.buffer_size as f64).ln(), // 3: buffer_size (log) + self.max_position_absolute, // 4: max_position_absolute + self.huber_delta.ln(), // 5: huber_delta (log) + self.entropy_coefficient, // 6: entropy_coefficient + self.transaction_cost_multiplier, // 7: transaction_cost_multiplier + self.per_alpha, // 8: per_alpha + self.per_beta_start, // 9: per_beta_start + self.dueling_hidden_dim as f64, // 10: dueling_hidden_dim + self.n_steps as f64, // 11: n_steps + self.tau.ln(), // 12: tau (log) + self.num_atoms as f64, // 13: num_atoms + self.v_max, // 14: v_max (v_min = -v_max) + self.noisy_sigma_init, // 15: noisy_sigma_init + self.minimum_profit_factor, // 16: minimum_profit_factor + self.weight_decay.ln(), // 17: weight_decay (log) + self.hidden_dim_base as f64, // 18: hidden_dim_base + self.min_hold_bars as f64, // 19: min_hold_bars + self.c51_warmup_epochs as f64, // 20: c51_warmup_epochs + self.iqn_lambda, // 21: iqn_lambda ] } fn param_names() -> Vec<&'static str> { - // 46 tuned parameters (Task 11: +4 new dims at indices 41-44, Task 4: +1 dim at index 45) + // 22 tuned parameters (consolidated from 46D) vec![ - "learning_rate", // 0 - "batch_size", // 1 - "gamma", // 2 - "buffer_size", // 3 - "max_position_absolute", // 4 - "huber_delta", // 5 - "entropy_coefficient", // 6 + "learning_rate", // 0 + "batch_size", // 1 + "gamma", // 2 + "buffer_size", // 3 + "max_position_absolute", // 4 + "huber_delta", // 5 + "entropy_coefficient", // 6 "transaction_cost_multiplier", // 7 - "per_alpha", // 8 - "per_beta_start", // 9 - "v_range", // 10: symmetric support ±v_range - "use_branching_fixed", // 11: always true (slot kept for index compat) - "noisy_sigma_init", // 12 - "dueling_hidden_dim", // 13 - "n_steps", // 14 - "num_atoms", // 15 - "weight_decay", // 16 - "kelly_fractional", // 17 - "kelly_max_fraction", // 18 - "volatility_window", // 19 - "tau", // 20 - "hidden_dim_base", // 21 - "cql_alpha", // 22 - "lr_decay_type", // 23 - "dsr_eta", // 24 - "minimum_profit_factor", // 25 - "count_bonus_coefficient", // 26 - "sharpe_weight", // 27: C4 FIX - "branch_hidden_dim", // 28: C5 branching head capacity - "gradient_accumulation_steps", // 29: C6 effective batch multiplier - "iqn_lambda", // 30: C7 IQN dual-head loss weight - // C8: Composite reward weights - "w_dsr", // 31 - "w_pnl", // 32 - "w_dd", // 33 - "w_idle", // 34 - "dd_threshold", // 35 - "loss_aversion", // 36 - "time_decay_rate", // 37 - "q_gap_threshold", // 38 - "spectral_norm_sigma_max", // 39 - "c51_warmup_epochs", // 40 - // Task 11: New hyperopt-tunable parameters - "her_ratio", // 41 - "curiosity_weight", // 42 - "use_cvar_action_selection", // 43 - "cvar_alpha", // 44 - "min_hold_bars", // 45 + "per_alpha", // 8 + "per_beta_start", // 9 + "dueling_hidden_dim", // 10 + "n_steps", // 11 + "tau", // 12 + "num_atoms", // 13 + "v_max", // 14 + "noisy_sigma_init", // 15 + "minimum_profit_factor", // 16 + "weight_decay", // 17 + "hidden_dim_base", // 18 + "min_hold_bars", // 19 + "c51_warmup_epochs", // 20 + "iqn_lambda", // 21 ] } @@ -1058,13 +740,7 @@ impl ParameterSpace for DQNParams { let mut bounds = Self::continuous_bounds(); let vram_mb = budget.gpu_memory_mb as f64; - // Derive upper bounds from VRAM via HardwareBudget methods (same estimators - // used by AutoBatchSizer and the DQN constructor). No hardcoded tiers. - // Skip fixed bounds (lo==hi from Phase Fast). - - // batch_size: constrain upper bound to VRAM capacity, but never exceed - // the TOML-configured upper bound (b.1). The TOML is the source of truth; - // VRAM sizing only REDUCES the upper bound, never expands it. + // batch_size (idx 1): constrain upper bound to VRAM capacity if let Some(b) = bounds.get_mut(1) { if (b.0 - b.1).abs() > 1e-6 { let toml_max = b.1; @@ -1074,34 +750,30 @@ impl ParameterSpace for DQNParams { } } - // buffer_size (ln-scale): 15% of VRAM / ~120 bytes per replay entry - // Cap to VRAM capacity but never exceed TOML upper bound + // buffer_size (idx 3, ln-scale): 15% of VRAM / ~120 bytes per replay entry let max_buffer = (vram_mb * 0.15 * 1024.0 * 1024.0 / 120.0).clamp(5000.0, 1_000_000.0); if let Some(b) = bounds.get_mut(3) { if (b.0 - b.1).abs() > 1e-6 { b.1 = b.1.min(max_buffer.ln()); } } - // hidden_dim_base + related dims: delegate to budget.max_hidden_dim_base_full() + // hidden_dim_base (idx 18) + dueling_hidden_dim (idx 10): cap to VRAM let worst_batch = bounds.get(1).map_or(512, |b| b.1 as usize); - let worst_atoms = bounds.get(15).map_or(51, |b| b.1 as usize); + let worst_atoms = bounds.get(13).map_or(101, |b| b.1 as usize); let worst_buffer = (vram_mb * 0.15 * 1024.0 * 1024.0 / 120.0) as usize; let max_base = budget.max_hidden_dim_base_full( 1, worst_batch, 56, 5, worst_atoms, true, true, worst_buffer, ); - // Cap hidden, dueling, branch proportionally from max_base let max_hidden = max_base as f64; - let max_dueling = (max_hidden * 0.5).clamp(128.0, 768.0); - let max_branch = (max_hidden * 0.25).clamp(64.0, 512.0); - // num_atoms: proportional to VRAM (5% budget / per-atom cost) + let max_dueling = (max_hidden * 0.5).clamp(64.0, 256.0); + // num_atoms (idx 13): proportional to VRAM let max_atoms = (vram_mb * 0.05 * 1024.0 * 1024.0 - / (worst_batch as f64 * 11.0 * 4.0 * 2.0)).clamp(11.0, 301.0); - // gradient_accum: proportional to compute budget (no extra VRAM) - let max_accum = (vram_mb / 10240.0 * 8.0).clamp(1.0, 8.0); + / (worst_batch as f64 * 11.0 * 4.0 * 2.0)).clamp(21.0, 101.0); - // Apply VRAM-derived caps (skip fixed Phase Fast bounds) + // Apply VRAM-derived caps (new 22D indices) for &(idx, cap) in &[ - (15, max_atoms), (21, max_hidden), (13, max_dueling), - (28, max_branch), (29, max_accum), + (13, max_atoms), // num_atoms + (18, max_hidden), // hidden_dim_base + (10, max_dueling), // dueling_hidden_dim ] { if let Some(b) = bounds.get_mut(idx) { if (b.0 - b.1).abs() > 1e-6 { b.1 = b.1.min(cap); } @@ -2187,6 +1859,7 @@ impl DQNTrainer { .map(|m| m.max_drawdown as f64) .fold(0.0_f64, f64::max); let mean_sortino = metrics.iter().map(|m| m.sortino as f64).sum::() / n; + let mean_calmar = metrics.iter().map(|m| m.calmar as f64).sum::() / n; let mean_wr = metrics.iter().map(|m| m.win_rate as f64).sum::() / n; let total_trades = metrics.iter().map(|m| m.total_trades as f64).sum::(); // Extended metrics from the GPU kernel (indices 6-9) @@ -2209,14 +1882,10 @@ impl DQNTrainer { Ok(Some(BacktestMetrics { sharpe_ratio: mean_sharpe, - total_return_pct: mean_pnl * 100.0, - max_drawdown_pct: worst_dd * 100.0, + total_return_pct: (mean_pnl * 100.0).clamp(-100_000.0, 100_000.0), + max_drawdown_pct: (worst_dd * 100.0).clamp(0.0, 100.0), sortino_ratio: mean_sortino, - calmar_ratio: if worst_dd > 1e-8 { - mean_pnl / worst_dd - } else { - 0.0 - }, + calmar_ratio: mean_calmar, win_rate: mean_wr, total_trades: total_trades as usize, var_95: mean_var_95, @@ -2315,12 +1984,6 @@ fn write_trial_result_dqn( /// - Stability (10%): Prevents Q-value collapse and gradient explosion /// - Completion (5%): Encourages full training runs /// - Hard constraints (35% implicit): Rejects catastrophically bad trials -fn normalize_reward(reward: f64) -> f64 { - // Target range: [-1.0, 1.0] corresponds to [-10.0, 10.0] raw reward - // Clamp to prevent outliers from dominating - (reward / 10.0).clamp(-1.0, 1.0) -} - /// Calculate Shannon entropy of action distribution /// /// Returns entropy in range [0, log2(3)] where: @@ -2597,19 +2260,6 @@ fn calculate_stability_penalty(gradient_norm: f64, q_value_std: f64) -> f64 { /// - Wave 3-A2: Multi-objective function design /// - Wave 4-A5: Completion penalty implementation /// - Smooth linear penalty replaces cliff (1000/500) for PSO gradient signal -fn calculate_completion_penalty( - epochs_completed: u32, - min_epochs: u32, - _early_stop_triggered: bool, -) -> f64 { - if min_epochs == 0 { - return 0.0; - } - // Linear scale: 50.0 * (1 - completed/min_epochs), clamped to [0, 50] - let completion_ratio = (epochs_completed as f64 / min_epochs as f64).clamp(0.0, 1.0); - 50.0 * (1.0 - completion_ratio) -} - impl HyperparameterOptimizable for DQNTrainer { type Params = DQNParams; type Metrics = DQNMetrics; @@ -2650,7 +2300,7 @@ impl HyperparameterOptimizable for DQNTrainer { // Fix 1: Clamp buffer size to max (4GB GPU constraint) let clamped_buffer_size = params.buffer_size.min(self.buffer_size_max); - info!("Training DQN with parameters (46D search space):"); + info!("Training DQN with parameters (22D search space):"); info!(" Learning rate: {:.6}", params.learning_rate); info!(" Batch size: {}", params.batch_size); info!(" Gamma: {:.3}", params.gamma); @@ -2937,7 +2587,7 @@ impl HyperparameterOptimizable for DQNTrainer { gradient_collapse_multiplier: 100.0, // Adaptive threshold (LR × 100) gradient_collapse_patience: 5, // 5 consecutive epochs before early stop - // LR scheduling — tunable via hyperopt search space (idx 23 in 39D) + // LR scheduling — fixed to constant (lr_decay_type=0.0, removed from 22D search space) // 0=constant, 1=linear decay, 2=cosine annealing // BUG FIX: total_steps = self.epochs (scheduler steps once per epoch, // NOT per training step). Previous calc used epochs × steps_per_epoch @@ -3224,8 +2874,6 @@ impl HyperparameterOptimizable for DQNTrainer { }, }; - tracing::debug!("Training completed successfully, extracting metrics"); - tracing::debug!("Checking validation data - internal_trainer.get_val_data().len() = {}", internal_trainer.get_val_data().len()); // Extract metrics from TrainingMetrics struct // Note: TrainingMetrics.loss is a single f64, not a Vec @@ -3453,7 +3101,7 @@ impl HyperparameterOptimizable for DQNTrainer { } Err(e) => { return Err(MLError::ModelError(format!( - "GPU backtest FAILED (no CPU fallback): {e}" + "GPU backtest FAILED: {e}" ))); } } @@ -3466,7 +3114,7 @@ impl HyperparameterOptimizable for DQNTrainer { Some(metrics) } else { return Err(MLError::TrainingError( - "GPU backtest returned no metrics — GPU evaluation is mandatory".to_owned() + "GPU backtest returned no metrics — evaluation is mandatory".to_owned() )); } } @@ -3514,11 +3162,11 @@ impl HyperparameterOptimizable for DQNTrainer { // VERBOSE: If backtest metrics are available, log comprehensive evaluation metrics if let Some(ref backtest) = metrics.backtest_metrics { info!( - "EVAL_METRICS: trial={}, sharpe={:.4}, win_rate={:.2}%, max_dd={:.2}%, \ + "EVAL_METRICS: trial={}, sharpe={:.4}, win_rate={:.1}%, max_dd={:.2}%, \ total_return={:.2}%, trades={}, sortino={:.4}, calmar={:.4}, omega={:.4}", current_trial, backtest.sharpe_ratio, - backtest.win_rate, + backtest.win_rate * 100.0, backtest.max_drawdown_pct, backtest.total_return_pct, backtest.total_trades, @@ -3626,7 +3274,7 @@ impl HyperparameterOptimizable for DQNTrainer { // VERBOSE: Comprehensive trial summary let objective_value = trial_result.objective; let best_sharpe = metrics.backtest_metrics.as_ref().map(|b| b.sharpe_ratio).unwrap_or(0.0); - let best_win_rate = metrics.backtest_metrics.as_ref().map(|b| b.win_rate).unwrap_or(0.0); + let best_win_rate = metrics.backtest_metrics.as_ref().map(|b| b.win_rate * 100.0).unwrap_or(0.0); let best_max_dd = metrics.backtest_metrics.as_ref().map(|b| b.max_drawdown_pct).unwrap_or(0.0); let best_epoch = metrics.epochs_completed; @@ -3774,16 +3422,20 @@ impl HyperparameterOptimizable for DQNTrainer { // CVaR = -0.005 → penalty ≈ 2.8 (elevated) // CVaR = -0.010 → penalty ≈ 9.8 (max, dangerous) let cvar_threshold = 0.05 / (common::thresholds::time::BARS_PER_DAY).sqrt(); - let cvar_penalty = ((-backtest.cvar_95 - cvar_threshold).max(0.0) * 1400.0).min(10.0); + // Softer ramp: cap at 3.0 instead of 10.0 to prevent CVaR + // from dominating the objective for under-trained models. + // Linear: 0 at threshold, 3.0 at threshold + 0.03 (~3% per-bar tail loss). + let cvar_penalty = ((-backtest.cvar_95 - cvar_threshold).max(0.0) * 100.0).min(3.0); // Component 2: HFT activity score (25% weight) let hft_activity = calculate_hft_activity_score_wave10(buy_pct, sell_pct, hold_pct); - // Mild diversity bonus: 5/5 unique → 0.0, 1/5 → 0.8 penalty. + // Mild diversity bonus: 9/9 unique → 0.0, 1/9 → 0.8 penalty. // Soft signal for TPE, never dominates real backtest metrics. - let mid_diversity_penalty = if backtest.unique_actions < 5 { - let ratio = backtest.unique_actions as f64 / 5.0; - 0.8 * (1.0 - ratio) // 0.64 for 1 action, 0.32 for 3, 0.0 for 5 + let num_exposure_actions = 9.0_f64; // 9 exposure levels (S100..L100) + let mid_diversity_penalty = if (backtest.unique_actions as f64) < num_exposure_actions { + let ratio = backtest.unique_actions as f64 / num_exposure_actions; + 0.8 * (1.0 - ratio) } else { 0.0 }; @@ -3794,9 +3446,11 @@ impl HyperparameterOptimizable for DQNTrainer { let objective = base_objective + trade_penalty + mid_diversity_penalty; - info!( - "OBJECTIVE: {:.4} = base {:.4} + trade_penalty {:.4} + diversity_penalty {:.4} | trades={} unique={}/5 composite={:.4}", - objective, base_objective, trade_penalty, mid_diversity_penalty, backtest.total_trades, backtest.unique_actions, composite_score + debug!( + "OBJECTIVE: {:.4} = base {:.4} + trade_penalty {:.4} + diversity {:.4} | Sharpe={:.2} Sortino={:.2} Calmar={:.2} Omega={:.2} CVaR={:.4} trades={} unique={}/9", + objective, base_objective, trade_penalty, mid_diversity_penalty, + backtest.sharpe_ratio, backtest.sortino_ratio, backtest.calmar_ratio, backtest.omega_ratio, + backtest.cvar_95, backtest.total_trades, backtest.unique_actions ); debug!( @@ -3826,22 +3480,12 @@ impl HyperparameterOptimizable for DQNTrainer { objective } } else { - // No backtest metrics — training-only fallback. - // In production hyperopt on GPU, upstream trial runner MUST produce - // backtest_metrics. This path is for unit tests and edge cases. - { - debug!("extract_objective without backtest_metrics — using training-only fallback"); - let reward_component = normalize_reward(metrics.avg_episode_reward); - let reward_weighted = 0.40 * reward_component; - let min_epochs = 5; - let completion_penalty = calculate_completion_penalty( - metrics.epochs_completed as u32, - min_epochs, - metrics.epochs_completed < (min_epochs as usize), - ); - let hft_activity = calculate_hft_activity_score_wave10(buy_pct, sell_pct, hold_pct); - reward_weighted + hft_activity + 0.15 * stability_penalty_raw + completion_penalty - } + // Backtest metrics MUST be present — GPU evaluator is mandatory. + // If we reach here, a bug prevented the evaluator from running. + // Return maximum penalty so the optimizer avoids this configuration + // and the error is visible in the trial log. + eprintln!("BUG: extract_objective called without backtest_metrics — returning 1e6 penalty"); + 1e6 }; objective_total @@ -3881,207 +3525,122 @@ mod tests { #[test] fn test_dqn_params_roundtrip() { - // Roundtrip test for the 46D search space (Task 11: +4 new dims, Task 4: +1 dim) - // Only the tuned parameters roundtrip; fixed params get defaults from from_continuous + // Roundtrip test for the 22D search space (consolidated from 46D) + // Only the 22 tuned parameters roundtrip; fixed params get defaults from from_continuous let params = DQNParams { learning_rate: 3.37e-05, batch_size: 92, - gamma: 0.92, - buffer_size: 97_273, + gamma: 0.97, + buffer_size: 50_000, max_position_absolute: 2.5, - huber_delta: 24.77, - entropy_coefficient: 0.03, // C2: narrowed range [0.005, 0.05] + huber_delta: 5.0, + entropy_coefficient: 0.1, transaction_cost_multiplier: 1.0, per_alpha: 0.6, per_beta_start: 0.4, dueling_hidden_dim: 128, n_steps: 3, - tau: 0.007, + tau: 0.005, num_atoms: 51, - // v_range = 10/(1-0.92)*1.2 = 150, clamped [20,300] → ±150 - v_min: -(10.0_f64 / (1.0 - 0.92) * 1.2).clamp(20.0, 300.0), - v_max: (10.0_f64 / (1.0 - 0.92) * 1.2).clamp(20.0, 300.0), + v_min: -150.0, + v_max: 150.0, noisy_sigma_init: 0.5, minimum_profit_factor: 1.5, - warmup_ratio: 0.0, - lr_decay_type: 2.0, - dsr_eta: 0.02, weight_decay: 1e-4, - kelly_fractional: 0.5, - kelly_max_fraction: 0.25, - kelly_min_trades: 20, - volatility_window: 20, - use_ensemble_uncertainty: false, - ensemble_size: 5.0, - beta_variance: 0.5, beta_disagreement: 0.5, beta_entropy: 0.2, - curiosity_weight: 0.0, // C2: fixed to 0.0 - td_error_clamp_max: 10.0, - batch_diversity_cooldown: 50.0, - sharpe_weight: 0.0, - gae_lambda: 0.95, - noisy_sigma_initial: 0.5, - noisy_sigma_final: 0.3, - use_spectral_norm: false, - use_attention: false, - use_residual: false, - norm_type: 1.0, - activation_type: 1.0, - num_quantiles: 64, - qr_kappa: 1.0, - iqn_lambda: 0.5, - spectral_norm_sigma_max: 3.0, - hidden_dim_base: 512, - noisy_epsilon_floor: 0.10, // Fixed: 10% floor to prevent action collapse - count_bonus_coefficient: 0.0, // C2: fixed to 0.0 - cql_alpha: 0.05, - eval_softmax_temp: 0.8, - branch_hidden_dim: 128, - gradient_accumulation_steps: 4, - // C8: Composite reward weights - w_dsr: 1.5, - w_pnl: 0.5, - w_dd: 2.0, - w_idle: 0.05, - dd_threshold: 0.04, - loss_aversion: 2.0, - time_decay_rate: 0.001, - q_gap_threshold: 0.2, - c51_warmup_epochs: 5, - dt_pretrain_epochs: 0, - // Task 11: New hyperopt-tunable parameters - her_ratio: 0.3, - curiosity_weight_tunable: 0.15, - use_cvar_action_selection: 1.0, - cvar_alpha: 0.08, - // Task 4: min_hold_bars in search space + hidden_dim_base: 256, min_hold_bars: 8, + c51_warmup_epochs: 5, + iqn_lambda: 0.5, + ..DQNParams::default() }; let continuous = params.to_continuous(); - assert_eq!(continuous.len(), 46, "to_continuous must return 46D vector"); + assert_eq!(continuous.len(), 22, "to_continuous must return 22D vector"); let recovered = DQNParams::from_continuous(&continuous).unwrap(); - // Tuned parameters must roundtrip exactly + // The 22 tuned parameters must roundtrip exactly assert!((recovered.learning_rate - params.learning_rate).abs() < 1e-6); assert_eq!(recovered.batch_size, params.batch_size); assert!((recovered.gamma - params.gamma).abs() < 1e-6); assert_eq!(recovered.buffer_size, params.buffer_size); assert!((recovered.max_position_absolute - params.max_position_absolute).abs() < 1e-6); + assert!((recovered.huber_delta - params.huber_delta).abs() < 0.1); + assert!((recovered.entropy_coefficient - params.entropy_coefficient).abs() < 1e-6); + assert!((recovered.transaction_cost_multiplier - params.transaction_cost_multiplier).abs() < 1e-6); assert!((recovered.per_alpha - params.per_alpha).abs() < 1e-6); assert!((recovered.per_beta_start - params.per_beta_start).abs() < 1e-6); - assert!((recovered.weight_decay - params.weight_decay).abs() < 1e-8); - assert!((recovered.kelly_fractional - params.kelly_fractional).abs() < 1e-6); - assert!((recovered.kelly_max_fraction - params.kelly_max_fraction).abs() < 1e-6); - assert_eq!(recovered.volatility_window, params.volatility_window); + assert_eq!(recovered.dueling_hidden_dim, params.dueling_hidden_dim); + assert_eq!(recovered.n_steps, params.n_steps); assert!((recovered.tau - params.tau).abs() < 1e-6); - assert_eq!(recovered.hidden_dim_base, params.hidden_dim_base); - assert!((recovered.cql_alpha - params.cql_alpha).abs() < 1e-6); - assert!((recovered.lr_decay_type - params.lr_decay_type).abs() < 1e-6); - assert!((recovered.dsr_eta - params.dsr_eta).abs() < 1e-4); + assert_eq!(recovered.num_atoms, params.num_atoms); + assert!((recovered.v_max - params.v_max).abs() < 1e-6); + assert!((recovered.v_min + recovered.v_max).abs() < 1e-6, "v_min must equal -v_max"); + assert!((recovered.noisy_sigma_init - params.noisy_sigma_init).abs() < 1e-6); assert!((recovered.minimum_profit_factor - params.minimum_profit_factor).abs() < 1e-6); + assert!((recovered.weight_decay - params.weight_decay).abs() < 1e-8); + assert_eq!(recovered.hidden_dim_base, params.hidden_dim_base); + assert_eq!(recovered.min_hold_bars, params.min_hold_bars); + assert_eq!(recovered.c51_warmup_epochs, params.c51_warmup_epochs); + assert!((recovered.iqn_lambda - params.iqn_lambda).abs() < 1e-6); - // C2: curiosity_weight, count_bonus_coefficient fixed; noisy_epsilon_floor fixed to 0.10 - assert!((recovered.curiosity_weight - 0.0).abs() < 1e-6); - assert!((recovered.noisy_epsilon_floor - 0.10).abs() < 1e-6); - assert!((recovered.count_bonus_coefficient - 0.0).abs() < 1e-6); - - // Other fixed parameters - assert!((recovered.warmup_ratio - 0.0).abs() < 1e-6); - assert!((recovered.eval_softmax_temp - 1.0).abs() < 1e-3); + // Fixed parameters get their defaults from from_continuous + assert!((recovered.w_dsr - 1.0).abs() < 1e-6); + assert!((recovered.w_pnl - 0.3).abs() < 1e-6); + assert!((recovered.w_dd - 1.0).abs() < 1e-6); + assert!((recovered.w_idle - 0.01).abs() < 1e-6); + assert!((recovered.dd_threshold - 0.01).abs() < 1e-6); + assert!((recovered.loss_aversion - 1.5).abs() < 1e-6); + assert!((recovered.time_decay_rate - 0.0005).abs() < 1e-6); + assert!((recovered.kelly_fractional - 0.5).abs() < 1e-6); assert_eq!(recovered.kelly_min_trades, 20); + assert_eq!(recovered.volatility_window, 23); + assert!(recovered.use_ensemble_uncertainty); assert!((recovered.ensemble_size - 5.0).abs() < 1e-6); + assert!((recovered.warmup_ratio - 0.0).abs() < 1e-6); + assert!((recovered.curiosity_weight - 0.0).abs() < 1e-6); assert!((recovered.td_error_clamp_max - 10.0).abs() < 1e-6); assert!((recovered.batch_diversity_cooldown - 50.0).abs() < 1e-6); + assert!((recovered.lr_decay_type - 0.0).abs() < 1e-6); + assert!((recovered.sharpe_weight - 0.1).abs() < 1e-6); + assert!((recovered.gae_lambda - 0.95).abs() < 1e-6); assert!((recovered.norm_type - 1.0).abs() < 1e-6); assert!((recovered.activation_type - 1.0).abs() < 1e-6); assert!((recovered.qr_kappa - 1.0).abs() < 1e-6); - - // C6: use_branching roundtrips via index 11 - // C6: gradient_accumulation_steps roundtrips via index 29 - assert_eq!(recovered.gradient_accumulation_steps, params.gradient_accumulation_steps); - - // C8: Composite reward weights roundtrip (indices 31-38) - assert!((recovered.w_dsr - params.w_dsr).abs() < 1e-6); - assert!((recovered.w_pnl - params.w_pnl).abs() < 1e-6); - assert!((recovered.w_dd - params.w_dd).abs() < 1e-6); - assert!((recovered.w_idle - params.w_idle).abs() < 1e-6); - assert!((recovered.dd_threshold - params.dd_threshold).abs() < 1e-6); - assert!((recovered.loss_aversion - params.loss_aversion).abs() < 1e-6); - assert!((recovered.time_decay_rate - params.time_decay_rate).abs() < 1e-6); - - // Task 11: New hyperopt parameters roundtrip (indices 41-44) - assert!((recovered.her_ratio - params.her_ratio).abs() < 1e-6); - assert!((recovered.curiosity_weight_tunable - params.curiosity_weight_tunable).abs() < 1e-6); - assert!((recovered.use_cvar_action_selection - params.use_cvar_action_selection).abs() < 1e-6); - assert!((recovered.cvar_alpha - params.cvar_alpha).abs() < 1e-6); - // Task 4: min_hold_bars roundtrip (index 45) - assert_eq!(recovered.min_hold_bars, params.min_hold_bars); + assert_eq!(recovered.gradient_accumulation_steps, 1); + assert!((recovered.eval_softmax_temp - 1.0).abs() < 1e-3); } #[test] fn test_dqn_params_bounds() { let bounds = DQNParams::continuous_bounds(); - assert_eq!(bounds.len(), 46); // 46D: 39 base + spectral_norm_sigma_max + c51_warmup_epochs + her_ratio + curiosity_weight + use_cvar + cvar_alpha + min_hold_bars + assert_eq!(bounds.len(), 22); // 22D consolidated search space // Check log-scale bounds are reasonable - assert!(bounds[0].0 < bounds[0].1); // learning_rate - assert!(bounds[3].0 < bounds[3].1); // buffer_size - assert!(bounds[5].0 < bounds[5].1); // huber_delta - assert!(bounds[12].0 < bounds[12].1); // noisy_sigma_init - assert!(bounds[16].0 < bounds[16].1); // weight_decay (log scale) - assert!(bounds[20].0 < bounds[20].1); // tau (log scale, was 21) + assert!(bounds[0].0 < bounds[0].1); // learning_rate (log) + assert!(bounds[3].0 < bounds[3].1); // buffer_size (log) + assert!(bounds[5].0 < bounds[5].1); // huber_delta (log) + assert!(bounds[12].0 < bounds[12].1); // tau (log) + assert!(bounds[17].0 < bounds[17].1); // weight_decay (log) // Check learning rate range includes production default 1e-4 let lr_min = bounds[0].0.exp(); let lr_max = bounds[0].1.exp(); - assert!(lr_min <= 1e-4 && 1e-4 <= lr_max, "Learning rate range [{}, {}] must include production default 1e-4", lr_min, lr_max); - assert!((lr_min - 1e-5).abs() < 1e-7, "Learning rate lower bound should be 1e-5, got {}", lr_min); - assert!((lr_max - 3e-4).abs() < 1e-6, "Learning rate upper bound should be 3e-4, got {}", lr_max); + assert!(lr_min <= 1e-4 && 1e-4 <= lr_max, + "Learning rate range [{:.2e}, {:.2e}] must include production default 1e-4", lr_min, lr_max); + assert!(lr_min > 0.0 && lr_min < 1e-3, "Learning rate lower bound should be small, got {:.2e}", lr_min); + assert!(lr_max >= 1e-4 && lr_max <= 1e-2, "Learning rate upper bound should be reasonable, got {:.2e}", lr_max); - // Check linear bounds - assert_eq!(bounds[1], (64.0, 512.0)); // batch_size (capped) - assert_eq!(bounds[2], (0.90, 0.99)); // gamma (widened) - assert_eq!(bounds[4], (1.0, 4.0)); // max_position_absolute - assert_eq!(bounds[6], (0.05, 0.5)); // entropy_coefficient (C1: 10x wider) - assert_eq!(bounds[7], (0.5, 2.0)); // transaction_cost_multiplier - assert_eq!(bounds[8], (0.4, 0.8)); // per_alpha - assert_eq!(bounds[9], (0.2, 0.6)); // per_beta_start - // Default phase is Fast — architecture dims are fixed to [phase_fast] TOML values. - assert_eq!(bounds[10], (1.0, 1.0)); // v_range (FIXED — now computed from gamma) - assert_eq!(bounds[11], (1.0, 1.0)); // use_branching (FIXED: always true) - assert_eq!(bounds[13], (128.0, 128.0)); // dueling_hidden_dim (FIXED in Fast phase) - assert_eq!(bounds[14], (3.0, 5.0)); // n_steps (searchable in all phases) - assert_eq!(bounds[15], (101.0, 101.0)); // num_atoms (FIXED in Fast phase to 101) - - // Kelly risk parameters - assert_eq!(bounds[17], (0.25, 0.75)); // kelly_fractional - assert_eq!(bounds[18], (0.1, 0.5)); // kelly_max_fraction - assert_eq!(bounds[19], (10.0, 30.0)); // volatility_window - - // C2: curiosity_weight fixed to 0.0, noisy_epsilon_floor fixed to 0.10, not in search space - assert_eq!(bounds[21], (128.0, 128.0)); // hidden_dim_base (FIXED in Fast phase) - assert_eq!(bounds[22], (0.0, 1.0)); // cql_alpha (full offline-RL range) - - // C5: branch_hidden_dim - assert_eq!(bounds[28], (64.0, 64.0)); // branch_hidden_dim (FIXED in Fast phase) - // C6: gradient_accumulation_steps (fixed at 1 for default GPUs) - assert_eq!(bounds[29], (1.0, 1.0)); // gradient_accumulation_steps - - // C8: Composite reward weights (FIXED in Fast phase to defaults) - assert_eq!(bounds[31], (1.0, 1.0)); // w_dsr (FIXED in Fast phase) - assert_eq!(bounds[32], (0.3, 0.3)); // w_pnl (FIXED in Fast phase) - assert_eq!(bounds[33], (1.0, 1.0)); // w_dd (FIXED in Fast phase) - assert_eq!(bounds[34], (0.01, 0.01)); // w_idle (FIXED in Fast phase) - assert_eq!(bounds[35], (0.01, 0.01)); // dd_threshold (FIXED in Fast phase) - assert_eq!(bounds[36], (1.5, 1.5)); // loss_aversion (FIXED in Fast phase) - assert_eq!(bounds[37], (0.0005, 0.0005)); // time_decay_rate (FIXED in Fast phase) - assert_eq!(bounds[38], (0.1, 0.1)); // q_gap_threshold (FIXED in Fast phase) + // Check all bounds have valid ranges (lower < upper) + // Exact values may be overridden by HyperoptProfile YAML + for (i, (lo, hi)) in bounds.iter().enumerate() { + assert!(lo < hi, "Bound {i}: lower ({lo}) must be < upper ({hi})"); + } } #[test] fn test_param_names() { let names = DQNParams::param_names(); - assert_eq!(names.len(), 46); // 46D: 39 base + spectral_norm_sigma_max + c51_warmup_epochs + her_ratio + curiosity_weight + use_cvar + cvar_alpha + min_hold_bars + assert_eq!(names.len(), 22); // 22D consolidated search space assert_eq!(names[0], "learning_rate"); assert_eq!(names[1], "batch_size"); assert_eq!(names[2], "gamma"); @@ -4092,154 +3651,84 @@ mod tests { assert_eq!(names[7], "transaction_cost_multiplier"); assert_eq!(names[8], "per_alpha"); assert_eq!(names[9], "per_beta_start"); - assert_eq!(names[10], "v_range"); - assert_eq!(names[11], "use_branching_fixed"); - assert_eq!(names[12], "noisy_sigma_init"); - assert_eq!(names[13], "dueling_hidden_dim"); - assert_eq!(names[14], "n_steps"); - assert_eq!(names[15], "num_atoms"); - assert_eq!(names[16], "weight_decay"); - assert_eq!(names[17], "kelly_fractional"); - assert_eq!(names[18], "kelly_max_fraction"); - assert_eq!(names[19], "volatility_window"); - assert_eq!(names[20], "tau"); - assert_eq!(names[21], "hidden_dim_base"); - assert_eq!(names[22], "cql_alpha"); - assert_eq!(names[23], "lr_decay_type"); - assert_eq!(names[24], "dsr_eta"); - assert_eq!(names[25], "minimum_profit_factor"); - assert_eq!(names[26], "count_bonus_coefficient"); - assert_eq!(names[27], "sharpe_weight"); // C4 - assert_eq!(names[28], "branch_hidden_dim"); // C5 - assert_eq!(names[29], "gradient_accumulation_steps"); // C6 - assert_eq!(names[39], "spectral_norm_sigma_max"); // 39 - assert_eq!(names[40], "c51_warmup_epochs"); // 40 - // Task 11: New hyperopt-tunable parameters - assert_eq!(names[41], "her_ratio"); // 41 - assert_eq!(names[42], "curiosity_weight"); // 42 - assert_eq!(names[43], "use_cvar_action_selection"); // 43 - assert_eq!(names[44], "cvar_alpha"); // 44 - assert_eq!(names[45], "min_hold_bars"); // 45 + assert_eq!(names[10], "dueling_hidden_dim"); + assert_eq!(names[11], "n_steps"); + assert_eq!(names[12], "tau"); + assert_eq!(names[13], "num_atoms"); + assert_eq!(names[14], "v_max"); + assert_eq!(names[15], "noisy_sigma_init"); + assert_eq!(names[16], "minimum_profit_factor"); + assert_eq!(names[17], "weight_decay"); + assert_eq!(names[18], "hidden_dim_base"); + assert_eq!(names[19], "min_hold_bars"); + assert_eq!(names[20], "c51_warmup_epochs"); + assert_eq!(names[21], "iqn_lambda"); } #[test] fn test_per_params_always_enabled() { // Test that PER is always enabled with tunable alpha/beta parameters - // 46D search space (Task 11: +4 new dims, Task 4: +1 min_hold_bars dim) + // 22D search space (consolidated from 46D) let continuous = vec![ - 3e-5_f64.ln(), 92.0, 0.92, 97_273_f64.ln(), 4.0, 24.77_f64.ln(), 0.03, 1.0, - 0.6, 0.4, // 8-9: per_alpha, per_beta_start - 25.0, 25.0, 0.5_f64.ln(), // 10-12: v_range, v_range_mirror, noisy_sigma_init - 256.0, 3.0, 101.0, // 13-15: dueling_hidden_dim, n_steps, num_atoms - 1e-4_f64.ln(), // 16: weight_decay - 0.5, 0.25, // 17-18: kelly_fractional, kelly_max_fraction - 20.0, // 19: volatility_window - 0.007_f64.ln(), // 20: tau - 512.0, // 21: hidden_dim_base - 0.1, // 22: cql_alpha - 2.0, // 23: lr_decay_type - 0.02_f64.ln(), // 24: dsr_eta (log scale) - 1.5, // 25: minimum_profit_factor - 0.1, // 26: count_bonus_coefficient - 0.3, // 27: sharpe_weight (C4) - 128.0, // 28: branch_hidden_dim (C5) - 1.0, // 29: gradient_accumulation_steps (C6) - 0.25, // 30: iqn_lambda (C7) - // C8: Composite reward weights (indices 31-38) - 1.0, // 31: w_dsr - 0.3, // 32: w_pnl - 1.0, // 33: w_dd - 0.01, // 34: w_idle - 0.02, // 35: dd_threshold - 1.5, // 36: loss_aversion - 0.0005, // 37: time_decay_rate - 0.1, // 38: q_gap_threshold + 3e-5_f64.ln(), // 0: learning_rate (log) + 92.0, // 1: batch_size + 0.97, // 2: gamma + 50_000_f64.ln(), // 3: buffer_size (log) + 2.0, // 4: max_position_absolute + 5.0_f64.ln(), // 5: huber_delta (log) + 0.1, // 6: entropy_coefficient + 1.0, // 7: transaction_cost_multiplier + 0.6, // 8: per_alpha + 0.4, // 9: per_beta_start + 128.0, // 10: dueling_hidden_dim + 3.0, // 11: n_steps + 0.005_f64.ln(), // 12: tau (log) + 51.0, // 13: num_atoms + 150.0, // 14: v_max + 0.5, // 15: noisy_sigma_init + 1.5, // 16: minimum_profit_factor + 1e-4_f64.ln(), // 17: weight_decay (log) + 256.0, // 18: hidden_dim_base + 5.0, // 19: min_hold_bars + 5.0, // 20: c51_warmup_epochs + 0.25, // 21: iqn_lambda ]; let params = DQNParams::from_continuous(&continuous).unwrap(); assert!((params.per_alpha - 0.6).abs() < 1e-6); assert!((params.per_beta_start - 0.4).abs() < 1e-6); assert!((params.warmup_ratio - 0.0).abs() < 1e-6); - assert!((params.curiosity_weight - 0.0).abs() < 1e-6); // C2: fixed - assert!((params.noisy_epsilon_floor - 0.10).abs() < 1e-6); // Fixed: 10% floor - // Verify symmetric support — v_range is dynamic from gamma + reward_scale. - // With gamma=0.92, reward_scale=10: v_range = 10 / (1-0.92) * 1.2 = 150 + assert!((params.curiosity_weight - 0.0).abs() < 1e-6); + assert!((params.noisy_epsilon_floor - 0.10).abs() < 1e-6); + // Verify symmetric support from v_max assert!(params.v_min < 0.0, "v_min should be negative"); assert!(params.v_max > 0.0, "v_max should be positive"); assert!((params.v_min + params.v_max).abs() < 1e-6, "v_min/v_max should be symmetric"); - // Test PER parameter bounds (min values) -- 39D + // Test PER parameter bounds (min values) -- 22D let continuous_min = vec![ - 2e-5_f64.ln(), 64.0, 0.88, 50_000_f64.ln(), 1.0, 10.0_f64.ln(), 0.05, 0.5, - 0.4, 0.2, // per_alpha min, per_beta_start min - 10.0, 10.0, 0.1_f64.ln(), // v_range min, v_range_mirror, noisy_sigma_init - 128.0, 3.0, 51.0, // dueling_hidden_dim, n_steps (min=3), num_atoms - 1e-4_f64.ln(), // weight_decay min - 0.25, 0.1, // kelly_fractional, kelly_max_fraction - 10.0, // volatility_window - 0.005_f64.ln(), // 20: tau min - 256.0, // 21: hidden_dim_base min - 0.0, // 22: cql_alpha min - 0.0, // 23: lr_decay_type (constant) - 0.001_f64.ln(), // 24: dsr_eta min (log scale) - 1.1, // 25: minimum_profit_factor min - 0.0, // 26: count_bonus_coefficient min - 0.0, // 27: sharpe_weight min (C4) - 64.0, // 28: branch_hidden_dim min (C5) - 1.0, // 29: gradient_accumulation_steps min (C6) - 0.0, // 30: iqn_lambda min (C7) - // C8: Composite reward weights min - 0.1, // 31: w_dsr min - 0.0, // 32: w_pnl min - 0.0, // 33: w_dd min - 0.0, // 34: w_idle min - 0.01, // 35: dd_threshold min - 1.0, // 36: loss_aversion min - 0.0001, // 37: time_decay_rate min - 0.0, // 38: q_gap_threshold min + 1e-5_f64.ln(), 64.0, 0.95, 10_000_f64.ln(), 1.0, 0.1_f64.ln(), 0.05, 0.5, + 0.4, 0.2, + 64.0, 1.0, 0.0001_f64.ln(), 21.0, 5.0, 0.1, + 1.1, 1e-5_f64.ln(), 64.0, 2.0, 3.0, 0.0, ]; let params_min = DQNParams::from_continuous(&continuous_min).unwrap(); assert!((params_min.per_alpha - 0.4).abs() < 1e-6); assert!((params_min.per_beta_start - 0.2).abs() < 1e-6); - // v_range now dynamic from gamma — just verify symmetric and positive assert!(params_min.v_min < 0.0, "v_min should be negative at min"); assert!(params_min.v_max > 0.0, "v_max should be positive at min"); assert!((params_min.v_min + params_min.v_max).abs() < 1e-6, "symmetric at min"); - // Test PER parameter bounds (max values) -- 39D + // Test PER parameter bounds (max values) -- 22D let continuous_max = vec![ - 8e-5_f64.ln(), 512.0, 0.99, 100_000_f64.ln(), 4.0, 40.0_f64.ln(), 0.5, 2.0, - 0.8, 0.6, // per_alpha max, per_beta_start max - 50.0, 50.0, 1.0_f64.ln(), // v_range max, v_range_mirror, noisy_sigma_init - 512.0, 5.0, 201.0, // dueling_hidden_dim, n_steps, num_atoms - 1e-2_f64.ln(), // weight_decay max - 0.75, 0.5, // kelly_fractional, kelly_max_fraction - 30.0, // volatility_window - 0.01_f64.ln(), // 20: tau max - 1024.0, // 21: hidden_dim_base max - 1.0, // 22: cql_alpha max - 2.0, // 23: lr_decay_type (cosine) - 0.05_f64.ln(), // 24: dsr_eta max (log scale) - 2.0, // 25: minimum_profit_factor max - 0.3, // 26: count_bonus_coefficient max - 0.5, // 27: sharpe_weight max (C4) - 256.0, // 28: branch_hidden_dim max (C5) - 1.0, // 29: gradient_accumulation_steps max (C6) - 2.0, // 30: iqn_lambda max (C7) - // C8: Composite reward weights max - 2.0, // 31: w_dsr max - 1.0, // 32: w_pnl max - 5.0, // 33: w_dd max - 0.1, // 34: w_idle max - 0.10, // 35: dd_threshold max - 3.0, // 36: loss_aversion max - 0.005, // 37: time_decay_rate max - 0.5, // 38: q_gap_threshold max + 1e-3_f64.ln(), 512.0, 0.999, 100_000_f64.ln(), 4.0, 20.0_f64.ln(), 0.5, 2.0, + 0.8, 0.6, + 256.0, 10.0, 0.01_f64.ln(), 101.0, 300.0, 1.0, + 2.0, 1e-2_f64.ln(), 256.0, 20.0, 15.0, 1.0, ]; let params_max = DQNParams::from_continuous(&continuous_max).unwrap(); assert!((params_max.per_alpha - 0.8).abs() < 1e-6); assert!((params_max.per_beta_start - 0.6).abs() < 1e-6); - // v_range dynamic from gamma — verify symmetric assert!(params_max.v_min < 0.0, "v_min should be negative at max"); assert!(params_max.v_max > 0.0, "v_max should be positive at max"); assert!((params_max.v_min + params_max.v_max).abs() < 1e-6, "symmetric at max"); @@ -4373,56 +3862,115 @@ mod tests { #[test] fn test_objective_function_maximizes_reward() { - // Test that objective function uses normalization + clamping + HFT activity scoring - // Note: Expected values updated for new multi-objective formula with HFT activity component + // Test that extract_objective correctly uses backtest_metrics for the multi-objective formula. + // The optimizer MINIMIZES the objective, so: + // - Good Sharpe → negative objective (good) + // - Bad Sharpe → positive objective (bad) + // + // Formula (when backtest_metrics is Some and total_trades >= 10): + // composite = 0.40*tanh(sharpe/5) + 0.25*tanh(sortino/8) + 0.20*tanh(calmar/5) + 0.15*tanh(omega/2) + // cvar_thresh = 0.05 / sqrt(BARS_PER_DAY) (~0.00253 for 1-min bars) + // cvar_penalty = ((-cvar_95 - cvar_thresh).max(0) * 100).min(3.0) + // hft_activity = calculate_hft_activity_score_wave10(buy%, sell%, hold%) + // diversity = 0.8 * (1 - unique_actions/9) if unique_actions < 9, else 0 + // base_obj = -0.60*composite + cvar_penalty - 0.05*hft_activity + 0.15*stability + // objective = base_obj + trade_penalty + diversity + // + // gradient_norm=2.0 (<50) and q_value_std=1.5 (<15) → stability_penalty_raw=0.0 - // Scenario 1: Positive reward with good action distribution (30/30/40) + // Scenario 1: Good Sharpe (2.0) with balanced action distribution (30/30/40). + // Approx: composite≈0.489, cvar_penalty=0 (cvar=-0.002 is above threshold), + // hft_activity=1.0, diversity≈0.178 (7/9 unique), trade_penalty=0 + // → objective ≈ -0.60*0.489 + 0 - 0.05*1.0 + 0.178 ≈ -0.165 (negative = good) let metrics_positive = DQNMetrics { train_loss: 0.5, val_loss: 0.4, avg_q_value: 10.0, final_epsilon: 0.01, - epochs_completed: 100, // >= min_epochs (5) → no completion penalty - avg_episode_reward: 100.0, // Clamped: (100/10).clamp(-1, 1) = 1.0 + epochs_completed: 100, + avg_episode_reward: 100.0, buy_action_pct: 0.3, sell_action_pct: 0.3, hold_action_pct: 0.4, gradient_norm: 2.0, q_value_std: 1.5, - backtest_metrics: None, + backtest_metrics: Some(BacktestMetrics { + sharpe_ratio: 2.0, + sortino_ratio: 3.0, + calmar_ratio: 5.0, + omega_ratio: 1.5, + win_rate: 55.0, + max_drawdown_pct: 15.0, + total_return_pct: 20.0, + total_trades: 500, + var_95: -0.0015, + cvar_95: -0.002, + beta: 0.0, + alpha: 0.0, + information_ratio: 0.0, + unique_actions: 7, + buy_action_pct: 0.3, + sell_action_pct: 0.3, + hold_action_pct: 0.4, + }), }; let objective_positive = DQNTrainer::extract_objective(&metrics_positive); - // Expected: 0.40 * 1.0 (reward) + ~1.0 (HFT activity, 60% BUY+SELL) + 0.0 (stability) + 0.0 (completion) ≈ 1.4 + // Good Sharpe → composite is positive → -0.60*composite pulls objective negative assert!( - objective_positive > 1.0, - "Objective should be > 1.0 with positive reward + HFT activity: {}", + objective_positive < 0.0, + "Scenario 1 (Sharpe=2.0): good composite should give negative objective (optimizer minimizes), got: {}", objective_positive ); - // Scenario 2: Negative reward (clamped to -1.0) + // Scenario 2: Bad Sharpe (-1.0) with balanced action distribution. + // Approx: composite≈-0.114, cvar_penalty≈0.747 (cvar=-0.01 below threshold), + // hft_activity=1.0, diversity≈0.356 (5/9 unique), trade_penalty=0 + // → objective ≈ -0.60*(-0.114) + 0.747 - 0.05*1.0 + 0.356 ≈ 1.12 (positive = bad) let metrics_negative = DQNMetrics { train_loss: 0.5, val_loss: 0.4, avg_q_value: 10.0, final_epsilon: 0.01, epochs_completed: 100, - avg_episode_reward: -50.0, // Clamped: (-50/10).clamp(-1, 1) = -1.0 + avg_episode_reward: -50.0, buy_action_pct: 0.3, sell_action_pct: 0.3, hold_action_pct: 0.4, gradient_norm: 2.0, q_value_std: 1.5, - backtest_metrics: None, + backtest_metrics: Some(BacktestMetrics { + sharpe_ratio: -1.0, + sortino_ratio: -0.5, + calmar_ratio: -2.0, + omega_ratio: 0.8, + win_rate: 35.0, + max_drawdown_pct: 40.0, + total_return_pct: -15.0, + total_trades: 300, + var_95: -0.008, + cvar_95: -0.01, + beta: 0.0, + alpha: 0.0, + information_ratio: 0.0, + unique_actions: 5, + buy_action_pct: 0.3, + sell_action_pct: 0.3, + hold_action_pct: 0.4, + }), }; let objective_negative = DQNTrainer::extract_objective(&metrics_negative); - // Expected: 0.40 * -1.0 + ~1.0 (HFT activity) + 0.0 + 0.0 ≈ 0.6 + // Bad Sharpe → composite is negative → -0.60*composite pushes objective positive + // plus CVaR penalty for deep tail loss assert!( objective_negative > 0.0, - "Negative reward should still have positive objective due to HFT activity: {}", + "Scenario 2 (Sharpe=-1.0): bad composite + CVaR penalty should give positive objective, got: {}", objective_negative ); - // Scenario 3: Zero reward with good action distribution + // Scenario 3: Near-zero Sharpe (0.1) — barely active model. + // Approx: composite≈0.104, cvar_penalty≈0.047 (cvar=-0.003 just below threshold), + // hft_activity=1.0, diversity≈0.267 (6/9 unique), trade_penalty=0 + // → objective ≈ -0.60*0.104 + 0.047 - 0.05*1.0 + 0.267 ≈ 0.20 let metrics_zero = DQNMetrics { train_loss: 0.5, val_loss: 0.4, @@ -4435,57 +3983,93 @@ mod tests { hold_action_pct: 0.4, gradient_norm: 2.0, q_value_std: 1.5, - backtest_metrics: None, + backtest_metrics: Some(BacktestMetrics { + sharpe_ratio: 0.1, + sortino_ratio: 0.2, + calmar_ratio: 0.5, + omega_ratio: 1.0, + win_rate: 50.0, + max_drawdown_pct: 10.0, + total_return_pct: 1.0, + total_trades: 200, + var_95: -0.002, + cvar_95: -0.003, + beta: 0.0, + alpha: 0.0, + information_ratio: 0.0, + unique_actions: 6, + buy_action_pct: 0.3, + sell_action_pct: 0.3, + hold_action_pct: 0.4, + }), }; let objective_zero = DQNTrainer::extract_objective(&metrics_zero); - // Expected: 0.0 (reward) + ~1.0 (HFT activity) + 0.0 + 0.0 ≈ 1.0 + // Near-zero Sharpe → objective near zero (small positive due to diversity penalty) assert!( - objective_zero > 0.5, - "Zero reward should still have positive objective due to HFT activity: {}", + objective_zero > -0.5 && objective_zero < 1.0, + "Scenario 3 (Sharpe=0.1): near-zero Sharpe should give near-zero objective, got: {}", objective_zero ); - // Scenario 4: Verify higher reward gives higher objective (with same action distribution) + // Scenario 4: Excellent Sharpe (3.0) — best model configuration. + // Using same CVaR as Scenario 1 (-0.002, above threshold) so CVaR penalty=0 for both. + // Approx: composite≈0.629, cvar_penalty=0 (cvar=-0.002 above threshold), + // hft_activity=1.0, diversity≈0.089 (8/9 unique), trade_penalty=0 + // → base_obj ≈ -0.60*0.629 + 0 - 0.05*1.0 ≈ -0.427 + // → objective ≈ -0.427 + 0 + 0.089 ≈ -0.338 (more negative than Scenario 1's ~-0.166) let high_reward = DQNMetrics { train_loss: 0.5, val_loss: 0.4, avg_q_value: 10.0, final_epsilon: 0.01, epochs_completed: 100, - avg_episode_reward: 200.0, // Clamped: (200/10).clamp(-1, 1) = 1.0 + avg_episode_reward: 200.0, buy_action_pct: 0.3, sell_action_pct: 0.3, hold_action_pct: 0.4, gradient_norm: 2.0, q_value_std: 1.5, - backtest_metrics: None, - }; - let low_reward = DQNMetrics { - train_loss: 0.5, - val_loss: 0.4, - avg_q_value: 10.0, - final_epsilon: 0.01, - epochs_completed: 100, - avg_episode_reward: 5.0, // (5/10).clamp(-1, 1) = 0.5 - buy_action_pct: 0.3, - sell_action_pct: 0.3, - hold_action_pct: 0.4, - gradient_norm: 2.0, - q_value_std: 1.5, - backtest_metrics: None, + backtest_metrics: Some(BacktestMetrics { + sharpe_ratio: 3.0, + sortino_ratio: 4.0, + calmar_ratio: 8.0, + omega_ratio: 2.0, + win_rate: 60.0, + max_drawdown_pct: 20.0, + total_return_pct: 30.0, + total_trades: 400, + var_95: -0.0015, + cvar_95: -0.002, // Same as Scenario 1 so CVaR penalty=0 for fair comparison + beta: 0.0, + alpha: 0.0, + information_ratio: 0.0, + unique_actions: 8, + buy_action_pct: 0.3, + sell_action_pct: 0.3, + hold_action_pct: 0.4, + }), }; let obj_high = DQNTrainer::extract_objective(&high_reward); - let obj_low = DQNTrainer::extract_objective(&low_reward); - - // Higher reward should give higher objective (HFT activity component is same) + // Very good Sharpe (3.0) → composite is large positive → objective should be negative assert!( - obj_high > obj_low, - "Higher reward (1.0) should give higher objective than lower reward (0.5): {} > {}", + obj_high < 0.0, + "Scenario 4 (Sharpe=3.0): excellent composite should give negative objective, got: {}", + obj_high + ); + // With same CVaR, Sharpe=3.0 should give lower (better) objective than Sharpe=2.0 + assert!( + obj_high < objective_positive, + "Sharpe=3.0 should give lower (better) objective than Sharpe=2.0 (same CVaR): {} < {}", obj_high, - obj_low + objective_positive ); - // Scenario 5: Test HFT activity penalty for low BUY/SELL actions + // Scenario 5: Low trading activity (5% BUY, 5% SELL, 90% HOLD) with low diversity. + // unique_actions=2 → diversity_penalty = 0.8*(7/9) ≈ 0.622 + // hft_activity = -5*(15-5)/15 = -3.333 (both < 15% threshold) → penalty + // Approx: composite≈0.222, cvar_penalty≈0.147, hft≈-3.333, diversity≈0.622, trade_penalty=0 + // → base_obj ≈ -0.60*0.222 + 0.147 - 0.05*(-3.333) ≈ -0.133+0.147+0.167 = 0.181 + // → objective ≈ 0.181 + 0 + 0.622 ≈ 0.803 let low_activity = DQNMetrics { train_loss: 0.5, val_loss: 0.4, @@ -4493,22 +4077,46 @@ mod tests { final_epsilon: 0.01, epochs_completed: 100, avg_episode_reward: 100.0, - buy_action_pct: 0.05, // Only 5% BUY (below 15% threshold) - sell_action_pct: 0.05, // Only 5% SELL (below 15% threshold) - hold_action_pct: 0.90, // 90% HOLD (passive behavior) + buy_action_pct: 0.05, + sell_action_pct: 0.05, + hold_action_pct: 0.90, gradient_norm: 2.0, q_value_std: 1.5, - backtest_metrics: None, + backtest_metrics: Some(BacktestMetrics { + sharpe_ratio: 0.5, + sortino_ratio: 1.0, + calmar_ratio: 2.0, + omega_ratio: 1.1, + win_rate: 45.0, + max_drawdown_pct: 25.0, + total_return_pct: 5.0, + total_trades: 50, + var_95: -0.003, + cvar_95: -0.004, + beta: 0.0, + alpha: 0.0, + information_ratio: 0.0, + unique_actions: 2, + buy_action_pct: 0.05, + sell_action_pct: 0.05, + hold_action_pct: 0.90, + }), }; let obj_low_activity = DQNTrainer::extract_objective(&low_activity); - // Low activity should have lower objective than balanced distribution + // Low activity + low diversity should produce a worse (higher) objective than Scenario 1 assert!( - obj_low_activity < objective_positive, - "Low BUY/SELL activity should be penalized: {} < {}", + obj_low_activity > objective_positive, + "Scenario 5 (low activity + unique_actions=2) should have higher objective than Scenario 1: {} > {}", obj_low_activity, objective_positive ); + // Diversity penalty is visible: unique_actions=2 → penalty ≈ 0.622 + assert!( + obj_low_activity > 0.0, + "Scenario 5: low activity + diversity penalty should give positive objective, got: {}", + obj_low_activity + ); } #[test] @@ -4522,7 +4130,7 @@ mod tests { fn test_qr_dqn_roundtrip_continuous() { let params = DQNParams::default(); let continuous = params.to_continuous(); - assert_eq!(continuous.len(), 46, "Should have 46 continuous dimensions"); + assert_eq!(continuous.len(), 22, "Should have 22 continuous dimensions"); let roundtrip = DQNParams::from_continuous(&continuous).unwrap(); // Default num_quantiles=32 (IQN default, used alongside C51) assert_eq!(roundtrip.num_quantiles, 32); @@ -4532,48 +4140,38 @@ mod tests { #[test] fn test_qr_dqn_activation_threshold() { let bounds = DQNParams::continuous_bounds(); - let dim = DQNParams::continuous_bounds().len(); + let dim = bounds.len(); let mut params = vec![0.0_f64; dim]; - // Fill with valid defaults + // Fill with valid defaults (22D) params[0] = (1e-4_f64).ln(); // learning_rate params[1] = 128.0; // batch_size - params[2] = 0.92; // gamma - params[3] = (100_000.0_f64).ln(); // buffer_size + params[2] = 0.97; // gamma + params[3] = (50_000.0_f64).ln(); // buffer_size params[4] = 3.0; // max_position_absolute - params[5] = (25.0_f64).ln(); // huber_delta - params[6] = 0.02; // entropy_coefficient (C2: narrowed) + params[5] = (5.0_f64).ln(); // huber_delta + params[6] = 0.1; // entropy_coefficient params[7] = 1.0; // transaction_cost_multiplier params[8] = 0.6; // per_alpha params[9] = 0.4; // per_beta_start // Fill remaining with midpoint of bounds - for i in 10..41 { - if i < bounds.len() { - params[i] = (bounds[i].0 + bounds[i].1) / 2.0; - } + for i in 10..dim { + params[i] = (bounds[i].0 + bounds[i].1) / 2.0; } - // num_atoms index = 15: rounding is (x.round() / 50.0).round() * 50.0, clamped to [51, 201] - // 51.0 -> (51/50).round()*50 = 1*50 = 50, clamped to 51 - params[15] = 51.0; + // num_atoms index = 13 in 22D: direct rounding, min 21 + params[13] = 51.0; let low_atoms = DQNParams::from_continuous(¶ms).unwrap(); - assert_eq!(low_atoms.num_atoms, 51, "51.0 rounds to 51 (clamped min)"); + assert_eq!(low_atoms.num_atoms, 51, "51.0 rounds to 51"); assert_eq!(low_atoms.num_quantiles, 32, "IQN default quantiles"); - // 100.0 -> (100/50).round()*50 = 2*50 = 100 (at boundary, should still be C51) - params[15] = 100.0; + params[13] = 100.0; let at_boundary = DQNParams::from_continuous(¶ms).unwrap(); assert_eq!(at_boundary.num_atoms, 100); - // 151.0 -> 151 (below 200 threshold, still C51) - params[15] = 151.0; - let high_atoms = DQNParams::from_continuous(¶ms).unwrap(); - assert_eq!(high_atoms.num_atoms, 151); - - // 201.0 -> 201 (above 200, QR-DQN activates) - params[15] = 201.0; - let max_atoms = DQNParams::from_continuous(¶ms).unwrap(); - assert_eq!(max_atoms.num_atoms, 201); - assert_eq!(max_atoms.num_quantiles, 32, "IQN quantile count is fixed at 32"); + params[13] = 21.0; + let min_atoms = DQNParams::from_continuous(¶ms).unwrap(); + assert_eq!(min_atoms.num_atoms, 21, "min 21 atoms"); + assert_eq!(min_atoms.num_quantiles, 32, "IQN quantile count is fixed at 32"); } #[test] @@ -4595,21 +4193,21 @@ mod tests { // batch_size is clamped to [64, 1024] in from_continuous — VRAM guard. // Even if PSO suggests 2048, from_continuous caps at 1024. let bounds = DQNParams::continuous_bounds(); - let dim = DQNParams::continuous_bounds().len(); + let dim = bounds.len(); let mut params = vec![0.0_f64; dim]; params[1] = 2048.0; // batch_size (index 1) — over the clamp // Fill other required params with valid defaults params[0] = (1e-4_f64).ln(); // learning_rate - params[2] = 0.92; // gamma - params[3] = (100_000.0_f64).ln(); // buffer_size + params[2] = 0.97; // gamma + params[3] = (50_000.0_f64).ln(); // buffer_size params[4] = 3.0; // max_position_absolute - params[5] = (25.0_f64).ln(); // huber_delta - params[6] = 0.02; // entropy_coefficient - params[7] = 0.5; // transaction_cost_multiplier + params[5] = (5.0_f64).ln(); // huber_delta + params[6] = 0.1; // entropy_coefficient + params[7] = 1.0; // transaction_cost_multiplier params[8] = 0.6; // per_alpha params[9] = 0.4; // per_beta_start - // Remaining params (indices 10-38) -- use midpoint of bounds - for i in 10..39 { + // Remaining params (indices 10-21) -- use midpoint of bounds + for i in 10..dim { params[i] = (bounds[i].0 + bounds[i].1) / 2.0; } let result = DQNParams::from_continuous(¶ms).unwrap(); @@ -4640,24 +4238,6 @@ mod tests { assert!(penalty_low > penalty_zero); } - #[test] - fn test_completion_penalty_smooth() { - let p0 = calculate_completion_penalty(0, 10, true); - assert!((p0 - 50.0).abs() < 0.1, "0 epochs = 50.0: {}", p0); - - let p5 = calculate_completion_penalty(5, 10, false); - assert!((p5 - 25.0).abs() < 0.1, "5/10 = 25.0: {}", p5); - - let p10 = calculate_completion_penalty(10, 10, false); - assert!(p10.abs() < 0.01, "Full = 0: {}", p10); - - let p15 = calculate_completion_penalty(15, 10, false); - assert!(p15.abs() < 0.01, "Over min = 0: {}", p15); - - assert!(p0 > p5); - assert!(p5 > p10); - } - #[test] fn test_trade_insufficiency_penalty_zero_trades() { let penalty = calculate_trade_insufficiency_penalty(0); @@ -4719,8 +4299,9 @@ mod tests { // eval_softmax_temp removed from search space (backtest uses greedy argmax). // Verify from_continuous always sets it to fixed 1.0 regardless of input. let bounds = DQNParams::continuous_bounds(); - let mut vec = vec![0.0_f64; 39]; // C8: 39D - for i in 0..38 { + let dim = bounds.len(); + let mut vec = vec![0.0_f64; dim]; // 22D + for i in 0..dim { vec[i] = (bounds[i].0 + bounds[i].1) / 2.0; } let params = DQNParams::from_continuous(&vec).unwrap(); @@ -4729,11 +4310,11 @@ mod tests { "eval_softmax_temp should be fixed to 1.0, got {}", params.eval_softmax_temp ); - // C2: verify exploration stacking params are fixed + // Verify fixed exploration params assert!((params.curiosity_weight - 0.0).abs() < f64::EPSILON, "curiosity_weight should be 0.0"); assert!((params.noisy_epsilon_floor - 0.10).abs() < 1e-6, "noisy_epsilon_floor should be 0.10"); - // C3: count_bonus_coefficient is now tuned (index 26), so midpoint of (0.0, 0.3) = 0.15 - assert!(params.count_bonus_coefficient >= 0.0 && params.count_bonus_coefficient <= 0.3, "count_bonus_coefficient should be in [0.0, 0.3], got {}", params.count_bonus_coefficient); + // count_bonus_coefficient is now fixed at 0.2 + assert!((params.count_bonus_coefficient - 0.2).abs() < 1e-6, "count_bonus_coefficient should be fixed to 0.2, got {}", params.count_bonus_coefficient); } #[test] diff --git a/crates/ml/src/hyperopt/campaign.rs b/crates/ml/src/hyperopt/campaign.rs index 788a160d0..a47f7390f 100644 --- a/crates/ml/src/hyperopt/campaign.rs +++ b/crates/ml/src/hyperopt/campaign.rs @@ -310,6 +310,12 @@ mod tests { #[test] #[ignore] // GPU + real data, ~30 min on RTX 3050 fn test_local_hyperopt() { + // Initialize tracing so info!/warn!/error! output is visible in test runs + let _ = tracing_subscriber::fmt() + .with_env_filter("info") + .with_test_writer() + .try_init(); + let data_dir = std::env::var("FOXHUNT_TEST_DATA") .unwrap_or_else(|_| "test_data/futures-baseline".to_owned()); let data_path = std::path::Path::new(&data_dir); diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs index 65321d558..34b3c635f 100644 --- a/crates/ml/src/trainers/dqn/fused_training.rs +++ b/crates/ml/src/trainers/dqn/fused_training.rs @@ -759,7 +759,7 @@ impl FusedTrainingCtx { // All auxiliary gradients (IQN, attention, ensemble, CQL) are now in grad_buf. // The single Adam update sees: C51 + iqn_lambda*IQN + attention + diversity + CQL. let fused_result = self.trainer.replay_adam_and_readback() - .map_err(|e| anyhow::anyhow!("graph_adam replay: {e}"))?; + .map_err(|e| { eprintln!("!!! ADAM REPLAY FAILED: {e}"); anyhow::anyhow!("graph_adam replay: {e}") })?; // ── Step 5f: Regime-adaptive PER scaling ────────────────────── // Kernel reads target ADX/CUSUM from first sample in states_buf. diff --git a/crates/ml/src/trainers/dqn/trainer/tests.rs b/crates/ml/src/trainers/dqn/trainer/tests.rs index ee1adebea..ddb459af9 100644 --- a/crates/ml/src/trainers/dqn/trainer/tests.rs +++ b/crates/ml/src/trainers/dqn/trainer/tests.rs @@ -551,24 +551,27 @@ fn test_c2_five_actions_produce_distinct_exposures() { assert_eq!(unique.len(), 5, "All 5 exposure levels must be distinct"); } -/// Verify hyperopt search space is 45D (Task 11: +4 new dims). +/// Verify hyperopt search space is 22D (consolidated from 46D). #[test] -fn test_c3_search_space_is_40d() { +fn test_c3_search_space_is_22d() { let bounds = crate::hyperopt::adapters::dqn::DQNParams::continuous_bounds(); - assert_eq!(bounds.len(), 46, "Search space must be 46D (45 base + min_hold_bars)"); + assert_eq!(bounds.len(), 22, "Search space must be 22D (consolidated from 46D)"); let names = crate::hyperopt::adapters::dqn::DQNParams::param_names(); - assert_eq!(names.len(), 46); - assert!(names.contains(&"count_bonus_coefficient"), "count_bonus_coefficient must be in search space (C3)"); - assert!(names.contains(&"sharpe_weight"), "sharpe_weight must be in search space (C4)"); - assert!(names.contains(&"branch_hidden_dim"), "branch_hidden_dim must be in search space (L2)"); + assert_eq!(names.len(), 22); assert!(names.contains(&"c51_warmup_epochs"), "c51_warmup_epochs must be in search space"); - // Task 11: curiosity_weight is now in the search space - assert!(names.contains(&"curiosity_weight"), "curiosity_weight must be in search space (Task 11)"); - assert!(names.contains(&"her_ratio"), "her_ratio must be in search space (Task 11)"); - assert!(names.contains(&"use_cvar_action_selection"), "use_cvar_action_selection must be in search space (Task 11)"); - assert!(names.contains(&"cvar_alpha"), "cvar_alpha must be in search space (Task 11)"); + assert!(names.contains(&"iqn_lambda"), "iqn_lambda must be in search space"); + assert!(names.contains(&"min_hold_bars"), "min_hold_bars must be in search space"); + assert!(names.contains(&"hidden_dim_base"), "hidden_dim_base must be in search space"); assert!(!names.contains(&"noisy_epsilon_floor"), "noisy_epsilon_floor must not be in search space"); + // These are now FIXED (not in search space): + assert!(!names.contains(&"count_bonus_coefficient"), "count_bonus_coefficient should be fixed, not in search space"); + assert!(!names.contains(&"sharpe_weight"), "sharpe_weight should be fixed, not in search space"); + assert!(!names.contains(&"branch_hidden_dim"), "branch_hidden_dim should be fixed, not in search space"); + assert!(!names.contains(&"curiosity_weight"), "curiosity_weight should be fixed, not in search space"); + assert!(!names.contains(&"her_ratio"), "her_ratio should be fixed, not in search space"); + assert!(!names.contains(&"use_cvar_action_selection"), "use_cvar_action_selection should be fixed, not in search space"); + assert!(!names.contains(&"cvar_alpha"), "cvar_alpha should be fixed, not in search space"); } /// Verify noisy_epsilon_floor is fixed to 0.10 (prevents action collapse). diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 71307ebc8..49459e13a 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -1070,7 +1070,7 @@ impl DQNTrainer { // cross-stream deadlock or data corruption. if let Some(ref stream) = self.cuda_stream { stream.synchronize() - .map_err(|e| anyhow::anyhow!("Stream sync before PER insert: {e}"))?; + .map_err(|e| { eprintln!("!!! SYNC FAILED: {e}"); anyhow::anyhow!("Stream sync before PER insert: {e}") })?; } if let Some(ref mut mon) = self.gpu_monitoring { @@ -1243,9 +1243,12 @@ impl DQNTrainer { let batch_data = explicit_batch.as_ref().ok_or_else(|| { anyhow::anyhow!("No batch data for fused training step") })?; - fused.run_full_step(batch_data, &mut *agent, &self.device) + let result = fused.run_full_step(batch_data, &mut *agent, &self.device) .map_err(|e| { eprintln!("!!! FUSED STEP ERROR: {:#}", e); e }) - .context("Fused CUDA training step failed")? + .context("Fused CUDA training step failed")?; + if train_step_count % 50 == 0 { + } + result } else { agent.train_step(explicit_batch) .map_err(|e| anyhow::anyhow!("Train step FAILED: {e}"))?