diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 9124eac21..7b67411f6 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -821,27 +821,52 @@ extern "C" __global__ void experience_action_select( (void)min_hold_bars; (void)portfolio_states; - /* Branch 0: direction — BOLTZMANN SOFTMAX selection. + /* Branch 0: direction — training uses BOLTZMANN SOFTMAX; eval uses strict + * argmax with Philox tie-break (Task 2.2 H10 fix). * - * C51's expected Q structurally favors Flat (zero PnL variance → tightest - * distribution → highest softmax expected Q). Argmax amplifies this into - * irrecoverable Flat dominance. Boltzmann samples direction PROPORTIONALLY - * to exp(Q/tau), ensuring Short/Long get explored even when Q-gaps are small. + * Training Boltzmann (C51's expected Q structurally favors Flat → zero PnL + * variance → tightest distribution → highest softmax expected Q; argmax + * amplifies into irrecoverable Flat dominance; Boltzmann samples direction + * PROPORTIONALLY to exp(Q/tau) so Short/Long get explored even when Q-gaps + * are small). Temperature tau = Q_range. * - * The Q-gap conviction filter is applied AFTER Boltzmann: if the sampled - * direction's Q is too close to Flat's Q, override to Flat. This preserves - * conviction gating while allowing diverse exploration. - * - * Temperature tau = Q_range: max e:1 ratio between best and worst direction. */ + * Eval strict-argmax + tie-break (Task 2.2): on ties within 1e-6, sample + * uniformly among tied indices via the existing Philox RNG state (same + * (i, timestep) seed used elsewhere in this kernel — no new atomics, no + * new RNG). Eval must be deterministic per (sample, epoch) so Boltzmann + * was replaced: with Q-range ≪ tau floor (0.01) the softmax spread was + * effectively uniform at eval even when Q-values meaningfully differed, + * which hid training-time argmax bias. Strict-argmax here honours the + * learned preference where it exists; tie-break only applies when + * |q_a − q_b| < 1e-6 (truly indistinguishable). */ if (!eval_mode && philox_uniform(i, timestep, rng_ctr++) < eps_dir) { int r = (int)(philox_uniform(i, timestep, rng_ctr++) * (float)b0_size); dir_idx = (r >= b0_size) ? b0_size - 1 : r; + } else if (eval_mode) { + /* Eval: strict argmax + uniform-among-tied Philox tie-break. */ + float q_max_d = q_sign * q_b0[0]; + for (int a = 1; a < b0_size; a++) { + float qv = q_sign * q_b0[a]; + if (qv > q_max_d + 1e-6f) q_max_d = qv; + } + int tie_count = 0; + for (int a = 0; a < b0_size; a++) { + float qv = q_sign * q_b0[a]; + if (fabsf(qv - q_max_d) <= 1e-6f) tie_count += 1; + } + int tie_pick = (int)(philox_uniform(i, timestep, rng_ctr++) * (float)tie_count); + if (tie_pick >= tie_count) tie_pick = tie_count - 1; + dir_idx = b0_size - 1; + int seen = 0; + for (int a = 0; a < b0_size; a++) { + float qv = q_sign * q_b0[a]; + if (fabsf(qv - q_max_d) <= 1e-6f) { + if (seen == tie_pick) { dir_idx = a; break; } + seen += 1; + } + } } else { - /* Both eval and training: Boltzmann softmax over direction Q-values. - * Eval uses the same Boltzmann — no RNG needed because tau→0 when - * Q-values differentiate (converges to argmax deterministically). - * When Q-values are flat, spreads evenly instead of flipping on noise. */ - /* Training mode: Boltzmann softmax over direction Q-values */ + /* Training mode: Boltzmann softmax over direction Q-values. */ float q_max_d = q_sign * (q_b0[0]); float q_min_d = q_max_d; for (int a = 1; a < b0_size; a++) { @@ -858,11 +883,6 @@ extern "C" __global__ void experience_action_select( exps_d[a] = expf((qv - q_max_d) / tau_d); sum_e += exps_d[a]; } - /* Both eval and training: Boltzmann sampling from softmax probs. - * Eval determinism comes from Philox seed (same i+timestep → same - * sample), NOT from argmax. argmax(softmax) = argmax(Q) which - * picks the same action for every state → val_Sharpe=0.00 freeze. - * Boltzmann sampling differentiates even when Q-gap is small. */ float r = philox_uniform(i, timestep, rng_ctr++) * sum_e; float cum = 0.0f; dir_idx = b0_size - 1; @@ -878,26 +898,55 @@ extern "C" __global__ void experience_action_select( if (dir_idx == 1 || dir_idx == 3) { mag_idx = 0; /* Hold/Flat: magnitude masked, set to 0 for deterministic encoding */ } else { - /* Branch 1: magnitude — BOLTZMANN SOFTMAX selection (variance-neutral). + /* Branch 1: magnitude — training uses BOLTZMANN SOFTMAX; eval uses + * strict argmax + Philox tie-break (Task 2.2 H10 fix). * - * C51's expected Q for magnitude has a structural distributional bias: - * tight return distributions (Small positions) get higher expected Q under - * the C51 softmax, regardless of actual expected returns. Argmax amplifies - * this into an irrecoverable collapse (target network feedback loop). + * Training Boltzmann rationale: C51's expected Q for magnitude has a + * structural distributional bias — tight return distributions (Small + * positions) get higher expected Q under the C51 softmax regardless + * of actual expected returns. Argmax amplifies this into irrecoverable + * collapse (target network feedback loop). Boltzmann samples magnitude + * PROPORTIONALLY to exp(Q/tau) so all magnitudes get chance even when + * Q-gaps are small. tau = Q-range (floor 0.01). IQN as primary loss + * (variance-neutral Huber) lets the model learn state-dependent + * magnitude preferences without collapse. * - * Boltzmann softmax samples magnitude PROPORTIONALLY to exp(Q/tau), - * giving all magnitudes a chance even when Q-gaps are small. Combined - * with IQN as primary training loss (variance-neutral Huber), the model - * can learn state-dependent magnitude preferences without collapse. - * - * Temperature tau = Q-range / 3 (adaptive, prevents division by zero). - * At tau → 0: greedy argmax. At tau → ∞: uniform random. */ + * Eval strict-argmax + tie-break: at eval the Boltzmann floor (0.01) + * was structurally collapsing to Quarter (bin 0) — Q-ranges during + * smoke training stay well below tau floor so softmax was effectively + * uniform, but cumulative sampling with Philox deterministically + * picked the first index over ties once eps_mag=0 (no random + * pre-sampling). Replace with strict argmax honouring learned + * preference, uniform Philox pick only on genuine ties < 1e-6. */ if (!eval_mode && philox_uniform(i, timestep, rng_ctr++) < eps_mag) { int r = (int)(philox_uniform(i, timestep, rng_ctr++) * (float)b1_size); mag_idx = (r >= b1_size) ? b1_size - 1 : r; + } else if (eval_mode) { + /* Eval: strict argmax + uniform-among-tied Philox tie-break. */ + float q_max_m = q_sign * q_b1[0]; + for (int a = 1; a < b1_size; a++) { + float qv = q_sign * q_b1[a]; + if (qv > q_max_m + 1e-6f) q_max_m = qv; + } + int tie_count = 0; + for (int a = 0; a < b1_size; a++) { + float qv = q_sign * q_b1[a]; + if (fabsf(qv - q_max_m) <= 1e-6f) tie_count += 1; + } + int tie_pick = (int)(philox_uniform(i, timestep, rng_ctr++) * (float)tie_count); + if (tie_pick >= tie_count) tie_pick = tie_count - 1; + mag_idx = b1_size - 1; + int seen = 0; + for (int a = 0; a < b1_size; a++) { + float qv = q_sign * q_b1[a]; + if (fabsf(qv - q_max_m) <= 1e-6f) { + if (seen == tie_pick) { mag_idx = a; break; } + seen += 1; + } + } } else { - /* Adaptive temperature: scale by Q-range so Boltzmann is meaningful - * regardless of absolute Q magnitude. Floor at 0.01 to avoid div/0. */ + /* Training: Boltzmann softmax with adaptive temperature, floor + * 0.01 to avoid div/0. */ float q_max_m = q_sign * (q_b1[0]); float q_min_m = q_max_m; for (int a = 1; a < b1_size; a++) { @@ -928,12 +977,36 @@ extern "C" __global__ void experience_action_select( } } } - /* Branch 2: order type — Boltzmann for both eval and training */ + /* Branch 2: order type — training Boltzmann; eval strict-argmax + Philox + * tie-break (Task 2.2 H10 fix). */ if (!eval_mode && philox_uniform(i, timestep, rng_ctr++) < eps_ord) { int r = (int)(philox_uniform(i, timestep, rng_ctr++) * (float)b2_size); a2 = (r >= b2_size) ? b2_size - 1 : r; + } else if (eval_mode) { + /* Eval: strict argmax + uniform-among-tied Philox tie-break. */ + float q_max_ord = q_sign * q_b2[0]; + for (int a = 1; a < b2_size; a++) { + float qv_ord = q_sign * q_b2[a]; + if (qv_ord > q_max_ord + 1e-6f) q_max_ord = qv_ord; + } + int tie_count = 0; + for (int a = 0; a < b2_size; a++) { + float qv_ord = q_sign * q_b2[a]; + if (fabsf(qv_ord - q_max_ord) <= 1e-6f) tie_count += 1; + } + int tie_pick = (int)(philox_uniform(i, timestep, rng_ctr++) * (float)tie_count); + if (tie_pick >= tie_count) tie_pick = tie_count - 1; + a2 = b2_size - 1; + int seen = 0; + for (int a = 0; a < b2_size; a++) { + float qv_ord = q_sign * q_b2[a]; + if (fabsf(qv_ord - q_max_ord) <= 1e-6f) { + if (seen == tie_pick) { a2 = a; break; } + seen += 1; + } + } } else { - /* Boltzmann softmax over order Q-values */ + /* Training: Boltzmann softmax over order Q-values. */ float q_max_ord = q_sign * q_b2[0], q_min_ord = q_max_ord; for (int a = 1; a < b2_size; a++) { float qv_ord = q_sign * q_b2[a]; @@ -961,12 +1034,36 @@ extern "C" __global__ void experience_action_select( } } - /* Branch 3: urgency — Boltzmann for both eval and training */ + /* Branch 3: urgency — training Boltzmann; eval strict-argmax + Philox + * tie-break (Task 2.2 H10 fix). */ if (!eval_mode && philox_uniform(i, timestep, rng_ctr++) < eps_urg) { int r = (int)(philox_uniform(i, timestep, rng_ctr++) * (float)b3_size); a3 = (r >= b3_size) ? b3_size - 1 : r; + } else if (eval_mode) { + /* Eval: strict argmax + uniform-among-tied Philox tie-break. */ + float q_max_urg = q_sign * q_b3[0]; + for (int a = 1; a < b3_size; a++) { + float qv_urg = q_sign * q_b3[a]; + if (qv_urg > q_max_urg + 1e-6f) q_max_urg = qv_urg; + } + int tie_count = 0; + for (int a = 0; a < b3_size; a++) { + float qv_urg = q_sign * q_b3[a]; + if (fabsf(qv_urg - q_max_urg) <= 1e-6f) tie_count += 1; + } + int tie_pick = (int)(philox_uniform(i, timestep, rng_ctr++) * (float)tie_count); + if (tie_pick >= tie_count) tie_pick = tie_count - 1; + a3 = b3_size - 1; + int seen = 0; + for (int a = 0; a < b3_size; a++) { + float qv_urg = q_sign * q_b3[a]; + if (fabsf(qv_urg - q_max_urg) <= 1e-6f) { + if (seen == tie_pick) { a3 = a; break; } + seen += 1; + } + } } else { - /* Boltzmann softmax over urgency Q-values */ + /* Training: Boltzmann softmax over urgency Q-values. */ float q_max_urg = q_sign * q_b3[0], q_min_urg = q_max_urg; for (int a = 1; a < b3_size; a++) { float qv_urg = q_sign * q_b3[a]; diff --git a/crates/ml/src/trainers/dqn/smoke_tests/magnitude_distribution.rs b/crates/ml/src/trainers/dqn/smoke_tests/magnitude_distribution.rs index 6d5b4bd2b..a5d094acd 100644 --- a/crates/ml/src/trainers/dqn/smoke_tests/magnitude_distribution.rs +++ b/crates/ml/src/trainers/dqn/smoke_tests/magnitude_distribution.rs @@ -52,5 +52,23 @@ fn test_magnitude_distribution() -> Result<()> { Quarter={:.3} Half={:.3}", f, q, h ); + + // Task 2.2 H10 regression — eval-mode distribution must include Half + Full. + // Pre-fix baseline (HEAD at 7e78bf4f8) was eval_dist [eq=1.000 eh=0.000 ef=0.000]; + // after the Philox tie-break fix the eval argmax samples tied indices + // uniformly, so eh + ef should approach training-time magnitude spread. + // Threshold 0.30 is the user-specified gate from Task 2.9 (`eval_dist eh + ef ≥ 0.3`). + // Do NOT relax — failure here means the H10 fix alone is insufficient and + // the H9 fallback (delete magnitude branch) may need to promote to primary. + let [eq, eh, ef] = trainer.last_eval_magnitude_dist(); + println!("[EVAL_DIST] Quarter={:.3} Half={:.3} Full={:.3}", eq, eh, ef); + assert!( + eh + ef >= 0.30, + "H10 regression: eval Half + Full share {:.3} < 0.30 (eq={:.3} eh={:.3} ef={:.3}) \ + — argmax tie-break did not prevent collapse. Train-mode dist_mag should show \ + similar spread; if training dist is also low, the H9 fallback (delete magnitude \ + branch) may be needed.", + eh + ef, eq, eh, ef + ); Ok(()) } diff --git a/crates/ml/src/trainers/dqn/trainer/mod.rs b/crates/ml/src/trainers/dqn/trainer/mod.rs index a76abf27f..6d839cb2d 100644 --- a/crates/ml/src/trainers/dqn/trainer/mod.rs +++ b/crates/ml/src/trainers/dqn/trainer/mod.rs @@ -1494,6 +1494,15 @@ impl DQNTrainer { self.last_magnitude_dist } + /// Per-magnitude action distribution from the most recent VALIDATION + /// backtest (eval-mode rollout, distinct from training-mode action_dist). + /// Layout: [Quarter, Half, Full] in [0, 1]. Used by the magnitude_distribution + /// smoke test for the H10 regression assertion (Task 2.2): eval-time + /// argmax over near-tied Q-values must not collapse to bin 0. + pub fn last_eval_magnitude_dist(&self) -> [f32; 3] { + self.last_eval_magnitude_dist + } + /// Per-epoch magnitude-branch action entropy history. /// Each entry is (epoch_0_indexed, normalized_entropy ∈ [0, 1]). /// Used by exploration_coverage smoke test.