# SP13 Redefine Success for Predictive Skill — Implementation Plan (v2)

> **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development (recommended) or superpowers:executing-plans to implement this plan task-by-task. Steps use checkbox (`- [ ]`) syntax for tracking.

**Goal:** Test the user's hypothesis that the persistent ~46% win rate is caused by Hold action hiding directional signal (not by absence of signal in data) by removing Hold cleanly in Phase 0a. If confirmed, build out direction-skill bonus + correct-call asymmetry (calibrated to prevent reward gaming) and aux head binary classification. v1 design dropped alpha-vs-benchmark reward (mathematically broken — see spec v1→v2 changes).

**Architecture:** 8 new ISV slots `[372..380)`, 1 new GPU producer kernel (`aux_dir_acc_reduce_kernel.cu`), corrected aux_w controller with stagnation detector (Phase 0b only), full action-space refactor (4 → 3, eliminating Hold), aux head refactor (regression → binary classification), reward composition extension (skill bonus + luck discount on top of SP12). Phases: P0a (Hold-only, decisive gate, 1 commit), P0b (aux amplification, only if P0a partial/red, 1 commit), B (aux classification, 1 commit), C (reward composition, 1 atomic commit), D (30-epoch validation + close-out).

**Tech Stack:** Rust 1.85+, CUDA 12.4, cudarc, mapped-pinned memory, block tree-reduce (no atomicAdd), GPU-only Pearls A+D via `apply_pearls_ad_kernel`, L40S smoke validation via Argo Workflows, RTX 3050 Ti for local oracle tests.

**Spec source:** `docs/superpowers/specs/2026-05-04-sp13-redefine-success-for-predictive-skill.md` (v2)

**Branch:** `sp11-reward-as-controlled-subsystem` (greenfield refactor — replay buffer compatibility intentionally broken in P0a)

**Operating principles** (from MEMORY.md):
- `feedback_no_cpu_compute_strict` — every producer GPU; Pearls A+D via `apply_pearls_ad_kernel`
- `feedback_no_atomicadd` — block tree-reduce only
- `feedback_no_partial_refactor` — P0a / P0b / B / C each single atomic commit
- `feedback_isv_for_adaptive_bounds` — all 8 knobs ISV-driven
- `feedback_no_htod_htoh_only_mapped_pinned` — mapped-pinned only
- `feedback_no_cpu_test_fallbacks` — GPU oracle testing only
- `feedback_no_stubs` / `feedback_no_todo_fixme` / `feedback_no_feature_flags`
- End-to-end training: NO pre-training of aux head
- `pearl_audit_unboundedness_for_implicit_asymmetry` — SP12 asymmetric cap stays
- `pearl_event_driven_reward_density_alignment` — reward fires only at trade events
- **NEW reward-quadrant audit rule** — every reward change must come with a 4-quadrant
  worked-example table (correct/wrong × winner/loser); enforced in Layer C tests

**Hypothesis under test** (from `project_hold_hides_directional_signal_hypothesis.md`):
- P0a green (10-epoch smoke shows dir_acc and WR climbing) → CONFIRMED, proceed
- P0a partial → P0b adds aux amplification probe
- Both red → DISCONFIRMED, halt SP13, plan SP14 data-augmentation work separately

---

## File structure

### New files (Phase 0a)

| File | Phase | Responsibility |
|---|---|---|
| `crates/ml/src/cuda_pipeline/sp13_isv_slots.rs` | P0a | 8 SP13 slot index constants `[372..380)` + defaults |
| `crates/ml/src/cuda_pipeline/aux_dir_acc_reduce_kernel.cu` | P0a | Single-block tree-reduce: 4 shared arrays (correct, pos_pred, pos_label, valid) → `(dir_acc, pos_pred_frac, pos_label_frac)` |
| `crates/ml/tests/sp13_phase0_oracle_tests.rs` | P0a | 6+ GPU oracle tests for dir_acc kernel + 4 action-selection tests for 3-direction argmax |

### Modified files (Phase 0a)

| File | Change |
|---|---|
| `crates/ml/src/cuda_pipeline/sp5_isv_slots.rs` | Bump `SP5_SLOT_END = 380`, `ISV_TOTAL_DIM = 380`, layout fingerprint seed |
| `crates/ml/src/cuda_pipeline/state_layout.cuh` | `NUM_DIRECTIONS: 4 → 3`; `DIRECTION_SHORT=0, DIRECTION_LONG=1, DIRECTION_FLAT=2` (DIRECTION_HOLD deleted); `ISV_TOTAL_DIM` mirror; SP13 slot defines |
| `crates/ml/src/trainers/dqn/state_reset_registry.rs` | 2 new entries (slots 373, 374) per-fold sentinel 0.5; line :974 `TRAIL_DIST_PER_DIR_BASE=270..274` comment update |
| `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` | Cubin include + launcher for `aux_dir_acc_reduce`; `aux_dir_acc_buf: MappedF32Buffer` (3 floats); static-init slot 372 = 0.55 |
| `crates/ml/build.rs` | Cubin registration; line :490 stale `Short=0, Hold=1, Long=2, Flat=3` doc-comment update |
| **`crates/ml-core/src/common/action.rs:124-132`** | Canonical `DirectionAction` Rust enum — eliminate `Hold` variant; `Short=0, Long=1, Flat=2` |
| **`crates/ml-core/src/state_layout.rs:48-49`** | `pub const NUM_DIRECTIONS: usize = 3` (was 4); doc-comment update |
| `crates/ml/src/cuda_pipeline/experience_kernels.cu` | Multi-site: line :1013 `experience_action_select` direction Thompson 4 → 3; line :4059 direction accumulation loop; line :5159 **direction-bias signal `[1.0, 0.5, 1.0, 0.0]` → `[1.0, 1.0, 0.0]`** (Hold's 0.5 softening gate removed; see P0a.T3 step 4 below); line :5306 conviction EMA loop; lines :6371-6397 per-atom direction loops (4 occurrences) |
| `crates/ml/src/cuda_pipeline/c51_loss_kernel.cu` | Lines :394 (doc), :460 (`for k<4` argmax), :1088 (argmax across direction bins) |
| `crates/ml/src/cuda_pipeline/c51_grad_kernel.cu` | Implementer AUDIT for `for.*<4` direction loops (Explore agent flagged for verification) |
| `crates/ml/src/cuda_pipeline/iqn_loss_kernel.cu` / `iqn_dual_head_kernel.cu` | Lines :245, :1199, :200 — `actions[sample * 4 + 0]` direction index UNCHANGED dimensionally (action layout still [direction, magnitude, order, urgency]); only direction VALUE range shrinks to 0..2 |
| `crates/ml/src/cuda_pipeline/dqn_inference.cu` | Direction Q-head output dim 4 → 3 |
| `crates/ml/src/cuda_pipeline/dqn_param_layout.rs` | Direction Q-head weight/bias shapes; layout fingerprint seed bump |
| `crates/ml/src/trainers/dqn/financials.rs:27,197` | `exp_idx = dir*3 + mag` packing — **range shrinks 12 → 9** (was 4 dirs × 3 mags = 12; now 3 dirs × 3 mags = 9); doc-comment :27 update |
| `crates/ml/src/trainers/dqn/monitoring.rs:13,341` | `action_counts[9..12]` Flat-mag indexing recomputes; doc-comment :13 stale `S=Short(0), H=Hold(1), L=Long(2), F=Flat(3)` update |
| `crates/ml/src/cuda_pipeline/sp5_isv_slots.rs:272` | `TRAIL_DIST_PER_DIR_BASE` doc-comment `[4 dirs: Short, Hold, Long, Flat]` update — slot range KEPT at 270..274 (4 entries) per spec; slot 273 becomes dead-space (consumers read by direction index 0..2; never touched) |
| `crates/ml/src/cuda_pipeline/pearl_8_trail_kernel.cu` | Producer kernel writes per-direction trail dist at NEW indices `{Short→0, Long→1, Flat→2, dead→3}`; consumer reads unchanged |
| `crates/ml/src/trainers/dqn/trainer/training_loop.rs:4046-4048` | Pearl 8 `apply_pearls_ad_kernel` calls — 4 calls become 3 (dead slot 273 not smoothed); also dir_acc launch + ISV write + HEALTH_DIAG `aux_dir_acc` + val_dir_dist 3-bucket emit |
| `crates/ml/src/trainers/dqn/trainer/metrics.rs:949` | `val_dir_dist [short=… hold=… long=… flat=…]` → `[short=… long=… flat=…]` |
| `crates/ml/src/trainers/dqn/smoke_tests/magnitude_distribution.rs:88,217` | `EVAL_DIR_DIST` 4-tuple → 3-tuple emit |
| `crates/ml/src/trainers/dqn/distributional_q_tests.rs:790,888` | `σ_C51 per direction` test output 4 → 3; `PROD_B0=4` test constant → 3 |
| `crates/ml/src/cuda_pipeline/kl_divergence_kernel.cu` (if exists) | 3-bucket update — Explore agent did NOT find this file; verify nonexistence or update if found |
| `crates/ml/src/trainers/dqn/trainer/tests.rs:542` | `assert_eq!(config.num_actions, 4, …)` → `3` |
| `crates/ml/src/trainers/dqn/trainer/mod.rs:1966` | Doc-comment `Layout: [Short, Hold, Long, Flat]` → `[Short, Long, Flat]` |
| `crates/ml/src/cuda_pipeline/thompson_test_kernel.cu` | Test kernel direction count 4 → 3 |
| 20+ stale doc lines across kernel headers (`.cu` files) | Implementer audit: `grep -rn "Short=0\|Hold=1\|Long=2\|Flat=3" crates/ml/src/` and update narrative comments to `Short=0, Long=1, Flat=2` |

**Note on slot 375 (`AUX_DIR_PREDICTION_INDEX`) in P0a:** the existing aux head still emits a regression scalar. Phase 0a writes `tanh(scalar_pred)` into slot 375 to fit the [-1, +1] interface so Layer B can refactor without changing wire format. This is NOT a stub — it produces a real prediction signal, just with the legacy regression head. Layer B replaces with softmax-derived logit_diff.

### Modified files (Phase 0b — only if P0a partial/red)

| File | Change |
|---|---|
| `crates/ml/src/trainers/dqn/trainer/training_loop.rs:4103-4104` | Replace inverted SP11 formula with deficit + stagnation controller |
| `crates/ml/src/cuda_pipeline/sp13_isv_slots.rs` | Add `AUX_W_BASE_PHASE0B = 0.5`, `AUX_W_STAGNATION_DECAY_K = 0.7` constants |

### Modified files (Layer B)

| File | Change |
|---|---|
| `crates/ml/src/cuda_pipeline/aux_head_kernel.cu` | Output dim 1 → 2; activation linear → softmax; loss MSE → CrossEntropy |
| `crates/ml/src/cuda_pipeline/aux_label_gen_kernel.cu` | Label = `sign(price[t+30] - price[t])` as 0/1 class index |
| `crates/ml/src/cuda_pipeline/aux_dir_acc_reduce_kernel.cu` | Read 2-class softmax instead of regression scalar; pred_pos = (softmax[1] > softmax[0]) |
| `crates/ml/src/cuda_pipeline/dqn_inference.cu` | Direction Q-head input concat: `[trunk_features, ISV[AUX_DIR_PREDICTION_INDEX]]`; in_dim grows by 1 |
| `crates/ml/src/cuda_pipeline/dqn_param_layout.rs` | Direction Q-head input dim +1 |
| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | Replace `tanh(scalar)` ISV[375] write with `softmax[1] - softmax[0]` |
| `crates/ml/tests/sp13_phase0_oracle_tests.rs` | +3 cross-entropy classification tests |

### New files (Layer C)

| File | Phase | Responsibility |
|---|---|---|
| `crates/ml/src/cuda_pipeline/sp13_reward_math_test_kernel.cu` | C | GPU oracle wrapper for 2 device-inline reward funcs (skill bonus, luck discount); mirrors SP12 sp12_reward_math_test_kernel.cu |
| `crates/ml/tests/sp13_reward_math_tests.rs` | C | 16 GPU oracle tests covering full 8-quadrant matrix from spec |

### Modified files (Layer C)

| File | Change |
|---|---|
| `crates/ml/src/cuda_pipeline/trade_physics.cuh` | Add 2 device-inline funcs: `compute_direction_skill_bonus_bounded`, `compute_luck_win_discount` |
| `crates/ml/src/cuda_pipeline/experience_kernels.cu` (~line 2788, exit-event reward) | After SP12 asymmetric cap, ADD skill bonus then ADD luck discount; min_hold_factor unchanged |
| `crates/ml/src/trainers/dqn/state_reset_registry.rs` | 4 new entries (slots 376, 377, 378, 379) per-fold |
| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | At fold-end HEALTH_DIAG: emit `per_fold_alpha_vs_buy_and_hold = total_pnl - (price[fold_end] - price[fold_start])` as METRIC ONLY |

### Documentation files (Layer D)

| File | Phase | Content |
|---|---|---|
| `docs/dqn-wire-up-audit.md` | D | SP13 entry: hypothesis, P0a/P0b results, Layer B/C deltas, 30-epoch metrics |
| `~/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/pearl_redefine_success_for_predictive_skill.md` | D | NEW pearl |
| `~/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/pearl_skill_bonus_must_be_alpha_bounded.md` | D | NEW pearl — calibration lesson from v1 review |
| `~/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/pearl_reward_quadrant_audit_required.md` | D | NEW meta-pearl from v1 review |
| `~/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/MEMORY.md` | D | Index entries |
| `~/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/project_hold_hides_directional_signal_hypothesis.md` | D | Status → CONFIRMED / DISCONFIRMED with smoke results |

---

## Phase 0a — Hold-pricing + dir_acc instrumentation (1 atomic commit, additive)

> **v3 NOTE**: P0a.T3 was v2's atomic Hold elimination (would have cascaded across 77 files / 200+ sites including 32 test files). After the T3 implementer's audit found `DirectionAction` enum doesn't exist (codebase uses 8-variant fused `ExposureLevel`), v3 reframed as **Hold-PRICING**: ~120 LOC, additive, no contract change, no cross-crate cascade. The File Structure table above was authored for v2 — entries that conflict with v3 (line 5159 direction-bias re-derivation, financials.rs/monitoring.rs action-packing 12→9, val_dir_dist 4→3, SP5 Pearl 8 trail-dist migration, 30+ stale-doc cleanup) are **NO LONGER IN SCOPE**. The authoritative T3 spec is the v3 P0a.T3 section below.

**P0a is additive, atomic per `feedback_no_partial_refactor`** — but "atomic" here means single commit, NOT contract change. Replay buffer / fxcache compatibility preserved.

### Task P0a.T1: ISV slot constants + state-reset registry

**Files:** Create `sp13_isv_slots.rs`; modify `sp5_isv_slots.rs`, `state_layout.cuh`, `state_reset_registry.rs`, `gpu_dqn_trainer.rs`

- [ ] **Step 1:** Create `crates/ml/src/cuda_pipeline/sp13_isv_slots.rs`:

```rust
//! SP13 ISV slot constants — see spec docs/superpowers/specs/2026-05-04-sp13-redefine-success-for-predictive-skill.md

pub const TARGET_DIR_ACC_INDEX: usize = 372;
pub const AUX_DIR_ACC_SHORT_EMA_INDEX: usize = 373;
pub const AUX_DIR_ACC_LONG_EMA_INDEX: usize = 374;
pub const AUX_DIR_PREDICTION_INDEX: usize = 375;
pub const DIR_SKILL_BONUS_ALPHA_INDEX: usize = 376;
pub const DIR_SKILL_BONUS_BETA_INDEX: usize = 377;
pub const LUCK_WIN_DISCOUNT_INDEX: usize = 378;
pub const SKILL_BONUS_CAP_RATIO_INDEX: usize = 379;

pub const TARGET_DIR_ACC_DEFAULT: f32 = 0.55;
pub const DIR_SKILL_BONUS_ALPHA_DEFAULT: f32 = 1.0;
pub const DIR_SKILL_BONUS_BETA_DEFAULT: f32 = 1.0;
pub const LUCK_WIN_DISCOUNT_DEFAULT: f32 = 0.3;
pub const SKILL_BONUS_CAP_RATIO_DEFAULT: f32 = 0.3;

/// Sentinel for both dir_acc EMAs at fold-start: 0.5 (random baseline)
/// — first observation replaces directly per pearl_first_observation_bootstrap
pub const DIR_ACC_EMA_SENTINEL: f32 = 0.5;
```

- [ ] **Step 2:** Bump SP5_SLOT_END in `sp5_isv_slots.rs`:

```rust
pub const SP5_SLOT_END: usize = 380;  // was 371 post-SP12; +8 for SP13
pub const ISV_TOTAL_DIM: usize = SP5_SLOT_END;
```

Add `pub use crate::cuda_pipeline::sp13_isv_slots::*;` re-export. Bump layout fingerprint seed.

- [ ] **Step 3:** Mirror in `state_layout.cuh` — bump `ISV_TOTAL_DIM` to 380; add `// === SP13 SLOT INDICES ===` block with all 8 defines.

- [ ] **Step 4:** Add reset registry entries:

```rust
FoldResetEntry { name: "aux_dir_acc_short_ema", slot_index: AUX_DIR_ACC_SHORT_EMA_INDEX, reset_value: DIR_ACC_EMA_SENTINEL, reset_kind: ResetKind::PerFold },
FoldResetEntry { name: "aux_dir_acc_long_ema",  slot_index: AUX_DIR_ACC_LONG_EMA_INDEX,  reset_value: DIR_ACC_EMA_SENTINEL, reset_kind: ResetKind::PerFold },
```

(Slots 372, 376-379 are static-initialised in trainer constructor — no reset entry needed; slot 375 is per-bar overwritten — no reset needed.)

- [ ] **Step 5:** In `GpuDqnTrainer::new` (or equivalent constructor), add static-init for the 6 static slots:

```rust
self.write_isv_slot(TARGET_DIR_ACC_INDEX, TARGET_DIR_ACC_DEFAULT)?;
self.write_isv_slot(DIR_SKILL_BONUS_ALPHA_INDEX, DIR_SKILL_BONUS_ALPHA_DEFAULT)?;
self.write_isv_slot(DIR_SKILL_BONUS_BETA_INDEX, DIR_SKILL_BONUS_BETA_DEFAULT)?;
self.write_isv_slot(LUCK_WIN_DISCOUNT_INDEX, LUCK_WIN_DISCOUNT_DEFAULT)?;
self.write_isv_slot(SKILL_BONUS_CAP_RATIO_INDEX, SKILL_BONUS_CAP_RATIO_DEFAULT)?;
```

(The α/β/luck/cap-ratio slots aren't *consumed* until Layer C, but writing them in P0a is harmless and avoids a follow-up registry update in C.)

- [ ] **Step 6:** Verify compile: `SQLX_OFFLINE=true cargo check -p ml --lib`

- [ ] **Step 7:** Stage (no commit yet — atomic at end of P0a).

---

### Task P0a.T2: aux_dir_acc_reduce kernel + GPU oracle tests

**Goal:** Single-block tree-reduce kernel with 4 shared-mem arrays (correct, pos_pred, pos_label, valid). Reads regression-mode aux scalar in P0a (Layer B refactors to softmax). 6+ GPU oracle tests on RTX 3050 Ti.

- [ ] **Step 1:** Write the kernel (`aux_dir_acc_reduce_kernel.cu`):

```cuda
#include "state_layout.cuh"

extern "C" __global__ void aux_dir_acc_reduce_kernel(
    const float* __restrict__ aux_pred,        // [B] regression scalar (Phase 0a); Layer B will swap to softmax read
    const float* __restrict__ next_bar_label,  // [B] sign-encoded: +1 / 0 / -1
    int batch_size,
    float* __restrict__ out_3                  // [3] = {dir_acc, pos_pred_frac, pos_label_frac}
) {
    const int tid = threadIdx.x;
    const int bdim = blockDim.x;

    extern __shared__ int shared[];
    int* sh_correct   = &shared[0 * bdim];
    int* sh_pos_pred  = &shared[1 * bdim];
    int* sh_pos_label = &shared[2 * bdim];
    int* sh_valid     = &shared[3 * bdim];

    int local_correct = 0, local_pos_pred = 0, local_pos_label = 0, local_valid = 0;
    for (int i = tid; i < batch_size; i += bdim) {
        const float p = aux_pred[i];
        const float l = next_bar_label[i];
        if (l == 0.0f) continue;
        local_valid += 1;
        const int pred_pos = (p > 0.0f) ? 1 : 0;
        const int label_pos = (l > 0.0f) ? 1 : 0;
        if (pred_pos == label_pos) local_correct += 1;
        local_pos_pred  += pred_pos;
        local_pos_label += label_pos;
    }
    sh_correct[tid]   = local_correct;
    sh_pos_pred[tid]  = local_pos_pred;
    sh_pos_label[tid] = local_pos_label;
    sh_valid[tid]     = local_valid;
    __syncthreads();

    for (int s = bdim / 2; s > 0; s >>= 1) {
        if (tid < s) {
            sh_correct[tid]   += sh_correct[tid + s];
            sh_pos_pred[tid]  += sh_pos_pred[tid + s];
            sh_pos_label[tid] += sh_pos_label[tid + s];
            sh_valid[tid]     += sh_valid[tid + s];
        }
        __syncthreads();
    }

    if (tid == 0) {
        const float denom = (float)sh_valid[0];
        const float sentinel = 0.5f;
        out_3[0] = (denom > 0.0f) ? ((float)sh_correct[0]   / denom) : sentinel;
        out_3[1] = (denom > 0.0f) ? ((float)sh_pos_pred[0]  / denom) : sentinel;
        out_3[2] = (denom > 0.0f) ? ((float)sh_pos_label[0] / denom) : sentinel;
    }
}
```

Block dim: 256, shared mem: `4 × 256 × sizeof(int) = 4 KB`.

- [ ] **Step 2:** Register cubin in `build.rs`.

- [ ] **Step 3:** Add launcher in `gpu_dqn_trainer.rs` (mirrors existing reduction-kernel launcher patterns, e.g. `launch_q_stats_reduce`):

```rust
pub fn launch_aux_dir_acc_reduce(
    &self,
    aux_pred: &CudaSlice<f32>,
    next_bar_label: &CudaSlice<f32>,
    batch_size: i32,
    out_3: &mut CudaSlice<f32>,
) -> Result<()> {
    let bdim: u32 = 256;
    let cfg = LaunchConfig {
        grid_dim: (1, 1, 1),
        block_dim: (bdim, 1, 1),
        shared_mem_bytes: 4 * bdim * std::mem::size_of::<i32>() as u32,
    };
    unsafe {
        self.aux_dir_acc_reduce
            .clone()
            .launch(cfg, (aux_pred, next_bar_label, batch_size, out_3))?;
    }
    Ok(())
}
```

Add `aux_dir_acc_reduce: CudaFunction` field + `aux_dir_acc_buf: MappedF32Buffer` (3 floats).

- [ ] **Step 4:** Write `crates/ml/tests/sp13_phase0_oracle_tests.rs`:

```rust
#![cfg(feature = "gpu-tests")]

use ml::cuda_pipeline::gpu_dqn_trainer::GpuDqnTrainer;

fn test_trainer() -> GpuDqnTrainer { /* mirror sp12_reward_math_tests scaffold */ }
fn run_dir_acc(trainer: &GpuDqnTrainer, pred: &[f32], label: &[f32]) -> [f32; 3] { /* ... */ }

#[test]
fn dir_acc_all_correct_returns_one() {
    let pred  = vec![1.0_f32, 1.0, -1.0, -1.0];
    let label = vec![1.0_f32, 1.0, -1.0, -1.0];
    let out = run_dir_acc(&test_trainer(), &pred, &label);
    assert!((out[0] - 1.0).abs() < 1e-6);
}

#[test]
fn dir_acc_all_wrong_returns_zero() {
    let pred  = vec![1.0_f32, 1.0, -1.0, -1.0];
    let label = vec![-1.0_f32, -1.0, 1.0, 1.0];
    let out = run_dir_acc(&test_trainer(), &pred, &label);
    assert!(out[0].abs() < 1e-6);
}

#[test]
fn dir_acc_half_correct() {
    let pred  = vec![1.0_f32, 1.0,  1.0,  1.0];
    let label = vec![1.0_f32, 1.0, -1.0, -1.0];
    let out = run_dir_acc(&test_trainer(), &pred, &label);
    assert!((out[0] - 0.5).abs() < 1e-6);
    assert!((out[1] - 1.0).abs() < 1e-6);
    assert!((out[2] - 0.5).abs() < 1e-6);
}

#[test]
fn dir_acc_skips_zero_labels() {
    let pred  = vec![1.0_f32, 1.0, -1.0];
    let label = vec![1.0_f32, 0.0, -1.0];
    let out = run_dir_acc(&test_trainer(), &pred, &label);
    assert!((out[0] - 1.0).abs() < 1e-6);
}

#[test]
fn dir_acc_zero_pred_is_negative() {
    // (p > 0.0f) ? 1 : 0 — zero predicts class 0 (negative)
    let pred  = vec![0.0_f32, -0.0];
    let label = vec![-1.0_f32, -1.0];
    let out = run_dir_acc(&test_trainer(), &pred, &label);
    assert!((out[0] - 1.0).abs() < 1e-6);
}

#[test]
fn dir_acc_empty_batch_returns_sentinel() {
    let pred: Vec<f32>  = vec![];
    let label: Vec<f32> = vec![];
    let out = run_dir_acc(&test_trainer(), &pred, &label);
    assert!((out[0] - 0.5).abs() < 1e-6);
}
```

- [ ] **Step 5:** Run on RTX 3050 Ti: `SQLX_OFFLINE=true cargo test -p ml --features gpu-tests --test sp13_phase0_oracle_tests`. Expected: 6/6 passing.

- [ ] **Step 6:** Stage.

---

### Task P0a.T3 (v3): Hold-pricing — adaptive cost controller (additive, no contract change)

**Goal:** Make Hold action cost something so the policy uses it deliberately. v2 was atomic Hold elimination; v3 prices Hold via ISV-driven adaptive controller targeting a Hold-rate (~20% — MFT-ish). Per `feedback_isv_for_adaptive_bounds` — base cost is the only constant; rate target and observed rate are ISV slots; cost output is ISV-driven.

**Architecture:**
- 3 new ISV slots (380-382): `HOLD_COST_INDEX`, `HOLD_RATE_TARGET_INDEX`, `HOLD_RATE_OBSERVED_EMA_INDEX`
- Hold-rate observer: small GPU kernel writing `(action == DIR_HOLD ? 1.0 : 0.0)` to scratch per-bar; Pearls A+D smoothing (α=0.05) into ISV[382]
- Hold-cost controller: 5-line formula in `training_loop.rs` reading observed rate and writing cost to ISV[380]
- Reward composition: per-bar `reward -= ISV[HOLD_COST_INDEX]` when action == Hold at the existing reward composition site in `experience_kernels.cu`
- 2 GPU oracle tests for the controller

**Per `pearl_event_driven_reward_density_alignment` tension** (acknowledged): per-bar Hold cost is per-bar reward shaping, which the pearl warns against. The cost is justified as: (a) economically realistic (broker fees, margin interest, opportunity cost), (b) ISV-bounded by the controller, (c) inverse-direction from the pearl's failure mode (pulls policy AWAY from Hold-default, doesn't create exposure-positive bias).

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/sp13_isv_slots.rs` (add 3 constants + defaults)
- Modify: `crates/ml/src/cuda_pipeline/sp5_isv_slots.rs` (`SP5_SLOT_END = 383`; layout fingerprint extension)
- Modify: `crates/ml/src/cuda_pipeline/state_layout.cuh` (3 new SP13 slot defines + `ISV_TOTAL_DIM = 383`)
- Modify: `crates/ml/src/trainers/dqn/state_reset_registry.rs` (1 new entry — slot 382 per-fold, sentinel 0.0)
- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (reset arm for slot 382 + per-step controller formula)
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (constructor static-init for slots 380, 381; bump `ISV_TOTAL_DIM = 383`; fingerprint extension)
- Create: `crates/ml/src/cuda_pipeline/hold_rate_observer_kernel.cu` (~30 LOC GPU kernel)
- Modify: `crates/ml/build.rs` (cubin registration)
- Modify: `crates/ml/src/cuda_pipeline/experience_kernels.cu` (locate reward composition site near SP12 alpha cap; add per-bar Hold cost subtraction when action == DIR_HOLD)
- Modify: `crates/ml/tests/sp13_phase0_oracle_tests.rs` (add 2 controller tests)

- [ ] **Step 1:** Add 3 ISV slot constants + defaults to `sp13_isv_slots.rs`:

```rust
pub const HOLD_COST_INDEX:                usize = 380;
pub const HOLD_RATE_TARGET_INDEX:         usize = 381;
pub const HOLD_RATE_OBSERVED_EMA_INDEX:   usize = 382;

/// Per-bar Hold cost base (controller can scale up to 5×, down to 0.5×).
/// Small per-bar (~10 ticks of price for ES.FUT at ~0.25 tick value) but cumulative
/// across long Hold runs becomes meaningful relative to capped trade reward (±5..10).
pub const HOLD_COST_BASE: f32 = 0.001;

pub const HOLD_RATE_TARGET_DEFAULT: f32 = 0.20;
pub const HOLD_COST_CONTROLLER_GAIN: f32 = 5.0;
pub const HOLD_COST_FLOOR_RATIO: f32 = 0.5;   /* never below 0.5 × base */
pub const HOLD_COST_CEIL_RATIO:  f32 = 5.0;   /* never above 5.0 × base */
```

- [ ] **Step 2:** Bump `SP5_SLOT_END = 383` (was 380 from T1) in `sp5_isv_slots.rs`. Extend `SP5_LAYOUT_FINGERPRINT_FRAGMENT` with the 3 new slot names. Update test expectations: `SP5_PRODUCER_COUNT = 209`.

- [ ] **Step 3:** Mirror in `state_layout.cuh`: `#define ISV_TOTAL_DIM 383`; add the 3 new SP13 slot defines.

- [ ] **Step 4:** Add reset registry entry in `state_reset_registry.rs` for slot 382 (per-fold, sentinel `0.0` — no observations at fold start). Slots 380 and 381 are static-init in constructor.

- [ ] **Step 5:** Add reset_named_state dispatch arm for slot 382 in `training_loop.rs` (writes 0.0).

- [ ] **Step 6:** Constructor static-init in `gpu_dqn_trainer.rs` for slots 380 (`HOLD_COST_BASE`) and 381 (`HOLD_RATE_TARGET_DEFAULT`). Extend `layout_fingerprint_seed()` byte-string with 3 new slot names + `ISV_TOTAL_DIM=383`. Bump `ISV_TOTAL_DIM` constant.

- [ ] **Step 7:** Write the Hold-rate observer kernel `crates/ml/src/cuda_pipeline/hold_rate_observer_kernel.cu`:

```cuda
// hold_rate_observer_kernel.cu
// Per-step: writes 1.0 if any bar in batch picked Hold, 0.0 otherwise (or aggregated
// fraction). Single-block reduce; mapped-pinned output for Pearls A+D smoothing.

#include "state_layout.cuh"

extern "C" __global__ void hold_rate_observer_kernel(
    const int* __restrict__ actions_dir,  /* [B] direction actions, values ∈ {0..3} */
    int batch_size,
    float* __restrict__ hold_rate_out      /* [1] = (count(Hold) / B) */
) {
    const int tid = threadIdx.x;
    const int bdim = blockDim.x;
    extern __shared__ int sh_hold[];

    int local_hold = 0;
    for (int i = tid; i < batch_size; i += bdim) {
        if (actions_dir[i] == DIR_HOLD) local_hold += 1;
    }
    sh_hold[tid] = local_hold;
    __syncthreads();

    for (int s = bdim / 2; s > 0; s >>= 1) {
        if (tid < s) sh_hold[tid] += sh_hold[tid + s];
        __syncthreads();
    }

    if (tid == 0) {
        hold_rate_out[0] = (batch_size > 0) ? ((float)sh_hold[0] / (float)batch_size) : 0.0f;
    }
}
```

(Single-block, 256 threads, 1 KB shared mem. Mirror the launch infrastructure of `aux_dir_acc_reduce` from T2.)

- [ ] **Step 8:** Register cubin in `build.rs` matching the existing pattern.

- [ ] **Step 9:** Add launcher in `gpu_dqn_trainer.rs` (mirroring `launch_aux_dir_acc_reduce`):

```rust
pub fn launch_hold_rate_observer(
    &self,
    actions_dir: &CudaSlice<i32>,
    batch_size: i32,
    out_1: &mut CudaSlice<f32>,
) -> Result<(), MLError> {
    let bdim: u32 = 256;
    let cfg = LaunchConfig {
        grid_dim: (1, 1, 1),
        block_dim: (bdim, 1, 1),
        shared_mem_bytes: bdim * std::mem::size_of::<i32>() as u32,
    };
    unsafe {
        self.hold_rate_observer
            .clone()
            .launch(cfg, (actions_dir, batch_size, out_1))?;
    }
    Ok(())
}
```

Add `hold_rate_observer: CudaFunction` and `hold_rate_buf: MappedF32Buffer` (1 float).

- [ ] **Step 10:** Wire the per-step launch + Pearls A+D smoothing + controller in `training_loop.rs` (after action selection, before next step):

```rust
// SP13 v3: Hold-rate observer + Hold-cost controller
self.launch_hold_rate_observer(
    &actions_dir,           /* [B] direction actions from policy this step */
    batch_size as i32,
    self.hold_rate_buf.device_slice_mut(),
)?;

// Pearls A+D smoothing (α=0.05) into ISV[HOLD_RATE_OBSERVED_EMA_INDEX]
self.launch_apply_pearls_ad(
    self.hold_rate_buf.device_slice(),
    1,
    HOLD_RATE_OBSERVED_EMA_INDEX,
    /* alpha */ 0.05,
)?;

// Controller (deficit-driven): when observed > target, raise cost
let observed = self.read_isv_slot(HOLD_RATE_OBSERVED_EMA_INDEX);
let target   = self.read_isv_slot(HOLD_RATE_TARGET_INDEX);
let excess   = (observed - target).max(0.0);
let hold_cost = HOLD_COST_BASE * (1.0 + HOLD_COST_CONTROLLER_GAIN * excess);
let hold_cost = hold_cost.clamp(
    HOLD_COST_BASE * HOLD_COST_FLOOR_RATIO,
    HOLD_COST_BASE * HOLD_COST_CEIL_RATIO,
);
self.write_isv_slot(HOLD_COST_INDEX, hold_cost)?;
```

(Match existing patterns for ISV reads/writes in this file. The `apply_pearls_ad_kernel` API may differ slightly; consult adjacent SP11/SP12 launches for the exact signature.)

- [ ] **Step 11:** Add per-bar Hold cost subtraction in `experience_kernels.cu`. Locate the reward composition site (SP12 asymmetric cap is around line 2788 per `pearl_symmetric_clamp_audit`):

```cuda
// SP13 v3: per-bar Hold cost (priced action — ISV-driven adaptive)
if (action_dir == DIR_HOLD) {
    base_reward -= isv[HOLD_COST_INDEX];
}

// (then existing SP12 asymmetric cap, etc.)
float r = fmaxf(-10.0f, fminf(base_reward, 5.0f));
```

(Adapt to actual variable names at the call site. Hold cost subtraction goes BEFORE the SP12 asymmetric cap so the cost participates in the bounded clamp — keeps total reward in the same range.)

- [ ] **Step 12:** Add 2 GPU oracle tests to `sp13_phase0_oracle_tests.rs`:

```rust
#[test]
fn hold_rate_observer_all_hold_returns_one() {
    // Batch where every action == DIR_HOLD (=1) → out = 1.0
    let actions = vec![1_i32; 10];
    let out = run_hold_rate_observer(&actions);
    assert!((out - 1.0).abs() < 1e-6);
}

#[test]
fn hold_rate_observer_no_hold_returns_zero() {
    // Batch with mix of Short(0)/Long(2)/Flat(3), no Hold(1) → out = 0.0
    let actions = vec![0_i32, 2, 3, 0, 2, 3];
    let out = run_hold_rate_observer(&actions);
    assert!(out.abs() < 1e-6);
}

// (Optional 3rd: half-hold case verifying out = 0.5 — implementer adds if scaffold permits)
```

- [ ] **Step 13:** Verify on RTX 3050 Ti:

```bash
SQLX_OFFLINE=true cargo check -p ml --lib  # clean
SQLX_OFFLINE=true cargo build -p ml --release  # cubin compiles
SQLX_OFFLINE=true cargo test -p ml --features cuda --test sp13_phase0_oracle_tests -- --ignored
# Expected: 8 passing (6 from T2 + 2 new)
SQLX_OFFLINE=true cargo test -p ml --features cuda --test sp12_reward_math_tests -- --ignored  # SP12 still green
```

- [ ] **Step 14:** Stage all P0a.T3 v3 changes (atomic with T1, T2, T4 in P0a.T5 commit):

```bash
git add crates/ml/src/cuda_pipeline/sp13_isv_slots.rs \
        crates/ml/src/cuda_pipeline/sp5_isv_slots.rs \
        crates/ml/src/cuda_pipeline/state_layout.cuh \
        crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \
        crates/ml/src/cuda_pipeline/hold_rate_observer_kernel.cu \
        crates/ml/src/cuda_pipeline/experience_kernels.cu \
        crates/ml/src/trainers/dqn/state_reset_registry.rs \
        crates/ml/src/trainers/dqn/trainer/training_loop.rs \
        crates/ml/build.rs \
        crates/ml/tests/sp13_phase0_oracle_tests.rs
```

---

### Task P0a.T4: Wire dir_acc launch + ISV write + HEALTH_DIAG

**Files:** Modify `training_loop.rs` (validation/eval step where aux head is forward'd)

- [ ] **Step 1:** Locate existing aux HEALTH_DIAG site:

```bash
grep -n "aux_nb_mse\|aux_rg_ce" crates/ml/src/trainers/dqn/trainer/training_loop.rs
```

- [ ] **Step 2:** After existing aux forward on validation batch, add:

```rust
// SP13 P0a: directional accuracy probe (regression-mode aux scalar in P0a; Layer B → softmax)
self.launch_aux_dir_acc_reduce(
    &fused.aux_pred,
    &fused.next_bar_label,
    batch_size as i32,
    self.aux_dir_acc_buf.device_slice_mut(),
)?;

// Pearls A+D smoothing on slot 0 (dir_acc) — short EMA (α=0.3) → slot 373
self.launch_apply_pearls_ad(
    self.aux_dir_acc_buf.device_slice_at(0, 1),
    1,
    AUX_DIR_ACC_SHORT_EMA_INDEX,
    /* alpha */ 0.3,
)?;
// Slow EMA (α=0.05) → slot 374 (stagnation detector)
self.launch_apply_pearls_ad(
    self.aux_dir_acc_buf.device_slice_at(0, 1),
    1,
    AUX_DIR_ACC_LONG_EMA_INDEX,
    /* alpha */ 0.05,
)?;

// Per-bar logit signal write for direction Q-head input (slot 375)
// In P0a: regression scalar squashed to [-1, +1] via tanh; Layer B replaces with softmax logit_diff
self.launch_aux_pred_to_isv_tanh(
    &fused.aux_pred,
    AUX_DIR_PREDICTION_INDEX,
)?;

let dir_acc_arr: [f32; 3] = self.aux_dir_acc_buf.read_pinned()?;
tracing::info!(
    "HEALTH_DIAG aux_dir_acc accuracy={:.4} pos_pred={:.4} pos_label={:.4}",
    dir_acc_arr[0], dir_acc_arr[1], dir_acc_arr[2],
);
```

(Adjust `apply_pearls_ad` signature to match project conventions; if it takes a single α with two slots needing different rates, write a small helper or two calls.)

The `launch_aux_pred_to_isv_tanh` is a 1-line per-bar kernel — write it inline if no
existing helper fits, ~20 LOC.

- [ ] **Step 3:** Verify compile + clippy clean.

- [ ] **Step 4:** Stage.

---

### Task P0a.T5: Atomic Phase 0a commit + smoke

- [ ] **Step 1:** Final compile + tests:

```bash
SQLX_OFFLINE=true cargo check --workspace 2>&1 | tail
SQLX_OFFLINE=true cargo build -p ml --release 2>&1 | grep -E "error|warning" | head
SQLX_OFFLINE=true cargo test -p ml --features gpu-tests --test sp13_phase0_oracle_tests
SQLX_OFFLINE=true cargo test -p ml --features gpu-tests --test sp12_reward_math_tests  # SP12 still green
```

- [ ] **Step 2:** Sanity grep for Hold leaks (Step 10 of T3) — must be zero.

- [ ] **Step 3:** Commit:

```bash
git commit -m "$(cat <<'EOF'
feat(sp13): P0a — Hold elimination + dir_acc instrumentation (atomic)

Tests user's hypothesis (Hold hides directional signal) by removing Hold from the
direction action space. 4 → 3-way: NUM_DIRECTIONS, DirectionAction enum {Long,
Short, Flat}, Q-head output dim, C51/IQN per-direction-bucket state, action
selection argmax, replay buffer action encoding, HEALTH_DIAG val_dir_dist 3-bucket.

Adds aux_dir_acc_reduce GPU kernel + 2 EMA slots (short α=0.3 / long α=0.05) for
dir_acc tracking, plus the aux logit-signal slot wired to ISV[375] via tanh of
the existing regression head (Layer B will replace with softmax logit_diff).

8 SP13 ISV slots [372..380); SP5_SLOT_END = 380. Per-direction slot 3 dead-space
documented; recompaction deferred to close-out.

Greenfield: replay buffer / checkpoints invalidated by design.

Tests: 10 GPU oracle tests in sp13_phase0_oracle_tests (6 dir_acc + 4 action select).

Spec: docs/superpowers/specs/2026-05-04-sp13-redefine-success-for-predictive-skill.md (v2)

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
EOF
)"
```

- [ ] **Step 4:** Push + 10-epoch L40S smoke:

```bash
git push --set-upstream origin sp11-reward-as-controlled-subsystem
./scripts/argo-train.sh --model dqn --commit $(git rev-parse HEAD) --epochs 10 --pool ci-training-l40s
```

Capture workflow name. Monitor early per `feedback_kill_runs_on_anomaly_quickly`:
- ep1 HEALTH_DIAG `aux_dir_acc` line emitting (slot 373 wired)
- val_dir_dist 3-bucket (no `hold=` field)
- aux_w stays at SP11 inverted-formula values (~0.09) — controller unchanged in P0a
- No NaN/inf

Kill on anomaly.

- [ ] **Step 5:** Phase 0a gate decision (post-completion):

```bash
argo logs <wf> -n foxhunt --tail=2000 | grep -E "HEALTH_DIAG.*epoch=9|aux_dir_acc|val_dir_dist|win_rate" | tail -30
```

Apply gate:

| `dir_acc` trajectory | `dir_acc[ep9]` | `win_rate[ep9]` | Decision |
|---|---|---|---|
| ↑ rising | > 0.52 | > 0.51 | ✅ **CONFIRMED** — proceed Layer B (skip P0b) |
| flat ~0.50 | < 0.51 | < 0.50 | proceed P0b (test aux amplification) |
| mixed (one ↑, one flat) | — | — | proceed P0b |

Update `project_hold_hides_directional_signal_hypothesis.md` with smoke ID, commit SHA, final dir_acc and WR.

If CONFIRMED → skip to **Layer B**. If proceeding to P0b → next section.

---

## Phase 0b — Aux amplification probe (1 commit, only on partial/red P0a)

### Task P0b.T1: Replace SP11 inverted aux_w with deficit + stagnation controller

**Files:** Modify `training_loop.rs:4103-4104`; modify `sp13_isv_slots.rs` to add base/decay constants

- [ ] **Step 1:** Add constants to `sp13_isv_slots.rs`:

```rust
pub const AUX_W_BASE: f32 = 0.5;
pub const AUX_W_DEFICIT_GAIN: f32 = 5.0;   // aux_w = base × (1 + 5 × deficit)
pub const AUX_W_STAGNATION_DECAY: f32 = 0.7;  // (1 - 0.7 × stagnation)
pub const AUX_W_HARD_FLOOR_RATIO: f32 = 0.3;  // never below 0.3 × base
pub const AUX_W_HARD_CEIL_RATIO: f32 = 3.0;   // never above 3.0 × base
```

- [ ] **Step 2:** Replace the SP11-era formula at `training_loop.rs:4103-4104`:

```rust
// SP13 P0b: deficit + stagnation aux_w controller. Replaces SP11 inverted
// `0.1 × health × (1 - sharpe)` formula which lowered aux_w when sharpe was
// high — exactly wrong for the failure mode (high sharpe + low WR).
let target = self.read_isv_slot(TARGET_DIR_ACC_INDEX);          // 0.55
let short_ema = self.read_isv_slot(AUX_DIR_ACC_SHORT_EMA_INDEX); // fast (α=0.3)
let long_ema  = self.read_isv_slot(AUX_DIR_ACC_LONG_EMA_INDEX);  // slow (α=0.05)

let deficit     = (target - short_ema).max(0.0);
let improvement = (short_ema - long_ema).max(0.0);
// stagnation = 1 means no improvement; 0 means improving fast
let stagnation = if deficit > 0.005 {
    (1.0 - improvement / deficit).clamp(0.0, 1.0)
} else {
    0.0  // already at target — not stagnant, just satisfied
};

let aux_w_raw = AUX_W_BASE
    * (1.0 + AUX_W_DEFICIT_GAIN * deficit)
    * (1.0 - AUX_W_STAGNATION_DECAY * stagnation);
let aux_w = aux_w_raw.clamp(
    AUX_W_BASE * AUX_W_HARD_FLOOR_RATIO,
    AUX_W_BASE * AUX_W_HARD_CEIL_RATIO,
);
fused.set_aux_weight(aux_w);

tracing::debug!(
    "aux_w controller: target={:.3} short={:.3} long={:.3} deficit={:.4} stag={:.3} aux_w={:.3}",
    target, short_ema, long_ema, deficit, stagnation, aux_w,
);
```

Imports needed:
```rust
use crate::cuda_pipeline::sp13_isv_slots::{
    TARGET_DIR_ACC_INDEX, AUX_DIR_ACC_SHORT_EMA_INDEX, AUX_DIR_ACC_LONG_EMA_INDEX,
    AUX_W_BASE, AUX_W_DEFICIT_GAIN, AUX_W_STAGNATION_DECAY,
    AUX_W_HARD_FLOOR_RATIO, AUX_W_HARD_CEIL_RATIO,
};
```

- [ ] **Step 3:** Verify compile + 10-bar simulation in head:
  - At fold start: short=long=0.5, deficit=0.05, improvement=0, stagnation=1.0, aux_w = 0.5 × 1.25 × 0.3 = 0.1875 → clamped to 0.15 (floor) — gentle start, controller waits for data
  - After improvement (short=0.55, long=0.50): deficit=0, improvement=0.05, aux_w = 0.5 × 1.0 × 1.0 = 0.5 — back to base, satisfied
  - After stagnation (short=0.50, long=0.50, target=0.55): deficit=0.05, improvement=0, stagnation=1.0, aux_w = floor 0.15 — backed off, no destabilisation

- [ ] **Step 4:** Commit:

```bash
git commit -m "$(cat <<'EOF'
feat(sp13): P0b — aux_w deficit + stagnation controller (replaces SP11 inverted)

Replaces SP11-era `aux_w = 0.1 × health × (1 - sharpe_tanh).clamp(0.05, 0.3)` at
training_loop.rs:4103-4104 with deficit-driven amplification + stagnation decay.

Mechanism:
  deficit    = max(0, target_dir_acc - short_ema)
  improvement = max(0, short_ema - long_ema)
  stagnation  = (deficit > 0) ? clamp(1 - improvement/deficit, 0, 1) : 0
  aux_w = base × (1 + 5 × deficit) × (1 - 0.7 × stagnation), bounded [0.3×base, 3×base]

Prevents the v1 failure mode where aux_w would push to hard cap forever in the
data-limited case (no signal → no improvement → controller stuck at 1.5× base
destabilising Q-head). Stagnation decay backs off when controller can't make progress.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
EOF
)"
git push
```

- [ ] **Step 5:** 10-epoch L40S smoke. Same gate as P0a. If still red → halt SP13, update hypothesis memory to DISCONFIRMED, plan SP14 separately.

---

## Layer B — Aux head regression → binary classification (1 atomic commit, only on green Phase 0)

### Implementation notes (discovered during P0a, 2026-05-04)

Concrete codebase locations the Layer B implementer needs (saves re-investigation):

- **Aux head forward**: `crates/ml/src/cuda_pipeline/aux_heads_kernel.cu` — `aux_heads_fwd.forward_next_bar(...)` produces `aux_nb_pred_buf [B]` (regression scalar). Layer B changes output dim 1 → 2 (softmax logits).
- **Aux head loss**: `crates/ml/src/cuda_pipeline/aux_heads_loss_ema_kernel.cu` — `next_bar_loss_reduce` computes MSE against `aux_nb_label_buf` and reduces. Layer B changes loss to CrossEntropy.
- **Label generation**: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs:15249-15280` — `aux_nb_label_buf` is gathered as **column 0 of next_states** (= `log_return` per pipeline.rs, first feature in MARKET group). Continuous signed value. Uses `vsn_logit_gather_kernel` with `col=0`. Layer B replaces with **sign-class** label generation: `class = (price[t+30] - price[t] > 0) ? 1 : 0`. Either modify the gather to apply sign, OR introduce a new `aux_nb_label_classify_kernel` reading next_states or price tensor directly.
- **WARNING from history** (`gpu_dqn_trainer.rs:15252`): "Use the DEDICATED `aux_nb_label_buf` (NOT aliased over a backward partial — the previous WIP aliased `aux_partial_nb_b2` here and triggered the F1/F2 regression)." Layer B must NOT alias the new 2-class label buffer over any backward-partial scratch.
- **Aux pred consumer** (P0a placeholder): `crates/ml/src/cuda_pipeline/aux_pred_to_isv_tanh_kernel.cu` reads scalar `aux_pred [B]`, computes `mean(tanh(aux_pred))` → ISV[375]. Layer B replaces with `mean(softmax[1] - softmax[0])` reduction (logit-diff ∈ [−1, +1]). Either modify in place OR create `aux_pred_softmax_to_isv_logit_diff_kernel.cu` and update the launcher to point to the new cubin. Cleanest: rewrite `aux_pred_to_isv_tanh_kernel.cu` since it's a P0a placeholder per its own header comment.
- **dir_acc consumer**: `crates/ml/src/cuda_pipeline/aux_dir_acc_reduce_kernel.cu` reads scalar `aux_pred[i]` and tests `(p > 0.0f) ? 1 : 0`. Layer B changes to read 2-class softmax: `pred_pos = (softmax[i*2+1] > softmax[i*2+0]) ? 1 : 0`. Update the 6 oracle tests in `sp13_phase0_oracle_tests.rs` to construct softmax pairs instead of scalar predictions.
- **Direction Q-head input concat**: `crates/ml/src/cuda_pipeline/dqn_inference.cu` direction Q-head input layer currently consumes `[trunk_features]`. Layer B concatenates `ISV[AUX_DIR_PREDICTION_INDEX]` (single scalar = mean logit-diff) → input dim grows by 1. Update `dqn_param_layout.rs` direction Q-head weight shape `(in_dim, NUM_DIRECTIONS)` → `(in_dim + 1, NUM_DIRECTIONS)`. Bump layout fingerprint seed to invalidate cached params (greenfield acceptable).
- **No information leakage**: aux uses state s_t at inference; price[t+30] only at training-time (hindsight labels). Same pattern as the existing aux MSE head — no new leakage path introduced.

### Why this is the right next move (rationale post-P0a)

P0a empirics (`train-67gqb`):
- aux_dir_acc climbed 0.149 → 0.483 in 10 epochs at aux_w=0.05 (suppressed)
- The trajectory has decaying deltas (0.107 → 0.069 → 0.054 → ... → 0.005) — signature of bias-term decay toward zero, NOT learning the regression task
- val_win_rate stuck at 0.4638 — the Q-head can't extract directional signal from a regression head that isn't actually learning

Layer B + P0b together address this:
- P0b: aux_w 0.05 → 0.5 gives the head 10× more gradient → it can actually learn the task
- Layer B: 2-class classification with cross-entropy is a STRONGER training signal than MSE on a continuous target (especially when the target sign is what we care about — not its magnitude)
- Layer B: Q-head input now sees the head's confidence directly (logit-diff ∈ [−1, +1]) instead of a `tanh(scalar)` of a noisy regression — explicit predictive feature

Combined effect: Q-head should start consuming the aux signal within ~5-10 epochs after Layer B lands, surfacing as WR > 0.50.

### Task B.T1: Aux head softmax + cross-entropy

**Files:** `aux_head_kernel.cu`, `dqn_param_layout.rs`

- [ ] **Step 1:** Update aux head forward:

```cuda
// Output dim 1 (regression scalar) → 2 (softmax logits over [neg, pos])
extern "C" __global__ void aux_head_forward_kernel(
    /* ... existing inputs ... */
    float* __restrict__ aux_softmax_out  // [B, 2]
) {
    /* ... existing trunk → linear → 2 logits ... */
    float logit_neg = /* ... */;
    float logit_pos = /* ... */;
    float m = fmaxf(logit_neg, logit_pos);
    float exp_neg = expf(logit_neg - m);
    float exp_pos = expf(logit_pos - m);
    float Z = exp_neg + exp_pos;
    aux_softmax_out[b * 2 + 0] = exp_neg / Z;
    aux_softmax_out[b * 2 + 1] = exp_pos / Z;
}
```

- [ ] **Step 2:** Update aux head backward — CrossEntropy gradient is `softmax - one_hot(label)`:

```cuda
// d_logit[c] = softmax[c] - (label == c ? 1 : 0)
```

- [ ] **Step 3:** Update param shapes: aux head output features 1 → 2 in `dqn_param_layout.rs`. Bump layout fingerprint seed.

- [ ] **Step 4:** Stage.

---

### Task B.T2: Label generation kernel (sign of next-30-bar return)

**Files:** `aux_label_gen_kernel.cu` (modify or create)

- [ ] **Step 1:** Update / write kernel:

```cuda
extern "C" __global__ void aux_label_gen_kernel(
    const float* __restrict__ price,     // [B, T]
    int t_index,
    int horizon,                          // 30 bars from spec
    int t_max,                            // T (sequence length)
    int* __restrict__ label_class_out,    // [B] {0=negative, 1=positive}; bars too close to end set to -1 (skip)
    int* __restrict__ label_valid_out     // [B] 0/1 mask
) {
    int b = blockIdx.x * blockDim.x + threadIdx.x;
    if (b >= /* B */) return;
    int t_target = t_index + horizon;
    if (t_target >= t_max) {
        label_class_out[b] = -1;
        label_valid_out[b] = 0;
        return;
    }
    float ret = price[b * t_max + t_target] - price[b * t_max + t_index];
    if (ret > 0.0f) {
        label_class_out[b] = 1;
        label_valid_out[b] = 1;
    } else if (ret < 0.0f) {
        label_class_out[b] = 0;
        label_valid_out[b] = 1;
    } else {
        label_class_out[b] = -1;  // skip flat bars
        label_valid_out[b] = 0;
    }
}
```

CrossEntropy loss masks invalid bars (multiply by `label_valid_out`).

**No information leakage**: aux head sees state s_t at inference; only training-time
labels use price[t+30]. Same hindsight pattern as existing next_bar_mse.

- [ ] **Step 2:** Wire label gen launch in training loop (replaces existing next-bar regression label gen).

- [ ] **Step 3:** Adapt `aux_dir_acc_reduce_kernel` to read 2-class softmax:

```cuda
// Layer B: softmax read instead of regression scalar
const float p_pos = aux_softmax[i * 2 + 1];
const float p_neg = aux_softmax[i * 2 + 0];
const int pred_pos = (p_pos > p_neg) ? 1 : 0;
// label comparison stays the same (sign-encoded label still works)
```

- [ ] **Step 4:** Stage.

---

### Task B.T3: Wire aux logit_diff → ISV[375] → direction Q-head input

**Files:** `training_loop.rs`, `dqn_inference.cu`, `dqn_param_layout.rs`

- [ ] **Step 1:** Replace P0a `tanh(scalar)` write with softmax logit_diff:

```rust
self.launch_aux_softmax_to_logit_diff_isv(
    &fused.aux_softmax,  // [B, 2]
    AUX_DIR_PREDICTION_INDEX,
)?;
```

Tiny kernel:
```cuda
extern "C" __global__ void aux_softmax_to_logit_diff_kernel(
    const float* __restrict__ aux_softmax,  // [B, 2]
    int batch_size,
    float* __restrict__ isv,
    int isv_stride
) {
    int b = blockIdx.x * blockDim.x + threadIdx.x;
    if (b >= batch_size) return;
    float diff = aux_softmax[b * 2 + 1] - aux_softmax[b * 2 + 0];  // ∈ [-1, +1]
    isv[b * isv_stride + AUX_DIR_PREDICTION_INDEX] = diff;
}
```

- [ ] **Step 2:** Concat `ISV[AUX_DIR_PREDICTION_INDEX]` into direction Q-head input:

In `dqn_inference.cu`, locate the direction Q-head forward — input layer was `[trunk_features]`, becomes `[trunk_features, aux_logit_diff]`. In_dim grows by 1.

In `dqn_param_layout.rs`, update direction Q-head weight shape: `(in_dim, 3)` → `(in_dim + 1, 3)`. Re-init.

- [ ] **Step 3:** Add 3 GPU oracle classification tests:

```rust
#[test] fn aux_softmax_logit_diff_positive_class() { /* logits=[-1, 1] → diff>0 */ }
#[test] fn aux_softmax_logit_diff_negative_class() { /* logits=[1, -1] → diff<0 */ }
#[test] fn aux_cross_entropy_grad_sign_correct() { /* softmax-onehot sign per class */ }
```

- [ ] **Step 4:** Atomic Layer B commit:

```bash
git commit -m "$(cat <<'EOF'
feat(sp13): Layer B — aux head regression → 30-bar binary classification

Replaces aux MSE on next-bar return with CrossEntropy on next-30-bar return sign.
Aux output dim 1 → 2 + softmax. logit_diff = softmax[1] - softmax[0] ∈ [-1, +1]
written to ISV[AUX_DIR_PREDICTION_INDEX] each bar, concatenated into direction
Q-head input layer (in_dim grows by 1). aux_dir_acc_reduce_kernel adapted to read
softmax instead of regression scalar.

No information leakage: aux uses state s_t at inference; price[t+30] only at
training time (hindsight labels). Same pattern as the existing aux MSE head.

Tests: 13 in sp13_phase0_oracle_tests (10 P0a + 3 classification).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
EOF
)"
git push
```

- [ ] **Step 5:** 5-epoch L40S smoke verifying:
  - Aux CrossEntropy loss decreasing
  - val_dir_acc continues at or above Phase 0 level
  - WR not regressing
  - aux signal feeding Q-head: perturbation test (zero out ISV[375] for 1 bar in eval, observe direction action distribution shift)

If clean → proceed Layer C.

---

## Layer C — Direction-skill bonus + luck discount (1 atomic commit)

### Task C.T1: Device-inline reward functions in trade_physics.cuh

**Files:** `trade_physics.cuh`

- [ ] **Step 1:** Add 2 device-inline functions:

```cuda
__device__ __forceinline__ float compute_direction_skill_bonus_bounded(
    int action_dir,                  // 0=Long, 1=Short, 2=Flat
    float realized_30_bar_return,    // sign-encoded
    float alpha_capped,              // SP12 asymmetric-capped trade pnl, used for cap calc
    float bonus_alpha_param,         // ISV[376], default 1.0
    float bonus_beta_param,          // ISV[377], default 1.0
    float cap_ratio                  // ISV[379], default 0.3
) {
    if (action_dir == DIRECTION_FLAT) return 0.0f;
    int trade_sign = (action_dir == DIRECTION_LONG) ? +1 : -1;
    int label_sign = (realized_30_bar_return > 0.0f) ? +1
                   : ((realized_30_bar_return < 0.0f) ? -1 : 0);
    if (label_sign == 0) return 0.0f;  // no signal to score
    float skill_cap = cap_ratio * fabsf(alpha_capped);
    if (trade_sign == label_sign) {
        return +fminf(bonus_alpha_param, skill_cap);
    } else {
        return -fminf(bonus_beta_param, skill_cap);
    }
}

__device__ __forceinline__ float compute_luck_win_discount(
    float reward_pre_discount,
    float realized_trade_pnl,
    bool direction_was_correct,
    float discount_factor    // ISV[378], default 0.3
) {
    bool trade_won = realized_trade_pnl > 0.0f;
    if (trade_won && !direction_was_correct) {
        return reward_pre_discount * discount_factor;
    }
    return reward_pre_discount;
}
```

- [ ] **Step 2:** Stage.

---

### Task C.T2: Replace exit-event reward composition in experience_kernels.cu

**Files:** `experience_kernels.cu` (~line 2788, SP12 exit-event site)

- [ ] **Step 1:** Locate the SP12 reward composition site:

```bash
grep -n "fmaxf(-10\|asymmetric_capped\|min_hold_factor\|compute_min_hold_penalty" \
    crates/ml/src/cuda_pipeline/experience_kernels.cu
```

- [ ] **Step 2:** Layer in skill bonus + luck discount AFTER SP12's asymmetric cap, BEFORE min_hold_factor multiplication:

```cuda
// SP12 asymmetric cap (preserved per pearl_audit_unboundedness_for_implicit_asymmetry)
float alpha_capped = compute_asymmetric_capped_pnl(realized_trade_pnl, -10.0f, +5.0f);

// SP13 Layer C: direction-skill bonus (bounded relative to alpha)
int trade_sign  = (trade_direction == DIRECTION_LONG) ? +1
                : ((trade_direction == DIRECTION_SHORT) ? -1 : 0);
int label_sign_for_30bar = (realized_30_bar_return > 0.0f) ? +1
                         : ((realized_30_bar_return < 0.0f) ? -1 : 0);
bool dir_correct = (trade_sign != 0) && (trade_sign == label_sign_for_30bar);

float skill_bonus = compute_direction_skill_bonus_bounded(
    trade_direction,
    realized_30_bar_return,
    alpha_capped,
    isv[DIR_SKILL_BONUS_ALPHA_INDEX],
    isv[DIR_SKILL_BONUS_BETA_INDEX],
    isv[SKILL_BONUS_CAP_RATIO_INDEX]
);

float r_pre_discount = alpha_capped + skill_bonus;

// SP13 Layer C: luck-win discount
float r_post_discount = compute_luck_win_discount(
    r_pre_discount,
    realized_trade_pnl,
    dir_correct,
    isv[LUCK_WIN_DISCOUNT_INDEX]
);

// SP12 min-hold penalty (preserved — orthogonal to skill incentive)
float r = r_post_discount * min_hold_factor;
```

(Implementer adapts variable names to actuals at the call site.)

- [ ] **Step 3:** Stage.

---

### Task C.T3: GPU oracle tests — full quadrant matrix

**Files:** Create `sp13_reward_math_test_kernel.cu` (mirrors SP12 pattern), Create `sp13_reward_math_tests.rs`, Modify `build.rs`

- [ ] **Step 1:** Test wrapper kernel exposing the 2 device-inline funcs to host (mirror `sp12_reward_math_test_kernel.cu`).

- [ ] **Step 2:** 16 oracle tests covering the 8-quadrant matrix from spec PLUS edge cases:

```rust
// All defaults: α=β=1.0, cap_ratio=0.3, luck=0.3, SP12 cap [-10,+5]
// Direction encoding (post-P0a): Short=0, Long=1, Flat=2

#[test]
fn correct_long_winner_big() {
    // alpha=+5, dir=Long(1), label=+1 → bonus=+min(1.0, 1.5)=+1.0
    let r = run_skill(/*alpha=*/5.0, /*dir=*/1, /*label=*/1.0);
    assert!((r - 6.0).abs() < 1e-6);
}

#[test]
fn correct_long_winner_small() {
    // alpha=+0.5, dir=Long(1), bonus_cap = 0.3 × 0.5 = 0.15
    let r = run_skill(0.5, 1, 1.0);
    assert!((r - 0.65).abs() < 1e-6);
}

#[test]
fn correct_long_loser_small_still_negative() {
    // alpha=-0.5, dir=Long(1), bonus=+0.15 → r=-0.35 (NOT positive — verify NO gaming)
    let r = run_skill(-0.5, 1, 1.0);
    assert!(r < 0.0, "correct-long-loser must stay negative; got {}", r);
}

#[test]
fn correct_long_loser_big_dominated_by_loss() {
    // alpha=-10, dir=Long(1), bonus_cap=3.0 → bonus=+1.0 → r=-9.0
    let r = run_skill(-10.0, 1, 1.0);
    assert!((r - (-9.0)).abs() < 1e-6);
}

#[test]
fn wrong_long_winner_lucky_discounted() {
    // alpha=+5, dir=Long(1), label=-1 → bonus=-1.0; lucky win → × 0.3
    let r = run_full(/*pnl_realized=*/2.0, 5.0, 1, -1.0);  // realized > 0, dir wrong
    let pre = 5.0 - 1.0;  // 4.0
    let expected = pre * 0.3;  // 1.2
    assert!((r - expected).abs() < 1e-6);
}

#[test]
fn wrong_long_loser() {
    // alpha=-10, dir=Long(1), label=-1 → bonus=-1.0 → r=-11.0 (no luck discount on loss)
    let r = run_full(-5.0, -10.0, 1, -1.0);
    assert!((r - (-11.0)).abs() < 1e-6);
}

#[test]
fn flat_zero_skill_bonus() {
    let r = run_skill(2.0, 2 /* Flat */, 1.0);
    assert!((r - 2.0).abs() < 1e-6);  // alpha + 0 bonus
}

#[test]
fn zero_label_zero_skill_bonus() {
    let r = run_skill(2.0, 1 /* Long */, 0.0);  // label=0 → no signal to score
    assert!((r - 2.0).abs() < 1e-6);
}

#[test]
fn correct_short_winner() {
    let r = run_skill(5.0, 0 /* Short */, -1.0);  // short + falling = correct
    assert!((r - 6.0).abs() < 1e-6);
}

#[test]
fn wrong_short_winner_lucky_discounted() {
    // dir=Short(0), label=+1 → wrong; pnl_realized > 0 → lucky → × 0.3
    let r = run_full(2.0, 5.0, 0 /* Short */, 1.0);
    let pre = 5.0 - 1.0;
    let expected = pre * 0.3;
    assert!((r - expected).abs() < 1e-6);
}

#[test]
fn no_negative_ev_gaming_invariant() {
    // Across many small-alpha correct-direction losers, reward must stay < 0
    for alpha_loss in (-10..=-1).map(|i| i as f32 * 0.01) {  // -0.1, -0.09, ...
        let r = run_skill(alpha_loss, 1 /* Long */, 1.0);
        assert!(r < 0.0,
            "negative-EV gaming invariant violated: alpha={} reward={}",
            alpha_loss, r);
    }
}

// + 6 more covering: bonus magnitude calibration, cap ratio extremes,
//   luck discount boundary (pnl exactly 0), Flat bypass invariant, etc.
```

- [ ] **Step 3:** Run on RTX 3050 Ti, expect 16/16 passing.

- [ ] **Step 4:** Add 4 ISV reset entries (slots 376, 377, 378, 379) to `state_reset_registry.rs` (per-fold).

- [ ] **Step 5:** Add per-fold HEALTH_DIAG metric:

```rust
// At fold-end, after all per-fold metrics:
let bh_pnl = price_at_fold_end - price_at_fold_start;
let alpha_vs_bh = total_pnl - bh_pnl;
tracing::info!(
    "HEALTH_DIAG fold_alpha_vs_buy_and_hold fold={} alpha={:.4} bh_pnl={:.4} our_pnl={:.4}",
    fold_idx, alpha_vs_bh, bh_pnl, total_pnl,
);
```

(Metric only — never feeds into reward. Per Major-1 fix.)

- [ ] **Step 6:** Atomic Layer C commit:

```bash
git commit -m "$(cat <<'EOF'
feat(sp13): Layer C — direction-skill bonus + luck-win discount (atomic)

Adds bounded skill bonus and luck-win discount on top of SP12 asymmetric-capped
P&L. Per spec v2 Major-3 fix, skill bonus is bounded relative to |alpha| via
cap_ratio (default 0.3) to prevent reward gaming on small-alpha correct-direction
losers.

  alpha_capped = SP12 cap [-10, +5]
  skill_bonus  = sign × min(α_param, cap_ratio × |alpha|)  if action ∈ {L,S} & label ≠ 0 else 0
  luck_disc    = (won AND wrong_dir) ? × 0.3 : 1.0
  reward       = (alpha_capped + skill_bonus) × luck_disc × min_hold_factor

4 new ISV slots [376, 377, 378, 379] populated in state-reset registry. Tests:
16-test quadrant matrix in sp13_reward_math_tests verifying no negative-EV gaming.

Per-fold buy-and-hold delta added to HEALTH_DIAG as METRIC ONLY (per spec v2
Major-1: alpha-vs-benchmark dropped as reward signal due to mathematical
tautology against always-long benchmark).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
EOF
)"
git push
```

---

## Layer D — 30-epoch validation + close-out

### Task D.T1: 30-epoch L40S validation

- [ ] **Step 1:** Push + deploy:

```bash
./scripts/argo-train.sh --model dqn --commit $(git rev-parse HEAD) --epochs 30 --pool ci-training-l40s
```

- [ ] **Step 2:** Optional 3-action SP12 control run (~€1) for fair before/after baseline:

Check out a parallel branch from pre-Layer-B SP12 + 3-action commit, deploy 30-epoch, compare.

- [ ] **Step 3:** Monitor early; kill on anomaly. Capture final-epoch HEALTH_DIAG metrics.

---

### Task D.T2: Close-out commit (audit doc + 3 new pearls + memory updates)

**Files:** `docs/dqn-wire-up-audit.md`, 3 new pearl files, `MEMORY.md`, `project_hold_hides_directional_signal_hypothesis.md`

- [ ] **Step 1:** Audit doc — SP13 entry: hypothesis, P0a result, P0b result (if run), Layer B/C deltas, 30-epoch metrics, control comparison (if run).

- [ ] **Step 2:** Write 3 pearls:

- `pearl_redefine_success_for_predictive_skill.md` — when policy plateaus on a structural-extraction objective across many tunes (12 across SP1-SP11 didn't move 46% WR), the fix is redefining the success criterion, not tuning. Canonical SP13.

- `pearl_skill_bonus_must_be_alpha_bounded.md` — direction-skill bonuses must be bounded relative to alpha magnitude (e.g. `min(α_param, ratio × |alpha|)`) to prevent reward gaming on small-alpha correct-direction losers. Worked example from SP13 v1 review (where unbounded α=1.0 + 0.5 alpha = -0.5 trade gave +0.5 reward, model rationally entered expected losers).

- `pearl_reward_quadrant_audit_required.md` — meta-pearl: any reward-composition change must include a 4-quadrant worked-example table (correct/wrong × winner/loser) showing all signs and the no-negative-EV-gaming invariant. Caught the SP13 v1 alpha-vs-benchmark math tautology and the unbounded skill bonus gaming.

- [ ] **Step 3:** Update `MEMORY.md` index: 3 new pearls under "Architectural constraints".

- [ ] **Step 4:** Update `project_hold_hides_directional_signal_hypothesis.md`:
- Status: UNVERIFIED → CONFIRMED / DISCONFIRMED / PARTIAL
- Phase 0a smoke ID, commit SHA, dir_acc, WR, val_dir_dist
- Phase 0b smoke if run

- [ ] **Step 5:** Commit + push.

---

## Estimated cost

| Phase | LOC | Time | Cost |
|---|---|---|---|
| P0a | ~300 | 4-5h | local |
| P0a smoke | — | 30min | ~€0.60 (10-epoch) |
| P0b (conditional) | ~80 | 1h | local |
| P0b smoke (conditional) | — | 30min | ~€0.60 |
| Layer B | ~120 | 2-3h | + ~€0.30 (5-epoch) |
| Layer C | ~150 | 2-3h | local |
| Layer D | — | 1h | ~€1.00 (30-epoch) + ~€1.00 control (recommended) |
| **Best case** (P0a green, no P0b) | **~570** | **~10h** | **~€2.90** |
| **Worst confirm** (P0a partial → P0b green) | **~650** | **~12h** | **~€3.50** |
| **Disconfirm** (both red) | **~380** | **~6h** | **~€1.20** |

Cheap-fail at ~€1.20 if hypothesis disconfirmed. Confirm cost ~€3 paid against the
chance of finding 55%+ WR (the prize across the entire SP1-SP12 lineage).
