From 2fb30f098e9aee6630ca2ada4c19a06e582e635a Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Wed, 22 Apr 2026 00:55:25 +0200 Subject: [PATCH] =?UTF-8?q?fix(monitoring):=20cascade=209=E2=86=9212=20act?= =?UTF-8?q?ion=20bins=20for=204-direction=20action=20space?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The runtime DQN action space is 4×3×3×3 = 108 (dir×mag×order×urgency) per the kernel's b0_size=4 (Short/Hold/Long/Flat) layout, but the monitoring stack was still coded for the legacy 3×3=9-bin exposure layout. Flat actions (dir=3) had exp_idx = 9 >= num_actions=9 and were silently dropped by the monitoring_reduce kernel's bounds check. Cascade: - monitoring_kernel.cu: num_exposure_bins=12 (EXP_MAX constant), summary layout 27 floats with bins [5..17), order [17..20), urgency [20..23), N [23], trades [24], pad [25..27). - gpu_monitoring.rs: MonitoringSummary.action_counts [usize; 12]; summary_buf 27 f32; reduce() now takes b0_size and validates num_exposure_bins ≤ EXP_MAX. - trainers/dqn/monitoring.rs: action_counts/q_value_sums/q_value_counts all [_; 12], factored_action_counts [_; 108], EXPOSURE_NAMES with 4 dirs × 3 mags ("S_Small".."F_Full"); track_action now maps the 8-level ExposureLevel enum onto the kernel's dir*3+mag layout by going through direction()/magnitude() accessors (was indexing by raw `exposure as usize` which is an entirely different scheme). - trainer/training_loop.rs: total_action_counts [_; 12], total_factored_action_counts [_; 108]; GPU monitoring reduce call now passes b0 from agent.branch_sizes(); direction-entropy gains a 4th dir bin; per-action Q diagnostics names array extended to 12. - trainer/metrics.rs: create_final_metrics arrays widened; q_diagnostics per_action_avgs [_; 12]; combo_idx = d*3 + m (matching kernel layout); action_space_size 108/12 and buy/sell/hold uses the new exp_idx boundaries. - financials.rs: action_counts [_; 12]; buy/sell/hold classification updated — BUY = Long [6..9], SELL = Short [0..3], HOLD = Hold + Flat ([3..6] ∪ [9..12]). Tests updated for 12-bin layout. Also fixes several collateral bugs exposed by the audit: 1. magnitude_action_dist_final conflated Hold (dir=1, mag-forced-to-0) with "Quarter magnitude". New formula restricts the denominator to tradable directions (Short + Long) and computes Q/H/F fractions only over those, producing the real magnitude preference signal. 2. training_loop.rs:3012 flat_count indexed [3,4,5] (which is Hold in the new layout, but was already the wrong bucket — "Flat" was never at that index even under the legacy 9-bin reading, where [3,4,5] was the Flat bin but mag was forced to 1 not 0). New code correctly separates flat_count [9..12] from hold_count [3..6] and excludes both from the directional denominator and diversity cell threshold. 3. monitoring.rs had q_value_sums: [f64; 7] but q_value_counts: [usize; 9] — different sizes for what should be the same exposure axis. Both now [_; 12]. 4. log_action_distribution iterated over 7 indices when printing per- action Q-values but the array is logically 12-wide — now uses `0..12` with the 12-entry EXPOSURE_NAMES lookup. Tests: all 7 financials tests pass. Monitoring cubin loads + default summary sanity checks pass. Smoke test magnitude_distribution now emits correct 12-bin breakdown — GPU summary shows (example epoch 20): actions=[786, 157, 280, 377, 2, 7, 806, 156, 317, 303, 5, 4] which decodes as Short (S) = 786/157/280 (~38%), Hold (H) = 377/2/7 (~12%, mag-forced), Long (L) = 806/156/317 (~40%), Flat (F) = 303/5/4 (~10%, mag-forced). Previously the Flat bucket (303+5+4 = ~10% of all actions) was silently dropped — policy was picking Flat ~10% of the time and nobody could see it. No atomic ops added; all reductions remain via warp-scratch + lane-0 sequential merge. No feature flags, no stubs, no quickfixes. --- crates/ml/src/cuda_pipeline/gpu_monitoring.rs | 68 +++++++--- .../ml/src/cuda_pipeline/monitoring_kernel.cu | 50 +++++--- crates/ml/src/trainers/dqn/financials.rs | 65 ++++++---- crates/ml/src/trainers/dqn/monitoring.rs | 112 ++++++++++------ crates/ml/src/trainers/dqn/trainer/metrics.rs | 65 +++++----- .../src/trainers/dqn/trainer/training_loop.rs | 121 ++++++++++++------ 6 files changed, 307 insertions(+), 174 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/gpu_monitoring.rs b/crates/ml/src/cuda_pipeline/gpu_monitoring.rs index 7b0f603e8..2eb503618 100644 --- a/crates/ml/src/cuda_pipeline/gpu_monitoring.rs +++ b/crates/ml/src/cuda_pipeline/gpu_monitoring.rs @@ -11,11 +11,17 @@ use crate::MLError; /// Precompiled monitoring kernel cubin, embedded at compile time by build.rs. static MONITORING_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/monitoring_kernel.cubin")); -/// Compact monitoring summary from GPU reduction (48-byte GPU transfer, 24 × f32). +/// Compact monitoring summary from GPU reduction (108-byte GPU transfer, 27 × f32). /// /// Reward statistics (mean, std, min, max, sharpe) are computed **per trade** /// (non-zero rewards only), not per bar. With sparse trade-completion rewards /// (98% of bars have reward=0), per-bar mean is always ~0 and uninformative. +/// +/// Exposure bins cover all `b0 * b1` = 4 * 3 = 12 direction×magnitude combos used +/// by the runtime DQN. Layout `[dir*3 + mag]` with dir ∈ {Short=0, Hold=1, Long=2, +/// Flat=3} and mag ∈ {Small=0, Half=1, Full=2}. Hold (dir=1) and Flat (dir=3) +/// force mag=0 in the kernel, so bins 4, 5, 10, 11 are empty by construction — +/// they still occupy slots for diagnostic cleanliness. #[derive(Debug, Clone, Default)] pub struct MonitoringSummary { /// Mean reward per trade completion (non-zero rewards only). @@ -26,8 +32,8 @@ pub struct MonitoringSummary { pub max_reward: f32, /// Sharpe estimate: mean_reward / reward_std (per-trade). pub sharpe_estimate: f32, - /// Per dir×mag level counts (9 bins: dir(3) × mag(3), e.g. ShortSmall=0 .. LongFull=8). - pub action_counts: [usize; 9], + /// Per dir×mag level counts (12 bins: dir(4) × mag(3) laid out as `dir*3 + mag`). + pub action_counts: [usize; 12], /// Per-order-type counts (3: Market, LimitMaker, IoC). pub order_counts: [usize; 3], /// Per-urgency counts (3: Patient, Normal, Aggressive). @@ -43,7 +49,7 @@ pub struct MonitoringSummary { pub struct GpuMonitoringReducer { stream: Arc, kernel_func: CudaFunction, - summary_buf: CudaSlice, // [24] + summary_buf: CudaSlice, // [27] } impl GpuMonitoringReducer { @@ -53,7 +59,8 @@ impl GpuMonitoringReducer { .map_err(|e| MLError::ModelError(format!("monitoring module load: {e}")))?; let kernel_func = module.load_function("monitoring_reduce") .map_err(|e| MLError::ModelError(format!("monitoring_reduce load: {e}")))?; - let summary_buf = stream.alloc_zeros::(24) // 5 stats + 7 exp + 3 ord + 3 urg + 1 total + 5 pad + // Layout: 5 stats + 12 exp + 3 ord + 3 urg + total + trades + 2 pad = 27 + let summary_buf = stream.alloc_zeros::(27) .map_err(|e| MLError::ModelError(format!("monitoring summary alloc: {e}")))?; Ok(Self { stream: Arc::clone(stream), kernel_func, summary_buf }) @@ -61,11 +68,16 @@ impl GpuMonitoringReducer { /// Launch reduction over rewards/actions buffers already on GPU. /// Does NOT synchronize — caller must sync before reading result. + /// + /// `b0_size` is the direction branch count (4 for the 4-branch DQN). It is + /// multiplied by `b1_size` = 3 (magnitude) to derive the 12-bin exposure + /// dimensionality that the kernel writes into `summary[5..17]`. pub fn reduce( &mut self, rewards: &CudaSlice, actions: &CudaSlice, n: usize, + b0_size: usize, ) -> Result<(), MLError> { let _nvtx = NvtxRange::new("gpu_monitoring"); if n == 0 { @@ -74,14 +86,25 @@ impl GpuMonitoringReducer { if n > i32::MAX as usize { return Err(MLError::ModelError(format!("monitoring n={n} exceeds i32::MAX"))); } + // b1_size (magnitude branch) is always 3 in the 4-branch DQN. The + // kernel has a static EXP_MAX=12 that matches b0_size*3 for b0∈{3,4}. + const B1_SIZE: usize = 3; + const EXP_MAX: usize = 12; + let num_exposure_bins = b0_size.saturating_mul(B1_SIZE); + if num_exposure_bins == 0 || num_exposure_bins > EXP_MAX { + return Err(MLError::ModelError(format!( + "monitoring num_exposure_bins={num_exposure_bins} (b0={b0_size} × b1={B1_SIZE}) \ + out of range [1, {EXP_MAX}] — kernel EXP_MAX mismatch" + ))); + } let config = LaunchConfig { grid_dim: (1, 1, 1), block_dim: (256, 1, 1), shared_mem_bytes: 0, // kernel uses static __shared__ only — no dynamic shmem needed }; // Safety: rewards and actions are valid GPU slices with at least n elements. - // summary_buf is pre-allocated 12-float device buffer on the same stream. - let num_actions_i32 = 12_i32; // dir(4) × mag(3) = 12 combos + // summary_buf is pre-allocated 27-float device buffer on the same stream. + let num_exposure_bins_i32 = num_exposure_bins as i32; let order_actions_i32 = 3_i32; // DQN_ORDER_ACTIONS let urgency_actions_i32 = 3_i32; // DQN_URGENCY_ACTIONS unsafe { @@ -91,7 +114,7 @@ impl GpuMonitoringReducer { .arg(actions) .arg(&self.summary_buf) .arg(&(n as i32)) - .arg(&num_actions_i32) + .arg(&num_exposure_bins_i32) .arg(&order_actions_i32) .arg(&urgency_actions_i32) .launch(config) @@ -100,27 +123,34 @@ impl GpuMonitoringReducer { Ok(()) } - /// Download summary from GPU (single 48-byte transfer). + /// Download summary from GPU (single 108-byte transfer). pub fn download_summary(&self) -> Result { - let mut raw = vec![0.0_f32; 24]; + let mut raw = vec![0.0_f32; 27]; super::dtoh_f32(&self.stream, &self.summary_buf, &mut raw)?; + // Kernel layout: + // summary[0..5) = stats (mean, std, min, max, sharpe) + // summary[5..17) = exposure bins [12] (dir*3 + mag, dir∈0..4, mag∈0..3) + // summary[17..20) = order counts [3] + // summary[20..23) = urgency counts [3] + // summary[23] = total bars N + // summary[24] = total trades (non-zero rewards) + // summary[25..27) = pad Ok(MonitoringSummary { mean_reward: raw[0], reward_std: raw[1], min_reward: raw[2], max_reward: raw[3], sharpe_estimate: raw[4], - // Kernel layout: summary[5..14]=exp[9], summary[14..17]=order[3], - // summary[17..20]=urgency[3], summary[20]=total, summary[21]=trades action_counts: [ - raw[5] as usize, raw[6] as usize, raw[7] as usize, - raw[8] as usize, raw[9] as usize, raw[10] as usize, + raw[5] as usize, raw[6] as usize, raw[7] as usize, + raw[8] as usize, raw[9] as usize, raw[10] as usize, raw[11] as usize, raw[12] as usize, raw[13] as usize, + raw[14] as usize, raw[15] as usize, raw[16] as usize, ], - order_counts: [raw[14] as usize, raw[15] as usize, raw[16] as usize], - urgency_counts: [raw[17] as usize, raw[18] as usize, raw[19] as usize], - total_experiences: raw[20] as usize, - total_trades: raw[21] as usize, + order_counts: [raw[17] as usize, raw[18] as usize, raw[19] as usize], + urgency_counts: [raw[20] as usize, raw[21] as usize, raw[22] as usize], + total_experiences: raw[23] as usize, + total_trades: raw[24] as usize, }) } } @@ -133,7 +163,7 @@ mod tests { fn test_monitoring_summary_default() { let s = MonitoringSummary::default(); assert_eq!(s.total_experiences, 0); - assert_eq!(s.action_counts, [0; 9]); + assert_eq!(s.action_counts, [0; 12]); } #[test] diff --git a/crates/ml/src/cuda_pipeline/monitoring_kernel.cu b/crates/ml/src/cuda_pipeline/monitoring_kernel.cu index 893eb61d4..9a30c0e45 100644 --- a/crates/ml/src/cuda_pipeline/monitoring_kernel.cu +++ b/crates/ml/src/cuda_pipeline/monitoring_kernel.cu @@ -7,23 +7,26 @@ // order. Was atomicAdd-into-shared-scalar — atomic is order-of-arrival, so // for floats the result varied across runs by a few ULPs. Now bit-stable. // -// Shared memory: ~2.6 KB for MAX_WARPS=32 warp scratch (well under 48 KB). +// Shared memory: ~3.0 KB for MAX_WARPS=32 warp scratch (well under 48 KB). extern "C" __global__ void monitoring_reduce( const float* __restrict__ rewards, // [N] f32 const int* __restrict__ actions, // [N] - float* summary, // [24]: mean, std, min, max, sharpe, exp[9], order[3], urgency[3], total, _pad2 + float* summary, // [27]: mean, std, min, max, sharpe, exp[12], order[3], urgency[3], total, trades, _pad2 int N, - int num_actions, // runtime: dir*mag combos (9 = 3*3) + int num_exposure_bins, // runtime: b0*b1 = dir*mag combos (12 = 4*3) int order_actions, // runtime: DQN_ORDER_ACTIONS (3) int urgency_actions // runtime: DQN_URGENCY_ACTIONS (3) ) { - /* Per-warp scratch arrays — one slot per warp (max 32 warps for 1024-thread block). */ + /* Per-warp scratch arrays — one slot per warp (max 32 warps for 1024-thread block). + * EXP_MAX = 12 matches the 4-direction × 3-magnitude layout used by the runtime + * DQN (Short/Hold/Long/Flat × Small/Half/Full). */ + constexpr int EXP_MAX = 12; __shared__ float w_sum[32]; __shared__ float w_sq_sum[32]; __shared__ float w_min[32]; __shared__ float w_max[32]; __shared__ int w_nonzero[32]; - __shared__ int w_exp[32][9]; + __shared__ int w_exp[32][EXP_MAX]; __shared__ int w_ord[32][3]; __shared__ int w_urg[32][3]; @@ -38,7 +41,7 @@ extern "C" __global__ void monitoring_reduce( float local_min = 1e30f; float local_max = -1e30f; int local_nonzero = 0; - int lc_exp[9] = {0,0,0,0,0,0,0,0,0}; + int lc_exp[EXP_MAX] = {0,0,0,0,0,0,0,0,0,0,0,0}; int lc_ord[3] = {0,0,0}; int lc_urg[3] = {0,0,0}; @@ -56,13 +59,15 @@ extern "C" __global__ void monitoring_reduce( local_nonzero++; } /* Decode dir*mag + order + urgency from factored action: - * action = exp * (b1*b2) + ord * b2 + urg */ + * action = exp * (b1*b2) + ord * b2 + urg + * where exp = dir * b1 + mag, with b0=4 (dir: Short/Hold/Long/Flat) + * and b1=3 (mag: Small/Half/Full). exp range = [0, 12). */ int a = actions[i]; int exp_idx = a / b1b2; int rem = a % b1b2; int ord_idx = rem / urgency_actions; int urg_idx = rem % urgency_actions; - if (exp_idx >= 0 && exp_idx < num_actions) lc_exp[exp_idx]++; + if (exp_idx >= 0 && exp_idx < num_exposure_bins) lc_exp[exp_idx]++; if (ord_idx >= 0 && ord_idx < order_actions) lc_ord[ord_idx]++; if (urg_idx >= 0 && urg_idx < urgency_actions) lc_urg[urg_idx]++; } @@ -75,7 +80,7 @@ extern "C" __global__ void monitoring_reduce( local_min = fminf(local_min, __shfl_xor_sync(0xFFFFFFFF, local_min, mask)); local_max = fmaxf(local_max, __shfl_xor_sync(0xFFFFFFFF, local_max, mask)); local_nonzero += __shfl_xor_sync(0xFFFFFFFF, local_nonzero, mask); - for (int i = 0; i < 9; i++) + for (int i = 0; i < EXP_MAX; i++) lc_exp[i] += __shfl_xor_sync(0xFFFFFFFF, lc_exp[i], mask); for (int i = 0; i < 3; i++) { lc_ord[i] += __shfl_xor_sync(0xFFFFFFFF, lc_ord[i], mask); @@ -91,7 +96,7 @@ extern "C" __global__ void monitoring_reduce( w_min[warp_id] = local_min; w_max[warp_id] = local_max; w_nonzero[warp_id] = local_nonzero; - for (int i = 0; i < 9; i++) w_exp[warp_id][i] = lc_exp[i]; + for (int i = 0; i < EXP_MAX; i++) w_exp[warp_id][i] = lc_exp[i]; for (int i = 0; i < 3; i++) { w_ord[warp_id][i] = lc_ord[i]; w_urg[warp_id][i] = lc_urg[i]; @@ -105,7 +110,7 @@ extern "C" __global__ void monitoring_reduce( float t_sum = 0.0f, t_sq = 0.0f; float t_min = 1e30f, t_max = -1e30f; int t_nonzero = 0; - int t_exp[9] = {0,0,0,0,0,0,0,0,0}; + int t_exp[EXP_MAX] = {0,0,0,0,0,0,0,0,0,0,0,0}; int t_ord[3] = {0,0,0}; int t_urg[3] = {0,0,0}; for (int w = 0; w < num_warps; w++) { @@ -114,7 +119,7 @@ extern "C" __global__ void monitoring_reduce( t_min = fminf(t_min, w_min[w]); t_max = fmaxf(t_max, w_max[w]); t_nonzero += w_nonzero[w]; - for (int i = 0; i < 9; i++) t_exp[i] += w_exp[w][i]; + for (int i = 0; i < EXP_MAX; i++) t_exp[i] += w_exp[w][i]; for (int i = 0; i < 3; i++) { t_ord[i] += w_ord[w][i]; t_urg[i] += w_urg[w][i]; @@ -131,16 +136,25 @@ extern "C" __global__ void monitoring_reduce( ? t_sq / denom - mean * mean : 0.0f; float std_val = sqrtf(fmaxf(var, 0.0f)); + /* Summary layout (27 floats): + * [0..5) : mean, std, min, max, sharpe + * [5..17) : exposure bins (12 = b0*b1 = 4*3) + * [17..20): order type counts (3) + * [20..23): urgency counts (3) + * [23] : total bars N + * [24] : total trades (non-zero rewards) + * [25..27): pad */ summary[0] = mean; summary[1] = std_val; summary[2] = (n_trades > 0) ? t_min : 0.0f; summary[3] = (n_trades > 0) ? t_max : 0.0f; summary[4] = (std_val > 1e-8f) ? mean / std_val : 0.0f; - for (int i = 0; i < 9; i++) summary[5 + i] = (float)t_exp[i]; - for (int i = 0; i < 3; i++) summary[14 + i] = (float)t_ord[i]; - for (int i = 0; i < 3; i++) summary[17 + i] = (float)t_urg[i]; - summary[20] = (float)N; - summary[21] = (float)n_trades; - for (int i = 22; i < 24; i++) summary[i] = 0.0f; + for (int i = 0; i < num_exposure_bins; i++) summary[5 + i] = (float)t_exp[i]; + for (int i = 0; i < 3; i++) summary[17 + i] = (float)t_ord[i]; + for (int i = 0; i < 3; i++) summary[20 + i] = (float)t_urg[i]; + summary[23] = (float)N; + summary[24] = (float)n_trades; + summary[25] = 0.0f; + summary[26] = 0.0f; } } diff --git a/crates/ml/src/trainers/dqn/financials.rs b/crates/ml/src/trainers/dqn/financials.rs index 44ce2c75b..eeb037ead 100644 --- a/crates/ml/src/trainers/dqn/financials.rs +++ b/crates/ml/src/trainers/dqn/financials.rs @@ -23,14 +23,15 @@ pub(crate) struct EpochFinancials { /// Compute financial metrics from GPU-sourced trade statistics and action counts. /// /// - `trade_stats`: per-epoch trade statistics from `collector.collect_trade_stats()` -/// - `action_counts`: 7-element array (exposure levels), grouped into BUY/SELL/HOLD -/// Indices: 0=ShortSmall, 1=ShortHalf, 2=ShortFull, 3=Flat, 4=LongSmall, 5=LongHalf, 6=LongFull +/// - `action_counts`: 12-bin `dir*3 + mag` layout matching the monitoring kernel. +/// dir ∈ {Short=0, Hold=1, Long=2, Flat=3}, mag ∈ {Small=0, Half=1, Full=2}. +/// Grouped into BUY/SELL/HOLD (see below). /// - `initial_capital`: starting equity for return calculation (default 100_000) /// - `bars_per_day`: number of bars in one trading day (390 for 1-min, 78 for 5-min, etc.) /// Used for Sharpe/Sortino annualization and return scaling. Pass from data pipeline's BarSize. pub(crate) fn compute_epoch_financials( trade_stats: &TradeStats, - action_counts: &[usize; 9], + action_counts: &[usize; 12], initial_capital: f64, bars_per_day: f64, ) -> EpochFinancials { @@ -168,15 +169,18 @@ pub(crate) fn compute_epoch_financials( // Cap at 100% — drawdown cannot exceed total liquidation let max_dd = max_dd.min(1.0); - // Action distribution: 7 exposure actions -> BUY/SELL/HOLD - // [ShortSmall, ShortHalf, ShortFull, Flat, LongSmall, LongHalf, LongFull] - // 0 1 2 3 4 5 6 - // Sell = indices 0-2, Hold = index 3, Buy = indices 4-6 + // Action distribution: 12-bin exposure layout → BUY/SELL/HOLD. + // Layout: exp_idx = dir*3 + mag with dir ∈ {Short=0, Hold=1, Long=2, Flat=3} + // and mag ∈ {Small=0, Half=1, Full=2}. + // Sell = Short dir = indices 0..3 + // Hold = Hold + Flat dirs = indices 3..6 ∪ 9..12 (both are non-directional positions) + // Buy = Long dir = indices 6..9 let total_actions: usize = action_counts.iter().sum(); let (buy_pct, sell_pct, hold_pct) = if total_actions > 0 { - let sell: usize = action_counts.iter().take(3).sum(); // ShortSmall+ShortHalf+ShortFull - let hold = *action_counts.get(3).unwrap_or(&0); // Flat - let buy: usize = action_counts.iter().skip(4).sum(); // LongSmall+LongHalf+LongFull + let sell: usize = action_counts[0..3].iter().sum(); + let hold: usize = action_counts[3..6].iter().sum::() + + action_counts[9..12].iter().sum::(); + let buy: usize = action_counts[6..9].iter().sum(); let t = total_actions as f64; (buy as f64 / t, sell as f64 / t, hold as f64 / t) } else { @@ -207,7 +211,7 @@ mod tests { #[test] fn test_empty_trade_stats() { let ts = TradeStats::default(); - let f = compute_epoch_financials(&ts, &[0; 9], 100_000.0, 390.0); + let f = compute_epoch_financials(&ts, &[0; 12], 100_000.0, 390.0); assert_eq!(f.total_trades, 0); assert_eq!(f.sharpe, 0.0); } @@ -225,7 +229,7 @@ mod tests { step_returns: vec![0.01, 0.02, 0.03, 0.015, 0.025], ..Default::default() }; - let f = compute_epoch_financials(&ts, &[0; 9], 100_000.0, 390.0); + let f = compute_epoch_financials(&ts, &[0; 12], 100_000.0, 390.0); assert_eq!(f.win_rate, 1.0); assert_eq!(f.total_trades, 5); assert!(f.sharpe > 0.0, "sharpe={}", f.sharpe); @@ -246,7 +250,7 @@ mod tests { step_returns: vec![0.01, -0.005, 0.008, -0.003, 0.005], ..Default::default() }; - let f = compute_epoch_financials(&ts, &[0; 9], 100_000.0, 390.0); + let f = compute_epoch_financials(&ts, &[0; 12], 100_000.0, 390.0); assert_eq!(f.total_trades, 5); assert!((f.win_rate - 0.6).abs() < 1e-10, "win_rate={}", f.win_rate); assert!(f.total_return > 0.0, "total_return={}", f.total_return); @@ -256,13 +260,15 @@ mod tests { } #[test] - fn test_action_distribution_7_actions() { - let mut actions = [0usize; 9]; - actions[0] = 10; // ShortSmall -> SELL - actions[1] = 10; // ShortHalf -> SELL - actions[3] = 20; // Flat -> HOLD - actions[5] = 30; // LongHalf -> BUY - actions[6] = 30; // LongFull -> BUY + fn test_action_distribution_12_bins() { + // 12-bin layout (dir*3 + mag): Short [0..3], Hold [3..6], Long [6..9], Flat [9..12]. + let mut actions = [0usize; 12]; + actions[0] = 10; // Short_Small -> SELL + actions[1] = 10; // Short_Half -> SELL + actions[3] = 15; // Hold_Small -> HOLD + actions[7] = 30; // Long_Half -> BUY + actions[8] = 30; // Long_Full -> BUY + actions[9] = 5; // Flat_Small -> HOLD let ts = TradeStats { total_trades: 1, winning_trades: 1, @@ -275,16 +281,19 @@ mod tests { ..Default::default() }; let f = compute_epoch_financials(&ts, &actions, 100_000.0, 390.0); - // SELL=20/100=0.2, HOLD=20/100=0.2, BUY=60/100=0.6 - assert!((f.sell_pct - 0.2).abs() < 1e-10, "sell_pct={}", f.sell_pct); - assert!((f.hold_pct - 0.2).abs() < 1e-10, "hold_pct={}", f.hold_pct); - assert!((f.buy_pct - 0.6).abs() < 1e-10, "buy_pct={}", f.buy_pct); + // Total = 10+10+15+30+30+5 = 100. + // SELL = 20/100 = 0.20, HOLD = (15+5)/100 = 0.20, BUY = 60/100 = 0.60. + assert!((f.sell_pct - 0.20).abs() < 1e-10, "sell_pct={}", f.sell_pct); + assert!((f.hold_pct - 0.20).abs() < 1e-10, "hold_pct={}", f.hold_pct); + assert!((f.buy_pct - 0.60).abs() < 1e-10, "buy_pct={}", f.buy_pct); } #[test] fn test_action_distribution_all_indices_counted() { - let mut actions = [0usize; 9]; - for i in 0..7 { actions[i] = 10; } + // Fill every bin with a non-zero count to prove all 12 indices contribute + // and BUY+SELL+HOLD partitions the total. + let mut actions = [0usize; 12]; + for i in 0..12 { actions[i] = 10; } let ts = TradeStats { total_trades: 1, winning_trades: 1, @@ -328,7 +337,7 @@ mod tests { step_returns, done_flags, }; - let f = compute_epoch_financials(&ts, &[0; 9], 100_000.0, 390.0); + let f = compute_epoch_financials(&ts, &[0; 12], 100_000.0, 390.0); // Episode 1 drawdown: (1-0.8)*(1-0.15)*(1-0.10)*(1-0.05) ≈ 0.4131 → ~42% DD // Episode 2 starts fresh, drawdown is smaller. @@ -360,7 +369,7 @@ mod tests { step_returns: vec![0.01, -0.005, 0.008, -0.003], done_flags: vec![], // empty — no episode info }; - let f = compute_epoch_financials(&ts, &[0; 9], 100_000.0, 390.0); + let f = compute_epoch_financials(&ts, &[0; 12], 100_000.0, 390.0); // Should not panic and should produce a valid number assert!(f.max_drawdown >= 0.0, "max_dd={}", f.max_drawdown); assert!(f.max_drawdown <= 1.0, "max_dd={}", f.max_drawdown); diff --git a/crates/ml/src/trainers/dqn/monitoring.rs b/crates/ml/src/trainers/dqn/monitoring.rs index f6ecfcc92..d424f47cb 100644 --- a/crates/ml/src/trainers/dqn/monitoring.rs +++ b/crates/ml/src/trainers/dqn/monitoring.rs @@ -8,18 +8,35 @@ use tracing::{debug, warn}; use crate::dqn::action_space::FactoredAction; +/// 12-entry exposure-bin labels matching the monitoring kernel's `dir*3 + mag` +/// layout. Direction prefix: +/// S = Short (dir=0), H = Hold (dir=1), L = Long (dir=2), F = Flat (dir=3). +/// Magnitude suffix: _Small (mag=0), _Half (mag=1), _Full (mag=2). +/// Hold and Flat force mag=0 in the kernel, so the _Half / _Full slots for +/// H and F are empty by construction but kept here so index → name is a +/// straight `dir*3 + mag` lookup everywhere. +const EXPOSURE_NAMES: [&str; 12] = [ + "S_Small", "S_Half", "S_Full", // dir=0 Short + "H_Small", "H_Half", "H_Full", // dir=1 Hold (Half/Full unreachable) + "L_Small", "L_Half", "L_Full", // dir=2 Long + "F_Small", "F_Half", "F_Full", // dir=3 Flat (Half/Full unreachable) +]; + /// Training monitor for per-epoch tracking and validation pub(crate) struct TrainingMonitor { pub(crate) epoch: usize, pub(crate) reward_history: Vec, - pub(crate) action_counts: [usize; 9], // 7 exposure levels (ShortSmall..LongFull) - pub(crate) q_value_sums: [f64; 7], // Sum of Q-values per exposure level - pub(crate) q_value_counts: [usize; 9], // Count of Q-values per exposure level + /// Per dir×mag exposure counts (12 = b0*b1 = 4*3, laid out as dir*3+mag). + pub(crate) action_counts: [usize; 12], + /// Sum of Q-values per exposure bin (same 12-bin layout). + pub(crate) q_value_sums: [f64; 12], + /// Count of Q-values per exposure bin (same 12-bin layout). + pub(crate) q_value_counts: [usize; 12], pub(crate) order_type_counts: [usize; 3], // Market, LimitMaker, IoC pub(crate) urgency_counts: [usize; 3], // Patient, Normal, Aggressive - /// Factored action counts: 7 exposure * 3 order * 3 urgency = 63 actions. - /// Index = exp_idx * 9 + order * 3 + urgency (0-62). - pub(crate) factored_action_counts: [usize; 63], + /// Factored action counts: 12 exposure × 3 order × 3 urgency = 108 actions. + /// Index = exp_idx * 9 + order * 3 + urgency (0-107), where exp_idx = dir*3 + mag. + pub(crate) factored_action_counts: [usize; 108], /// Per-branch Q-value tracking: direction[3], magnitude[3], order[3], urgency[3] pub(crate) branch_q_sums: [[f64; 3]; 4], pub(crate) branch_q_counts: [[u64; 3]; 4], @@ -40,12 +57,12 @@ impl TrainingMonitor { Self { epoch, reward_history: Vec::new(), - action_counts: [0; 9], - q_value_sums: [0.0; 7], - q_value_counts: [0; 9], + action_counts: [0; 12], + q_value_sums: [0.0; 12], + q_value_counts: [0; 12], order_type_counts: [0; 3], urgency_counts: [0; 3], - factored_action_counts: [0; 63], + factored_action_counts: [0; 108], branch_q_sums: [[0.0; 3]; 4], branch_q_counts: [[0; 3]; 4], consecutive_constant_epochs: 0, @@ -71,24 +88,35 @@ impl TrainingMonitor { } } - /// Add action to tracking by exposure index (0-6) + /// Add action to tracking by exposure index (0..12, = dir*3 + mag). pub(crate) fn track_action_by_exposure(&mut self, exposure_idx: usize) { - if exposure_idx < 7 { + if exposure_idx < 12 { self.action_counts[exposure_idx] += 1; } } - /// Add action to tracking (extracts exposure from FactoredAction) + /// Add action to tracking (extracts exposure from FactoredAction). + /// + /// FactoredAction.exposure is the 8-level `ExposureLevel` enum. This function + /// maps it onto the kernel's 12-bin `dir*3 + mag` layout. Hold and Flat + /// collapse to mag=0 (same convention the kernel uses). pub(crate) fn track_action(&mut self, action: &FactoredAction) { - let exp_idx = action.exposure as usize; // 0-4 - self.action_counts[exp_idx] += 1; + let dir = action.exposure.direction() as usize; // 0..4 + let mag_raw = action.exposure.magnitude() as usize; // 0..3 (3 = N/A for Hold/Flat) + let mag = if mag_raw >= 3 { 0 } else { mag_raw }; // Hold/Flat collapse to mag=0 + let exp_idx = dir * 3 + mag; // 0..12 + if exp_idx < 12 { + self.action_counts[exp_idx] += 1; + } let order_idx = action.order as usize; // 0-2 self.order_type_counts[order_idx] += 1; let urgency_idx = action.urgency as usize; // 0-2 self.urgency_counts[urgency_idx] += 1; - // Track composite factored action index (0-62): exp(0-6) * 9 + order(0-2) * 3 + urgency(0-2) + // Track composite factored action index (0..108): + // factored_idx = exp_idx * 9 + order * 3 + urgency + // Max = 11*9 + 2*3 + 2 = 107. let factored_idx = exp_idx * 9 + order_idx * 3 + urgency_idx; - if factored_idx < 63 { + if factored_idx < 108 { self.factored_action_counts[factored_idx] += 1; } } @@ -115,19 +143,25 @@ impl TrainingMonitor { } } - /// Add Q-value to tracking by exposure index (0-6) + /// Add Q-value to tracking by exposure index (0..12, = dir*3 + mag). pub(crate) fn track_q_value_by_exposure(&mut self, exposure_idx: usize, q_value: f64) { - if exposure_idx < 7 { + if exposure_idx < 12 { self.q_value_sums[exposure_idx] += q_value; self.q_value_counts[exposure_idx] += 1; } } - /// Add Q-value to tracking (extracts exposure from FactoredAction) + /// Add Q-value to tracking (extracts exposure from FactoredAction). + /// Uses the same 12-bin `dir*3 + mag` layout as the kernel. pub(crate) fn track_q_value(&mut self, action: &FactoredAction, q_value: f64) { - let idx = action.exposure as usize; // 0-4 - self.q_value_sums[idx] += q_value; - self.q_value_counts[idx] += 1; + let dir = action.exposure.direction() as usize; + let mag_raw = action.exposure.magnitude() as usize; + let mag = if mag_raw >= 3 { 0 } else { mag_raw }; + let idx = dir * 3 + mag; + if idx < 12 { + self.q_value_sums[idx] += q_value; + self.q_value_counts[idx] += 1; + } } /// Track Q-value range for monitoring (WAVE 9-11 production) @@ -203,13 +237,12 @@ impl TrainingMonitor { return Ok(()); } - let exposure_names = ["S_Small", "S_Half", "S_Full", "F_Small", "F_Half", "F_Full", "L_Small", "L_Half", "L_Full"]; for (i, &count) in self.action_counts.iter().enumerate() { let pct = (count as f64 / total_exposure as f64) * 100.0; if pct < 5.0 { warn!( "LOW EXPOSURE DIVERSITY at epoch {}: {} only {:.1}% ({}/{})", - self.epoch, exposure_names[i], pct, count, total_exposure + self.epoch, EXPOSURE_NAMES[i], pct, count, total_exposure ); } } @@ -233,9 +266,9 @@ impl TrainingMonitor { /// Validate Q-value balance across exposure actions pub(crate) fn validate_q_value_balance(&self) -> Result<()> { - // Calculate average Q-value per exposure level - let mut avg_q_values = [0.0_f64; 7]; - for i in 0..7 { + // Calculate average Q-value per exposure bin (12 = dir×mag) + let mut avg_q_values = [0.0_f64; 12]; + for i in 0..12 { if self.q_value_counts[i] > 0 { avg_q_values[i] = self.q_value_sums[i] / self.q_value_counts[i] as f64; } @@ -246,12 +279,11 @@ impl TrainingMonitor { let min_q = avg_q_values.iter().cloned().fold(f64::INFINITY, f64::min); if (max_q - min_q).abs() > 1000.0 { - let exposure_names = ["S_Small", "S_Half", "S_Full", "F_Small", "F_Half", "F_Full", "L_Small", "L_Half", "L_Full"]; warn!( "Q-VALUE DIVERGENCE at epoch {}: {}", self.epoch, avg_q_values.iter().enumerate() - .map(|(i, q)| format!("{}={:.2}", exposure_names[i], q)) + .map(|(i, q)| format!("{}={:.2}", EXPOSURE_NAMES[i], q)) .collect::>() .join(", ") ); @@ -265,16 +297,14 @@ impl TrainingMonitor { if self.epoch % 10 == 0 { let total_actions: usize = self.action_counts.iter().sum(); if total_actions > 0 { - let exposure_names = ["S_Small", "S_Half", "S_Full", "F_Small", "F_Half", "F_Full", "L_Small", "L_Half", "L_Full"]; - - // Log all 7 exposure level actions + // Log all 12 exposure level actions (dir×mag = 4×3) debug!( - "Action Distribution [Epoch {}] - 7 Exposure Level Actions:", + "Action Distribution [Epoch {}] - 12 Exposure Level Actions (dir×mag):", self.epoch ); for (idx, &count) in self.action_counts.iter().enumerate() { let pct = (count as f64 / total_actions as f64) * 100.0; - debug!(" [{}] {}: {} ({:.1}%)", idx, exposure_names[idx], count, pct); + debug!(" [{}] {}: {} ({:.1}%)", idx, EXPOSURE_NAMES[idx], count, pct); } // Log order type distribution (branching DQN) @@ -305,10 +335,10 @@ impl TrainingMonitor { } } - // Per-branch Q-value averages + // Per-branch Q-value averages (4 branches: direction×4, magnitude×3, order×3, urgency×3) let branch_names = ["Direction", "Magnitude", "Order", "Urgency"]; let action_names = [ - ["Short", "Flat", "Long"], + ["Short", "Hold", "Long"], // Direction slot 0..3 — the 4th (Flat) is tracked via action_counts[9..12] ["Small", "Half", "Full"], ["Market", "LimitMaker", "IoC"], ["Patient", "Normal", "Aggressive"], @@ -330,16 +360,16 @@ impl TrainingMonitor { let unique_ord = self.order_type_counts.iter().filter(|&&c| c > 0).count(); let unique_urg = self.urgency_counts.iter().filter(|&&c| c > 0).count(); debug!( - "Branch Diversity [Epoch {}]: exp={}/7 order={}/3 urgency={}/3", + "Branch Diversity [Epoch {}]: exp={}/12 order={}/3 urgency={}/3", self.epoch, unique_exp, unique_ord, unique_urg ); - // Log average Q-values per exposure level + // Log average Q-values per exposure bin (12 = dir×mag) debug!("Average Q-values [Epoch {}]:", self.epoch); - for idx in 0..7 { + for idx in 0..12 { if self.q_value_counts[idx] > 0 { let avg_q = self.q_value_sums[idx] / self.q_value_counts[idx] as f64; - debug!(" [{}] {}: Q={:.4}", idx, exposure_names[idx], avg_q); + debug!(" [{}] {}: Q={:.4}", idx, EXPOSURE_NAMES[idx], avg_q); } } } diff --git a/crates/ml/src/trainers/dqn/trainer/metrics.rs b/crates/ml/src/trainers/dqn/trainer/metrics.rs index 7f1ea7734..ca3c43d70 100644 --- a/crates/ml/src/trainers/dqn/trainer/metrics.rs +++ b/crates/ml/src/trainers/dqn/trainer/metrics.rs @@ -99,8 +99,8 @@ impl DQNTrainer { num_epochs: usize, training_duration: std::time::Duration, early_stopped: bool, - total_action_counts: [usize; 9], // 7 exposure levels - total_factored_action_counts: [usize; 63], // 63 factored actions (7 exp * 3 ord * 3 urg) + total_action_counts: [usize; 12], // 12 exposure bins (dir×mag = 4×3) + total_factored_action_counts: [usize; 108], // 108 factored actions (12 exp * 3 ord * 3 urg) ) -> Result { let final_loss = total_loss / num_epochs as f64; let avg_q_value_final = total_q_value / num_epochs as f64; @@ -141,26 +141,26 @@ impl DQNTrainer { metrics.add_metric("first_sharpe", first_sharpe); metrics.add_metric("final_val_loss", final_val_loss); - // Action metrics: use factored 63-action space if branching, else 7 exposure levels + // Action metrics: use factored 108-action space if branching, else 12 exposure bins let total_factored: usize = total_factored_action_counts.iter().sum(); let total_exposure: usize = total_action_counts.iter().sum(); let total_actions = total_factored.max(total_exposure); if total_actions > 0 { - // Factored 63-action diversity (primary metric when branching) + // Factored 108-action diversity (primary metric when branching) if total_factored > 0 { let unique_factored = total_factored_action_counts.iter() .filter(|&&count| count > 0).count(); - let factored_diversity = (unique_factored as f64 / 63.0) * 100.0; + let factored_diversity = (unique_factored as f64 / 108.0) * 100.0; metrics.add_metric("action_diversity", factored_diversity); metrics.add_metric("factored_unique_actions", unique_factored as f64); - metrics.add_metric("action_space_size", 63.0); + metrics.add_metric("action_space_size", 108.0); // Active factored actions (used >0.5% of the time) let active_threshold = (total_factored as f64 * 0.005).max(1.0); let active_count = total_factored_action_counts.iter() .filter(|&&count| count as f64 >= active_threshold) .count(); - let active_diversity_pct = (active_count as f64 / 63.0) * 100.0; + let active_diversity_pct = (active_count as f64 / 108.0) * 100.0; metrics.add_metric("active_actions_count", active_count as f64); metrics.add_metric("active_diversity_pct", active_diversity_pct); @@ -182,18 +182,18 @@ impl DQNTrainer { let top5_coverage_pct = (top5_count as f64 / total_factored as f64) * 100.0; metrics.add_metric("top5_coverage_pct", top5_coverage_pct); } else { - // Fallback: exposure-only (7 levels) + // Fallback: exposure-only (12 bins: dir×mag = 4×3) let unique_actions = total_action_counts.iter() .filter(|&&count| count > 0).count(); - let action_diversity = (unique_actions as f64 / 7.0) * 100.0; + let action_diversity = (unique_actions as f64 / 12.0) * 100.0; metrics.add_metric("action_diversity", action_diversity); - metrics.add_metric("action_space_size", 7.0); + metrics.add_metric("action_space_size", 12.0); let active_threshold = (total_exposure as f64 * 0.005).max(1.0); let active_count = total_action_counts.iter() .filter(|&&count| count as f64 >= active_threshold) .count(); - let active_diversity_pct = (active_count as f64 / 7.0) * 100.0; + let active_diversity_pct = (active_count as f64 / 12.0) * 100.0; metrics.add_metric("active_actions_count", active_count as f64); metrics.add_metric("active_diversity_pct", active_diversity_pct); @@ -217,15 +217,17 @@ impl DQNTrainer { metrics.add_metric("total_actions", total_actions as f64); - // Buy/sell/hold from 7-exposure space (meaningful for P&L) - // 0=ShortSmall, 1=ShortHalf, 2=ShortFull, 3=Flat, 4=LongSmall, 5=LongHalf, 6=LongFull - let sell_count: usize = total_action_counts.get(0).copied().unwrap_or(0) - + total_action_counts.get(1).copied().unwrap_or(0) - + total_action_counts.get(2).copied().unwrap_or(0); - let hold_count: usize = total_action_counts.get(3).copied().unwrap_or(0); - let buy_count: usize = total_action_counts.get(4).copied().unwrap_or(0) - + total_action_counts.get(5).copied().unwrap_or(0) - + total_action_counts.get(6).copied().unwrap_or(0); + // Buy/sell/hold from 12-bin exposure space (meaningful for P&L). + // Layout: exp_idx = dir*3 + mag, dir ∈ {Short=0, Hold=1, Long=2, Flat=3}. + // Sell = Short directions = indices 0..3 + // Hold = Hold + Flat bins = indices 3..6 ∪ 9..12 + // Buy = Long directions = indices 6..9 + // Flat is "close to zero" which from a trading-flow perspective is neither + // a long nor a short — we classify it as Hold (no new directional bet). + let sell_count: usize = total_action_counts[0..3].iter().sum(); + let hold_count: usize = total_action_counts[3..6].iter().sum::() + + total_action_counts[9..12].iter().sum::(); + let buy_count: usize = total_action_counts[6..9].iter().sum(); metrics.add_metric("buy_count", buy_count as f64); metrics.add_metric("sell_count", sell_count as f64); metrics.add_metric("hold_count", hold_count as f64); @@ -307,10 +309,10 @@ impl DQNTrainer { /// /// Returns (gap_stats, per_action_avgs) where: /// - gap_stats: (mean_gap, min_gap, max_gap) of Q_best - Q_second_best - /// - per_action_avgs: `[f64; 9]` averages (one per exposure action, capped at 5) + /// - per_action_avgs: `[f64; 12]` averages indexed by `dir*3 + mag` (4 dirs × 3 mags) pub(crate) async fn compute_epoch_q_diagnostics(&mut self, agent: &mut DQNAgentType) -> Option<( (f64, f64, f64), - [f64; 9], + [f64; 12], )> { let buffer = agent.memory_mut(); @@ -361,7 +363,9 @@ impl DQNTrainer { ag.branch_sizes() }; let mut gaps = Vec::with_capacity(rows); - let mut col_sums = [0.0_f64; 9]; + // 12 bins = b0 * b1 for the 4-branch DQN (dir×mag = 4×3). Aligns with + // the monitoring kernel's exposure layout and per_action_avgs consumer. + let mut col_sums = [0.0_f64; 12]; for r in 0..rows { let row_offset = r * cols; // Direction gap: best vs second_best within first b0 columns only @@ -377,13 +381,16 @@ impl DQNTrainer { } } gaps.push(best_val - second_best); - // Per-action sums: map branch outputs to 9 dir×mag combos - // dir_d × mag_m → index d*b1 + m (for display only) - for d in 0..b0.min(3) { + // Per-action sums: map branch outputs to dir×mag combos + // dir_d × mag_m → index d*3 + m (matches monitoring kernel layout). + // Cap at 4 directions × 3 magnitudes = 12 slots. + let b0_cap = b0.min(4); + let b1_cap = b1.min(3); + for d in 0..b0_cap { let q_dir = host_q.get(row_offset + d).copied().unwrap_or(0.0) as f64; - for m in 0..b1.min(3) { + for m in 0..b1_cap { let q_mag = host_q.get(row_offset + b0 + m).copied().unwrap_or(0.0) as f64; - let combo_idx = d * b1.min(3) + m; + let combo_idx = d * 3 + m; // kernel layout: always multiply by 3 (full mag dim) if let Some(s) = col_sums.get_mut(combo_idx) { // Combined Q ≈ Q_dir + Q_mag (additive advantage) *s += q_dir + q_mag; @@ -396,7 +403,7 @@ impl DQNTrainer { let gap_mean = gaps.iter().sum::() / rows as f32; let gap_min = gaps.iter().copied().fold(f32::INFINITY, f32::min); let gap_max = gaps.iter().copied().fold(f32::NEG_INFINITY, f32::max); - let mut per_action_avgs = [0.0_f64; 9]; + let mut per_action_avgs = [0.0_f64; 12]; for (i, avg) in per_action_avgs.iter_mut().enumerate() { *avg = col_sums.get(i).copied().unwrap_or(0.0) / rows_f64; } diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index f4bed076e..c3c3a1f3c 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -93,8 +93,8 @@ impl DQNTrainer { let mut total_q_value = 0.0; let mut total_gradient_norm = 0.0; let mut total_reward = 0.0; - let mut total_action_counts = [0_usize; 9]; - let mut total_factored_action_counts = [0_usize; 63]; + let mut total_action_counts = [0_usize; 12]; + let mut total_factored_action_counts = [0_usize; 108]; self.log_training_config().await; @@ -537,8 +537,10 @@ impl DQNTrainer { } } - // M3: Q-value diagnostics — disabled, causes hang in single-threaded tokio - let q_diagnostics: Option<((f64, f64, f64), [f64; 7])> = None; + // M3: Q-value diagnostics — disabled, causes hang in single-threaded tokio. + // Tuple is (gap_stats, per_action_avgs) with per_action_avgs length = 12 + // matching the kernel's dir×mag exposure layout (see metrics.rs). + let q_diagnostics: Option<((f64, f64, f64), [f64; 12])> = None; // Three-Phase Training Pipeline let total_epochs = self.hyperparams.epochs; @@ -1379,7 +1381,16 @@ impl DQNTrainer { } if let Some(ref mut mon) = self.gpu_monitoring { - if let Err(e) = mon.reduce(collector.rewards_gpu(), collector.actions_gpu(), count) { + // GPU monitoring expects b0_size for the exposure-bin dimensionality. + // The 4-branch DQN uses b0=4 (direction); acquire it from the agent so + // the kernel's `num_exposure_bins` matches the actual branch layout. + let (b0, _, _, _) = { + let ag = self.agent.read().await; + ag.branch_sizes() + }; + if let Err(e) = mon.reduce( + collector.rewards_gpu(), collector.actions_gpu(), count, b0, + ) { debug!("GPU monitoring reduce failed: {e}"); } if let Some(ref stream) = self.cuda_stream { @@ -2284,15 +2295,25 @@ impl DQNTrainer { }; // Track 1 action distribution: per-magnitude usage across the epoch. - // action_counts[9] layout is dir*3 + mag — Quarter = mag 0 (indices 0,3,6), - // Half = mag 1 (indices 1,4,7), Full = mag 2 (indices 2,5,8). Computed - // from the same monitor struct the LOW EXPOSURE DIVERSITY check uses. + // Layout: action_counts[12], exp_idx = dir*3 + mag with + // dir ∈ {0=Short, 1=Hold, 2=Long, 3=Flat}. + // Quarter = mag 0 (indices 0, 3, 6, 9) + // Half = mag 1 (indices 1, 4, 7, 10) + // Full = mag 2 (indices 2, 5, 8, 11) + // BUT: Hold (dir=1) and Flat (dir=3) force mag=0 in the kernel — + // their non-zero mag slots are always empty by construction. To + // measure "magnitude preference when the policy actually picks a + // tradable direction", exclude Hold and Flat from the denominator: + // denom = Σ (Short + Long) bins = indices [0..3] ∪ [6..9] + // This is what the smoke test actually wants to measure. let action_total_epoch: usize = monitor.action_counts.iter().sum(); - let (dist_q, dist_h, dist_f) = if action_total_epoch > 0 { - let t = action_total_epoch as f32; - let q = (monitor.action_counts[0] + monitor.action_counts[3] + monitor.action_counts[6]) as f32 / t; - let h = (monitor.action_counts[1] + monitor.action_counts[4] + monitor.action_counts[7]) as f32 / t; - let f = (monitor.action_counts[2] + monitor.action_counts[5] + monitor.action_counts[8]) as f32 / t; + let trade_total: usize = monitor.action_counts[0..3].iter().sum::() + + monitor.action_counts[6..9].iter().sum::(); + let (dist_q, dist_h, dist_f) = if trade_total > 0 { + let t = trade_total as f32; + let q = (monitor.action_counts[0] + monitor.action_counts[6]) as f32 / t; + let h = (monitor.action_counts[1] + monitor.action_counts[7]) as f32 / t; + let f = (monitor.action_counts[2] + monitor.action_counts[8]) as f32 / t; (q, h, f) } else { (0.0_f32, 0.0_f32, 0.0_f32) @@ -2303,18 +2324,20 @@ impl DQNTrainer { // Track 4 exploration entropy — per-branch action entropy normalized // to [0, 1]. Magnitude branch uses the [dist_q, dist_h, dist_f] above; // direction is summed across magnitudes per direction bin. - // Layout (dir*3 + mag): dir 0 = indices 0..3, dir 1 (Flat) = 3..6, dir 2 = 6..9. + // Layout (dir*3 + mag, 4 directions): dir 0=Short [0..3], dir 1=Hold + // [3..6], dir 2=Long [6..9], dir 3=Flat [9..12]. let ent_mag = shannon_entropy_normalized(&[dist_q, dist_h, dist_f]); - let (dir0, dir1, dir2) = if action_total_epoch > 0 { + let (dir0, dir1, dir2, dir3) = if action_total_epoch > 0 { let t = action_total_epoch as f32; - let d0 = (monitor.action_counts[0] + monitor.action_counts[1] + monitor.action_counts[2]) as f32 / t; - let d1 = (monitor.action_counts[3] + monitor.action_counts[4] + monitor.action_counts[5]) as f32 / t; - let d2 = (monitor.action_counts[6] + monitor.action_counts[7] + monitor.action_counts[8]) as f32 / t; - (d0, d1, d2) + let d0 = monitor.action_counts[0..3].iter().sum::() as f32 / t; + let d1 = monitor.action_counts[3..6].iter().sum::() as f32 / t; + let d2 = monitor.action_counts[6..9].iter().sum::() as f32 / t; + let d3 = monitor.action_counts[9..12].iter().sum::() as f32 / t; + (d0, d1, d2, d3) } else { - (0.0_f32, 0.0_f32, 0.0_f32) + (0.0_f32, 0.0_f32, 0.0_f32, 0.0_f32) }; - let ent_dir = shannon_entropy_normalized(&[dir0, dir1, dir2]); + let ent_dir = shannon_entropy_normalized(&[dir0, dir1, dir2, dir3]); // Task 0.14: append to exploration entropy history for smoke-test readback. self.explore_entropy_mag_history.push((epoch as u32, ent_mag)); @@ -2573,9 +2596,9 @@ impl DQNTrainer { boundary: &Option, monitor: &mut TrainingMonitor, epoch_duration: std::time::Duration, - total_action_counts: &mut [usize; 9], - total_factored_action_counts: &mut [usize; 63], - mut q_diagnostics: Option<((f64, f64, f64), [f64; 7])>, + total_action_counts: &mut [usize; 12], + total_factored_action_counts: &mut [usize; 108], + mut q_diagnostics: Option<((f64, f64, f64), [f64; 12])>, ) -> Result { // Calculate epoch metrics (average over training steps) let (epoch_loss, epoch_q_value, epoch_gradient_norm) = match boundary { @@ -2914,7 +2937,16 @@ impl DQNTrainer { gap_mean, gap_min, gap_max ); - let names = ["S100", "S75", "S50", "S25", "Flat", "L25", "L50", "L75", "L100"]; + // 12 labels matching the dir*3 + mag layout (dir: Short/Hold/Long/Flat). + // Hold and Flat magnitude slots are masked to mag=0 by the kernel, so + // H_Half/H_Full/F_Half/F_Full are structurally empty — listed here so + // index → name is a straight lookup everywhere in monitoring. + let names = [ + "S_Small", "S_Half", "S_Full", // dir=0 Short + "H_Small", "H_Half", "H_Full", // dir=1 Hold + "L_Small", "L_Half", "L_Full", // dir=2 Long + "F_Small", "F_Half", "F_Full", // dir=3 Flat + ]; let parts: Vec = names.iter() .zip(per_action_avgs.iter()) .map(|(n, q)| format!("{}={:.4}", n, q)) @@ -3003,20 +3035,29 @@ impl DQNTrainer { let total_factored_space = b0.max(1) * b1.max(1) * b2.max(1) * b3.max(1); let epoch_total: usize = monitor.action_counts.iter().sum(); - // Diversity threshold: 0.5% of DIRECTIONAL actions (excluding Flat). - // With Flat-dominant policies, using total inflates the threshold and - // mechanically kills diversity when Flat > 80%. The metric should measure - // magnitude diversity WITHIN directional positions, not overall share. - // 9-bin layout: dir(3)×mag(3). Flat cells = indices 3,4,5 (dir=1 × mag=0,1,2). - // In practice only Flat×Half (index 4) is reachable (Flat forces mag=1). - let flat_count: usize = monitor.action_counts[3] + monitor.action_counts[4] + monitor.action_counts[5]; - let directional_total = epoch_total.saturating_sub(flat_count).max(1); + // Diversity threshold: 1% of DIRECTIONAL actions (Short + Long only). + // With Flat/Hold-dominant policies, using total inflates the threshold and + // mechanically kills diversity when non-trading > 80%. The metric should + // measure magnitude diversity WITHIN directional (tradable) positions, + // not overall share. + // + // 12-bin layout: dir(4)×mag(3) with dir ∈ {Short=0, Hold=1, Long=2, Flat=3} + // and exp_idx = dir*3 + mag. Non-directional cells: + // Hold = indices 3,4,5 (mag forced to 0 by kernel → only index 3 populated) + // Flat = indices 9,10,11 (mag forced to 0 → only index 9 populated) + // Both buckets are excluded from the directional denominator and given the + // "any actions" pass-through rule below. + let flat_count: usize = monitor.action_counts[9..12].iter().sum(); + let hold_count: usize = monitor.action_counts[3..6].iter().sum(); + let directional_total = epoch_total.saturating_sub(flat_count + hold_count).max(1); let active_threshold = (directional_total as f64 * 0.01).max(1.0); - // action_counts[9]: dir×mag combos. Flat cells (3,4,5) always count if non-zero. + // action_counts[12]: dir×mag combos. Hold (3..6) and Flat (9..12) cells always + // count if non-zero — they're rare by the forced-mag=0 kernel convention, so + // any presence is meaningful. let active_dirmag = monitor.action_counts.iter().enumerate() .filter(|&(i, &c)| { - if (3..=5).contains(&i) { - c > 0 // Flat cells: count if any actions + if (3..6).contains(&i) || (9..12).contains(&i) { + c > 0 // Hold / Flat cells: count if any actions } else { c as f64 >= active_threshold // Directional cells: threshold } @@ -3028,8 +3069,10 @@ impl DQNTrainer { let active_factored = active_dirmag * active_ord * active_urg; let diversity_pct = (active_factored as f64 / total_factored_space as f64) * 100.0; - // Max dir*mag = b0 * b1 = 9 (full combinatorial space). - // Epsilon exploration can hit any combo including Flat×{Small,Full}. + // Max dir*mag = b0 * b1 = 12 for the 4-branch DQN (dir=4, mag=3). + // Epsilon exploration can hit any combo. Hold/Flat collapse their + // magnitude slots to mag=0 in the kernel, so the effective reachable + // set is 8 unique cells (Short×{S,H,F}, Long×{S,H,F}, Hold, Flat). let max_dirmag = b0 * b1; info!( "Epoch {}/{}: Action diversity={}/{} ({:.1}%) — dir*mag={}/{} order={}/{} urgency={}/{}", @@ -3048,7 +3091,7 @@ impl DQNTrainer { h -= p * p.ln(); } } - h / ((b0 * b1) as f64).ln() // normalize by log(9) for dir*mag combos + h / ((b0 * b1) as f64).ln() // normalize by log(b0*b1) = log(12) for dir*mag combos } else { 0.0 };