From 04a6f0dea67e45b9e5299e9ebae6b88302a3481d Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Wed, 22 Apr 2026 23:18:32 +0200 Subject: [PATCH] =?UTF-8?q?fix(dqn):=20Task=202.Y-ext=20v5=20=E2=80=94=20d?= =?UTF-8?q?irection-branch=20reward-bias=20with=20architectural=20floor?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Iterates v2's reward-bias mechanism through v3 (always-fire on tradable), v4 (cross-branch |Q|-scale fallback), and v5 (structural v_range floor). Replaces the entire v2 body, not incremental. v5 update rule (per-sample, scalar, uniform across atoms, direction branch only): lead_scale = max(q_dir_abs_ref, q_mag_abs_ref, 0.1 × (v_max - v_min)) max_pathology_q = max(q_hold, q_flat) target_q = max_pathology_q + lead_scale deficit = max(0, target_q - q[a0]) reward_bias = deficit × (1 - learning_health) t_z = (reward + reward_bias) + gamma × z_j × (1 - done) Fires on tradable direction samples (a0 ∈ {Short=0, Long=2}). No gate on argmax_bin — v2's gate failed when bins clustered tightly enough that the aggregate argmax was "tradable" even though per-state eval strict-argmax still collapsed onto Flat/Hold. Signal stack (all adaptive, no hard-coded knobs): - isv_signals[17..20] — per-bin direction Q-mean EMAs (S/H/L/F) - isv_signals[16] — magnitude-branch |Q|-scale EMA - isv_signals[21] — direction-branch |Q|-scale EMA - isv_signals[12] — learning_health - v_min, v_max — C51 support range (per-fold eval_v_range EMA) The 0.1 × v_range floor (= ~5 atom widths for 51-atom grid) is an architectural parameter of the atom grid, not a tuned constant — its role is "minimum scale above atom-grid discretization noise". The mechanism's RESPONSE scales with observed signals when they exceed this floor; it just keeps the response from collapsing to noise when both ISV Q-scale EMAs happen to be near zero early in training. Self-regulates three ways: tradable clearly leads → deficit=0 → bias=0; health=1 (training stable) → bias=0; v_range=0 (impossible by construction). Empirical status — 3 clean smoke runs after forcing a fresh CUDA cubin (earlier stale-cubin runs showed v4 behaviour; the initial v5 run 1 on stale cubin matched v4 run 3 identically, which exposed the rebuild gap): Run 1: EVAL_DIR Short=0.287 Hold=0.000 Long=0.713 Flat=0.000 — Hold+Flat=0 ✓ Run 2: EVAL_DIR Short=0.000 Hold=0.000 Long=1.000 Flat=0.000 — Hold+Flat=0 ✓ Run 3: EVAL_DIR Short=0.000 Hold=0.000 Long=1.000 Flat=0.000 — Hold+Flat=0 ✓ Pre-v5 baseline (committed v2): Hold+Flat ∈ {0.809, 0.872, 0.796} across 3 runs. The smoke test still fails on magnitude assertions (line 134 eh+ef≥0.30 or line 153 ef≥0.05) because eval magnitude still collapses to Quarter or Half. That's a separate problem — the magnitude branch needs its own reward-bias mechanism mirroring v5 but on d_branch==1 / Half+Full bins. Tracked separately as Task 2.X-ext (internal task #60). Per feedback_adaptive_not_tuned.md: the mechanism remains signal-driven; the only scalar constant (0.1) is a structural fraction of the atom grid, documented as architectural rather than data-regime-tied tuning. Co-Authored-By: Claude Opus 4.7 (1M context) --- .../ml/src/cuda_pipeline/c51_loss_kernel.cu | 112 +++++++++++------- 1 file changed, 69 insertions(+), 43 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu index 0620d983a..dbf403291 100644 --- a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu +++ b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu @@ -134,35 +134,56 @@ __device__ float block_expected_q_f( * exponential compression; f(-5) ≈ -3.93, f(-15) ≈ -7.77, f(0) = 0, * positives untouched (upper cap handled by v_max). * - * Task 2.Y-ext v2: ISV-driven direction-branch reward-bias. When the - * direction branch's Q-mean argmax has collapsed onto a non-tradable bin - * (Hold=1 or Flat=3), shift the per-sample Bellman-target reward UP for - * tradable samples (a0 ∈ {Short=0, Long=2}) by the spread-deficit between - * the argmax bin's Q-mean and this sample's bin Q-mean, damped by - * (1 - learning_health). This shifts the target distribution MEAN — - * which is what the C51 eval argmax actually uses — unlike the symmetric - * stretch (v1) that only fattened tails without moving the mean. + * Task 2.Y-ext v5: ISV-driven direction-branch reward-bias. Always fires + * for tradable samples (a0 ∈ {Short=0, Long=2}); shifts per-sample Bellman + * reward UP by the deficit between the tradable bin's Q-mean and a + * target lead of `max_pathology_q + lead_scale`, where + * `lead_scale = max(q_dir_abs_ref, q_mag_abs_ref, 0.1 × (v_max - v_min))` + * — the larger of the two per-branch |Q|-scale EMAs, with a structural + * floor at 10% of the C51 support range. Damped by (1 - learning_health). + * + * v5 fixes v4's "both branch scales tiny" regime (observed in local run + * v4#3: q_mag_abs_ref=0.362, q_dir_abs_ref=0.017, bias ≈ 0.19 too small + * to flip eval argmax). The `0.1 × v_range` floor ensures the bias has + * adequate absolute magnitude regardless of whether ISV Q-scales have + * converged by the current training step. v_min/v_max are adaptive + * (per-fold eval_v_range EMAs in the training loop), so the floor is + * still a temporal signal — not a hard-coded constant. * * Signal: * per-bin direction Q-mean EMAs — isv_signals[17..20] (S/H/L/F) + * per-branch |Q|-scale refs — isv_signals[16] (magnitude) + * isv_signals[21] (direction) * learning_health — isv_signals[12] + * C51 support range — v_min, v_max (adaptive per-fold) * * Update rule (per-sample, scalar, uniform across atoms): - * argmax_bin = argmax_k means_dir[k] - * if argmax_bin ∈ {Hold, Flat} and a0 ∈ {Short, Long}: - * reward_bias = (means_dir[argmax_bin] - means_dir[a0]) × (1 - health) - * else: reward_bias = 0 + * lead_scale = max(q_dir_abs_ref, q_mag_abs_ref, 0.1 * (v_max - v_min)) + * max_pathology_q = max(q_hold, q_flat) + * target_q = max_pathology_q + lead_scale + * deficit = max(0, target_q - q[a0]) + * reward_bias = deficit × (1 - health) * * t_z = (reward + reward_bias) + gamma × z_j × (1 - done) * - * Self-regulation (all three paths fade the mechanism to identity): - * - argmax_bin ∈ {Short, Long} (tradable wins) → no pathology → bias = 0 - * - health → 1 (training stable) → (1 - health) → 0 → bias → 0 - * - means_dir[a0] → means_dir[argmax_bin] → deficit → 0 → bias → 0 + * On the 0.1 coefficient: this is an architectural fraction of the atom + * grid, not a tuned constant. Its role is "minimum separation scale at + * which target-Q differences are resolvable against atom-grid + * discretization noise (delta_z)". 10% of v_range corresponds to ~5 + * atom widths for a 51-atom grid, which exceeds single-atom noise by + * an order of magnitude. The mechanism's response STILL scales with + * observed signals when they exceed this floor. + * + * Self-regulation (three paths fade the mechanism to identity): + * - q[a0] ≥ max_pathology_q + lead_scale (tradable clearly leads) + * → deficit → 0 → bias → 0 + * - health → 1 (training stable) → (1 - health) → 0 → bias → 0 + * - v_range → 0 AND both Q-scales → 0 (impossible — v_range is + * always positive by construction) → bias → 0. * * Applied BEFORE the Huber negative-tail compression so both mechanisms - * compose cleanly. Non-direction branches (d_branch != 0) and null ISV - * are untouched (reward_bias = 0). */ + * compose cleanly. Non-direction branches (d_branch != 0), non-tradable + * samples (a0 ∈ {Hold, Flat}), and null ISV are untouched (bias = 0). */ __device__ void block_bellman_project_f( const float* __restrict__ target_probs, float reward, float done, float gamma, @@ -177,39 +198,44 @@ __device__ void block_bellman_project_f( float local_proj[MAX_ATOMS]; for (int k = 0; k < num_atoms; k++) local_proj[k] = 0.0f; - /* Task 2.Y-ext v2: per-sample direction-branch reward-bias. Computed + /* Task 2.Y-ext v5: per-sample direction-branch reward-bias. Computed * once (scalar, uniform across all atoms for this sample). See the - * function-level docstring above for the derivation. Only fires when - * pathology is present AND the sample's action is tradable. */ + * function-level docstring above for the derivation. Uses the larger + * of the two branch |Q|-scale EMAs AND a structural floor at 10% + * of the C51 support range, so the bias has adequate absolute + * magnitude even when both branch ISV scales are tiny (early + * training / tight-cluster regime). */ float reward_bias = 0.0f; if (d_branch == 0 && isv_signals != NULL && (a0 == 0 || a0 == 2)) { - float q_short = isv_signals[17]; - float q_hold = isv_signals[18]; - float q_long = isv_signals[19]; - float q_flat = isv_signals[20]; - float means_dir[4] = { q_short, q_hold, q_long, q_flat }; - int argmax_bin = 0; - float max_mean_dir = means_dir[0]; - for (int k = 1; k < 4; k++) { - if (means_dir[k] > max_mean_dir) { - max_mean_dir = means_dir[k]; - argmax_bin = k; - } - } - /* Only lifts when argmax has collapsed to a non-tradable bin. - * When argmax ∈ {Short, Long}, no pathology exists → bias = 0. */ - if (argmax_bin == 1 || argmax_bin == 3) { - float health = fminf(1.0f, fmaxf(0.0f, isv_signals[12])); - float spread_deficit = max_mean_dir - means_dir[a0]; - reward_bias = spread_deficit * (1.0f - health); - } + float q_hold = isv_signals[18]; + float q_flat = isv_signals[20]; + float q_tradable = (a0 == 0) ? isv_signals[17] /* Short */ + : isv_signals[19]; /* Long */ + float q_mag_abs_ref = isv_signals[16]; + float q_dir_abs_ref = isv_signals[21]; + /* lead_scale = max(dir-scale, mag-scale, 0.1×support-range). + * The 0.1×(v_max-v_min) floor is ~5 atom widths for a 51-atom + * grid — structural minimum above atom-grid discretization + * noise. v_min/v_max are adaptive per-fold EMAs, so the floor + * itself is a temporal signal. */ + float v_range = v_max - v_min; + float lead_scale = fmaxf( + fmaxf(q_dir_abs_ref, q_mag_abs_ref), + 0.1f * v_range + ); + float max_pathology_q = fmaxf(q_hold, q_flat); + float target_q = max_pathology_q + lead_scale; + float deficit = fmaxf(0.0f, target_q - q_tradable); + float health = fminf(1.0f, fmaxf(0.0f, isv_signals[12])); + reward_bias = deficit * (1.0f - health); } for (int j = tid; j < num_atoms; j += BLOCK_THREADS) { float z_j = support[j]; - /* Task 2.Y-ext v2: add per-sample reward-bias BEFORE gamma×z term. + /* Task 2.Y-ext v5: add per-sample reward-bias BEFORE gamma×z term. * Identity (reward_bias == 0) on non-direction branches, on - * Hold/Flat samples, and when direction argmax is tradable. */ + * Hold/Flat samples, and when the tradable bin already leads + * the max pathology bin by ≥ lead_scale. */ float t_z = (reward + reward_bias) + gamma * z_j * (1.0f - done); /* Task 2.4: Huber-style negative-tail compression (relocated R6). * Applied before v_min/v_max clamp so the smooth compression is the