diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu
index cc944ade1..0c4fdde30 100644
--- a/crates/ml/src/cuda_pipeline/experience_kernels.cu
+++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu
@@ -817,6 +817,49 @@ extern "C" __global__ void experience_action_select(
         eps_mag = fmaxf(eps_mag, EPS_FLOOR);
         eps_ord = fmaxf(eps_ord, EPS_FLOOR);
         eps_urg = fmaxf(eps_urg, EPS_FLOOR);
+
+        /* C51-bias breakout: when the policy collapses toward Hold/Flat (no
+         * trade attempts), boost the direction-branch eps floor proportional
+         * to how far below TRADE_TARGET_RATE the current attempt EMA sits.
+         *
+         * Why this is needed (not redundant with the static EPS_FLOOR above):
+         * C51's expected Q makes Flat = δ(0) = 0 attractive against directional
+         * actions whose distributions concentrate slightly below 0 from
+         * tx_cost without compensating edge. Once Q values lock into this
+         * preference (within ~1 epoch of training), Boltzmann correctly
+         * samples Flat/Hold most of the time. A 2% random floor is too small
+         * to break the lock — Long/Short get 0.5% sampling each, generating
+         * insufficient experience for Q-values to discover real edge. The
+         * adaptive boost scales eps_dir up to half (0.5) when the policy is
+         * fully passive (zero trade attempts), preserving learned preferences
+         * while forcing enough exploration to escape the attractor.
+         *
+         * passive_pressure = clamp(0, 1, 1 − attempt_ema/target_rate). At
+         * target rate or above, no boost (passive_pressure=0, eps_dir at
+         * baseline floor). At zero attempts, full boost (eps_dir floor 0.5
+         * = half random sampling). The 0.5 ceiling is a structural blend
+         * point — half random / half policy — not a tuned magnitude; it's
+         * the maximum exploration that still preserves directional Q-signal
+         * propagation through the replay buffer.
+         *
+         * Cold-start safety: until ISV[72] freezes from the measured EMA
+         * (epoch 5 per project_reward), target_raw is 0 and the formula
+         * no-ops (passive_pressure clamps to 0 via the target floor). Once
+         * frozen, the feedback loop activates.
+         *
+         * Direction branch only — magnitude/order/urgency don't have the
+         * Flat-attractor problem (Kelly cap already shapes magnitude
+         * realisation; order/urgency operate on top of an already-decided
+         * directional pick). */
+        if (isv_signals_ptr != NULL) {
+            float attempt_ema = isv_signals_ptr[71];
+            float target_raw  = isv_signals_ptr[72];
+            float target_rate = fmaxf(target_raw, 1e-4f);  /* cold-start floor */
+            float ratio = attempt_ema / target_rate;
+            float passive_pressure = fmaxf(0.0f, fminf(1.0f, 1.0f - ratio));
+            float adaptive_floor = 0.5f * passive_pressure;
+            eps_dir = fmaxf(eps_dir, adaptive_floor);
+        }
     }
 
     int q_stride = b0_size + b1_size + b2_size + b3_size;
diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md
index a4d2fac06..6ee764ff0 100644
--- a/docs/dqn-wire-up-audit.md
+++ b/docs/dqn-wire-up-audit.md
@@ -184,7 +184,8 @@ P5T5 Phase E (2026-04-26): per-fold reset for `MetricBandsRegistry` regression-d
 | C.3 B.1/B.2 KL-amp consumers (`experience_kernels.cu` Flat opp_cost + entering_trade B.2 bonus) | `experience_env_step` kernel: B.1 site multiplies `reward = -shaping_scale × holding_cost_rate × conviction_core × vol_proxy_flat × q_abs_ref × kl_amp_b1`; B.2 site multiplies `bonus_scaled = shaping_scale × bonus × kl_amp_b2`. Both consumers use `fmaxf(1.0, isv_signals_ptr[ISV_STATE_KL_AMP_IDX])` to no-op against cold-start. `ISV_STATE_KL_AMP_IDX=79` macro defined in `state_layout.cuh`. | Wired | Plan 3 Task 7 C.3 — bounded amp ∈ [1, 2] stacks safely with B.1's existing `q_abs_ref` unbounded multiplicand per `pearl_one_unbounded_signal_per_reward.md`. | — |
 | `trainers/dqn/monitors/state_kl_monitor.rs` | Read-only observer for `state_kl_moment_match` kernel output (ISV slots 78 + 79); consumers: HEALTH_DIAG `state_kl.train_val_ema` / `state_kl.amp` + `controller_activity` smoke fire-rate tracking | Wired | Plan 3 Task 7 C.3 | — |
 | `cuda_pipeline/scripted_policy_kernel.cu` | `GpuExperienceCollector::launch_timestep_loop` (per-timestep dispatch when `seed_phase_active_cache=true`); 4 policies (UNIFORM 40% / MOMENTUM 20% / MEAN_REV 20% / VWAP_DEV 20%) deterministically mixed by `i % 5`. Per-thread one sample. Reads `batch_states[i, MARKET_START]` for current close + `portfolio_states[i, PS_PREV_CLOSE]` for prev close; writes `batch_actions[i]` (factored `dir*27 + mag*9 + ord*3 + urg`) and `conviction_buf[i] ∈ [0, 1]` — same outputs the network's `experience_action_select` would have written. Direction-flip thresholds (`±0.0001f` momentum/reversal, `±5bp` VWAP band) are noise-floor cutoffs, not scaling coefficients. | Wired | Plan 3 Task 8 B.3 — GPU-only seeded warm-start. CPU only orchestrates per-epoch dispatch (cold-path read of ISV[SEED_STEPS_DONE/TARGET]); the action computation itself is 100% GPU. No CPU physics mirror — existing `experience_env_step` runs unchanged. | — |
-| C51 bias C.1 ISV-adaptive direction-Boltzmann tau floor (`experience_kernels.cu` `experience_action_select` direction branch) | `experience_action_select` Branch 0 direction softmax: `tau_d = max(q_range, max(ISV[Q_DIR_ABS_REF_INDEX=21], 0.01))`. Replaces static `0.01` floor (tuned constant, dangerous when Q-magnitudes grow during training: spread that's small in absolute terms becomes deterministic argmax even though it represents no real edge relative to scale). Now scales with the network's actual Q magnitude EMA, preserving relative spread for sampling and preventing the C51 expected-Q Hold/Flat attractor that the val-Flat-collapse Kelly fix exposed (val_dir_dist S=.23/H=.17/L=.42/F=.18 epoch 1 → S=.14/H=.35/L=.15/F=.37 epoch 2 = 35→72% Hold+Flat shift in one epoch). Cold-start fallback retains 0.01 minimum so kernel doesn't divide by zero pre-first-update. Same ISV[21] reference used by conviction below — coherent adaptive mechanism. | Wired | val-collapse follow-up — C51 expected-Q bias fix. ISV-driven, no tuned constants per `feedback_isv_for_adaptive_bounds.md` and `feedback_adaptive_not_tuned.md`. | — |
+| C51 bias C.1 ISV-adaptive direction-Boltzmann tau floor (`experience_kernels.cu` `experience_action_select` direction branch) | `experience_action_select` Branch 0 direction softmax: `tau_d = max(q_range, max(ISV[Q_DIR_ABS_REF_INDEX=21], 0.01))`. Replaces static `0.01` floor (tuned constant, dangerous when Q-magnitudes grow during training: spread that's small in absolute terms becomes deterministic argmax even though it represents no real edge relative to scale). Now scales with the network's actual Q magnitude EMA, preserving relative spread for sampling and preventing the C51 expected-Q Hold/Flat attractor that the val-Flat-collapse Kelly fix exposed (val_dir_dist S=.23/H=.17/L=.42/F=.18 epoch 1 → S=.14/H=.35/L=.15/F=.37 epoch 2 = 35→72% Hold+Flat shift in one epoch). Cold-start fallback retains 0.01 minimum so kernel doesn't divide by zero pre-first-update. Same ISV[21] reference used by conviction below — coherent adaptive mechanism. NOTE: empirical verification (train-bscl2) showed this tau-floor change had no measurable effect on the post-training Hold/Flat collapse — once Q-values reflect tx_cost-driven aversion, sampling stays biased regardless of tau. Retained as cold-start protection; bias breakout handled by the eps_dir adaptive floor below. | Wired | val-collapse follow-up — C51 expected-Q bias fix (cold-start protection only). ISV-driven, no tuned constants per `feedback_isv_for_adaptive_bounds.md` and `feedback_adaptive_not_tuned.md`. | — |
+| C51 bias C.2 trade-attempt-rate-driven adaptive eps_dir floor (`experience_kernels.cu` `experience_action_select` eps floor block) | `experience_action_select` Branch 0 direction floor: `eps_dir = max(eps_dir, 0.5 × passive_pressure)` where `passive_pressure = clamp(0, 1, 1 − ISV[71]/max(ISV[72], 1e-4))`. Reads `TRADE_ATTEMPT_RATE_EMA_INDEX=71` (current Flat→Positioned rate) and `TRADE_TARGET_RATE_INDEX=72` (target rate frozen at epoch 5 from measured EMA). When the policy collapses to passive (zero trade attempts), eps_dir floor rises to 0.5 — half random direction sampling — forcing enough Long/Short experiences into the replay buffer for Q-values to discover real edge. At target rate or above, passive_pressure=0, eps_dir at baseline 0.02 floor. Direction branch only: magnitude/order/urgency don't have the Flat-attractor problem (Kelly cap shapes magnitude realisation; order/urgency operate on top of already-decided directional picks). The 0.5 ceiling is a structural blend point (half random / half policy), not a tuned magnitude — maximum exploration that still preserves directional Q-signal propagation through the replay buffer. Active in training only (eval_mode skips eps logic entirely so val remains deterministic). | Wired | val-collapse follow-up — C51 expected-Q bias breakout. ISV-driven feedback control via existing trade-rate EMA (B.2 producer); no new ISV slots. | — |
 | `cuda_pipeline/seed_step_counter_update_kernel.cu` | `GpuExperienceCollector::launch_seed_step_counter_update_inplace` → `training_loop.rs` (called alongside the other Plan 3 EMA producers each epoch); single-block single-thread cold-path. Increments ISV[SEED_STEPS_DONE_INDEX=83] by `n_samples`, capped at ISV[SEED_STEPS_TARGET_INDEX=82], and EMAs the derived `max(0, 1 - DONE/TARGET)` into ISV[SEED_FRAC_EMA_INDEX=84]. Adaptive α=α_base × (1+0.5×\|clamp(sharpe,−2,2)\|), α_base=0.05. | Wired | Plan 3 Task 8 B.3 — cold-path (per-epoch). No atomicAdd. SEED_FRAC_EMA cold-start 1.0 ensures Task 9's CQL ramp sees `target=0` until the seed phase actually decays. | — |
 | `cuda_pipeline/cql_alpha_seed_update_kernel.cu` | `GpuExperienceCollector::launch_cql_alpha_seed_update_inplace` → `training_loop.rs` (called alongside `launch_seed_step_counter_update_inplace` each epoch); single-block single-thread cold-path. EMAs ISV[CQL_ALPHA_INDEX=48] toward `cql_alpha_final × max(0, 1 - ISV[SEED_FRAC_EMA_INDEX=84])` where `cql_alpha_final = config.cql_alpha`. Adaptive α matches Task 8 convention. | Wired | Plan 3 Task 9 C.5 — producer upgrade for ISV[CQL_ALPHA_INDEX=48]: SchemaContract → FoldReset. CQL gradient kernel consumer (already reading slot 48 per Plan 1 Task 12) unchanged. | — |
 | `trainers/dqn/monitors/seed_monitor.rs` | Read-only observer for `seed_step_counter_update` kernel output (ISV slots 82 / 83 / 84); consumers: HEALTH_DIAG `seed.steps_target` / `seed.steps_done` / `seed.frac_ema` + `controller_activity` smoke fire-rate tracking | Wired | Plan 3 Task 8 B.3 | — |