diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 7370308b4..60b1ca9d7 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -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]; diff --git a/crates/ml/src/cuda_pipeline/mod.rs b/crates/ml/src/cuda_pipeline/mod.rs index 715292889..d4052fe46 100644 --- a/crates/ml/src/cuda_pipeline/mod.rs +++ b/crates/ml/src/cuda_pipeline/mod.rs @@ -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::(batch).expect("alloc actions"); + let mut q_gaps_buf = stream.alloc_zeros::(batch).expect("alloc q_gaps"); + let mut intent_buf = stream.alloc_zeros::(batch).expect("alloc intent"); + let mut conv_buf = stream.alloc_zeros::(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}", + ); + } } diff --git a/crates/ml/src/trainers/dqn/trainer/metrics.rs b/crates/ml/src/trainers/dqn/trainer/metrics.rs index acec40720..31dafe38c 100644 --- a/crates/ml/src/trainers/dqn/trainer/metrics.rs +++ b/crates/ml/src/trainers/dqn/trainer/metrics.rs @@ -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) diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index e99d11d37..5bd1174c7 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -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