gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

feat(dqn-v2): C.4 temporal timing bonus on trade exit — peak-bar proximity

Browse Source 
Plan 3 Task 5.

Portfolio-state tail-append (shared-contract migration, all in same commit):
- PS_PEAK_PNL_BAR = 38 (hold_time snapshotted when MAX_PNL updates)
- PS_STRIDE 38 -> 39 in state_layout.cuh and ml-core/state_layout.rs
- PORTFOLIO_STRIDE 38 -> 39 in trade_stats_kernel.cu (hardcoded copy)
- PORTFOLIO_STRIDE 38 -> 39 in gpu_experience_collector.rs allocator
- ps_stride 38 -> 39 in gpu_dqn_trainer.rs launch_kelly_cap_update

Producer (experience_kernels.cu):
- Peak bar snapshotted alongside every MAX_PNL update (uses local
  hold_time, not ps[PS_HOLD_TIME], because the portfolio-state commit
  block runs later in the kernel).
- Peak bar reset to 0 at every MAX_PNL reset site: plan-entry (1856),
  entering_trade (2014), reversing_trade (2019), fold hard-reset (2736),
  trade-complete soft-reset (2751).

Consumer (experience_kernels.cu segment_complete block):
- bars_early = max(0, segment_hold_time - PS_PEAK_PNL_BAR)
- timing_bonus = shaping_scale x (bars_early / segment_hold_time)
                 x |final_pnl| x conviction_core
- reward += timing_bonus; rc[5] += timing_bonus
  (accumulates with Task 3 B.2 entry bonus — different (i,t) slots).

No new ISV slot — rc[5] bonus semantics unchanged; B.2 and C.4 share it
via += accumulate semantics (defensively idempotent, but the two sites
fire at distinct (i,t) by construction: entry vs exit).

Self-scaling: shaping_scale x conviction_core x |pnl| keeps the bonus
proportional to trade magnitude, no tuned coefficients.

Smoke multi_fold_convergence (RTX 3050 Ti): all 3 folds complete,
fold-2 best Sharpe 84.44 at epoch 1 (expected ~85 range).
cargo check --workspace clean at 11 warnings baseline.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-24 23:34:27 +02:00
parent b5d19c1004
commit a0abc3da35
7 changed files with 68 additions and 21 deletions
Download Patch File Download Diff File Expand all files Collapse all files

3
crates/ml-core/src/state_layout.rs
View File

@@ -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)).

56
crates/ml/src/cuda_pipeline/experience_kernels.cu
View File

@@ -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 */

2
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
View File

@@ -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)

12
crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
View File

@@ -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;

5
crates/ml/src/cuda_pipeline/state_layout.cuh
View File

@@ -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)).

4
crates/ml/src/cuda_pipeline/trade_stats_kernel.cu
View File

@@ -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(

7
docs/dqn-wire-up-audit.md
View File

@@ -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.