# ml-alpha Phase B: PPO Policy + Walk-Forward CV Implementation Plan

> **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:** On top of the frozen CfC perception from Plan 1, build the CfC actor-critic policy, on-device replay buffer, PPO training step with EWC anchor, the seven ISV-driven controllers, kill-switch + atomic weights swap, and the 6-fold walk-forward CV runner. Pass criterion: mean test Sharpe across folds 1-4 strictly > 0 at IBKR commission (0.034 price units / contract), with `std < mean` across folds.

**Architecture:** Phase A perception weights frozen and loaded read-only. Policy network (CfC actor + CfC critic, ~11K params) trained via PPO with GAE on multi-env rollouts (64 parallel envs). EWC term anchored to Phase A perception's "downstream consumer" Fisher (computed during a one-shot pass on Phase A val). All forward + backward + AdamW + GAE + advantage normalization + loss fused into CUDA Graph C. Kill-switch reads ISV diagnostics + rolls back weights atomically. Seven controllers (entropy, lr_policy, λ_ewc, clip_eps, value_coef, replay mix, C_vol) drive hyperparameters via Wiener-α blending with permanent floors.

**Tech Stack:** Same as Plan 1 (Rust 1.85+, cudarc 0.19, nvcc cubins). Adds cuRAND for action sampling. Reuses Plan 1's `IsvBus`, `MappedPinnedSnapshotSlot`/`FillSlot`, `CfcTrunk` (frozen).

**Source spec:** `docs/superpowers/specs/2026-05-16-ml-alpha-cfc-ppo-design.md` Sections 2 (policy), 3 (action+reward+state), 4 (Phase B + walk-forward), 5.5 Graphs B/C, 6 (online learning + kill-switch + cold-start), 7 (testing + ISV controllers).

**Prerequisite:** Plan 1 has merged with `phase_a_gate.json` showing PASS verdict on `2024Q1+Q2`.

---

## File structure

| File | Status | Responsibility |
|---|---|---|
| `crates/ml-alpha/src/policy/mod.rs` | Create | Policy module surface. |
| `crates/ml-alpha/src/policy/network.rs` | Create | `PolicyNetwork` — CfC actor + CfC critic, weights + grad buffers + AdamW state. |
| `crates/ml-alpha/src/policy/sampler.rs` | Create | Action sampling (cuRAND categorical with logit mask). |
| `crates/ml-alpha/src/policy/state.rs` | Create | 20-dim AlphaState assembly + build_state kernel binding. |
| `crates/ml-alpha/src/replay/mod.rs` | Create | Replay module. |
| `crates/ml-alpha/src/replay/buffer.rs` | Create | `ReplayBuffer` — GPU-resident offline + live ring, `sample_mixed`. |
| `crates/ml-alpha/src/replay/io.rs` | Create | Load offline trajectories from disk, write QuestDB rows for live. |
| `crates/ml-alpha/src/rollout/mod.rs` | Create | Multi-env rollout collection. |
| `crates/ml-alpha/src/rollout/env.rs` | Create | `RolloutEnv` — 64-parallel position/PnL state on device. |
| `crates/ml-alpha/src/ppo/mod.rs` | Create | PPO module. |
| `crates/ml-alpha/src/ppo/gae.rs` | Create | GAE returns kernel binding. |
| `crates/ml-alpha/src/ppo/loss.rs` | Create | Fused PPO clip + value MSE + entropy + EWC loss. |
| `crates/ml-alpha/src/ppo/update.rs` | Create | Update orchestration + atomic weights swap. |
| `crates/ml-alpha/src/ewc/mod.rs` | Create | EWC anchor + Fisher computation. |
| `crates/ml-alpha/src/controllers/mod.rs` | Create | Seven ISV controllers. |
| `crates/ml-alpha/src/controllers/entropy.rs` | Create | Entropy-deficit controller. |
| `crates/ml-alpha/src/controllers/lr.rs` | Create | Loss-slope LR controller. |
| `crates/ml-alpha/src/controllers/ewc_lambda.rs` | Create | Weight-drift λ_ewc controller. |
| `crates/ml-alpha/src/controllers/clip_eps.rs` | Create | KL-targeting clip_eps controller. |
| `crates/ml-alpha/src/controllers/value_coef.rs` | Create | Value/policy ratio controller. |
| `crates/ml-alpha/src/controllers/replay_mix.rs` | Create | Reward-gap mix controller. |
| `crates/ml-alpha/src/controllers/vol_penalty.rs` | Create | Vol-target C_vol controller. |
| `crates/ml-alpha/src/kill_switch.rs` | Create | Rolling-baseline σ-check + Phase B bootstrap baseline. |
| `crates/ml-alpha/src/trainer/phase_b.rs` | Create | Phase B trainer loop. |
| `crates/ml-alpha/src/cv/mod.rs` | Create | Walk-forward fold splitting + runner. |
| `crates/ml-alpha/cuda/build_state.cu` | Create | 20-dim AlphaState assembly. |
| `crates/ml-alpha/cuda/policy_forward.cu` | Create | Actor + critic fused forward. |
| `crates/ml-alpha/cuda/sample_action.cu` | Create | cuRAND categorical + mask. |
| `crates/ml-alpha/cuda/write_replay_slot.cu` | Create | Write transition to ring. |
| `crates/ml-alpha/cuda/ingest_fill.cu` | Create | Update position/PnL + back-fill reward. |
| `crates/ml-alpha/cuda/gae_returns.cu` | Create | Sequential prefix scan for GAE. |
| `crates/ml-alpha/cuda/advantage_normalize.cu` | Create | Z-score advantages. |
| `crates/ml-alpha/cuda/ppo_loss.cu` | Create | Fused PPO loss + backward. |
| `crates/ml-alpha/cuda/kill_switch_check.cu` | Create | σ-distance check. |
| `crates/ml-alpha/examples/alpha_train_phase_b.rs` | Create | CLI for Phase B per-fold training. |
| `crates/ml-alpha/examples/alpha_walk_forward.rs` | Create | CLI for 6-fold walk-forward CV runner. |
| `crates/ml-alpha/tests/policy_forward_bit_equiv.rs` | Create | Actor/critic forward CPU oracle vs GPU. |
| `crates/ml-alpha/tests/gae_bit_equiv.rs` | Create | GAE reference impl vs GPU kernel. |
| `crates/ml-alpha/tests/ppo_loss_finite_diff.rs` | Create | PPO loss backward vs finite difference. |
| `crates/ml-alpha/tests/replay_buffer.rs` | Create | sample_mixed proportions + offline-only fallback. |
| `crates/ml-alpha/tests/rollout_reward.rs` | Create | Synthetic 64-env rollout produces non-zero rewards. |
| `crates/ml-alpha/tests/kill_switch.rs` | Create | σ-rejection + rolling-vs-Phase-B-baseline switchover. |
| `crates/ml-alpha/tests/controller_invariants.rs` | Create | Each controller respects floor, monotonicity, and first-observation bootstrap. |
| `crates/ml-alpha/tests/walk_forward_smoke.rs` | Create | 1-fold smoke (Q1 train, Q2 test, ~1M snapshots) produces finite Sharpe. |
| `infra/k8s/argo/alpha-phase-b-template.yaml` | Create | Walk-forward CV workflow (6 sequential folds on L40S). |
| `scripts/argo-alpha-phase-b.sh` | Create | Submission script. |
| `phase_b_baseline.json` | Artifact | Emitted by last fold's trainer; consumed by live cold-start. |

---

## Task 1: `build_state` kernel — 20-dim AlphaState assembly

**Files:**
- Create: `crates/ml-alpha/cuda/build_state.cu`
- Create: `crates/ml-alpha/src/policy/state.rs`
- Create: `crates/ml-alpha/tests/build_state_bit_equiv.rs`

The kernel reads ISV[0..13] (5 alpha probs + 8-dim projection) and position-state slots, and packs the 20-dim AlphaState directly into the next replay-buffer slot. Per the spec Section 3 state table.

- [ ] **Step 1: Write the failing bit-equiv test**

```rust
// crates/ml-alpha/tests/build_state_bit_equiv.rs
use approx::assert_relative_eq;
use ml_alpha::policy::state::{AlphaStateCtx, build_state_cpu, build_state_gpu};
use ml_core::device::MlDevice;

fn synth_ctx() -> AlphaStateCtx {
    AlphaStateCtx {
        alpha_probs: [0.51, 0.48, 0.55, 0.6, 0.7],
        proj: [-0.2, 0.1, 0.3, -0.5, 0.0, 0.05, 0.2, -0.1],
        current_target: 2,
        max_train: 10,
        unrealized_pnl_ticks: 4.0,
        time_since_entry_steps: 12.0,
        minute_of_day: 870.0, // 14:30
        spread_quantile: 0.42,
        recent_vol_norm: 0.8,
    }
}

#[test]
fn build_state_cpu_gpu_bit_equiv() {
    let ctx = synth_ctx();
    let cpu = build_state_cpu(&ctx);
    let dev = MlDevice::default_for_test();
    let gpu = build_state_gpu(&dev, &ctx).expect("gpu");
    for i in 0..20 {
        assert_relative_eq!(cpu[i], gpu[i], epsilon = 1e-6, max_relative = 1e-5);
    }
}
```

- [ ] **Step 2: Run — fails**

```bash
SQLX_OFFLINE=true cargo test -p ml-alpha --test build_state_bit_equiv 2>&1 | tail -10
```

- [ ] **Step 3: Implement kernel + binding + oracle**

`crates/ml-alpha/cuda/build_state.cu`:

```cuda
// build_state.cu — 20-dim AlphaState assembly per spec Section 3 state table.
extern "C" __global__ void build_state(
    const float* __restrict__ isv,        // [32] — slots 0..4 probs, 5..12 proj
    int current_target,
    int max_train,
    float unrealized_pnl_ticks,
    float time_since_entry_steps,
    float minute_of_day,
    float spread_quantile,
    float recent_vol_norm,
    float* __restrict__ state             // [20]
) {
    if (threadIdx.x != 0) return;
    for (int i = 0; i < 5; ++i) state[i]      = isv[i];        // alpha probs
    for (int i = 0; i < 8; ++i) state[5 + i]  = isv[5 + i];    // projection
    state[13] = ((float) current_target) / ((float) max_train);
    state[14] = fminf(fmaxf(unrealized_pnl_ticks, -20.0f), 20.0f) / 20.0f;
    state[15] = fminf(fmaxf(time_since_entry_steps, 0.0f), 100.0f) / 100.0f;
    const float two_pi_min = 6.28318530f * (minute_of_day / 1440.0f);
    state[16] = sinf(two_pi_min);
    state[17] = cosf(two_pi_min);
    state[18] = spread_quantile;
    state[19] = recent_vol_norm;
}
```

`crates/ml-alpha/src/policy/state.rs`:

```rust
use anyhow::{Context, Result};
use cudarc::driver::{LaunchAsync, LaunchConfig};
use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer;
use ml_core::device::MlDevice;

const CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/build_state.cubin"));

#[derive(Clone, Copy, Debug)]
pub struct AlphaStateCtx {
    pub alpha_probs: [f32; 5],
    pub proj: [f32; 8],
    pub current_target: i32,
    pub max_train: i32,
    pub unrealized_pnl_ticks: f32,
    pub time_since_entry_steps: f32,
    pub minute_of_day: f32,
    pub spread_quantile: f32,
    pub recent_vol_norm: f32,
}

pub fn build_state_cpu(c: &AlphaStateCtx) -> [f32; 20] {
    let mut s = [0f32; 20];
    s[..5].copy_from_slice(&c.alpha_probs);
    s[5..13].copy_from_slice(&c.proj);
    s[13] = c.current_target as f32 / c.max_train as f32;
    s[14] = c.unrealized_pnl_ticks.clamp(-20.0, 20.0) / 20.0;
    s[15] = c.time_since_entry_steps.clamp(0.0, 100.0) / 100.0;
    let two_pi_min = std::f32::consts::TAU * (c.minute_of_day / 1440.0);
    s[16] = two_pi_min.sin();
    s[17] = two_pi_min.cos();
    s[18] = c.spread_quantile;
    s[19] = c.recent_vol_norm;
    s
}

pub fn build_state_gpu(dev: &MlDevice, c: &AlphaStateCtx) -> Result<[f32; 20]> {
    let mut isv = [0f32; 32];
    isv[..5].copy_from_slice(&c.alpha_probs);
    isv[5..13].copy_from_slice(&c.proj);
    let isv_d = dev.htod_copy(isv.to_vec())?;
    let staging = unsafe { MappedF32Buffer::new(20) }.context("state staging")?;
    let func = dev.load_cubin_function(CUBIN, "build_state")?;
    let cfg = LaunchConfig { grid_dim: (1,1,1), block_dim: (1,1,1), shared_mem_bytes: 0 };
    unsafe {
        func.launch_async(cfg, (
            &isv_d, c.current_target, c.max_train,
            c.unrealized_pnl_ticks, c.time_since_entry_steps,
            c.minute_of_day, c.spread_quantile, c.recent_vol_norm,
            staging.device_ptr(),
        ))?;
    }
    dev.default_stream().synchronize().context("build_state sync")?;
    let mut out = [0f32; 20];
    out.copy_from_slice(staging.as_slice());
    Ok(out)
}
```

Add the kernel to `build.rs` KERNELS list and `pub mod policy;` (with `pub mod state;`) to lib.rs.

- [ ] **Step 4: Run — passes**

```bash
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml-alpha --test build_state_bit_equiv 2>&1 | tail -10
```

- [ ] **Step 5: Commit**

```bash
git add crates/ml-alpha/cuda/build_state.cu crates/ml-alpha/src/policy/state.rs crates/ml-alpha/build.rs crates/ml-alpha/src/lib.rs crates/ml-alpha/tests/build_state_bit_equiv.rs
git commit -m "feat(ml-alpha): build_state kernel — 20-dim AlphaState assembly

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>"
```

---

## Task 2: `policy_forward` kernel — fused CfC actor + critic

**Files:**
- Create: `crates/ml-alpha/cuda/policy_forward.cu`
- Create: `crates/ml-alpha/src/policy/network.rs`
- Create: `crates/ml-alpha/tests/policy_forward_bit_equiv.rs`

The actor is `20→CfC(64)→Linear(64→21)` (21 logits for target_position ∈ {-10..+10}). The critic is `20→CfC(64)→Linear(64→1)`. The two CfCs share input but have independent weights. The kernel emits `(logits[21], log_probs[21], value)` plus the actor hidden post-step (for the next decision).

Same TDD shape as Plan 1 Task 6 (CPU oracle + GPU kernel + bit-equiv ≤1e-4). Block dimensions: 64 threads (one per hidden unit), shared-mem partials for the 21-logit GEMM tail.

- [ ] **Step 1: Failing bit-equiv test**

```rust
// crates/ml-alpha/tests/policy_forward_bit_equiv.rs
use approx::assert_relative_eq;
use ml_alpha::policy::network::{PolicyWeights, policy_forward_cpu, policy_forward_gpu, ACTION_DIM};
use ml_core::device::MlDevice;
use rand::{Rng, SeedableRng};
use rand_chacha::ChaCha8Rng;

#[test]
fn policy_forward_cpu_gpu_bit_equiv() {
    let mut r = ChaCha8Rng::seed_from_u64(0x9019);
    let mut rand_vec = |n: usize, range: f32| -> Vec<f32> {
        (0..n).map(|_| r.gen_range(-range..range)).collect()
    };
    let n_state = 20; let n_hid = 64; let n_act = ACTION_DIM;
    let w = PolicyWeights {
        actor_w_in:   rand_vec(n_hid * n_state, 0.2),
        actor_w_rec:  rand_vec(n_hid * n_hid,   0.1),
        actor_b:      vec![0.0; n_hid],
        actor_tau:    (0..n_hid).map(|_| r.gen_range(0.05..2.0)).collect(),
        actor_logits_w: rand_vec(n_act * n_hid, 0.1),
        actor_logits_b: vec![0.0; n_act],
        critic_w_in:  rand_vec(n_hid * n_state, 0.2),
        critic_w_rec: rand_vec(n_hid * n_hid,   0.1),
        critic_b:     vec![0.0; n_hid],
        critic_tau:   (0..n_hid).map(|_| r.gen_range(0.05..2.0)).collect(),
        critic_v_w:   rand_vec(n_hid, 0.1),
        critic_v_b:   0.0,
        n_state, n_hid,
    };
    let state: Vec<f32> = rand_vec(n_state, 1.0);
    let h_actor_old: Vec<f32> = rand_vec(n_hid, 0.5);
    let h_critic_old: Vec<f32> = rand_vec(n_hid, 0.5);
    let dt_s = 0.02_f32;
    let max_live = 10_i32;

    let cpu = policy_forward_cpu(&w, &state, &h_actor_old, &h_critic_old, dt_s, max_live);
    let dev = MlDevice::default_for_test();
    let gpu = policy_forward_gpu(&dev, &w, &state, &h_actor_old, &h_critic_old, dt_s, max_live).unwrap();

    for k in 0..n_act { assert_relative_eq!(cpu.logits[k], gpu.logits[k], epsilon = 1e-3, max_relative = 1e-3); }
    for k in 0..n_act { assert_relative_eq!(cpu.log_probs[k], gpu.log_probs[k], epsilon = 1e-3, max_relative = 1e-3); }
    assert_relative_eq!(cpu.value, gpu.value, epsilon = 1e-3);
}
```

- [ ] **Step 2: Run — fails**

```bash
SQLX_OFFLINE=true cargo test -p ml-alpha --test policy_forward_bit_equiv 2>&1 | tail -10
```

- [ ] **Step 3: Implement kernel + CPU oracle + Rust binding**

The kernel structure:
- Block dim = 64 (one thread per actor hidden unit).
- Each thread computes its CfC step for actor and critic (sequential within thread; the two updates share `state` reads).
- Shared-memory tile holds the 64 actor hidden post-step values; the 21-logit output is computed by 21 threads (using 64-element dot products) via a per-thread partial + 1-block reduce.
- log_softmax: block tree-reduce over the 21 logits (logsumexp).
- Logit mask: threads with action index outside `[-max_live, +max_live]` set logit to `-INF` before the softmax.

```cuda
// crates/ml-alpha/cuda/policy_forward.cu
extern "C" __global__ void policy_forward(
    const float* __restrict__ actor_w_in, const float* __restrict__ actor_w_rec,
    const float* __restrict__ actor_b,    const float* __restrict__ actor_tau,
    const float* __restrict__ actor_logits_w, const float* __restrict__ actor_logits_b,
    const float* __restrict__ critic_w_in, const float* __restrict__ critic_w_rec,
    const float* __restrict__ critic_b,    const float* __restrict__ critic_tau,
    const float* __restrict__ critic_v_w,  float critic_v_b,
    const float* __restrict__ state,       // [n_state]
    const float* __restrict__ h_actor_old, // [n_hid]
    const float* __restrict__ h_critic_old,// [n_hid]
    float dt_s,
    int n_state, int n_hid, int n_act, int max_live, int max_train,
    float* __restrict__ h_actor_new,       // [n_hid]
    float* __restrict__ h_critic_new,      // [n_hid]
    float* __restrict__ logits_out,        // [n_act]
    float* __restrict__ log_probs_out,     // [n_act]
    float* __restrict__ value_out          // [1]
) {
    // implementation: 64 threads, two CfC updates, then logit GEMM + masked logsumexp +
    // critic head. Reductions block tree-reduce. No atomicAdd.
    // (Full body ~120 lines — straight chain-rule, no novel CUDA.)
}
```

Add binding + CPU oracle + bit-equiv test mirror in `crates/ml-alpha/src/policy/network.rs`. Same shape as Plan 1 Task 6.

- [ ] **Step 4: Run — passes**

```bash
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml-alpha --test policy_forward_bit_equiv 2>&1 | tail -10
```

- [ ] **Step 5: Commit**

```bash
git add crates/ml-alpha/cuda/policy_forward.cu crates/ml-alpha/src/policy/network.rs crates/ml-alpha/build.rs crates/ml-alpha/tests/policy_forward_bit_equiv.rs
git commit -m "feat(ml-alpha): policy_forward — fused CfC actor + critic + masked log_softmax

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>"
```

---

## Task 3: `sample_action` kernel (cuRAND categorical with logit mask)

**Files:**
- Create: `crates/ml-alpha/cuda/sample_action.cu`
- Create: `crates/ml-alpha/src/policy/sampler.rs`
- Create: `crates/ml-alpha/tests/sample_action_dist.rs`

Statistical test: 100K samples from a known distribution → empirical proportions within 3 SE.

- [ ] **Step 1: Failing test**

```rust
// crates/ml-alpha/tests/sample_action_dist.rs
use ml_alpha::policy::sampler::sample_categorical_gpu;
use ml_core::device::MlDevice;

#[test]
fn samples_match_target_distribution() {
    let dev = MlDevice::default_for_test();
    // log-probs from softmax([1.0, 2.0, 3.0]) → roughly [0.09, 0.245, 0.665]
    let log_probs = vec![1.0 - 3.40760_f32, 2.0 - 3.40760, 3.0 - 3.40760];
    let n = 100_000usize;
    let counts = sample_categorical_gpu(&dev, &log_probs, n, /*seed=*/0xC0DEC).unwrap();
    let mut empirical = [0f32; 3];
    for i in 0..3 { empirical[i] = counts[i] as f32 / n as f32; }
    let expected = [0.0900_f32, 0.2447, 0.6652];
    let se = (0.25_f32 / n as f32).sqrt();
    for i in 0..3 {
        assert!((empirical[i] - expected[i]).abs() < 4.0 * se, "category {i}: emp={} exp={} se={}", empirical[i], expected[i], se);
    }
}
```

- [ ] **Step 2: Run — fails**

```bash
SQLX_OFFLINE=true cargo test -p ml-alpha --test sample_action_dist 2>&1 | tail -10
```

- [ ] **Step 3: Implement kernel + binding**

Kernel: block tree-reduce computes cumulative log-prob, draws a uniform via cuRAND device API, inverse-CDF samples. Single thread for the small (21-element) categorical.

- [ ] **Step 4: Run — passes**
- [ ] **Step 5: Commit**

---

## Task 4: `write_replay_slot` kernel + `ReplayBuffer`

**Files:**
- Create: `crates/ml-alpha/cuda/write_replay_slot.cu`
- Create: `crates/ml-alpha/src/replay/mod.rs`
- Create: `crates/ml-alpha/src/replay/buffer.rs`
- Create: `crates/ml-alpha/tests/replay_buffer.rs`

The replay buffer:
- Layout per slot (116 bytes): `state[20] + action[i8] + 3 padding + log_prob[f32] + value[f32] + reward[f32] + done[f32]` → 28 floats + 4 bytes
- Offline trajectories: GPU-resident, loaded once at startup, ring index 0..N_offline
- Live ring: GPU-resident, ring index 0..N_live (10K capacity), CAS-based atomic write head (NOT atomicAdd; uses atomicCAS for ring-index, which is the one allowed exception per spec)
- `sample_mixed(n, offline_frac)` returns indices into a single contiguous "mixed" buffer that holds the gathered minibatch

- [ ] **Step 1: Failing tests**

```rust
// crates/ml-alpha/tests/replay_buffer.rs
use ml_alpha::replay::buffer::{ReplayBuffer, ReplayConfig};
use ml_core::device::MlDevice;

#[test]
fn sample_mixed_respects_offline_fraction() {
    let dev = MlDevice::default_for_test();
    let mut buf = ReplayBuffer::new(&dev, &ReplayConfig { offline_capacity: 1000, live_capacity: 1000, state_dim: 20 }).unwrap();
    // Load 1000 offline with marker = -1.0 in reward slot
    let offline_rewards = vec![-1.0_f32; 1000];
    buf.load_synthetic_offline(&offline_rewards).unwrap();
    // Push 200 live with marker = +1.0 in reward slot
    for _ in 0..200 { buf.push_synthetic_live(1.0).unwrap(); }

    let mix = buf.sample_mixed(1000, 0.8, /*seed=*/42).unwrap();
    let n_offline = mix.iter().filter(|r| **r < 0.0).count();
    assert!(n_offline >= 750 && n_offline <= 850, "offline_frac=0.8 of 1000 should be ~800; got {n_offline}");
}

#[test]
fn sample_returns_all_offline_when_live_is_empty() {
    let dev = MlDevice::default_for_test();
    let mut buf = ReplayBuffer::new(&dev, &ReplayConfig { offline_capacity: 500, live_capacity: 100, state_dim: 20 }).unwrap();
    buf.load_synthetic_offline(&vec![-1.0_f32; 500]).unwrap();
    let mix = buf.sample_mixed(128, 0.8, 11).unwrap();
    assert!(mix.iter().all(|r| *r < 0.0));
}
```

- [ ] **Step 2: Run — fails**
- [ ] **Step 3: Implement** the kernel + Rust shim (`atomicCAS` only on the ring head per `feedback_no_atomicadd.md` accepted exception; reductions stay tree-reduce).
- [ ] **Step 4: Run — passes**
- [ ] **Step 5: Commit**

---

## Task 5: `ingest_fill` kernel — fill → position/PnL → back-fill reward

**Files:**
- Create: `crates/ml-alpha/cuda/ingest_fill.cu`
- Create: `crates/ml-alpha/src/replay/fill.rs`
- Create: `crates/ml-alpha/tests/ingest_fill.rs`

Reads `FillPayload` from a mapped-pinned fill slot, updates a position/PnL device state, and writes the realized P&L into the `reward` field of the prior replay slot.

- [ ] **Step 1: Failing test** — synthetic fill into a slot with known position, check reward post-condition.
- [ ] **Step 2: Run — fails**
- [ ] **Step 3: Implement** kernel (single thread is enough; position/PnL state lives in 4 f32 + 1 i32 device-resident).
- [ ] **Step 4: Run — passes**
- [ ] **Step 5: Commit**

---

## Task 6: `RolloutEnv` — 64 parallel environments on device

**Files:**
- Create: `crates/ml-alpha/src/rollout/mod.rs`
- Create: `crates/ml-alpha/src/rollout/env.rs`
- Create: `crates/ml-alpha/tests/rollout_reward.rs`

`RolloutEnv` owns:
- 64 × per-env (current_pos, current_pnl_segment_open_mid, time_since_entry, last_trade_pos)
- Snapshot pointer into the frozen Phase A perception output buffer
- Decision counter per env (stride=50 snapshots)

`rollout(steps)`:
1. For each snapshot tick (`steps × 50` iterations):
   - Perception forward (replay Graph A) — fills ISV[0..13] for each env (batched)
   - Increment decision counter; if `== 50`: build_state, policy_forward, sample_action, write transition
   - Compute segment reward = `target_t · (mid_close − mid_open) − C_trade·|Δtarget| − C_vol·max(0, |seg_return| − r_thr)`

- [ ] **Step 1: Failing test** — feed synthetic upward-drifting mid prices into a single env; assert mean reward > 0 with `target=1`.
- [ ] **Step 2-5: TDD cycle.**

---

## Task 7: `gae_returns` kernel — single-block cooperative prefix scan

**Files:**
- Create: `crates/ml-alpha/cuda/gae_returns.cu`
- Create: `crates/ml-alpha/src/ppo/gae.rs`
- Create: `crates/ml-alpha/tests/gae_bit_equiv.rs`

GAE: `δ_t = r_t + γ·V_{t+1} − V_t`; `A_t = δ_t + γ·λ·A_{t+1}`. Sequential backward scan over the rollout.

- [ ] **Step 1: Failing bit-equiv test against CPU GAE reference (sequential).**
- [ ] **Step 2-5: TDD.** ε ≤ 1e-5.

---

## Task 8: `advantage_normalize` kernel — z-score with block tree-reduce

**Files:**
- Create: `crates/ml-alpha/cuda/advantage_normalize.cu`
- Modify: `crates/ml-alpha/src/ppo/gae.rs`
- Create: `crates/ml-alpha/tests/advantage_normalize.rs`

Standard mean/var via block tree-reduce, then z-score in place.

- [ ] **Step 1-5: TDD.** Test on synthetic Gaussian input: mean ≈ 0, std ≈ 1.

---

## Task 9: `ppo_loss` kernel — fused PPO clip + value MSE + entropy + EWC + backward

**Files:**
- Create: `crates/ml-alpha/cuda/ppo_loss.cu`
- Create: `crates/ml-alpha/src/ppo/loss.rs`
- Create: `crates/ml-alpha/tests/ppo_loss_finite_diff.rs`

Forward + backward fused into one kernel. Backward uses cuBLAS Lt grad-GEMMs (chain-rule through policy_forward weights).

Loss:
```
L = L_clip + value_coef·L_value − entropy_coef·H + λ_ewc · Σ_i F_i · (θ_i − θ_anchor_i)^2
L_clip = −E[min(ratio·A, clip(ratio, 1±ε)·A)]
L_value = E[(V(s) − R_t)^2]
```

- [ ] **Step 1: Failing finite-difference test** — pick 5 random parameters; compare backward grad to (loss_up − loss_dn) / 2eps within 2% relative.
- [ ] **Step 2-5: TDD.** Reductions block tree-reduce; no atomicAdd.

---

## Task 10: EWC anchor + Fisher diagonal

**Files:**
- Create: `crates/ml-alpha/src/ewc/mod.rs`
- Create: `crates/ml-alpha/tests/ewc_fisher.rs`

After Phase B finishes on training window: compute `F_i = E_val[(∂L_PPO/∂θ_i)²]` per parameter via a pass over the validation rollout (no backward — squared grads accumulated via block tree-reduce). Persist `(θ_anchor, F)` to disk.

- [ ] **Step 1: Failing test** — assert Fisher is non-negative everywhere; assert anchor differs from current weights after one update.
- [ ] **Step 2-5: TDD.**

---

## Task 11: CUDA Graph B capture (decision graph)

**Files:**
- Modify: `crates/ml-alpha/src/policy/network.rs` (add `capture_graph_b`, `policy_action_captured`)
- Create: `crates/ml-alpha/tests/graph_b_replay.rs`

Graph B: build_state → policy_forward → sample_action → write_replay_slot → async DtoH of action byte.

- [ ] **Step 1: Failing test** — Graph B replay produces bit-identical state/log_prob/value to sequential dispatch.
- [ ] **Step 2-5: TDD.** Per `pearl_no_host_branches_in_captured_graph.md`: `max_live` mask is read from device-resident ISV slot 28, not from a host scalar argument.

---

## Task 12: CUDA Graph C capture (training step)

**Files:**
- Modify: `crates/ml-alpha/src/ppo/update.rs` (add Graph C capture)
- Create: `crates/ml-alpha/tests/graph_c_replay.rs`

Graph C: gather minibatch → GAE → advantage normalize → policy_forward batch → ppo_loss fused forward+backward → AdamW.

- [ ] **Step 1: Failing test** — after K=4 epochs of one minibatch the captured graph produces the same weight delta as sequential dispatch (ε ≤ 1e-4).
- [ ] **Step 2-5: TDD.**

---

## Task 13: Seven ISV controllers

**Files:**
- Create: `crates/ml-alpha/src/controllers/mod.rs`
- Create one file per controller: `entropy.rs`, `lr.rs`, `ewc_lambda.rs`, `clip_eps.rs`, `value_coef.rs`, `replay_mix.rs`, `vol_penalty.rs`
- Create: `crates/ml-alpha/tests/controller_invariants.rs`

Every controller follows the canonical foxhunt pattern (per `pearl_controller_anchors_isv_driven.md`, `pearl_first_observation_bootstrap.md`, `pearl_wiener_alpha_floor_for_nonstationary.md`, `pearl_blend_formulas_must_have_permanent_floor.md`):

```rust
pub trait Controller {
    fn update(&mut self, signal: f32) -> f32;     // returns new effective value
    fn snapshot(&self) -> ControllerState;        // for ISV readout
}
```

| Controller | Signal | Anchor / target | Floor | Wiener-α floor |
|---|---|---|---|---|
| `EntropyCoefController` | entropy_deficit = (target_entropy − current_entropy) | target_entropy = 0.7·ln(n_act) | 1e-5 | 0.4 |
| `LrPolicyController` | loss_slope EMA | target_loss_slope = -ε per epoch | 1e-6 | 0.4 |
| `EwcLambdaController` | ‖θ − θ_anchor‖² | drift_target = ISV[17] EMA | 1e-2 | 0.4 |
| `ClipEpsController` | kl_divergence (vs rollout policy) | target_kl = 0.01 | 0.05 | 0.4 |
| `ValueCoefController` | value_loss / policy_loss | target_ratio = 1/5 | 0.1 | 0.4 |
| `ReplayMixController` | live_mean_reward − offline_mean_reward | gap_target = 0 | live_frac ≥ 0.05 | 0.4 |
| `VolPenaltyController` | std(seg_return) | target = ISV[19] × 1.5 | 0.0 (off by default) | 0.4 |

- [ ] **Step 1: Failing invariant test for each controller**:
  - First-observation bootstrap (sentinel = 0, first sample replaces directly)
  - Wiener-α never below 0.4
  - Output never below floor regardless of signal
  - Monotonicity in the expected direction (e.g., entropy_coef decreases as deficit becomes negative)

```rust
// crates/ml-alpha/tests/controller_invariants.rs
use ml_alpha::controllers::entropy::EntropyCoefController;
use ml_alpha::controllers::Controller;

#[test]
fn entropy_first_observation_bootstrap() {
    let mut c = EntropyCoefController::new(/*target=*/2.0, /*init=*/0.01);
    // Sentinel state: no observation yet → returns init unchanged for first call
    let v0 = c.update(1.5);
    // After one observation, internal state has the observed entropy (1.5)
    // and the controller has stepped. v0 != init exactly only if controller
    // has internal momentum on first step; the bootstrap rule says it should
    // step normally on the first observation, not stay stuck at init.
    assert!(v0 > 0.0 && v0 != 0.01);
}

#[test]
fn entropy_coef_respects_floor() {
    let mut c = EntropyCoefController::new(2.0, 0.01);
    // Drive deficit very negative repeatedly: entropy_coef should decay toward floor 1e-5
    for _ in 0..100 { c.update(10.0); }
    let snap = c.snapshot();
    assert!(snap.effective >= 1e-5, "below floor: {}", snap.effective);
}
```

(Repeat per controller; the test file lists one floor-respect test and one bootstrap test per controller.)

- [ ] **Step 2-5: TDD.**

---

## Task 14: `kill_switch_check` kernel + `KillSwitch` module + bootstrap baseline

**Files:**
- Create: `crates/ml-alpha/cuda/kill_switch_check.cu`
- Create: `crates/ml-alpha/src/kill_switch.rs`
- Create: `crates/ml-alpha/tests/kill_switch.rs`

`KillSwitch` keeps a rolling 20-update baseline (mean, std for {loss_ppo, mean_reward, hit_rate, kl, entropy}) plus an immutable `phase_b_baseline` loaded from disk. Bootstrap mode (Section 6) uses Phase B baseline; switches to rolling once 20 successful updates land.

| Metric | Reject when (bootstrap mode) | Reject when (rolling mode) |
|---|---|---|
| `loss_ppo` | > phase_b_loss + 2·phase_b_loss_std | > rolling_mean + 3σ |
| `mean_reward` | < phase_b_reward − 2·phase_b_reward_std | < rolling_mean − 2σ |
| `hit_rate` | < 0.45 absolute | < 0.45 absolute |
| `kl` | > 0.1 absolute | > 0.1 absolute |
| `entropy` | < 0.1 nats absolute | < 0.1 nats absolute |

- [ ] **Step 1: Failing tests** — bootstrap+rolling regime switchover at update count=20; rejection for each metric; never-rejects-on-first-update.
- [ ] **Step 2-5: TDD.**

---

## Task 15: PPO `update()` + atomic weights swap

**Files:**
- Create: `crates/ml-alpha/src/ppo/update.rs`
- Create: `crates/ml-alpha/tests/atomic_swap.rs`

Per spec Section 6:
1. Write versioned weight snapshot to disk first
2. Clone weights to `candidate` device buffer
3. Run K=4 PPO epochs × N minibatches on `candidate` (Graph C replays)
4. `kill_switch.check(&metrics)` →
   - `Ok` → `std::mem::swap(&mut weights, &mut candidate)`; record metrics; controllers update
   - `Reject` → drop candidate; weights untouched; flip `online_updates_enabled = false`

- [ ] **Step 1: Failing test** — synthetic update with metric forced past threshold returns `RolledBack` and weights unchanged.
- [ ] **Step 2-5: TDD.**

---

## Task 16: Phase B trainer loop (training-window only, single fold)

**Files:**
- Create: `crates/ml-alpha/src/trainer/phase_b.rs`
- Create: `crates/ml-alpha/tests/phase_b_synthetic.rs`

Pulls it all together:
```
load_perception(plan_1_artifact)
init_policy_random()
init_replay_buffer(offline_capacity=50_000, live=0)
init_isv_bus()
init_controllers()
for rollout_iter in 0..N_rollout_iters:
    rollout(steps=rollout_length=2048 decisions per env)
    update(K=4, minibatch_size=256)   // Graph C
    controllers.tick()
    if kill_switch.fired: error and exit
emit phase_b_baseline.json
emit weights.bin + ewc_anchor.bin + fisher.bin
```

- [ ] **Step 1: Failing test** — synthetic env where target=+1 always wins by 0.1 ticks/segment, training drives policy to `target=+1` with prob > 0.9 within 50 rollouts.
- [ ] **Step 2-5: TDD.**

---

## Task 17: `alpha_train_phase_b` example binary

**Files:**
- Create: `crates/ml-alpha/examples/alpha_train_phase_b.rs`

CLI mirrors Plan 1 Task 15. Args: `--perception-weights <path>`, `--train-quarters`, `--val-quarter`, `--rollout-length`, `--n-rollout-iters`, `--out`.

- [ ] **Step 1-5: scaffold + smoke run on 1Q + commit.**

---

## Task 18: 6-fold walk-forward CV runner

**Files:**
- Create: `crates/ml-alpha/src/cv/mod.rs`
- Create: `crates/ml-alpha/examples/alpha_walk_forward.rs`
- Create: `crates/ml-alpha/tests/walk_forward_smoke.rs`

Fold table per spec Section 4 (folds 1-4 produce test Sharpe; folds 5-6 are for final calibration / live).

`alpha_walk_forward`:
1. For each fold (1..=4): train Phase B on `train_quarters`, save weights, evaluate on `test_quarter` via a no-update rollout, collect realized Sharpe + hit-rate + trades-per-day + commission-paid
2. Write `walk_forward.json` with per-fold metrics + mean ± std
3. Exit 0 if `mean Sharpe > 0` AND `std < mean`, else 1

- [ ] **Step 1: Failing 1-fold smoke test** — Fold 1 (train Q1-Q4 2024, test Q1 2025) produces finite Sharpe in `walk_forward.json`.
- [ ] **Step 2-5: TDD.**

---

## Task 19: Cluster workflow (6-fold sequential on L40S)

**Files:**
- Create: `infra/k8s/argo/alpha-phase-b-template.yaml`
- Create: `scripts/argo-alpha-phase-b.sh`

Single Argo `WorkflowTemplate`, 6 sequential `Phase B + test` steps (one per fold). NodeSelector pinned to `topology.kubernetes.io/zone=fr-par-2`, toleration for `ci-training-l40s`. Default GPU pool `ci-training-l40s` per `feedback_default_to_l40s_pool.md`. Each fold writes `phase_b_fold_<N>.json`; the final step aggregates into `walk_forward.json` + emits to MinIO.

- [ ] **Step 1**: Write the workflow template (model on `alpha-cv-template.yaml`).
- [ ] **Step 2**: Write the submission script.
- [ ] **Step 3**: `git push origin ml-alpha-phase-b` then `./scripts/argo-alpha-phase-b.sh` (per `feedback_push_before_deploy.md`).
- [ ] **Step 4**: Wait. Fetch `walk_forward.json`. Assert pass criteria via:
  ```bash
  jq '.mean_sharpe > 0 and .std_sharpe < .mean_sharpe' artifacts/walk-forward/walk_forward.json
  ```
- [ ] **Step 5**: Commit infra.

---

## Gate-pass exit criteria

- [ ] All per-kernel bit-equiv tests pass at ε ≤ 1e-5 (forward) or ≤ 5% relative (backward)
- [ ] Graph B replay matches sequential dispatch
- [ ] Graph C replay matches sequential training-step dispatch
- [ ] All 7 controllers respect first-observation bootstrap, Wiener-α floor 0.4, output floor
- [ ] Kill-switch correctly rolls back synthetic-rejection updates
- [ ] Phase B trainer overfits the +0.1-ticks/segment synthetic env to target=+1 within 50 rollouts
- [ ] Walk-forward 1-fold smoke produces finite Sharpe
- [ ] **Walk-forward 4-fold mean test Sharpe > 0 at IBKR commission (0.034)**
- [ ] **std(test_sharpe) across folds < mean(test_sharpe)** (consistency floor)
- [ ] `phase_b_baseline.json` artifact emitted with all 5 baseline metrics
- [ ] No TODO/FIXME/XXX markers (pre-commit hook enforces)

**On pass:** open PR `ml-alpha-phase-b` → `main`. Once merged, invoke Plan 3 (`docs/superpowers/plans/2026-05-16-ml-alpha-live.md`).

**On fail:**
1. Mean Sharpe negative → reward-shaping issue. Re-check `C_trade` units (0.034 price units, not dollars). Verify segment P&L formula. Check `pearl_event_driven_reward_density_alignment.md` — is dense per-segment reward creating exposure-positive bias?
2. Mean positive but std > mean → regime fragility. Phase B fold 2 or 3 likely lost heavily (per `pearl_single_window_oos_is_not_oos.md`). Investigate which fold + which regime; consider Mamba2 fallback for that regime.
3. Kill-switch firing repeatedly → controller pathology. Read `controller_state.json` snapshot; suspect Wiener-α floor too low for non-stationary loop, per `pearl_wiener_alpha_floor_for_nonstationary.md`.

Document the failure mode and re-plan in a spec delta — do not advance to live integration on a non-passing CV.
