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