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

fix(dqn): unify eval action selection with training Boltzmann softmax

Browse Source 
The eval policy used strict-argmax with an ISV-tied tie-break across all four
factored heads (direction, magnitude, order, urgency). The tie threshold was
`0.01 × isv_signals[V_HALF_*_INDEX]` — i.e. 1% of the C51 atom support range
(~40 for direction). That threshold did not match the actual per-sample
Q-spread (~1.5 once the IQN trunk gradient unstuck), so tie-break never fired
and eval became pure strict-argmax over a peaked Q distribution → val argmax
glued to one direction → 1-25 trades per 214k-bar window across cluster runs
`vg2r9` and `vnwtn`.

Replace with the same Boltzmann softmax training already uses: `tau =
max(q_range, floor)` where `q_range` is computed per-sample. Softmax is
mathematically bounded to `P(best) ≤ 47.5%` for 4 actions with `tau=q_range`,
so eval can never collapse to pure-greedy regardless of how peaked the
Q-values become. State-adaptive without tuned constants — confident states
(large q_range) still favour the best direction near-deterministically;
ambiguous states (q_range at floor) sample uniformly. The Philox stream is
seeded by (i, timestep) so eval remains bit-reproducible across runs at the
same checkpoint.

Three additions:
  1. `experience_kernels.cu`: drop the four `else if (eval_mode)` strict-argmax
     blocks; eval falls through to the existing Boltzmann path. Net -149 lines.
  2. `cuda_pipeline/mod.rs`: add `test_eval_action_select_boltzmann_bounded`,
     a focused unit test that exercises the kernel directly with peaked
     synthetic Q-values and asserts the histogram matches Boltzmann theory
     (P(best) ≈ 0.366, ≤ 0.6, ≥ 0.25). Runs in 1.65s after build, replaces
     15-min smoke runs for kernel-level validation.
  3. `trainers/dqn/trainer/metrics.rs`: log per-direction eval distribution
     (`val_dir_dist [short hold long flat]`) to HEALTH_DIAG. The kernel-side
     `dir_entropy` collapses Hold+Flat into one bucket, masking whether the
     eval policy actually picks one direction or balances Hold/Flat.

Verified: unit test produces histogram short=0.146 hold=0.239 long=0.382
flat=0.233 — matches Boltzmann math, confirms the eval kernel produces
diverse picks for peaked Q-input.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-26 19:28:42 +02:00
parent 326c133782
commit 0a6a615d83
4 changed files with 207 additions and 192 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -841,74 +841,25 @@ extern "C" __global__ void experience_action_select(
(void)min_hold_bars;
(void)portfolio_states;
/* Branch 0: direction — training uses BOLTZMANN SOFTMAX; eval uses strict
* argmax with Philox tie-break (Task 2.2 H10 fix).
/* Branch 0: direction — Boltzmann softmax with adaptive tau (q_range,
* floor 0.01) for both training and eval.
*
* 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.
* State-adaptive without tuned constants: large per-sample q_range
* (confident state) gives large tau → softmax picks the best direction
* near-deterministically; small q_range (ambiguous state) sits at the
* 0.01 floor → softmax samples uniformly across directions. The Philox
* stream is seeded by (i, timestep) so eval is bit-reproducible across
* runs at the same checkpoint.
*
* 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). */
* C51's expected Q structurally favors Flat (zero PnL variance → tightest
* distribution → highest expected Q). Boltzmann samples PROPORTIONALLY to
* exp(Q/tau) so Short/Long retain probability mass even when their Q is
* mildly worse, preserving the diversity training relies on. */
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.
*
* Adaptive tie epsilon (val-Flat-collapse fix, 2026-04-23):
* Prior code used a fixed 1e-6 threshold, which was too tight when
* C51 atom utilisation collapsed (direction Q-values drifted by
* 1e-4 to 1e-5 from numerical precision alone, below the atom
* resolution but above 1e-6). That gave one bin a consistent but
* meaningless edge — typically Flat (bin 3) — producing ~22 val
* trades per epoch while training ran 30-200k trades via
* Boltzmann. Widen the tie threshold to a small fraction of the
* direction branch's Q half-width read from the ISV bus
* (V_HALF_DIR_INDEX=24). Q-differences smaller than 1% of the
* learned Q support are numerically indistinguishable from atom
* jitter → fire the tie-break. The 0.01 fraction is a
* numerical-safety bound per feedback_isv_for_adaptive_bounds.md.
* NULL ISV falls back to 1e-6 (pre-fix behaviour). */
float tie_eps = 1e-6f;
if (isv_signals_ptr != NULL) {
float half_dir = isv_signals_ptr[24]; /* V_HALF_DIR_INDEX */
if (half_dir > 1e-6f) {
tie_eps = fmaxf(1e-6f, half_dir * 0.01f);
}
}
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 + tie_eps) 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) <= tie_eps) 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) <= tie_eps) {
if (seen == tie_pick) { dir_idx = a; break; }
seen += 1;
}
}
} else {
/* Training mode: Boltzmann softmax over direction Q-values. */
/* Boltzmann softmax over direction Q-values (train + eval). */
float q_max_d = q_sign * (q_b0[0]);
float q_min_d = q_max_d;
for (int a = 1; a < b0_size; a++) {
@@ -940,65 +891,25 @@ extern "C" __global__ void experience_action_select(
if (dir_idx == DIR_HOLD || dir_idx == DIR_FLAT) {
mag_idx = 0; /* Hold/Flat: magnitude masked, set to 0 for deterministic encoding */
} else {
/* Branch 1: magnitude — training uses BOLTZMANN SOFTMAX; eval uses
* strict argmax + Philox tie-break (Task 2.2 H10 fix).
/* Branch 1: magnitude — Boltzmann softmax with adaptive tau (q_range,
* floor 0.01) for both training and eval. See Branch 0 above for the
* shared state-adaptive temperature rationale.
*
* 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.
* Magnitude-specific: C51's expected Q has a structural bias toward
* Small (tight return distributions → highest expected Q regardless of
* actual returns). Boltzmann samples PROPORTIONALLY to exp(Q/tau) so
* Half/Full retain probability mass even when their Q is mildly worse.
* IQN as primary loss (variance-neutral Huber) lets state-dependent
* magnitude preferences develop without the collapse argmax produces.
*
* 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. */
* Realised magnitude is still gated by Kelly cap downstream (see
* trade_physics.cuh::apply_kelly_cap); this kernel writes the policy's
* INTENT magnitude, which env_step then caps to the Kelly fraction. */
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.
* Adaptive tie_eps from ISV V_HALF_MAG_INDEX (26) — same
* rationale as direction branch. See the dir block above for
* the full explanation of the val-Flat-collapse fix. */
float tie_eps_m = 1e-6f;
if (isv_signals_ptr != NULL) {
float half_mag = isv_signals_ptr[26];
if (half_mag > 1e-6f) {
tie_eps_m = fmaxf(1e-6f, half_mag * 0.01f);
}
}
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 + tie_eps_m) 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) <= tie_eps_m) 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) <= tie_eps_m) {
if (seen == tie_pick) { mag_idx = a; break; }
seen += 1;
}
}
} else {
/* Training: Boltzmann softmax with adaptive temperature, floor
* 0.01 to avoid div/0. */
/* Boltzmann softmax over magnitude Q-values (train + eval). */
float q_max_m = q_sign * (q_b1[0]);
float q_min_m = q_max_m;
for (int a = 1; a < b1_size; a++) {
@@ -1029,55 +940,25 @@ extern "C" __global__ void experience_action_select(
}
}
}
/* Branch 2: order type — training Boltzmann; eval strict-argmax + Philox
* tie-break (Task 2.2 H10 fix). */
/* Branch 2: order type — Boltzmann softmax with adaptive tau (q_range,
* floor 0.5) for both training and eval.
*
* Order-specific: floor 0.5 (vs 0.01 for direction/magnitude) because
* order-type Q-values start nearly identical (same reward signal across
* Market/Limit/etc.). A 0.01 floor amplifies init noise 100× and locks
* in one order type before counterfactual training can differentiate
* them; 0.5 keeps meaningful stochasticity until genuine Q-gaps develop. */
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.
* Adaptive tie_eps from ISV V_HALF_ORD_INDEX (28). */
float tie_eps_ord = 1e-6f;
if (isv_signals_ptr != NULL) {
float half_ord = isv_signals_ptr[28];
if (half_ord > 1e-6f) {
tie_eps_ord = fmaxf(1e-6f, half_ord * 0.01f);
}
}
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 + tie_eps_ord) 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) <= tie_eps_ord) 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) <= tie_eps_ord) {
if (seen == tie_pick) { a2 = a; break; }
seen += 1;
}
}
} else {
/* Training: Boltzmann softmax over order Q-values. */
/* Boltzmann softmax over order Q-values (train + eval). */
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];
q_max_ord = fmaxf(q_max_ord, qv_ord);
q_min_ord = fminf(q_min_ord, qv_ord);
}
/* Higher temperature floor (0.5 vs 0.01) for order: Q-values for different
* order types start nearly identical (same reward signal). The low 0.01 floor
* amplifies initialization noise 100×, locking in a single order type before
* learning can differentiate them. 0.5 gives meaningful Boltzmann stochasticity
* until the order branch develops genuine Q-gaps from counterfactual training. */
float tau_ord = fmaxf(q_max_ord - q_min_ord, 0.5f);
float sum_e = 0.0f;
float exps_ord[3];
@@ -1094,44 +975,15 @@ extern "C" __global__ void experience_action_select(
}
}
/* Branch 3: urgency — training Boltzmann; eval strict-argmax + Philox
* tie-break (Task 2.2 H10 fix). */
/* Branch 3: urgency — Boltzmann softmax with adaptive tau (q_range,
* floor 0.5) for both training and eval. Same rationale as Branch 2:
* urgency Q-values start nearly identical; the higher floor preserves
* stochasticity until learning differentiates the urgency levels. */
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.
* Adaptive tie_eps from ISV V_HALF_URG_INDEX (30). */
float tie_eps_urg = 1e-6f;
if (isv_signals_ptr != NULL) {
float half_urg = isv_signals_ptr[30];
if (half_urg > 1e-6f) {
tie_eps_urg = fmaxf(1e-6f, half_urg * 0.01f);
}
}
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 + tie_eps_urg) 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) <= tie_eps_urg) 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) <= tie_eps_urg) {
if (seen == tie_pick) { a3 = a; break; }
seen += 1;
}
}
} else {
/* Training: Boltzmann softmax over urgency Q-values. */
/* Boltzmann softmax over urgency Q-values (train + eval). */
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];
@@ -1168,10 +1020,9 @@ extern "C" __global__ void experience_action_select(
/* Pure-policy magnitude intent — diagnostic only, bypasses the Hold/Flat
* dir_idx gate that forces mag_idx=0 and the downstream Kelly/margin
* clamps inside env_step. Computed from q_b1 with strict argmax and the
* same tie-break convention used elsewhere (prefer higher bin index on
* ties within 1e-6). Uses q_sign to stay consistent with contrarian
* mode. Written to out_intent_mag[i] if non-NULL. */
* clamps inside env_step. Computed from q_b1 with strict argmax (prefer
* higher bin index on ties within 1e-6). Uses q_sign to stay consistent
* with contrarian mode. Written to out_intent_mag[i] if non-NULL. */
if (out_intent_mag != NULL) {
int intent = 0;
float q_intent_max = q_sign * q_b1[0];

127
crates/ml/src/cuda_pipeline/mod.rs
View File

@@ -1061,4 +1061,131 @@ mod tests {
);
}
}
/// Direct kernel test: feed `experience_action_select` synthetic peaked Q-values
/// in eval mode (eps=0) and verify the returned direction histogram matches the
/// Boltzmann theoretical bound. With `tau = q_range` and 4 actions linearly
/// spaced in Q, the softmax has `P(best) ≤ 47.5%`. If the realised histogram
/// shows P(best) ≈ 1.0 the eval policy is silently strict-argmax; ≤0.6 confirms
/// Boltzmann is firing at eval.
///
/// Runs in seconds. No DQN training needed. Pulls the actual cubin compiled
/// from the current `experience_kernels.cu`.
#[test]
fn test_eval_action_select_boltzmann_bounded() {
use cudarc::driver::{CudaContext, LaunchConfig, PushKernelArg};
static EXP_CUBIN: &[u8] = include_bytes!(concat!(
env!("OUT_DIR"),
"/experience_kernels.cubin"
));
let ctx = CudaContext::new(0).expect("CUDA required");
let stream = ctx.default_stream();
let module = ctx
.load_cubin(EXP_CUBIN.to_vec())
.expect("load experience_kernels.cubin");
let kernel = module
.load_function("experience_action_select")
.expect("load experience_action_select");
// 4 dirs + 3 mags + 3 ords + 3 urgs = 13 floats per sample.
let b0: i32 = 4;
let b1: i32 = 3;
let b2: i32 = 3;
let b3: i32 = 3;
let q_stride: usize = (b0 + b1 + b2 + b3) as usize;
let batch: usize = 8192;
let dir_best: usize = 2; // Long
// Build Q with direction PEAKED at bin 2 (Long): linearly spaced
// [0.0, 0.5, 1.0, 0.5] — best gap = 0.5, max-min = 1.0 (so tau=1.0).
// Boltzmann theory: P(best=2) = e^0/(e^0 + 2*e^-0.5 + e^-1)
// = 1/(1 + 2*0.6065 + 0.3679) ≈ 0.366.
// Magnitude/order/urgency are uniform (Q=0) so all branches default
// toward random sampling under Boltzmann floor 0.01/0.5.
let mut q_host = vec![0.0_f32; batch * q_stride];
for i in 0..batch {
let base = i * q_stride;
q_host[base + 0] = 0.0; // Short
q_host[base + 1] = 0.5; // Hold
q_host[base + 2] = 1.0; // Long (best)
q_host[base + 3] = 0.5; // Flat
}
let q_buf = stream.memcpy_stod(&q_host).expect("upload q_values");
let mut actions_buf = stream.alloc_zeros::<i32>(batch).expect("alloc actions");
let mut q_gaps_buf = stream.alloc_zeros::<f32>(batch).expect("alloc q_gaps");
let mut intent_buf = stream.alloc_zeros::<i32>(batch).expect("alloc intent");
let mut conv_buf = stream.alloc_zeros::<f32>(batch).expect("alloc conviction");
let blocks = (batch as u32).div_ceil(256);
let cfg = LaunchConfig {
grid_dim: (blocks, 1, 1),
block_dim: (256, 1, 1),
shared_mem_bytes: 0,
};
let n_i32: i32 = batch as i32;
let null_ptr: u64 = 0;
unsafe {
stream
.launch_builder(&kernel)
.arg(&q_buf)
.arg(&mut actions_buf)
.arg(&mut q_gaps_buf)
.arg(&0.0_f32) // eps_start (eval)
.arg(&0.0_f32) // eps_end (eval)
.arg(&0_i32) // current_epoch
.arg(&1_i32) // total_epochs
.arg(&n_i32) // N
.arg(&b0).arg(&b1).arg(&b2).arg(&b3)
.arg(&0.0_f32) // q_gap_threshold
.arg(&null_ptr) // portfolio_states (NULL)
.arg(&0_i32) // min_hold_bars
.arg(&0.0_f32) // max_position
.arg(&1.0_f32) // eps_exp_mult
.arg(&1.0_f32) // eps_ord_mult
.arg(&1.0_f32) // eps_urg_mult
.arg(&0_i32) // timestep
.arg(&null_ptr) // per_sample_epsilon (NULL)
.arg(&null_ptr) // isv_signals_ptr (NULL)
.arg(&0_i32) // contrarian_active
.arg(&mut intent_buf)
.arg(&mut conv_buf)
.launch(cfg)
.expect("launch experience_action_select");
}
stream.synchronize().expect("sync");
let host_actions = stream.memcpy_dtov(&actions_buf).expect("readback actions");
let mut dir_hist = [0_u32; 4];
for &a in &host_actions {
let dir = (a / (b1 * b2 * b3)) as usize;
assert!(dir < 4, "direction {dir} out of range from action {a}");
dir_hist[dir] += 1;
}
let total = batch as f32;
let p_short = dir_hist[0] as f32 / total;
let p_hold = dir_hist[1] as f32 / total;
let p_long = dir_hist[2] as f32 / total;
let p_flat = dir_hist[3] as f32 / total;
eprintln!(
"eval_action_select boltzmann histogram: short={:.3} hold={:.3} long={:.3} flat={:.3} (best=Long expected ≈0.366)",
p_short, p_hold, p_long, p_flat,
);
let p_best = p_long;
// Boltzmann theory: P(best) ≈ 0.366 for Q=[0,0.5,1.0,0.5], tau=1.0.
// Allow ±10% tolerance on a finite-sample (8192) histogram.
// Strict argmax would give P(best)=1.0; degenerate uniform gives 0.25.
// Anything ≥ 0.7 indicates Boltzmann is NOT firing at eval.
assert!(
p_best <= 0.6,
"P(Long) = {p_best:.3} exceeds 0.6 — eval is effectively strict-argmax (Boltzmann not firing). Hist: short={p_short:.3} hold={p_hold:.3} long={p_long:.3} flat={p_flat:.3}",
);
assert!(
p_best >= 0.25,
"P(Long) = {p_best:.3} below 0.25 — Boltzmann should still favour the best bin. Hist: short={p_short:.3} hold={p_hold:.3} long={p_long:.3} flat={p_flat:.3}",
);
}
}

17
crates/ml/src/trainers/dqn/trainer/metrics.rs
View File

@@ -779,6 +779,23 @@ impl DQNTrainer {
"eval intent magnitude dist readback failed (non-fatal): {e}"
),
}
// Per-direction eval distribution — needed to disambiguate eval
// dir_entropy near zero. The kernel-side dir_entropy collapses
// Hold+Flat into one bucket (3 buckets), so a healthy 25/25/25/25
// direction split would still register as ~ln(3) entropy. A flat
// dir_dist confirms the eval policy actually picks one direction;
// a varied dir_dist with low entropy in the 3-bucket metric points
// at the Hold/Flat collapse instead of mono-direction.
let epoch = self.current_epoch;
match ev.read_eval_action_distribution_per_direction() {
Ok(dd) => tracing::info!(
"HEALTH_DIAG[{}]: val_dir_dist [short={:.4} hold={:.4} long={:.4} flat={:.4}]",
epoch, dd[0], dd[1], dd[2], dd[3],
),
Err(e) => tracing::warn!(
"eval direction dist readback failed (non-fatal): {e}"
),
}
}
Ok(-val_sharpe)

20
docs/dqn-wire-up-audit.md
View File

@@ -2,6 +2,26 @@
**Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7.
Eval-mode action selection unified with Boltzmann softmax (2026-04-26):
`experience_kernels.cu` — the four factored action heads (direction, magnitude,
order, urgency) all dropped their `else if (eval_mode)` strict-argmax + ISV-
tied tie-break blocks. The tie threshold was `0.01 × isv_signals[V_HALF_*_INDEX]`
(1% of the C51 atom support range, ~40 for direction); in practice the actual
per-sample direction Q-spread post-IQN is ~1.5, so tie-break never fired and
eval was pure strict-argmax over a peaked distribution → val argmax glued to
one direction → 1-25 trades / 214k-bar window across cluster runs. Replaced
with the same Boltzmann path training already used: `tau = max(q_range, floor)`
where `q_range` is per-sample. Mathematically `P(best) ≤ 47.5%` for 4 actions
with `tau=q_range`, so eval can never collapse to pure-greedy regardless of
how peaked the Q-values become. Confirmed by new unit test
`cuda_pipeline::tests::test_eval_action_select_boltzmann_bounded` — exercises
the kernel directly with peaked synthetic Q-values `[0, 0.5, 1.0, 0.5]` and
asserts the realised histogram matches Boltzmann theory (P(best) ≈ 0.366,
≤ 0.6, ≥ 0.25); the test runs in ~1.6 s after build, replacing 15-min smoke
runs for kernel-level validation. New `val_dir_dist` line in HEALTH_DIAG logs
per-direction eval distribution (short / hold / long / flat) — disambiguates
the kernel-side `dir_entropy` collapse that bucketed Hold+Flat into one slot.
IQN trunk-gradient readiness gate removed (2026-04-26): `gpu_dqn_trainer.rs:6461,6492`
SAXPY scale was `iqn_lambda × iqn_readiness × iqn_budget`. `iqn_readiness` initialises
to 0.0 and only ramps when `iqn_loss_ema` drops below `iqn_loss_initial`, but that