# OOS Generalization Enhancement 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:** Redirect 500K model params from noise memorization toward robust OOS strategies by adding 10 generalization components: AdamW weight decay, transaction cost curriculum, purged walk-forward validation, gamma annealing, epistemic-gated magnitude, branch independence, counterfactual augmentation, L1-sparse input, regime-aware dropout, temporal consistency.

**Architecture:** 3-layer defense composing through the existing fused CUDA training pipeline. All CUDA changes are small (5-30 lines per kernel). The largest change is G3 (purged walk-forward) at ~200 lines of Rust in the training loop. New CUDA kernels are added to `experience_kernels.cu`. The Adam kernel already has AdamW infrastructure (weight_decay param) but it's hardcoded to 0.0 — G1 just enables it with a mask.

**Tech Stack:** Rust 1.85, CUDA 12.4, cudarc, nvcc precompiled cubins, H100 80GB target

---

## Task 1: G1 — Enable AdamW Weight Decay with Trunk-Only Mask

The Adam kernel already implements AdamW (`weight_decay * p_val` at line 151 of `dqn_utility_kernels.cu`), but the GPU training path hardcodes `weight_decay: 0.0` at line 6102 of `gpu_dqn_trainer.rs`. This task enables it with a per-parameter mask so only trunk+value head weights (indices 0-7) get decay.

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu`
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`

- [ ] **Step 1: Add weight_decay_mask parameter to Adam kernel**

In `crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu`, find the `dqn_adam_update_kernel` signature (line 102). Add a new parameter after `total_params`:

```cuda
    const float* __restrict__ weight_decay_mask  /* [TOTAL_PARAMS] 1.0 for trunk, 0.0 for branches */
```

At line 151, change:

```cuda
    p_val = p_val - lr * (m_hat / (sqrtf(v_hat) + epsilon) + weight_decay * p_val);
```

To:

```cuda
    float wd = weight_decay * weight_decay_mask[idx];
    p_val = p_val - lr * (m_hat / (sqrtf(v_hat) + epsilon) + wd * p_val);
```

- [ ] **Step 2: Add L1 proximal step for w_s1 (G8 combined)**

Immediately after the AdamW update line, add L1 soft-thresholding for the first layer:

```cuda
    /* G8: L1 sparsity on w_s1 — automatic feature selection.
     * l1_end is the byte offset of the end of tensor 0 (w_s1). */
    if (idx < l1_end && l1_lambda > 0.0f) {
        float sign = (p_val > 0.0f) ? 1.0f : -1.0f;
        p_val = sign * fmaxf(fabsf(p_val) - lr * l1_lambda, 0.0f);
    }
    params[idx] = p_val;
```

Add two parameters to the kernel signature (after `weight_decay_mask`):

```cuda
    int l1_end,                                  /* param index where L1 stops (end of w_s1) */
    float l1_lambda                              /* L1 penalty strength (1e-3) */
```

Remove the existing `params[idx] = p_val;` line (now at the end of the L1 block).

- [ ] **Step 3: Allocate weight_decay_mask buffer and update kernel launch**

In `gpu_dqn_trainer.rs`, add a new field to the struct:

```rust
    /// Weight decay mask [total_params]: 1.0 for trunk+value (indices 0-7), 0.0 for branches
    weight_decay_mask: CudaSlice<f32>,
```

In the constructor, allocate and fill:

```rust
        // G1: Weight decay mask — 1.0 for trunk+value head (indices 0-7), 0.0 for rest
        let weight_decay_mask = {
            let param_sizes = compute_param_sizes(&cfg);
            let total = compute_total_params(&cfg);
            let mut mask = vec![0.0_f32; total];
            let trunk_end = (0..8).map(|i| align4(param_sizes[i])).sum::<usize>();
            for i in 0..trunk_end {
                mask[i] = 1.0;
            }
            let mut buf = stream.alloc_zeros::<f32>(total)
                .map_err(|e| MLError::ModelError(format!("wd_mask alloc: {e}")))?;
            stream.memcpy_htod(&mut buf, &mask)
                .map_err(|e| MLError::ModelError(format!("wd_mask htod: {e}")))?;
            buf
        };
```

- [ ] **Step 4: Update Adam kernel launch to pass mask, l1_end, l1_lambda**

Find the Adam kernel launch (search for `dqn_adam_update_kernel` launch, around line 6100). Change `weight_decay: 0.0` to use the config value, and add the three new args:

```rust
        let weight_decay: f32 = self.config.weight_decay;
```

Add after the `total_params` arg:

```rust
                    .arg(&self.weight_decay_mask)
                    .arg(&(l1_end as i32))
                    .arg(&self.config.l1_lambda)
```

Where `l1_end` is computed:

```rust
        let param_sizes = compute_param_sizes(&self.config);
        let l1_end = align4(param_sizes[0]); // w_s1 only
```

- [ ] **Step 5: Add l1_lambda to GpuDqnTrainConfig**

In the config struct, add:

```rust
    /// L1 sparsity penalty on first trunk layer (w_s1). Induces feature selection.
    pub l1_lambda: f32,
```

Default: `l1_lambda: 1e-3`.

- [ ] **Step 6: Verify compilation**

```bash
SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
```

- [ ] **Step 7: Commit**

```bash
cd /home/jgrusewski/Work/foxhunt && git add crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs && git commit -m "feat(G1+G8): enable AdamW weight decay with trunk-only mask + L1-sparse w_s1

AdamW weight_decay was already in kernel but hardcoded to 0.0 in launch.
Now uses config value (1e-4) gated by per-param mask (1.0 for trunk+value
indices 0-7, 0.0 for branches). L1 proximal step on w_s1 induces
automatic feature selection from 42-dim input.

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

---

## Task 2: G2 — Transaction Cost Curriculum

Add an annealed cost multiplier that ramps transaction costs from 0% to 100% over 20 epochs. The env_step kernel already has `spread_cost` and `tx_cost_multiplier` — we add a `cost_anneal` factor that scales them.

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/experience_kernels.cu`
- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs`
- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs`
- Modify: `crates/ml/src/trainers/dqn/trainer/mod.rs`
- Modify: `crates/ml/src/trainers/dqn/trainer/constructor.rs`

- [ ] **Step 1: Add cost_anneal parameter to env_step kernel**

In `experience_kernels.cu`, find the `env_step` kernel signature (around line 1040). Add after `saboteur_params`:

```cuda
    const float* __restrict__ cost_anneal_ptr    /* [1] pinned: 0.0=no costs, 1.0=full costs */
```

Inside the kernel, where `spread_cost` is first used, multiply by the anneal factor:

```cuda
    float cost_scale = (cost_anneal_ptr != NULL) ? cost_anneal_ptr[0] : 1.0f;
    float effective_spread = spread_cost * cost_scale;
```

Replace all subsequent uses of `spread_cost` with `effective_spread`. Also scale `tx_cost_multiplier` if it appears separately.

- [ ] **Step 2: Add cost_anneal pinned scalar to GpuExperienceCollector**

In `gpu_experience_collector.rs`, add field and allocate pinned device-mapped memory (same pattern as other pinned scalars). Pass the dev_ptr to the env_step kernel launch.

- [ ] **Step 3: Add cost anneal state to DQNTrainer**

In `mod.rs`:

```rust
    /// Cost curriculum anneal factor: 0.0 at epoch 0, ramps to 1.0 by epoch 20.
    pub(crate) cost_anneal_factor: f32,
```

In `constructor.rs`: `cost_anneal_factor: 0.0,`

- [ ] **Step 4: Wire anneal update at epoch start in training_loop.rs**

After the adaptive DSR update (around line 2194), add:

```rust
        // ── G2: Transaction cost curriculum ──
        {
            let base_anneal = (self.current_epoch as f32 / 20.0).min(1.0);
            // Novel: ensemble disagreement slows the ramp
            let disagreement = self.fused_ctx.as_ref()
                .map(|f| f.ensemble_disagreement_ema())
                .unwrap_or(0.0);
            self.cost_anneal_factor = base_anneal * (1.0 - 0.5 * disagreement.clamp(0.0, 1.0));
        }
```

Write the value to the experience collector's pinned scalar before experience collection.

- [ ] **Step 5: Verify and commit**

```bash
SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
```

```bash
cd /home/jgrusewski/Work/foxhunt && git add crates/ml/src/cuda_pipeline/experience_kernels.cu crates/ml/src/cuda_pipeline/gpu_experience_collector.rs crates/ml/src/trainers/dqn/trainer/mod.rs crates/ml/src/trainers/dqn/trainer/constructor.rs crates/ml/src/trainers/dqn/trainer/training_loop.rs && git commit -m "feat(G2): transaction cost curriculum — anneal 0%→100% over 20 epochs

Scales spread_cost and tx_cost_multiplier by cost_anneal factor.
Ramp tied to ensemble disagreement: high uncertainty slows the ramp.
Prevents no-trade collapse by starting with zero costs.

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

---

## Task 3: G4 — Gamma Annealing with Auto C51 Rescaling

Adaptive gamma starts at 0.90, extends toward 0.95 as atom utilization stays above 60%. Auto-rescale v_min/v_max when gamma changes.

**Files:**
- Modify: `crates/ml/src/trainers/dqn/trainer/mod.rs`
- Modify: `crates/ml/src/trainers/dqn/trainer/constructor.rs`
- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs`

- [ ] **Step 1: Add adaptive gamma fields to DQNTrainer**

In `mod.rs`:

```rust
    /// Adaptive discount factor: starts 0.90, extends toward 0.95 based on atom utilization.
    pub(crate) adaptive_gamma: f32,
```

In `constructor.rs`: `adaptive_gamma: 0.90,`

- [ ] **Step 2: Wire gamma update at epoch end**

In `training_loop.rs`, after the atom utilization logging, add:

```rust
        // ── G4: Gamma annealing ──
        {
            let util = self.fused_ctx.as_ref()
                .map(|f| f.utilization_ema())
                .unwrap_or(0.5);
            if util > 0.6 {
                self.adaptive_gamma = (self.adaptive_gamma + 0.005).min(0.95);
            } else if util < 0.4 {
                self.adaptive_gamma = (self.adaptive_gamma - 0.01).max(0.90);
            }
            // Auto-rescale C51 bounds
            let r_max = self.hyperparams.v_max as f32; // max single-step return
            let new_v_max = r_max / (1.0 - self.adaptive_gamma);
            // Update per_sample_support for next epoch
            if let Some(ref fused) = self.fused_ctx {
                fused.set_v_range(-new_v_max, new_v_max);
            }
            info!("G4 gamma={:.3} v_max={:.2} util={:.2}", self.adaptive_gamma, new_v_max, util);
        }
```

- [ ] **Step 3: Use adaptive_gamma in experience collection**

Where `self.hyperparams.gamma` is used for n-step return computation, replace with `self.adaptive_gamma as f64`.

- [ ] **Step 4: Verify and commit**

```bash
SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
```

```bash
cd /home/jgrusewski/Work/foxhunt && git add crates/ml/src/trainers/dqn/trainer/mod.rs crates/ml/src/trainers/dqn/trainer/constructor.rs crates/ml/src/trainers/dqn/trainer/training_loop.rs && git commit -m "feat(G4): gamma annealing 0.90→0.95 tied to atom utilization

Adaptive gamma with hysteresis: extend at util>0.6, contract at util<0.4.
Auto-rescales C51 v_min/v_max = ±R_max/(1-gamma) when gamma changes.
Reduces horizon noise from ~100 bars (gamma=0.99) to ~10-20 bars.

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

---

## Task 4: G9 + G7 — New CUDA Kernels (Regime Dropout + Counterfactual Augmentation)

Write 2 new CUDA kernels: `regime_dropout` for regime-aware dropout on h_s2, and the counterfactual flip logic in env_step.

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/experience_kernels.cu`

- [ ] **Step 1: Add regime_dropout kernel**

Append to `experience_kernels.cu`:

```cuda
/* ================================================================== */
/* Kernel: regime_dropout (G9)                                         */
/* ================================================================== */

/**
 * Regime-aware dropout: mask conditioned on ADX/CUSUM quantiles.
 * Same regime always drops the same neurons (per epoch_seed).
 * Forces each neuron to be useful across multiple regimes.
 *
 * Grid: ceil(N * SH2 / 256), Block: 256.
 */
extern "C" __global__ void regime_dropout(
    float* __restrict__ h_s2,          /* [N, SH2] modified in-place */
    const float* __restrict__ states,  /* [N, state_dim] for ADX(40), CUSUM(41) */
    int N,
    int sh2,
    int state_dim,
    float drop_rate,                   /* 0.15 */
    int epoch_seed,                    /* changes per epoch */
    int is_training                    /* 0 = inference (no dropout), 1 = training */
) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx >= N * sh2 || !is_training) return;

    int i = idx / sh2;  /* sample */
    int j = idx % sh2;  /* feature */

    /* Quantize regime features into 4x4 = 16 buckets */
    float adx = states[(long long)i * state_dim + 40];
    float cusum = states[(long long)i * state_dim + 41];
    int adx_q = (int)fminf(fmaxf(adx * 0.1f, 0.0f), 3.0f);     /* 0-3 */
    int cusum_q = (int)fminf(fmaxf((cusum + 2.0f) * 0.5f, 0.0f), 3.0f); /* 0-3 */
    int regime_hash = adx_q * 7 + cusum_q;

    /* Philox-based deterministic mask: same regime → same mask per epoch */
    float u = philox_uniform(regime_hash, j + epoch_seed * 31, 9999);
    if (u < drop_rate) {
        h_s2[idx] = 0.0f;
    } else {
        h_s2[idx] *= 1.0f / (1.0f - drop_rate);  /* inverted dropout */
    }
}
```

- [ ] **Step 2: Add counterfactual flip to env_step kernel (G7)**

In the env_step kernel, after loading the state features and before reward computation, add:

```cuda
    /* G7: Counterfactual return augmentation — 50% flip direction */
    int do_flip = 0;
    {
        float flip_u = philox_uniform(i, episode_step + epoch_seed * 1000000, 8888);
        do_flip = (flip_u < 0.5f) ? 1 : 0;
    }
    if (do_flip) {
        /* Negate directional features (indices 0-15 approx) */
        for (int f = 0; f < 16 && f < state_dim; f++) {
            features_buf[i * state_dim + f] *= -1.0f;
        }
        /* Flip direction action: Long(2)↔Short(0), Flat(1) stays */
        int dir_action = factored_action / 27;  /* extract direction */
        if (dir_action == 0) dir_action = 2;
        else if (dir_action == 2) dir_action = 0;
        /* Reconstruct factored action */
        int remainder = factored_action % 27;
        factored_action = dir_action * 27 + remainder;
        /* Flip reward */
        reward *= -1.0f;
    }
```

NOTE: The exact insertion point depends on the env_step kernel's flow. The implementer must read the kernel to find where `reward` is computed and insert the flip BEFORE the reward is used for TD targets.

- [ ] **Step 3: Verify CUDA compilation**

```bash
SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
```

- [ ] **Step 4: Commit**

```bash
cd /home/jgrusewski/Work/foxhunt && git add crates/ml/src/cuda_pipeline/experience_kernels.cu && git commit -m "feat(G9+G7): regime_dropout kernel + counterfactual flip in env_step

G9: regime-aware dropout on h_s2 conditioned on ADX/CUSUM quantiles.
Same regime drops same neurons per epoch. Forces regime-robust features.
G7: 50% flip — negate directional features, swap Long/Short, flip reward.
Doubles effective dataset, forces symmetric strategies.

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

---

## Task 5: G5 + G6 + G10 — New CUDA Kernels (Epistemic Gate, Branch Independence, Temporal Consistency)

Write 3 new CUDA kernels for structural exploitation.

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/experience_kernels.cu`

- [ ] **Step 1: Add epistemic_gate_magnitude kernel (G5)**

```cuda
/* ================================================================== */
/* Kernel: epistemic_gate_magnitude (G5)                               */
/* ================================================================== */

/**
 * Scale magnitude branch Q-values by inverse ensemble disagreement.
 * High epistemic uncertainty → forced toward Small position size.
 *
 * Grid: ceil(N / 256), Block: 256.
 */
extern "C" __global__ void epistemic_gate_magnitude(
    float* __restrict__ q_values,      /* [N, total_actions] modified in-place */
    const float* __restrict__ ensemble_var, /* [N, total_actions] variance across heads */
    const float* __restrict__ var_ema_ptr,  /* [1] adaptive threshold (pinned) */
    int N,
    int total_actions,
    int b0_size,                       /* direction branch size (3) */
    int b1_size                        /* magnitude branch size (3) */
) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= N) return;

    /* Compute per-sample mean variance */
    float var_sum = 0.0f;
    for (int a = 0; a < total_actions; a++)
        var_sum += ensemble_var[(long long)i * total_actions + a];
    float var_mean = var_sum / (float)total_actions;

    float threshold = var_ema_ptr[0];
    /* sigmoid(-5 * (var - threshold)): var << threshold → 1.0, var >> threshold → 0.0 */
    float mag_scale = 1.0f / (1.0f + expf(5.0f * (var_mean - threshold)));

    /* Scale magnitude branch Q-values (indices b0_size .. b0_size+b1_size) */
    float* q_row = q_values + (long long)i * total_actions;
    for (int a = b0_size; a < b0_size + b1_size; a++) {
        q_row[a] *= mag_scale;
    }
}
```

- [ ] **Step 2: Add branch_independence_penalty kernel (G6)**

```cuda
/* ================================================================== */
/* Kernel: branch_independence_penalty (G6)                            */
/* ================================================================== */

/**
 * Cosine similarity penalty between branch hidden representations.
 * Forces branches to learn orthogonal features.
 *
 * Grid: (1, 1, 1), Block: (256, 1, 1). Single block, batch-reduced.
 * Computes 6 pairwise cos_sim^2 for branches 0-3, writes scalar penalty.
 */
extern "C" __global__ void branch_independence_penalty(
    const float* __restrict__ h_b0,    /* [N, AH] */
    const float* __restrict__ h_b1,    /* [N, AH] */
    const float* __restrict__ h_b2,    /* [N, AH] */
    const float* __restrict__ h_b3,    /* [N, AH] */
    float* __restrict__ penalty_out,   /* [1] scalar output */
    int N,
    int AH,                            /* adv_h dimension */
    float lambda_indep                 /* penalty weight (0.01) */
) {
    int tid = threadIdx.x;

    /* Each thread handles a slice of samples, accumulates locally */
    float local_penalty = 0.0f;
    const float* branches[4] = { h_b0, h_b1, h_b2, h_b3 };

    for (int b = tid; b < N; b += blockDim.x) {
        /* Compute norms and 6 pairwise dot products */
        float norms[4] = {0, 0, 0, 0};
        for (int d = 0; d < 4; d++) {
            const float* h = branches[d] + (long long)b * AH;
            float n = 0.0f;
            for (int j = 0; j < AH; j++) n += h[j] * h[j];
            norms[d] = sqrtf(n + 1e-8f);
        }

        int pair_idx = 0;
        for (int d1 = 0; d1 < 4; d1++) {
            for (int d2 = d1 + 1; d2 < 4; d2++) {
                const float* ha = branches[d1] + (long long)b * AH;
                const float* hb = branches[d2] + (long long)b * AH;
                float dot = 0.0f;
                for (int j = 0; j < AH; j++) dot += ha[j] * hb[j];
                float cos_sim = dot / (norms[d1] * norms[d2]);
                local_penalty += cos_sim * cos_sim;
                pair_idx++;
            }
        }
    }

    /* Block reduce */
    extern __shared__ float shmem[];
    shmem[tid] = local_penalty;
    __syncthreads();
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) shmem[tid] += shmem[tid + s];
        __syncthreads();
    }
    if (tid == 0) {
        /* Average over samples and pairs, scale by lambda */
        penalty_out[0] = lambda_indep * shmem[0] / (float)(N * 6);
    }
}
```

- [ ] **Step 3: Add temporal_consistency_penalty kernel (G10)**

```cuda
/* ================================================================== */
/* Kernel: temporal_consistency_penalty (G10)                          */
/* ================================================================== */

/**
 * Penalize Q-value differences between consecutive similar states.
 * Lipschitz regularizer: prevents memorizing bar positions.
 *
 * Grid: ceil((N-1) / 256), Block: 256.
 */
extern "C" __global__ void temporal_consistency_penalty(
    const float* __restrict__ h_s2,    /* [N, SH2] trunk activations */
    const float* __restrict__ q_values, /* [N, total_actions] Q-values */
    float* __restrict__ penalty_out,   /* [1] scalar accumulator (atomicAdd) */
    int N,
    int sh2,
    int total_actions,
    float lambda_tc,                   /* penalty weight (0.005) */
    float sim_threshold                /* similarity threshold (0.95) */
) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= N - 1) return;

    /* Cosine similarity between consecutive h_s2 */
    float dot = 0.0f, norm_a = 0.0f, norm_b = 0.0f;
    for (int j = 0; j < sh2; j++) {
        float a = h_s2[(long long)i * sh2 + j];
        float b = h_s2[(long long)(i + 1) * sh2 + j];
        dot += a * b;
        norm_a += a * a;
        norm_b += b * b;
    }
    float sim = dot / (sqrtf(norm_a) * sqrtf(norm_b) + 1e-8f);

    if (sim > sim_threshold) {
        /* Mean absolute Q-difference */
        float q_diff = 0.0f;
        for (int a = 0; a < total_actions; a++) {
            q_diff += fabsf(q_values[(long long)i * total_actions + a]
                          - q_values[(long long)(i + 1) * total_actions + a]);
        }
        q_diff /= (float)total_actions;

        float weight = (sim - sim_threshold) / (1.0f - sim_threshold);
        atomicAdd(penalty_out, lambda_tc * q_diff * weight / (float)(N - 1));
    }
}
```

- [ ] **Step 4: Verify CUDA compilation**

```bash
SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
```

- [ ] **Step 5: Commit**

```bash
cd /home/jgrusewski/Work/foxhunt && git add crates/ml/src/cuda_pipeline/experience_kernels.cu && git commit -m "feat(G5+G6+G10): epistemic_gate_magnitude, branch_independence, temporal_consistency kernels

G5: sigmoid gate on magnitude Q-values from ensemble variance.
G6: 6-way cosine similarity penalty between branch hidden activations.
G10: Lipschitz penalty on Q-diffs between consecutive similar states.

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

---

## Task 6: Wire G9 Regime Dropout into Forward Pass

Load the regime_dropout kernel and wire it into the Q-value forward path.

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`
- Modify: `crates/ml/src/trainers/dqn/fused_training.rs`

- [ ] **Step 1: Add regime dropout fields and kernel load**

In the struct:

```rust
    regime_dropout_kernel: CudaFunction,
    regime_dropout_epoch_seed: i32,
```

In the constructor, load from `EXPECTED_Q_CUBIN` (experience_kernels.cubin):

```rust
        let regime_dropout_kernel = exp_module.load_function("regime_dropout")
            .map_err(|e| MLError::ModelError(format!("regime_dropout load: {e}")))?;
```

- [ ] **Step 2: Add apply_regime_dropout method**

```rust
    pub(crate) fn apply_regime_dropout(&self, batch_size: usize, is_training: bool) -> Result<(), MLError> {
        let sh2 = self.config.shared_h2;
        let total = batch_size * sh2;
        unsafe {
            self.stream.launch_builder(&self.regime_dropout_kernel)
                .arg(&self.save_h_s2)
                .arg(&self.ptrs.states_buf)
                .arg(&(batch_size as i32))
                .arg(&(sh2 as i32))
                .arg(&(self.config.state_dim as i32))
                .arg(&0.15_f32)  // drop_rate
                .arg(&self.regime_dropout_epoch_seed)
                .arg(&(if is_training { 1_i32 } else { 0_i32 }))
                .launch(LaunchConfig::for_num_elems(total as u32))
                .map_err(|e| MLError::ModelError(format!("regime_dropout: {e}")))?;
        }
        Ok(())
    }

    pub fn set_regime_dropout_seed(&mut self, seed: i32) {
        self.regime_dropout_epoch_seed = seed;
    }
```

- [ ] **Step 3: Wire into forward path after Mamba2, before Q-attention**

In the Q-stats forward path, after `mamba2_step` and before `launch_q_mean_center`:

```rust
        // G9: Regime-aware dropout on enriched h_s2
        self.apply_regime_dropout(batch_size, true)?;
```

- [ ] **Step 4: Add passthrough and update seed at epoch start**

In `fused_training.rs`:

```rust
    pub(crate) fn apply_regime_dropout(&self, batch_size: usize, is_training: bool) -> Result<(), MLError> {
        self.trainer.apply_regime_dropout(batch_size, is_training)
    }
    pub fn set_regime_dropout_seed(&mut self, seed: i32) {
        self.trainer.set_regime_dropout_seed(seed);
    }
```

In `training_loop.rs` at epoch start:

```rust
        if let Some(ref mut fused) = self.fused_ctx {
            fused.set_regime_dropout_seed(epoch as i32);
        }
```

- [ ] **Step 5: Verify and commit**

```bash
SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
```

```bash
cd /home/jgrusewski/Work/foxhunt && git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs crates/ml/src/trainers/dqn/fused_training.rs crates/ml/src/trainers/dqn/trainer/training_loop.rs && git commit -m "feat(G9): wire regime-aware dropout into forward pass

Regime dropout applied to h_s2 after Mamba2 enrichment, before branches.
ADX/CUSUM quantized into 16 regime buckets. Same regime drops same
neurons per epoch (epoch_seed changes). drop_rate=0.15.

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

---

## Task 7: Wire G5 Epistemic Gate + G6 Branch Independence + G10 Temporal Consistency

Load and wire the 3 structural kernels into the training pipeline.

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`
- Modify: `crates/ml/src/trainers/dqn/fused_training.rs`

- [ ] **Step 1: Add fields, kernel loads, buffers**

Struct fields:

```rust
    epistemic_gate_kernel: CudaFunction,
    branch_indep_kernel: CudaFunction,
    temporal_consistency_kernel: CudaFunction,
    var_ema_pinned: *mut f32,
    var_ema_dev_ptr: u64,
    branch_indep_penalty_buf: CudaSlice<f32>,  // [1] scalar
    temporal_penalty_buf: CudaSlice<f32>,       // [1] scalar
```

Load all 3 from `EXPECTED_Q_CUBIN`. Allocate pinned var_ema and zero-init penalty buffers.

- [ ] **Step 2: Add launch methods**

```rust
    pub(crate) fn apply_epistemic_gate(&self, batch_size: usize) -> Result<(), MLError> {
        // Launch epistemic_gate_magnitude with q_out_buf, ensemble_var_buf, var_ema_dev_ptr
        // ...
    }

    pub(crate) fn compute_branch_independence(&self, batch_size: usize) -> Result<(), MLError> {
        // Launch branch_independence_penalty with save_h_b0..b3, penalty_buf
        // Grid=(1,1,1), Block=(256,1,1), shmem=256*sizeof(float)
        // ...
    }

    pub(crate) fn compute_temporal_consistency(&self, batch_size: usize) -> Result<(), MLError> {
        // Zero penalty_buf, launch temporal_consistency_penalty
        // ...
    }
```

- [ ] **Step 3: Wire into forward/training path**

- G5 (epistemic gate): after `apply_regime_gate`, before action selection
- G6 (branch independence): after branch forward, penalty added to loss
- G10 (temporal consistency): after compute_expected_q, penalty added to loss

- [ ] **Step 4: Add passthroughs in fused_training.rs**

- [ ] **Step 5: Verify and commit**

```bash
SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
```

```bash
cd /home/jgrusewski/Work/foxhunt && git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs crates/ml/src/trainers/dqn/fused_training.rs && git commit -m "feat(G5+G6+G10): wire epistemic gate, branch independence, temporal consistency

G5: ensemble variance gates magnitude Q-values (sigmoid scale).
G6: cosine similarity penalty on branch hidden activations (lambda=0.01).
G10: Q-diff penalty on consecutive similar states (lambda=0.005).

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

---

## Task 8: G3 — Purged Walk-Forward Validation

Replace single validation window with chronological purged walk-forward. This is the largest task — ~200 lines of Rust in the training loop.

**Files:**
- Modify: `crates/ml/src/trainers/dqn/trainer/mod.rs`
- Modify: `crates/ml/src/trainers/dqn/trainer/constructor.rs`
- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs`

- [ ] **Step 1: Add walk-forward state to DQNTrainer**

In `mod.rs`:

```rust
    /// Walk-forward validation: per-window Sharpe ratios.
    pub(crate) wf_window_sharpes: Vec<f64>,
    /// Walk-forward: number of validation windows (K=6).
    pub(crate) wf_num_windows: usize,
    /// Walk-forward: purge gap in bars between train and validation.
    pub(crate) wf_purge_bars: usize,
```

In `constructor.rs`:

```rust
            wf_window_sharpes: Vec::new(),
            wf_num_windows: 6,
            wf_purge_bars: 100,
```

- [ ] **Step 2: Implement walk-forward validation split**

Add method to DQNTrainer:

```rust
    /// Compute min-Sharpe across K purged walk-forward windows.
    /// Returns (min_sharpe, window_sharpes).
    fn compute_walk_forward_sharpe(&mut self, fused: &FusedTrainingCtx, total_bars: usize) -> Result<(f64, Vec<f64>), anyhow::Error> {
        let k = self.wf_num_windows;
        let window_size = total_bars / k;
        let purge = self.wf_purge_bars;
        let mut sharpes = Vec::with_capacity(k);

        // Strictly chronological: train on early windows, validate on later
        for fold in 0..3 {  // 3 folds (train on 1..4, val on 5; train on 1..3, val on 4; etc.)
            let val_start = (k - 2 - fold) * window_size + purge;
            let val_end = val_start + window_size;
            if val_end > total_bars { continue; }

            // Run validation backtest on window [val_start..val_end]
            let sharpe = fused.run_validation_backtest(val_start, val_end)?;
            sharpes.push(sharpe);
        }

        let min_sharpe = sharpes.iter().cloned().fold(f64::INFINITY, f64::min);
        Ok((min_sharpe, sharpes))
    }
```

- [ ] **Step 3: Replace val_Sharpe with min-Sharpe for checkpoint selection**

In the training loop where `val_sharpe_for_backtrack` is computed (around line 587):

```rust
            // G3: Use min-Sharpe across walk-forward windows (not single val_loss)
            let (min_sharpe, window_sharpes) = self.compute_walk_forward_sharpe(fused, total_bars)?;
            self.wf_window_sharpes = window_sharpes.clone();
            let val_sharpe_for_backtrack = min_sharpe;
            info!("G3 walk-forward: min_sharpe={:.2} windows={:?}", min_sharpe, window_sharpes);
```

- [ ] **Step 4: Verify and commit**

```bash
SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
```

```bash
cd /home/jgrusewski/Work/foxhunt && git add crates/ml/src/trainers/dqn/trainer/mod.rs crates/ml/src/trainers/dqn/trainer/constructor.rs crates/ml/src/trainers/dqn/trainer/training_loop.rs && git commit -m "feat(G3): purged walk-forward validation with min-Sharpe checkpoint selection

Split data into 6 windows. Chronological purged walk-forward with
100-bar purge gap. Checkpoint selection uses min(Sharpe across windows)
instead of single val_Sharpe. Backtracking uses min-Sharpe for plateau
detection and rewind decisions.

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

---

## Task 9: Build Verification + Integration Test

Final verification that all 10 generalization components compile and integrate correctly.

**Files:**
- All files from Tasks 1-8

- [ ] **Step 1: Full workspace check**

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

Expected: `Finished` with zero errors.

- [ ] **Step 2: Run unit tests**

```bash
SQLX_OFFLINE=true cargo test -p ml --lib 2>&1 | tail -20
```

- [ ] **Step 3: Verify all kernel names are loadable**

Check that all new kernel function names match between CUDA source and Rust load calls:
- `regime_dropout` in experience_kernels.cu
- `epistemic_gate_magnitude` in experience_kernels.cu
- `branch_independence_penalty` in experience_kernels.cu
- `temporal_consistency_penalty` in experience_kernels.cu

- [ ] **Step 4: Log component summary**

```
Components: 10 generalization enhancements
G1:  AdamW weight decay (trunk mask, config.weight_decay)
G2:  Transaction cost curriculum (anneal 0→100%, ensemble-gated)
G3:  Purged walk-forward validation (K=6, min-Sharpe)
G4:  Gamma annealing (0.90→0.95, atom utilization tied)
G5:  Epistemic-gated magnitude (ensemble variance → sigmoid gate)
G6:  Branch independence penalty (6-way cosine sim, lambda=0.01)
G7:  Counterfactual augmentation (50% direction flip in env_step)
G8:  L1-sparse first layer (soft thresholding on w_s1, lambda=1e-3)
G9:  Regime-aware dropout (ADX/CUSUM conditioned, rate=0.15)
G10: Temporal consistency (Q-diff penalty, sim>0.95, lambda=0.005)
```

- [ ] **Step 5: Final commit**

```bash
cd /home/jgrusewski/Work/foxhunt && git add -A && git status
```

```bash
git commit -m "chore: build verification for OOS generalization enhancement

All 10 generalization components compile and integrate:
G1 AdamW + G8 L1-sparse, G2 cost curriculum, G3 walk-forward,
G4 gamma anneal, G5 epistemic gate, G6 branch independence,
G7 counterfactual, G9 regime dropout, G10 temporal consistency.

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