gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

fix(dqn): H10 — stable argmax tie-break at eval per Track 1 triage + Task 2.0 re-diagnosis

Browse Source 
Replaces eval-mode Boltzmann softmax with strict argmax + uniform-sample-
among-tied-indices (|q_a − q_b| < 1e-6). Applied to all 4 branches
(direction, magnitude, order, urgency) of experience_action_select.
Uses the existing Philox state (same (i, timestep) seed used elsewhere
in the kernel for CF-flip / exploration); eval mode is therefore
deterministic per (sample, epoch) — no new atomics, no new RNG.

Training mode keeps Boltzmann softmax unchanged (needed for exploration
+ gradient flow when C51 expected-Q structurally favors Flat/Quarter).

Root-cause re-diagnosis (commit 7e78bf4f8 + 41b0c559c): Task 2.0
instrumentation + bug fixes revealed Track 1 H4 was a measurement
artefact produced by two silent bugs in per_branch_grad_norms. Real
mag/dir gradient ratio is 50-400× (magnitude over-fed, NOT starved).
The genuine observable problem is H10 — argmax tie-break at eval,
which is what this commit closes.

Why the previous eval-mode Boltzmann collapsed: with Q-range ≪ tau
floor (0.01), softmax was effectively uniform at eval, but cumulative
sampling with a single Philox draw deterministically picked the first
index over every tied bin — hence eval_dist [eq=1.000 eh=0.000 ef=0.000]
across all 20 baseline epochs despite training-mode ent_mag ≈ 0.99.
Strict argmax honours the learned preference where it exists, tie-break
only applies on genuine ties < 1e-6.

Results from 20-epoch smoke (magnitude_distribution):

  Baseline (HEAD pre-fix):
    training dist [dist_q=0.577 dist_h=0.166 dist_f=0.257]
    eval     dist [eq=1.000 eh=0.000 ef=0.000]

  Post-fix (this commit):
    training dist [dist_q=0.577 dist_h=0.166 dist_f=0.257]  (unchanged — eval-only path)
    eval     dist [eq=0.580 eh=0.420 ef=0.000]

  eh + ef = 0.420 ≥ 0.30 regression gate → PASS.

Training-mode distribution is bit-identical as expected: the fix
only touches the eval_mode == 1 branch of experience_action_select.

Regression assertion in magnitude_distribution smoke requires
eh + ef ≥ 0.30 (the production threshold from Task 2.9 §2.9 gate).

Closes Track 1 H10 CONFIRMED verdict. With H4 REJECTED by Task 2.0's
real data, H10 is now Phase 2's primary fix. If future training
regresses this ratio, the fallback is the "delete magnitude branch"
path (H9 spec §5.1 proposed fix — 108 → 36 action space).
This commit is contained in:
jgrusewski
2026-04-22 11:24:41 +02:00
parent d1068de2b8
commit 8aef59f735
3 changed files with 162 additions and 38 deletions
Download Patch File Download Diff File Expand all files Collapse all files

173
crates/ml/src/cuda_pipeline/experience_kernels.cu
View File

@@ -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];

18
crates/ml/src/trainers/dqn/smoke_tests/magnitude_distribution.rs
View File

@@ -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(())
}

9
crates/ml/src/trainers/dqn/trainer/mod.rs
View File

@@ -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.