diff --git a/crates/ml-core/src/state_layout.rs b/crates/ml-core/src/state_layout.rs
index c3fbaeb9d..b0c1c33db 100644
--- a/crates/ml-core/src/state_layout.rs
+++ b/crates/ml-core/src/state_layout.rs
@@ -96,7 +96,8 @@ pub const PS_PLAN_CONVICTION:       usize = 27; // plan: conviction at entry [0,
 pub const PS_PLAN_ASYMMETRY:        usize = 28; // plan: profit/stop asymmetry
 pub const PS_PLAN_ENTRY_REGIME:     usize = 29; // plan: regime_stability at entry
 pub const PS_OFI_PREV_BASE:         usize = 30; // first OFI prev-bar slot (stride 8: [30..37])
-pub const PS_STRIDE:                usize = 38; // total portfolio state stride (== PORTFOLIO_STRIDE)
+pub const PS_PEAK_PNL_BAR:          usize = 38; // Plan 3 C.4: hold_time at which MAX_PNL was hit (temporal timing bonus)
+pub const PS_STRIDE:                usize = 39; // total portfolio state stride (== PORTFOLIO_STRIDE), grown 38→39 by Plan 3 C.4
 
 // ────────────────────────────────────────────────────────────────────────────
 // Mirrors of PLAN_ISV_* constants (plan_isv[0..PLAN_ISV_DIM=7)).
diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu
index df1da365e..d2abd9037 100644
--- a/crates/ml/src/cuda_pipeline/experience_kernels.cu
+++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu
@@ -20,7 +20,7 @@
  *   [market_dim .. market_dim+3)           : portfolio features (value_norm, position, cash_norm)
  *   [market_dim+3 .. state_dim)            : zero-pad for tensor-core alignment
  *
- * Portfolio state layout used by experience kernels ([N, PORTFOLIO_STRIDE=38]):
+ * Portfolio state layout used by experience kernels ([N, PORTFOLIO_STRIDE=39]):
  *   [0]  position        — current contract position (signed)
  *   [1]  cash            — cash balance
  *   [2]  portfolio_value — mark-to-market total value (cash + position * price)
@@ -59,10 +59,10 @@
  */
 
 /* ------------------------------------------------------------------ */
-/* Portfolio stride for experience kernels (38 floats per episode).    */
+/* Portfolio stride for experience kernels (39 floats per episode).    */
 /* portfolio_sim_kernel uses its own stride of 8 — do NOT change it.  */
 /* ------------------------------------------------------------------ */
-#define PORTFOLIO_STRIDE 38  /* was 30. Slots 30-37: prev-bar OFI for delta computation */
+#define PORTFOLIO_STRIDE 39  /* was 38. Slots 30-37: prev-bar OFI; slot 38: PS_PEAK_PNL_BAR (Plan 3 C.4) */
 #define DSR_A_SLOT 3
 #define DSR_B_SLOT 4
 #define DSR_TRADE_COUNT_SLOT 5
@@ -1292,7 +1292,7 @@ extern "C" __global__ void experience_action_select(
  *   [2] raw_close      — raw close price (for position cost + tx)
  *   [3] raw_next       — raw next-bar close (UNUSED in reward path)
  *
- * portfolio_states layout: [N, PORTFOLIO_STRIDE=38]  (read-write, float)
+ * portfolio_states layout: [N, PORTFOLIO_STRIDE=39]  (read-write, float)
  *   See file header for field definitions.
  */
 extern "C" __global__ void experience_env_step(
@@ -1596,7 +1596,7 @@ extern "C" __global__ void experience_env_step(
     /* Guard against degenerate prices from data gaps. */
     if (raw_close <= 0.0f) raw_close = 1.0f;
 
-    /* ---- Read full portfolio state (PORTFOLIO_STRIDE=38) ---- */
+    /* ---- Read full portfolio state (PORTFOLIO_STRIDE=39) ---- */
     /* Portfolio arithmetic uses float accumulators because trade physics
      * functions (execute_trade, apply_margin_cap, etc.) in trade_physics.cuh
      * are all float. Converting the entire physics engine to bf16 would
@@ -1853,6 +1853,7 @@ extern "C" __global__ void experience_env_step(
         /* Store entry regime_stability for drift detection (P12). */
         ps[PS_PLAN_ENTRY_REGIME] = (isv_signals_ptr != NULL) ? isv_signals_ptr[11] : 1.0f;
         ps[PS_INTRA_TRADE_MAX_PNL] = 0.0f;  /* reset max P&L for new trade */
+        ps[PS_PEAK_PNL_BAR]        = 0.0f;  /* Plan 3 C.4: reset peak-bar alongside MAX_PNL */
         /* Apply conviction to position size — scale by readiness for smooth ramp */
         position *= fmaxf(ps[PS_PLAN_CONVICTION] * readiness, 0.1f);
     }
@@ -2010,10 +2011,12 @@ extern "C" __global__ void experience_env_step(
     if (entering_trade) {
         intra_trade_max_dd = 0.0f;
         ps[PS_INTRA_TRADE_MAX_PNL] = 0.0f;  /* reset max P&L for new trade */
+        ps[PS_PEAK_PNL_BAR]        = 0.0f;  /* Plan 3 C.4: reset peak-bar alongside MAX_PNL */
     }
     if (reversing_trade) {
         intra_trade_max_dd = 0.0f;
         ps[PS_INTRA_TRADE_MAX_PNL] = 0.0f;  /* reset max P&L for reversed trade */
+        ps[PS_PEAK_PNL_BAR]        = 0.0f;  /* Plan 3 C.4: reset peak-bar alongside MAX_PNL */
     }
     if (fabsf(position) > 0.001f && entry_price > 0.0f) {
         float unrealized = (position > 0.0f)
@@ -2025,7 +2028,15 @@ extern "C" __global__ void experience_env_step(
         /* Track best unrealized P&L during trade (for hindsight plan labels) */
         float pnl_pct = (raw_close - entry_price) / fmaxf(fabsf(entry_price), 1.0f)
                          * ((position > 0.0f) ? 1.0f : -1.0f);
-        ps[PS_INTRA_TRADE_MAX_PNL] = fmaxf(ps[PS_INTRA_TRADE_MAX_PNL], pnl_pct);  /* max P&L during trade */
+        /* Plan 3 C.4: snapshot bar-of-peak alongside MAX_PNL so the temporal
+         * timing bonus at trade exit can measure bars_early = hold_time - peak_bar.
+         * Uses the local `hold_time` (already advanced by unified_env_step_core)
+         * — `ps[PS_HOLD_TIME]` isn't rewritten until the portfolio-state
+         * commit block further down. */
+        if (pnl_pct > ps[PS_INTRA_TRADE_MAX_PNL]) {
+            ps[PS_INTRA_TRADE_MAX_PNL] = pnl_pct;
+            ps[PS_PEAK_PNL_BAR]        = hold_time;
+        }
     }
 
     /* Idle counter (flat bars) — used for state features, no penalty */
@@ -2135,6 +2146,35 @@ extern "C" __global__ void experience_env_step(
          * before execute_trade) with health-coupled safety multiplier.
          * The network can't over-lever because the env won't let it —
          * not because the reward penalizes deviation from a formula. */
+
+        /* ── C.4 (Plan 3 Task 5): Temporal timing bonus ──
+         *
+         * Fires on trade exit only (inside segment_complete). Rewards exiting
+         * near the unrealized-PnL peak. `bars_early / hold_time` ∈ [0, 1] is
+         * 1 when the peak was at entry (degenerate — limit case only) and 0
+         * when the peak coincides with the exit bar. Scaled by |final_pnl|
+         * (trade outcome magnitude), conviction_core (the same self-scaling
+         * coefficient used by B.1/B.2), and the global shaping_scale gate —
+         * no tuned constants.
+         *
+         * Attribution: ACCUMULATES into rc[5] via `+=` rather than overwrite,
+         * because Task 3 B.2 may also touch rc[5] on entry. The two sites
+         * fire at different (i, t) slots (entry vs exit), so the accumulate
+         * semantics are defensively idempotent. `reward` holds the trade
+         * P&L at this point (Layer 2 result + Layer 4-9 credits), matching
+         * the comment at the "When segment_complete fired above" block. */
+        if (segment_hold_time > 0.0f) {
+            const float peak_bar   = ps[PS_PEAK_PNL_BAR];
+            const float bars_early = fmaxf(0.0f, segment_hold_time - peak_bar);
+            const float final_pnl  = reward;
+            const float timing_bonus = shaping_scale
+                                     * (bars_early / segment_hold_time)
+                                     * fabsf(final_pnl)
+                                     * conviction_core;
+            reward += timing_bonus;
+            /* C.2 attribution — accumulate into rc[5] bonus slot (see above). */
+            reward_components_per_sample[out_off * 6 + 5] += timing_bonus;
+        }
     }
 
     /* ── v10: Trade-level reward attribution ──────────────────────────────
@@ -2431,7 +2471,7 @@ extern "C" __global__ void experience_env_step(
     int data_end       = (next_bar >= total_bars) ? 1 : 0;
     int done           = trade_complete | capital_breach | data_end;
 
-    /* ---- Update full portfolio state (PORTFOLIO_STRIDE=38) ---- */
+    /* ---- Update full portfolio state (PORTFOLIO_STRIDE=39) ---- */
     ps[PS_POSITION]  = (position);
     ps[PS_CASH]  = (cash);
     ps[PS_PORTFOLIO_VALUE]  = (new_portfolio_value);
@@ -2725,6 +2765,7 @@ extern "C" __global__ void experience_env_step(
             ps[PS_KELLY_SUM_SQ_RETURNS] = 0.0f;
             ps[PS_INTRA_TRADE_MAX_DD] = 0.0f;           /* intra_trade_max_dd */
             ps[PS_INTRA_TRADE_MAX_PNL] = 0.0f;           /* intra_trade_max_pnl */
+            ps[PS_PEAK_PNL_BAR] = 0.0f;                  /* Plan 3 C.4: peak-bar alongside MAX_PNL */
             ps[PS_PLAN_TARGET_BARS] = 0.0f; ps[PS_PLAN_PROFIT_TARGET] = 0.0f; ps[PS_PLAN_STOP_LOSS] = 0.0f;
             ps[PS_PLAN_SCALE_AGGRESSION] = 0.0f; ps[PS_PLAN_CONVICTION] = 0.0f; ps[PS_PLAN_ASYMMETRY] = 0.0f;
             ps[PS_PLAN_ENTRY_REGIME] = 0.0f;           /* plan slots zeroed */
@@ -2739,6 +2780,7 @@ extern "C" __global__ void experience_env_step(
             ps[PS_TRADE_START_PNL] = (ps[PS_REALIZED_PNL]);       /* trade_start_pnl = current realized_pnl */
             ps[PS_INTRA_TRADE_MAX_DD] = 0.0f;           /* intra_trade_max_dd — reset for next trade */
             ps[PS_INTRA_TRADE_MAX_PNL] = 0.0f;           /* intra_trade_max_pnl — reset for next trade */
+            ps[PS_PEAK_PNL_BAR] = 0.0f;                  /* Plan 3 C.4: peak-bar reset for next trade */
             ps[PS_PLAN_TARGET_BARS] = 0.0f; ps[PS_PLAN_PROFIT_TARGET] = 0.0f; ps[PS_PLAN_STOP_LOSS] = 0.0f;
             ps[PS_PLAN_SCALE_AGGRESSION] = 0.0f; ps[PS_PLAN_CONVICTION] = 0.0f; ps[PS_PLAN_ASYMMETRY] = 0.0f;
             ps[PS_PLAN_ENTRY_REGIME] = 0.0f;           /* plan slots zeroed on trade complete */
diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
index 8d0af9c28..7d46d0edd 100644
--- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
@@ -10656,7 +10656,7 @@ impl GpuDqnTrainer {
     ) -> Result<(), MLError> {
         let isv_ptr = self.isv_signals_dev_ptr;
         let isv_out = isv_ptr; // kernel writes to the same ISV bus (offset 44)
-        let ps_stride: i32 = 38; // PS_STRIDE from state_layout.cuh
+        let ps_stride: i32 = 39; // PS_STRIDE from state_layout.cuh (grown 38→39 by Plan 3 C.4)
         unsafe {
             self.stream
                 .launch_builder(&self.kelly_cap_update_kernel)
diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
index 99906c273..72f4f1d62 100644
--- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
@@ -41,10 +41,10 @@ const MAX_EPISODES_LIMIT: usize = 0x8000;
 /// Absolute upper bound for validation — reject configs above this.
 const MAX_TIMESTEPS_LIMIT: usize = 1000;
 const PORTFOLIO_STATE_SIZE: usize = 8;
-/// Portfolio stride for DQN experience kernels (38 floats per episode).
-/// Matches PORTFOLIO_STRIDE in experience_kernels.cu.
+/// Portfolio stride for DQN experience kernels (39 floats per episode).
+/// Matches PORTFOLIO_STRIDE in experience_kernels.cu and trade_stats_kernel.cu.
 /// Do NOT change PORTFOLIO_STATE_SIZE above — it's for the PPO/legacy path.
-const PORTFOLIO_STRIDE: usize = 38;  // was 30. Slots 30-37: prev-bar OFI for delta
+const PORTFOLIO_STRIDE: usize = 39;  // was 38. Slots 30-37: prev-bar OFI; slot 38: PS_PEAK_PNL_BAR (Plan 3 C.4)
 
 /// Number of floats in the trade_stats_reduce output buffer.
 /// Layout: [win_count, loss_count, sum_wins, sum_losses, sum_returns, sum_sq_returns]
@@ -568,7 +568,7 @@ pub struct GpuExperienceCollector {
     trade_stats_buf: CudaSlice<f32>,
 
     // Per-episode state buffers [alloc_episodes, ...]
-    portfolio_states: CudaSlice<f32>,  // [alloc_episodes * PORTFOLIO_STRIDE] (38 floats per episode)
+    portfolio_states: CudaSlice<f32>,  // [alloc_episodes * PORTFOLIO_STRIDE] (39 floats per episode)
     episode_starts_buf: CudaSlice<i32>,// [alloc_episodes]
 
     // Output buffers [alloc_episodes * alloc_timesteps, ...]
@@ -981,7 +981,7 @@ impl GpuExperienceCollector {
         }
 
         // ── Step 6: Allocate per-episode buffers ────────────────────────
-        // Portfolio states for experience kernels: [N, PORTFOLIO_STRIDE=38]
+        // Portfolio states for experience kernels: [N, PORTFOLIO_STRIDE=39]
         let mut portfolio_states = stream
             .alloc_zeros::<f32>(alloc_episodes * PORTFOLIO_STRIDE)
             .map_err(|e| MLError::ModelError(format!("alloc portfolio_states: {e}")))?;
@@ -3284,7 +3284,7 @@ impl GpuExperienceCollector {
         _avg_spread: f32,
         _cash_reserve_pct: f32,
     ) -> Result<(), MLError> {
-        // Reset portfolio states [N, PORTFOLIO_STRIDE=38]
+        // Reset portfolio states [N, PORTFOLIO_STRIDE=39]
         let mut portfolio_init = vec![0.0_f32; self.alloc_episodes * PORTFOLIO_STRIDE];
         for i in 0..self.alloc_episodes {
             let off = i * PORTFOLIO_STRIDE;
diff --git a/crates/ml/src/cuda_pipeline/state_layout.cuh b/crates/ml/src/cuda_pipeline/state_layout.cuh
index 567a9658c..bdaf00023 100644
--- a/crates/ml/src/cuda_pipeline/state_layout.cuh
+++ b/crates/ml/src/cuda_pipeline/state_layout.cuh
@@ -30,7 +30,7 @@
 #define SL_STATE_DIM_PADDED   128
 
 // ────────────────────────────────────────────────────────────────────────────
-// Portfolio state slot indices (ps[0..PORTFOLIO_STRIDE=38)).
+// Portfolio state slot indices (ps[0..PORTFOLIO_STRIDE=39)).
 // Invariant 8: every slot has a named constant. Raw ps[N] access outside the
 // definition site in experience_kernels.cu is a lint violation.
 //
@@ -71,7 +71,8 @@
 #define PS_PLAN_ENTRY_REGIME        29   // plan: regime_stability at entry
 // Slots 30-37: prev-bar OFI scratch for delta computation (see experience_kernels.cu)
 #define PS_OFI_PREV_BASE            30   // first OFI prev-bar slot (stride 8: [30..37])
-#define PS_STRIDE                   38   // total portfolio state stride (== PORTFOLIO_STRIDE)
+#define PS_PEAK_PNL_BAR             38   // Plan 3 C.4: hold_time at which MAX_PNL was hit (temporal timing bonus)
+#define PS_STRIDE                   39   // total portfolio state stride (== PORTFOLIO_STRIDE), grown 38→39 by Plan 3 C.4
 
 // ────────────────────────────────────────────────────────────────────────────
 // Plan-ISV slot indices (plan_isv[0..PLAN_ISV_DIM=7)).
diff --git a/crates/ml/src/cuda_pipeline/trade_stats_kernel.cu b/crates/ml/src/cuda_pipeline/trade_stats_kernel.cu
index 1f3c0fa69..f45d43c86 100644
--- a/crates/ml/src/cuda_pipeline/trade_stats_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/trade_stats_kernel.cu
@@ -1,7 +1,7 @@
 /**
  * Trade stats reduction kernel.
  *
- * Reads portfolio_states[N, PORTFOLIO_STRIDE=38] and sums fields [14:19]
+ * Reads portfolio_states[N, PORTFOLIO_STRIDE=39] and sums fields [14:19]
  * (win_count, loss_count, sum_wins, sum_losses, sum_returns, sum_sq_returns)
  * across all N episodes into a 6-float output buffer.
  * Uses shared-memory warp reduction (same pattern as monitoring_reduce).
@@ -9,7 +9,7 @@
  * Launch config: grid=(1, 1, 1), block=(256, 1, 1).
  */
 
-#define PORTFOLIO_STRIDE 38
+#define PORTFOLIO_STRIDE 39  // was 38; grown by Plan 3 C.4 PS_PEAK_PNL_BAR append. Kelly slots [14..19] unchanged.
 #include "state_layout.cuh"
 
 extern "C" __global__ void trade_stats_reduce(
diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md
index 4ce62b63f..991d3d381 100644
--- a/docs/dqn-wire-up-audit.md
+++ b/docs/dqn-wire-up-audit.md
@@ -61,6 +61,7 @@
 | B.1 Flat opp-cost ISV scaling (`experience_kernels.cu` Flat branch) | `experience_env_step` kernel (Flat position=0, not segment_complete branch); writes `rc[4]` to `reward_components_per_sample`; consumed by `reward_component_ema` → ISV[67] | Wired | Plan 3 Task 2 B.1 — opp-cost multiplied by `isv_signals_ptr[21]` (Q_DIR_ABS_REF_INDEX, EMA of max(|Q_mean|) across direction bins). Self-scaling: as Q magnitudes drift during training, opp-cost tracks proportionally. Floor 1e-3 for cold-start. No tuned multiplier. | — |
 | `cuda_pipeline/trade_rate_ema_kernel.cu` | `GpuExperienceCollector::launch_trade_attempt_rate_ema_inplace` → `training_loop.rs` (called alongside `launch_reward_component_ema_inplace` each epoch); reads `flat_to_pos_per_sample [N*L]`, reduces count on-GPU, writes ISV[TRADE_ATTEMPT_RATE_EMA_INDEX=71] adaptive EMA (α=α_base × (1+0.5×\|sharpe\|), α_base=0.05) | Wired | Plan 3 Task 3 B.2 — cold-path (per-epoch); single-block 1-thread reduction + EMA. No atomicAdd. | — |
 | B.2 Flat→Positioned novelty bonus (`experience_kernels.cu` trade-lifecycle block) | `experience_env_step` kernel (entering_trade path, AFTER opp-cost branch, BEFORE drawdown penalty); writes `rc[5]` and adds `shaping_scale × bonus` to `reward`; consumed by `reward_component_ema` → ISV[68] and drives PopArt denominator via `total_reward_per_sample` | Wired | Plan 3 Task 3 B.2 — novelty = max(0, 1 − ISV[71]/max(1e-4, ISV[72])); bonus = conviction_core × vol_proxy × novelty; `TRADE_TARGET_RATE` frozen at epoch 5 from measured attempt-EMA (min 0.001). No tuned multiplier. | — |
+| C.4 Temporal timing bonus on trade exit (`experience_kernels.cu` segment_complete block) | `experience_env_step` kernel (inside `segment_complete && segment_hold_time > 0` block, AFTER Layer 2–9 credits land in `reward`); accumulates `+=` into `rc[5]` and adds `timing_bonus` to `reward`; consumed by `reward_component_ema` → ISV[68] (shares the bonus slot with B.2 — different (i,t) slots, so `+=` is idempotent) | Wired | Plan 3 Task 5 C.4 — bars_early = max(0, segment_hold_time − PS_PEAK_PNL_BAR); timing_bonus = shaping_scale × (bars_early / segment_hold_time) × \|final_pnl\| × conviction_core. Peak bar tracked in new portfolio-state slot PS_PEAK_PNL_BAR=38 (PS_STRIDE 38→39), snapshot alongside every MAX_PNL update, reset at every MAX_PNL reset site (entry/reverse/fold-reset). No tuned multiplier. | — |
 
 ## CUDA Pipeline — Rust Wrappers
 
@@ -243,18 +244,20 @@ Plan 3 Task 1 C.2 (2026-04-24): `reward_component_ema_kernel.cu` + `RewardCompon
 
 Plan 3 Task 3 B.2 (2026-04-24): `trade_rate_ema_kernel.cu` added. 2 new ISV slots [71] TRADE_ATTEMPT_RATE_EMA and [72] TRADE_TARGET_RATE; fingerprint shifted [69..71) → [73..75); `ISV_TOTAL_DIM` 71 → 75. `experience_kernels.cu` gains `flat_to_pos_per_sample [N*L]` output + 2 ISV slot-idx scalar parameters; a novelty-scaled bonus (conviction × vol_proxy × novelty) is added to reward and captured in `rc[5]` at every Flat→Positioned transition. TRADE_TARGET_RATE frozen in `training_loop.rs` at epoch 5 from measured attempt-EMA (min 0.001). 2 new Wired rows.
 
+Plan 3 Task 5 C.4 (2026-04-24): Temporal timing bonus on trade exit. No new ISV slot — accumulates into `rc[5]` bonus slot via `+=` (B.2 at entry, C.4 at exit: different (i,t) slots). New portfolio-state slot `PS_PEAK_PNL_BAR=38`; `PS_STRIDE` 38 → 39; `PORTFOLIO_STRIDE` 38 → 39 in all consumers (`experience_kernels.cu`, `trade_stats_kernel.cu`, `gpu_experience_collector.rs`, `gpu_dqn_trainer.rs::launch_kelly_cap_update`, `ml-core::state_layout.rs`). Peak bar snapshotted alongside every MAX_PNL update, reset at every MAX_PNL reset site (entry, reverse, fold hard-reset, trade-complete soft-reset). `timing_bonus = shaping_scale × (bars_early / hold_time) × |final_pnl| × conviction_core` where `bars_early = max(0, segment_hold_time − PS_PEAK_PNL_BAR)`. No tuned multiplier. 1 new Wired row.
+
 Plan 1 Tasks 12/15/16 + pre-allocation (2026-04-24): No new modules added. Changes are ISV slot allocation + consumer migration only. Task 15 confirmed no-op (`IQL_BRANCH_SCALE_FLOOR_INDEX` already serves conviction-floor role). Tasks 12 and 16 migrate `cql_alpha` and plan-threshold consumers from config fields / hardcoded literals to ISV slots. 8 new ISV slots allocated ([39..47)); fingerprint tail moves from [37..39) to [47..49); `ISV_TOTAL_DIM` 39 → 49. `GpuDqnTrainConfig` gains `total_epochs` field (written to `TOTAL_EPOCHS_INDEX` at construction). `write_isv_signal_at` bound extended from `ISV_DIM` to `ISV_TOTAL_DIM` to allow writes beyond slot 22.
 
 Plan 2 Task 6B D.3 (2026-04-24): IQL value head widened from 1 to 2 outputs (V_short + V_long). `v_out_buf` shape `[B]` → `[B*2]`. `gemm_fwd_v` M=1→2, `gemm_bwd_dw3` M=1→2, `gemm_bwd_dh2` K=1→2. `W3` param block `[H*1]` → `[H*2]`, `b3` `[1]` → `[2]`. `total_params` += H+1. `iql_expectile_loss` kernel extended with `num_heads` argument. 4 consumer kernels in `iql_value_kernel.cu` updated to read `v_out[b*2+0] + v_out[b*2+1]`. Checkpoint compat break — retrain required.
 
 | Classification | Count |
 |---|---|
-| Wired | 79 |
+| Wired | 80 |
 | Partial | 10 |
 | Orphan (held for follow-up) | 3 |
 | Ghost | 0 |
 | OUT-of-DQN-scope | 17 |
-| **Total** | **109** |
+| **Total** | **110** |
 
 The 3 remaining Orphan rows are:
 - `cuda_pipeline/gpu_statistics.rs` + `statistics_kernel.cu` — held for Plan 2 D.2 wire-or-delete decision.