# Introspective State Vector (ISV) — Design Spec

**Goal:** Replace the hardcoded Q-drift penalty with a self-awareness module that feeds training dynamics back into the model as learnable features. Instead of external penalties fighting the model, the model OBSERVES its own state (drift, uncertainty, gradient pressure, performance) and learns context-dependent responses. The model becomes meta-cognitive — it knows when it's drifting, uncertain, or degrading, and adapts its actions accordingly.

**Core Principle:** Penalty = blind force. Signal = learnable information. The model should understand it's drifting, not be punished for it.

---

## Architecture Overview

```
Market Features [B, 42]
        │
   ┌────┴────┐
   │  Trunk  │ (shared layers s1→s2, unchanged)
   └────┬────┘
        │
   h_s2 [B, SH2]
        │
   ┌────┴──────────────────────────────────────────┐
   │                                                │
   │  ISV Signals [8 pinned scalars, GPU-readable]: │
   │  ┌─ [0] q_drift (normalized)                   │
   │  ├─ [1] grad_norm_ema                           │
   │  ├─ [2] td_error_ema                            │
   │  ├─ [3] ensemble_var_ema                        │
   │  ├─ [4] ensemble_var_delta (velocity)           │
   │  ├─ [5] return_ema (performance)                │
   │  ├─ [6] atom_utilization                        │
   │  └─ [7] loss_ema                                │
   │         │                                       │
   │    ISV Encoder MLP                              │
   │    8 → 16 (ReLU) → 8                           │
   │         │                                       │
   │    ISV Embedding [8]                            │
   │         │                                       │
   │    ┌────┴──── Temporal ISV ────┐                │
   │    │ Rolling buffer [K=4, 8]   │                │
   │    │ Decay-weighted average    │                │
   │    └──────────┬────────────────┘                │
   │               │                                 │
   └───────────────┼─────────────────────────────────┘
                   │
     ┌─────────────┼──────────────┬──────────┬──────────┐
     │             │              │          │          │
  Direction    Magnitude       Order     Urgency     Risk
  (ISV gate)  (ISV gate)   (ISV gate) (ISV gate)  (h_s2 || ISV_raw)
     │             │              │          │          │
  confidence   confidence   confidence  confidence   │
   score         score        score      score       │
     │             │              │          │          │
     └──────── Branch Confidence Routing ──────────────┘
                        │
               Drift-Conditioned γ
               (ISV → effective_gamma)
                        │
                  Final Action
```

---

## Component Catalog

### PEARLS (established meta-RL)

**P1: TD-Error as Feature**
The temporal-difference error (predicted Q minus actual return) is the most fundamental learning signal in RL. RL² (Duan et al. 2016) showed that feeding TD-error back as an observation lets the model learn its own learning rate. High TD-error means "my predictions were wrong" — the model can learn to be cautious. We already compute TD-error in the C51 loss kernel. Route its EMA to ISV signal slot [2].

**P2: Action Entropy as Feature**
Policy entropy measures decisiveness. Low entropy = confident (one action dominates). High entropy = uncertain (uniform distribution). Computed from C51 atom probabilities: `H = -sum(p * log(p))` over 51 atoms. We track this as part of the entropy regularization already. Route the batch-mean entropy to ISV via `atom_utilization` (slot [6]) which is a proxy — high utilization correlates with low entropy (atoms are well-spread, distribution is peaked).

**P3: EMA of Own Returns**
Running average of the model's realized training returns. When return_ema is falling, the model sees "my strategy is degrading." This is what a human trader does — checks the equity curve. Route training Sharpe EMA (already computed in training_loop.rs) to ISV slot [5] via pinned memory.

### GEMS (less known, solid foundation)

**G1: Ensemble Disagreement Velocity**
Not just current ensemble variance, but its rate of change. Growing disagreement = entering unknown territory (distribution shift). Shrinking = converging on a learned pattern. The derivative of uncertainty is more informative than the level. Computed as `delta = (ensemble_var - ensemble_var_prev) / max(ensemble_var_prev, 1e-6)`. Positive delta = "I'm becoming more confused." ISV slot [4].

**G2: Gradient Norm as Feature**
The magnitude of gradients tells the model how aggressively the optimizer is correcting it. High gradient norm = "the optimizer thinks I'm very wrong." The model can learn: when gradient pressure is high, trust current weights less, be more conservative. We already compute gradient norms for clipping (`grad_norm_pinned`). EMA it and route to ISV slot [1].

**G3: Per-Branch Calibration (via Recursive Confidence — see N4)**
Track `predicted_Q - actual_return` per branch over a rolling window. A branch that's confidently wrong gets a negative calibration score. Folded into the Recursive Confidence novel (N4) below.

### NOVELS (genuinely new)

**N1: Introspective State Vector (ISV) — Unified Self-Awareness Module**
Instead of ad-hoc per-signal routing, ALL self-monitoring signals are packed into a single 8-dimensional vector that flows through a dedicated encoder MLP. The model develops a unified "how am I doing?" representation. No existing system provides comprehensive introspection — most meta-RL feeds back 1-2 signals, not a full state vector. The ISV is global (same for all samples in batch), broadcast via pinned device-mapped memory (zero memcpy).

**N2: Drift-Conditioned Discount (γ as Learned Function)**
Instead of fixed or externally-adapted gamma, the model's OWN drift assessment modulates the discount factor:

```
effective_gamma = base_gamma * sigmoid(w_gamma @ ISV_embedding + b_gamma)
```

High drift → model shortens its horizon ("my far-future Q-values are unreliable, focus near-term"). Low drift → model extends horizon ("Q-values are stable, I can trust longer predictions"). Gamma becomes endogenous — the model controls its own planning horizon based on self-assessed reliability. The sigmoid output ranges [0.5, 1.0] (clamped), so effective_gamma stays in [base_gamma*0.5, base_gamma].

**N3: Temporal ISV — Pattern Memory for Training Dynamics**
A rolling buffer of K=4 recent ISV snapshots with learned decay weights. The model sees not just "I'm drifting now" but "I've been drifting for 4 steps" or "drift just spiked after being stable." Computed as weighted average: `ISV_temporal[k] = sum(decay^t * ISV_history[t])` for each feature k. The decay weight is a single learned scalar (per ISV feature = 8 decay weights). Lightweight: 4×8=32 float buffer + 8 learned weights. No Mamba2 needed for 8 scalars.

**N4: Recursive Confidence — the Model Predicts Its Own Error**
An auxiliary head that predicts `|Q_predicted - Q_actual|` — the model's own future prediction error:

```
predicted_error = sigmoid(w_conf @ h_s2 + b_conf)  → [B, 1]
```

Supervised by actual absolute TD-error from the PREVIOUS step (lagged target, stored in pinned memory). The model learns to predict WHEN it will be wrong. When predicted_error is high, the risk branch automatically reduces exposure (predicted_error feeds into risk_budget_forward as an extra input). This creates a recursive self-improvement loop — the model's error prediction gets better each epoch, making it increasingly self-aware of its own limitations.

Loss: `L_conf = MSE(predicted_error, lagged_td_error_abs)`, weighted at 0.01× main loss (auxiliary, never dominates). Gradient flows through the shared trunk, teaching the trunk to encode error-predictive features.

**N5: Branch Confidence Routing — Selective Self-Doubt**
Each branch computes a confidence score from its Q-distribution:

```
confidence_d = max(softmax(Q_d)) / (1/A_d)    → ratio to uniform
```

When confidence is close to 1.0 (near uniform = no preference), the branch is uncertain. When the ISV shows instability (high drift, high grad_norm), uncertain branches get dampened:

```
effective_weight_d = confidence_d * sigmoid(ISV_gate_d)
Q_gated[b, d, a] = effective_weight_d * Q_advantage[b, d, a]
```

The model develops selective self-doubt — it can be confident about direction but uncertain about magnitude, and act accordingly. This prevents one uncertain branch from poisoning the action selection. Replaces the current `regime_branch_gate` with a strictly more powerful mechanism that incorporates both market regime AND model confidence.

---

## ISV Signal Computation (Pure GPU — Zero CPU Involvement)

All 8 ISV signals are stored in a **pinned device-mapped buffer** `isv_signals [8]`. A single-thread GPU kernel `isv_signal_update` reads from existing GPU-accessible pointers, computes all EMA updates, and writes the ISV vector in-place. **No CPU in the hot path.** The kernel runs after `graph_adam` replay, before the next `graph_forward`.

| Slot | Signal | GPU Source Pointer | Formula (computed in kernel) |
|------|--------|-------------------|------------------------------|
| 0 | `q_drift` | `q_mean_scratch_dev_ptr`, `q_mean_ema_dev_ptr` | `(q_mean - ema) / max(\|ema\|, 0.1)` |
| 1 | `grad_norm_ema` | `grad_norm_dev_ptr` | `0.95 * old + 0.05 * sqrt(grad_norm_sq)` |
| 2 | `td_error_ema` | `td_error_scratch_dev_ptr` (new pinned scalar) | `0.95 * old + 0.05 * mean_abs_td` |
| 3 | `ensemble_var_ema` | `ensemble_var_scratch_dev_ptr` (new pinned scalar) | `0.95 * old + 0.05 * mean_var` |
| 4 | `ensemble_var_delta` | Slot [3] internal | `(new_ema - old_ema) / max(old_ema, 1e-6)` |
| 5 | `reward_ema` | `reward_scratch_dev_ptr` (new pinned scalar) | `0.95 * old + 0.05 * batch_mean_reward` |
| 6 | `atom_util` | `atom_util_scratch_dev_ptr` (new pinned scalar) | EMA of atom utilization |
| 7 | `loss_ema` | `total_loss_dev_ptr` | `0.95 * old + 0.05 * total_loss` |

**New pinned scalars needed:** `td_error_scratch`, `ensemble_var_scratch`, `reward_scratch`, `atom_util_scratch` — each a single f32, pinned device-mapped. Written by existing GPU kernels (C51 loss writes mean |TD|, ensemble aggregate writes mean variance, etc.) via `atomicExch` or simple write at the end of those kernels. The `isv_signal_update` kernel then reads all of them.

### CUDA Kernel: `isv_signal_update`

Single-thread kernel that runs after each training step. Reads GPU-accessible scalars, computes EMA updates, writes ISV vector. Also shifts the temporal history buffer.

```cuda
extern "C" __global__ void isv_signal_update(
    float* __restrict__ isv_signals,          /* [8] pinned device-mapped — ISV vector */
    float* __restrict__ isv_history,          /* [K*8] pinned device-mapped — temporal buffer */
    const float* __restrict__ q_mean_ptr,     /* [1] pinned — current batch Q-mean */
    const float* __restrict__ q_mean_ema_ptr, /* [1] pinned — Q-mean EMA */
    const float* __restrict__ grad_norm_ptr,  /* [1] pinned — gradient norm² from Adam */
    const float* __restrict__ td_error_ptr,   /* [1] pinned — batch mean |TD-error| */
    const float* __restrict__ ens_var_ptr,    /* [1] pinned — batch mean ensemble var */
    const float* __restrict__ reward_ptr,     /* [1] pinned — batch mean reward */
    const float* __restrict__ atom_util_ptr,  /* [1] pinned — atom utilization */
    const float* __restrict__ loss_ptr,       /* [1] pinned — total loss */
    int K                                     /* temporal history length (4) */
) {
    if (threadIdx.x != 0) return;

    float alpha = 0.05f;  /* EMA smoothing */

    /* ── Shift temporal history: prepend current ISV, drop oldest ── */
    for (int t = K - 1; t > 0; t--)
        for (int k = 0; k < 8; k++)
            isv_history[t * 8 + k] = isv_history[(t - 1) * 8 + k];
    for (int k = 0; k < 8; k++)
        isv_history[k] = isv_signals[k];  /* save current before overwrite */

    /* ── Compute new ISV signals ── */

    /* [0] Q-drift: normalized deviation from EMA */
    float q_mean = q_mean_ptr[0];
    float q_ema  = q_mean_ema_ptr[0];
    isv_signals[0] = (q_mean - q_ema) / fmaxf(fabsf(q_ema), 0.1f);

    /* [1] Gradient norm EMA (grad_norm_ptr stores norm², take sqrt) */
    float gn = sqrtf(fmaxf(grad_norm_ptr[0], 0.0f));
    isv_signals[1] = (1.0f - alpha) * isv_signals[1] + alpha * gn;

    /* [2] TD-error EMA */
    isv_signals[2] = (1.0f - alpha) * isv_signals[2] + alpha * td_error_ptr[0];

    /* [3] Ensemble variance EMA (save old for delta) */
    float old_ens_ema = isv_signals[3];
    isv_signals[3] = (1.0f - alpha) * old_ens_ema + alpha * ens_var_ptr[0];

    /* [4] Ensemble variance velocity */
    isv_signals[4] = (isv_signals[3] - old_ens_ema) / fmaxf(old_ens_ema, 1e-6f);

    /* [5] Reward EMA (batch-level, not epoch-level) */
    isv_signals[5] = (1.0f - alpha) * isv_signals[5] + alpha * reward_ptr[0];

    /* [6] Atom utilization */
    isv_signals[6] = (1.0f - alpha) * isv_signals[6] + alpha * atom_util_ptr[0];

    /* [7] Loss EMA */
    isv_signals[7] = (1.0f - alpha) * isv_signals[7] + alpha * loss_ptr[0];
}
```

**Launch:** grid=(1,1,1), block=(1,1,1). Single thread, ~40 FLOPs. Negligible.

**Integration:** The existing kernels need small additions to write their batch-mean scalars to pinned scratch pointers:
- C51 loss kernel: at the end, one thread writes `mean(|td_errors|)` to `td_error_scratch_dev_ptr`
- Ensemble aggregate kernel: at the end, one thread writes `mean(ensemble_var)` to `ensemble_var_scratch_dev_ptr`
- Experience collector reward path: write batch mean reward to `reward_scratch_dev_ptr`
- Atom utilization kernel: already produces a scalar — route to `atom_util_scratch_dev_ptr`

These are all single-write additions to existing kernels, not new kernel launches.

---

## New Weight Tensors

NUM_WEIGHT_TENSORS: 68 → 78 (10 new tensors).

| Index | Name | Shape | Purpose |
|-------|------|-------|---------|
| 68 | `w_isv_fc1` | [16, 8] | ISV encoder layer 1 weights |
| 69 | `b_isv_fc1` | [16] | ISV encoder layer 1 bias |
| 70 | `w_isv_fc2` | [8, 16] | ISV encoder layer 2 weights |
| 71 | `b_isv_fc2` | [8] | ISV encoder layer 2 bias |
| 72 | `w_isv_gate` | [4, 8] | Branch confidence gating weights |
| 73 | `b_isv_gate` | [4] | Branch confidence gating bias |
| 74 | `w_isv_gamma` | [1, 8] | Drift-conditioned gamma weights |
| 75 | `b_isv_gamma` | [1] | Drift-conditioned gamma bias |
| 76 | `w_conf_fc` | [1, SH2] | Recursive confidence head weights |
| 77 | `b_conf_fc` | [1] | Recursive confidence head bias |

Plus 8 learned decay weights for temporal ISV (stored in `isv_decay_pinned [8]`, not in param buffer — updated by separate SGD like risk branch).

**Total new parameters in param buffer (indices 68-77):** 128 + 16 + 128 + 8 + 32 + 4 + 8 + 1 + SH2 + 1 = 326 + SH2. With SH2=256 (production default): **582 parameters**. Plus 8 decay weights in pinned memory (separate SGD). Negligible vs ~580K existing.

**Existing tensor size change:** `w_risk_fc` at index [64] grows from `AH * SH2 = 128 * 256 = 32768` to `AH * (SH2 + 9) = 128 * 265 = 33920` (+1152 elements). This is the ONLY existing tensor that changes — but it has blast radius: `risk_aligned_count`, `padded_byte_offset(param_sizes, 64)`, `risk_grad_buf`, and all risk SGD pointer arithmetic must be updated atomically with `compute_param_sizes`. The `risk_budget_backward` kernel stride changes from `j * SH2` to `j * (SH2 + 9)`.

---

## CUDA Kernel: `isv_forward`

Single kernel computes ISV encoder + branch gating + drift-conditioned gamma. One thread per sample.

```cuda
extern "C" __global__ void isv_forward(
    const float* __restrict__ isv_signals_dev_ptr,   /* [8] pinned device-mapped */
    const float* __restrict__ isv_history,            /* [K=4, 8] temporal buffer */
    const float* __restrict__ w_fc1,                  /* [16, 8] */
    const float* __restrict__ b_fc1,                  /* [16] */
    const float* __restrict__ w_fc2,                  /* [8, 16] */
    const float* __restrict__ b_fc2,                  /* [8] */
    const float* __restrict__ isv_decay,              /* [8] learned decay weights */
    const float* __restrict__ w_gate,                 /* [4, 8] */
    const float* __restrict__ b_gate,                 /* [4] */
    const float* __restrict__ w_gamma,                /* [1, 8] */
    const float* __restrict__ b_gamma,                /* [1] */
    float* __restrict__ isv_embedding_out,            /* [8] shared (single-instance) */
    float* __restrict__ branch_gate_out,              /* [4] shared (single-instance) */
    float* __restrict__ gamma_mod_out,                /* [1] shared (single-instance) */
    int K
) {
    /* Single thread — ISV is global, not per-sample */
    if (threadIdx.x != 0) return;

    /* Step 1: Temporal ISV — decay-weighted average of history + current */
    float isv_temporal[8];
    for (int k = 0; k < 8; k++) {
        float decay = 1.0f / (1.0f + expf(-isv_decay[k]));  /* sigmoid → (0,1) */
        float weighted_sum = isv_signals_dev_ptr[k];  /* current (weight=1) */
        float weight_sum = 1.0f;
        for (int t = 0; t < K; t++) {
            float w = powf(decay, (float)(t + 1));
            weighted_sum += w * isv_history[t * 8 + k];
            weight_sum += w;
        }
        isv_temporal[k] = weighted_sum / weight_sum;
    }

    /* Step 2: ISV encoder — FC1 (8→16, ReLU) → FC2 (16→8) */
    float hidden[16];
    for (int j = 0; j < 16; j++) {
        float val = b_fc1[j];
        for (int k = 0; k < 8; k++)
            val += w_fc1[j * 8 + k] * isv_temporal[k];
        hidden[j] = fmaxf(val, 0.0f);  /* ReLU */
    }

    float embedding[8];
    for (int k = 0; k < 8; k++) {
        float val = b_fc2[k];
        for (int j = 0; j < 16; j++)
            val += w_fc2[k * 16 + j] * hidden[j];
        embedding[k] = val;
    }

    /* Write embedding (single instance — downstream kernels read via pointer) */
    for (int k = 0; k < 8; k++)
        isv_embedding_out[k] = embedding[k];

    /* Step 3: Branch gating — softmax(W_gate @ embedding + b_gate) */
    float logits[4];
    float max_logit = -1e9f;
    for (int d = 0; d < 4; d++) {
        float val = b_gate[d];
        for (int k = 0; k < 8; k++)
            val += w_gate[d * 8 + k] * embedding[k];
        logits[d] = val;
        max_logit = fmaxf(max_logit, logits[d]);
    }
    float sum_exp = 0.0f;
    for (int d = 0; d < 4; d++) {
        logits[d] = expf(logits[d] - max_logit);
        sum_exp += logits[d];
    }
    for (int d = 0; d < 4; d++)
        branch_gate_out[d] = logits[d] / sum_exp;

    /* Step 4: Drift-conditioned gamma — sigmoid(W_gamma @ embedding + b_gamma) */
    float gamma_raw = b_gamma[0];
    for (int k = 0; k < 8; k++)
        gamma_raw += w_gamma[k] * embedding[k];
    /* Map to [0.5, 1.0] — 0.5 + 0.5*sigmoid(x). No clamp needed — formula
     * naturally produces this range. x→-∞: 0.5, x→+∞: 1.0. */
    gamma_mod_out[0] = 0.5f + 0.5f / (1.0f + expf(-gamma_raw));
}
```

**Launch:** grid=(1, 1, 1), block=(1, 1, 1). Single thread — ISV signals are global (same for all samples), so computing once and broadcasting is correct. The output buffers (`isv_embedding_out [8]`, `branch_gate_out [4]`, `gamma_mod_out [1]`) are single-instance, not per-sample. Downstream kernels read these shared buffers via pointer.

**Temporal ISV ring buffer convention:** `isv_history[0..8]` = most recent snapshot (1 step ago), `isv_history[8..16]` = 2 steps ago, etc. CPU prepends new snapshots by shifting existing entries down: `memmove(&history[8], &history[0], (K-1)*8*sizeof(float))`, then writes current ISV to `history[0..8]`.

---

## CUDA Kernel: `recursive_confidence_forward`

Auxiliary head that predicts the model's own TD-error.

```cuda
extern "C" __global__ void recursive_confidence_forward(
    const float* __restrict__ h_s2,        /* [B, SH2] */
    const float* __restrict__ w_conf,      /* [1, SH2] */
    const float* __restrict__ b_conf,      /* [1] */
    float* __restrict__ predicted_error,   /* [B] */
    int B, int SH2
) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= B) return;

    const float* h = h_s2 + (long long)i * SH2;
    float val = b_conf[0];
    for (int k = 0; k < SH2; k++)
        val += w_conf[k] * h[k];
    /* Sigmoid → [0, 1] range for error magnitude */
    predicted_error[i] = 1.0f / (1.0f + expf(-val));
}
```

**Supervision:** `L_conf = mean((predicted_error - lagged_td_error_abs)^2)` at 0.01× weight. `lagged_td_error_abs` is the batch-mean absolute TD-error from the PREVIOUS training step, stored in a pinned scalar.

---

## CUDA Kernel: `branch_confidence_routing`

Replaces `regime_branch_gate`. Combines ISV gating with per-branch confidence.

**Note on market regime:** The current `regime_branch_gate` uses ADX (state[40]) and CUSUM (state[41]) as market regime features. These are NOT lost — the trunk already processes them as input features, so branches see market regime through h_s2. ISV adds a NEW dimension: training dynamics awareness. `branch_confidence_routing` is strictly more powerful because it combines ISV (how the model is doing) with per-branch confidence (how certain each branch is).

```cuda
extern "C" __global__ void branch_confidence_routing(
    float* __restrict__ q_values,                    /* [B, total_actions] in-place */
    const float* __restrict__ branch_gate,           /* [4] shared from isv_forward */
    int B,
    int b0_size, int b1_size, int b2_size, int b3_size
) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= B) return;

    int branch_sizes[4] = { b0_size, b1_size, b2_size, b3_size };
    int total_actions = b0_size + b1_size + b2_size + b3_size;
    int branch_offset = 0;

    for (int d = 0; d < 4; d++) {
        int A_d = branch_sizes[d];

        /* Branch confidence from Q-value separation:
         * Q_max - Q_mean within the branch, normalized by Q_range.
         * High separation = confident (one action clearly preferred).
         * Low separation = uncertain (all actions similar).
         * Uses q_values directly (expected Q per action, already computed). */
        float q_max = -1e9f;
        float q_sum = 0.0f;
        for (int a = 0; a < A_d; a++) {
            float q = q_values[(long long)i * total_actions + branch_offset + a];
            q_max = fmaxf(q_max, q);
            q_sum += q;
        }
        float q_mean = q_sum / (float)A_d;
        float separation = q_max - q_mean;
        /* Normalize: sigmoid maps separation to (0, 1). Scale=5 gives
         * ~0.5 at separation=0 (uniform), ~0.99 at separation=1. */
        float confidence = 1.0f / (1.0f + expf(-5.0f * separation));

        /* Combined gate: ISV importance × branch confidence.
         * ISV gate from isv_forward (how training is going).
         * Confidence from Q-value separation (how decisive this branch is).
         * Floor at 0.3 — never fully mute a branch. */
        float isv_gate = branch_gate[d];  /* shared — same for all samples */
        float effective_weight = isv_gate * fmaxf(confidence, 0.3f);

        /* Apply to Q-values for this branch (in-place) */
        for (int a = 0; a < A_d; a++) {
            q_values[(long long)i * total_actions + branch_offset + a] *= effective_weight;
        }
        branch_offset += A_d;
    }
}
```

**Key change from regime_branch_gate:** No `q_distributions` parameter needed — confidence is computed directly from the Q-values that are already in the buffer. This avoids the incorrect indexing of packed advantage logits and is computationally cheaper (A_d additions per branch instead of A_d×NA exponentials).

---

## Risk Branch Integration

The risk branch currently takes `h_s2 [B, SH2]` as input. With ISV, it receives `h_s2 || ISV_raw || predicted_error`:

```
risk_input [B, SH2 + 8 + 1] = concat(h_s2, isv_signals, predicted_error)
```

This gives the risk branch the most detailed self-awareness. It sees:
- **h_s2**: market state representation
- **ISV raw signals**: 8 training dynamics indicators
- **predicted_error**: the model's own error prediction

The risk branch learns: "when I see high drift AND high predicted error AND high ensemble disagreement, reduce risk_budget → smaller positions, tighter CVaR."

**Change to `risk_budget_forward`:**

```cuda
extern "C" __global__ void risk_budget_forward(
    const float* __restrict__ h_s2,
    const float* __restrict__ isv_signals_dev_ptr,  /* [8] NEW: raw ISV */
    const float* __restrict__ predicted_error,      /* [B] NEW: recursive confidence */
    const float* __restrict__ w_risk_fc,            /* [AH, SH2+9] NEW: wider input */
    const float* __restrict__ b_risk_fc,
    const float* __restrict__ w_risk_out,
    const float* __restrict__ b_risk_out,
    float* __restrict__ risk_hidden,
    float* __restrict__ risk_budget_out,
    int B, int SH2, int AH
) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= B) return;

    int input_dim = SH2 + 9;  /* SH2 + 8 ISV + 1 predicted_error */
    const float* h = h_s2 + (long long)i * SH2;
    float* h_risk = risk_hidden + (long long)i * AH;

    for (int j = 0; j < AH; j++) {
        float val = b_risk_fc[j];
        /* h_s2 portion */
        for (int k = 0; k < SH2; k++)
            val += w_risk_fc[(long long)j * input_dim + k] * h[k];
        /* ISV signals (broadcast — same for all samples) */
        for (int k = 0; k < 8; k++)
            val += w_risk_fc[(long long)j * input_dim + SH2 + k] * isv_signals_dev_ptr[k];
        /* Predicted error (per-sample) */
        val += w_risk_fc[(long long)j * input_dim + SH2 + 8] * predicted_error[i];
        h_risk[j] = fmaxf(val, 0.0f);
    }

    float raw = b_risk_out[0];
    for (int j = 0; j < AH; j++)
        raw += w_risk_out[j] * h_risk[j];
    risk_budget_out[i] = 1.0f / (1.0f + expf(-raw));
}
```

**Weight tensor change:** `w_risk_fc` at index [64] grows from `AH * SH2` to `AH * (SH2 + 9)`. This is the ONLY existing tensor that changes size.

---

## Drift-Conditioned Gamma — Integration

The `gamma_mod_out [B]` from `isv_forward` modulates the Bellman target discount:

```
effective_gamma[b] = base_gamma * gamma_mod[b]
```

Where `gamma_mod` ∈ [0.5, 1.0], so effective_gamma ∈ [base_gamma*0.5, base_gamma].

**Integration point:** In `c51_loss_kernel.cu`, the `gamma` parameter becomes `gamma_buf [B]` — a per-sample buffer filled by `isv_forward`. The kernel reads `gamma_buf[b]` instead of a scalar `gamma`.

Currently `c51_loss_kernel` takes `float gamma` as a scalar. Change to:

```cuda
const float* __restrict__ gamma_buf   /* [B] per-sample effective gamma */
```

And in the Bellman projection:
```cuda
float g = gamma_buf[b];  /* was: float g = gamma; */
```

The per-sample gamma buffer is allocated once (`[max_batch_size]`) and filled by a `fill_gamma_buf` kernel that broadcasts `base_gamma * gamma_mod_out[0]` to all B entries. When ISV is not yet trained (early epochs), `gamma_mod_out[0]` starts at 0.75 (zero-init → sigmoid(0)=0.5 → 0.5+0.5*0.5=0.75), so gamma starts slightly conservative and adapts.

**CUDA Graph compatibility:** `c51_loss_kernel.cu` must be recompiled with the new signature. Set `graph_forward = None` after the kernel change to force recapture on next `train_step()`.

---

## Forward Pass Integration Order

```
1. trunk forward (cuBLAS)              → h_s2 [B, SH2]
2. mamba2_step()                       → h_enriched [B, SH2]
3. isv_forward()                       → isv_embedding [B, 8], branch_gate [B, 4], gamma_mod [B]
4. recursive_confidence_forward()      → predicted_error [B]
5. fill gamma_buf = base_gamma * gamma_mod
6. risk_budget_forward(h_s2, isv_signals, predicted_error) → risk_budget [B]
7. branch heads forward (cuBLAS)       → Q values [B, total_actions]
8. compute_expected_q()                → expected Q per action
9. branch_confidence_routing(Q, branch_gate) → gated Q (uses Q-value separation for confidence)
10. apply_risk_budget()                → scaled Q
11. action selection
```

Steps 3-5 are new. Step 6 changes (wider input). Step 9 replaces `regime_branch_gate` — needs Q-values already computed (step 8) to measure per-branch confidence from Q-value separation.

---

## Backward Pass

**ISV encoder backward:** The ISV encoder weights receive gradients from two paths:
1. Branch gating loss (through branch_confidence_routing)
2. Gamma modulation loss (through C51 Bellman target)

Both are small MLP backward passes (8→16→8). Computed analytically in a `isv_backward` kernel, gradients accumulated into the main `grad_buf` at offsets 68-75.

**Recursive confidence backward:** MSE loss between predicted_error and lagged_td_error_abs. Gradient flows through w_conf_fc into the trunk gradient (accumulated into existing `bw_d_h_s2`).

**Temporal ISV decay backward:** The 8 decay weights are updated by separate SGD (like risk branch), not part of main Adam. Learning rate: 0.001. Gradient: chain rule through the decay-weighted average formula.

**Risk branch backward:** Already exists but must be updated for wider input. The existing kernel at `experience_kernels.cu:4750` uses stride `j * SH2` for `d_w_risk_fc`. Must change to `j * (SH2 + 9)` AND add gradient accumulation for the 9 new ISV+error columns:

```cuda
/* Existing: gradient for h_s2 portion (columns 0..SH2) */
for (int k = 0; k < SH2; k++)
    atomicAdd(&d_w_risk_fc[(long long)j * input_dim + k], d_hj * h[k]);
/* NEW: gradient for ISV portion (columns SH2..SH2+8) */
for (int k = 0; k < 8; k++)
    atomicAdd(&d_w_risk_fc[(long long)j * input_dim + SH2 + k], d_hj * isv_signals_dev_ptr[k]);
/* NEW: gradient for predicted_error (column SH2+8) */
atomicAdd(&d_w_risk_fc[(long long)j * input_dim + SH2 + 8], d_hj * predicted_error[i]);
```

Where `input_dim = SH2 + 9`. The `risk_grad_buf` must be reallocated to match the new `risk_aligned_count`.

---

## Pinned Memory Layout

New pinned allocations (all `cuMemAllocHost` + `cuMemHostGetDevicePointer`):

| Name | Size | Purpose | Written by |
|------|------|---------|-----------|
| `isv_signals_pinned` | 8 × f32 = 32 B | ISV signal vector | `isv_signal_update` GPU kernel |
| `isv_history_pinned` | 4 × 8 × f32 = 128 B | Temporal ISV rolling buffer | `isv_signal_update` GPU kernel |
| `isv_decay_pinned` | 8 × f32 = 32 B | Temporal decay weights (learned) | ISV decay SGD GPU kernel |
| `td_error_scratch_pinned` | 1 × f32 = 4 B | Batch mean \|TD-error\| | C51 loss kernel (1 atomicExch) |
| `ensemble_var_scratch_pinned` | 1 × f32 = 4 B | Batch mean ensemble variance | Ensemble aggregate kernel |
| `reward_scratch_pinned` | 1 × f32 = 4 B | Batch mean reward | Reward aggregation |
| `atom_util_scratch_pinned` | 1 × f32 = 4 B | Atom utilization | Atom util kernel |
| `lagged_td_error_pinned` | 1 × f32 = 4 B | Previous step's mean \|TD-error\| (recursive confidence target) | `isv_signal_update` copies from slot [2] before overwrite |

Total: 212 bytes. All written by GPU kernels, read by GPU kernels. **Zero CPU involvement.**

---

## What ISV Replaces

| Old Component | Replacement | Why |
|---------------|-------------|-----|
| Hardcoded Q-drift penalty (removed in TEA) | ISV slot [0] q_drift → branch gating + risk input | Model learns context-dependent drift response |
| `regime_branch_gate` (4 inputs: ADX, CUSUM, Q_gap, atom_util) | `branch_confidence_routing` (ISV embedding + per-branch confidence) | Strictly more powerful — incorporates both market regime AND model self-awareness |
| Fixed gamma / externally-adapted gamma | Drift-conditioned gamma | Model controls its own planning horizon |
| External E1 Q-correction (additive bias) | ISV → risk_budget reduction when drift is high | Model learns WHEN correction is needed, not a blanket offset |

**E1 enrichment (Q-correction) is NOT removed** — it provides the external "ground truth" signal from validation. ISV provides the internal signal from training dynamics. They are complementary: E1 corrects from eval data (slow, epoch-level), ISV adapts from training dynamics (fast, per-step).

---

## Initialization

- ISV encoder weights: Xavier uniform (small MLP, standard init)
- Branch gate weights: zero-initialized biases, small Xavier weights → initial softmax ≈ uniform (no branch preferred)
- Gamma weights: zero-initialized → initial sigmoid = 0.5 → initial gamma_mod = 0.75 → slightly conservative (shorter horizon while ISV learns)
- Recursive confidence weights: zero-initialized → initial predicted_error ≈ 0.5 (moderate uncertainty)
- Decay weights: initialized to 0.0 → sigmoid(0) = 0.5 → moderate decay (geometric with ratio 0.5)
- Risk branch wider input: existing SH2 portion from current checkpoint, new ISV+error columns zero-initialized (no disruption to existing risk behavior until ISV trains)

---

## VRAM Budget

| Buffer | Size (B=4096, SH2=256, AH=128) |
|--------|------|
| `isv_embedding_out` [8] | 32 B |
| `branch_gate_out` [4] | 16 B |
| `gamma_mod_out` [1] | 4 B |
| `gamma_buf` [B] | 16 KB |
| `predicted_error` [B] | 16 KB |
| ISV weight tensors (582 params) | 2.3 KB |
| `w_risk_fc` growth (+1152 elements) | 4.5 KB |
| Pinned memory (signals + history + decay + lagged_td) | 196 B |
| **Total** | **~39 KB** |

Negligible. Less than 0.0001% of H100 80GB.

---

## Target Network

ISV weights (indices 68-77) are **online-only** — they are NOT copied to `target_params_buf`. Rationale: ISV conditioning reflects the model's current training dynamics (drift, gradient pressure, loss trajectory). The target network should use stable behavior independent of training dynamics:

- **Target network forward:** Use `base_gamma` directly (no drift-conditioned gamma). Fill `gamma_buf` with `base_gamma` for target Q computation.
- **Target network gating:** Use uniform branch gates (all 0.25) — no ISV-based branch suppression for Bellman targets.
- **Target network risk:** `risk_budget_forward` for target uses zeros for ISV slots and 0.5 for predicted_error — neutral conditioning.

This means `target_params_buf` stays at its current size (indices 0-67). The ISV encoder, gamma, and gate weights only exist in `params_buf` and `grad_buf`. No EMA sync needed for ISV weights.

---

## What the Model Learns

With ISV, the model develops context-dependent responses to its own state:

| Scenario | ISV Signals | Learned Response |
|----------|-------------|------------------|
| Stable training, low drift | q_drift≈0, grad_norm low, loss steady | Long gamma (trust far-future Q), full risk_budget |
| Q-drift high, trending market | q_drift>0, return_ema positive | Moderate gamma (drift is from real gains), maintain risk |
| Q-drift high, mean-reverting | q_drift>0, return_ema falling | Short gamma (overestimation danger), reduce risk |
| High uncertainty, exploring | ensemble_var high, ensemble_var_delta>0 | Reduce magnitude, prefer flat, low risk_budget |
| Loss spike, gradient pressure | loss_ema spiking, grad_norm high | Conservative actions, short gamma, low confidence routing |
| Model converging | td_error falling, loss falling | Extend gamma, increase risk_budget, trust branch confidence |

The key difference from hardcoded penalties: the model can learn that the SAME drift level means different things in different contexts. Q-drift=+3 during a strong trend is fine. Q-drift=+3 during consolidation is dangerous. ISV + market features let the model make this distinction.

---

## Training Step Flow (Pure GPU)

```
Per training step:
  1. replay graph_forward           → C51/MSE loss, gradients, td_errors
  2. inject auxiliary gradients     → IQN, ensemble, CQL into grad_buf
  3. replay graph_adam              → Adam update, grad_norm computed
  4. isv_signal_update()            → reads all GPU scalars, updates ISV [8] + history [K,8]
  5. (next step's forward):
     5a. isv_forward()              → encoder MLP, branch gate, gamma_mod (reads ISV from step 4)
     5b. recursive_confidence_fwd() → predicted_error [B]
     5c. fill gamma_buf             → base_gamma * gamma_mod[0], broadcast to [B]
     5d. risk_budget_forward()      → reads h_s2 + ISV signals + predicted_error
     5e. branch heads (cuBLAS)      → Q values
     5f. branch_confidence_routing()→ gated Q
     5g. apply_risk_budget()        → scaled Q
     5h. action selection
```

Step 4 (`isv_signal_update`) is a standalone kernel launch — NOT captured in graph_forward or graph_adam. It runs between graph replays on the training stream. Steps 5a-5c are new kernel launches inside graph_forward capture.

---

## Implementation Order

1. **Pinned scratch scalars** — allocate td_error_scratch, ensemble_var_scratch, reward_scratch, atom_util_scratch (4 pinned scalars). Add single-write lines to existing C51 loss, ensemble aggregate, and atom util kernels.
2. **ISV signal buffer + history** — allocate isv_signals [8], isv_history [K=4, 8], isv_decay [8] pinned. Write `isv_signal_update` kernel. Wire after graph_adam replay.
3. **ISV encoder kernel** (`isv_forward`) — single-thread MLP forward, branch gating output, gamma modulation output. Write kernel + Rust launch.
4. **Temporal ISV** — decay-weighted average in isv_forward (reads isv_history + isv_decay).
5. **Drift-conditioned gamma** — change c51_loss_kernel gamma from scalar `float` to `const float* gamma_buf [B]`. Recompile cubin. Add `fill_gamma_buf` kernel (broadcasts `base_gamma * gamma_mod[0]`). Set `graph_forward = None` for recapture.
6. **Branch confidence routing** — write `branch_confidence_routing` kernel. Replace `regime_branch_gate` at all call sites: `apply_regime_gate()`, `reduce_current_q_stats()`, `submit_forward_ops_ddqn()`. Drop `regime_gate_kernel` field.
7. **Risk branch wider input** — change `w_risk_fc` in `compute_param_sizes` from `AH * SH2` to `AH * (SH2 + 9)`. Update `risk_budget_forward` + `risk_budget_backward` kernels for wider stride. Reallocate `risk_grad_buf`. Update `risk_aligned_count` and all `padded_byte_offset(param_sizes, 64)` callers.
8. **Recursive confidence head** — write `recursive_confidence_forward` kernel + MSE backward. Route `lagged_td_error_pinned` from `isv_signal_update` (copies slot [2] before overwrite).
9. **Weight tensors** — NUM_WEIGHT_TENSORS 68→78, extend `compute_param_sizes`, `xavier_init_params_buf`, `fan_dims`. Target network: ISV weights online-only (indices 68-77 NOT in target_params_buf).
10. **Graph recapture + smoke tests (Phase 1)** — invalidate all graphs, verify 0 errors.
11. **Regime ISV extension** — expand ISV_DIM 8→12. Add 4 regime signals to `isv_signal_update` (needs states_buf pointer). Update `isv_forward` w_fc1 from [16,8]→[16,12]. Update isv_history K*8→K*12, isv_decay 8→12.
12. **Regime-conditioned Mamba2 decay** — multiply Mamba2 decay by regime_stability (ISV slot [11]). Modify `mamba2_temporal_scan` kernel.
13. **Risk branch regime extension** — widen risk input from SH2+9→SH2+13. Update forward/backward kernels and w_risk_fc size.
14. **Final smoke test + compute-sanitizer (Phase 2)** — invalidate all graphs (`graph_forward = None`, `graph_adam = None`, `graph_forward_ddqn = None`). Run smoke test: `FOXHUNT_TEST_DATA=test_data/futures-baseline cargo test -p ml --lib -- test_generalization_components_smoke --include-ignored --nocapture`. Run compute-sanitizer: `compute-sanitizer --tool memcheck --print-limit 5 target/debug/deps/ml-* "test_generalization_components_smoke" --test-threads=1 --include-ignored`. **Target: 0 errors.** Run full test suite: `SQLX_OFFLINE=true cargo test -p ml --lib`. **Target: 899+ passed, 0 failed.**

---

## Phase 2: Regime Awareness Extension

The core ISV (Phase 1, tasks 1-9) makes the model aware of its own training dynamics. Phase 2 extends ISV with **market regime awareness** — the model learns to detect regime transitions and adapt its behavior during the most dangerous periods.

### The Problem

Markets don't have clean regime boundaries. ADX=25 could mean:
- **Entering a trend** (from 15) → direction branch should be aggressive
- **Leaving a trend** (from 40) → direction branch should be cautious
- **Consolidating** → prefer flat

Same value, three opposite strategies. The model currently sees WHERE it is but not WHERE it's GOING. Regime transitions are the most dangerous periods — the model's strategy for the previous regime is maximally wrong.

### Extension Architecture

Expand ISV from 8→12 signals. The 4 new slots are computed in `isv_signal_update` alongside the existing 8 — same kernel, same pure-GPU path. ISV encoder MLP input grows from 8→12 (hidden stays 16, output stays 8).

### New ISV Signals

| Slot | Signal | Formula | Source |
|------|--------|---------|--------|
| 8 | `regime_velocity` | `\|ADX[t] - ADX[t-1]\| + \|CUSUM[t] - CUSUM[t-1]\|` | States buffer, indices 40-41 |
| 9 | `regime_disagreement` | `\|norm(ADX) - norm(CUSUM)\|` | States buffer, normalized |
| 10 | `regime_transition_ema` | EMA of regime_velocity — detects sustained transitions vs noise | Derived from slot [8] |
| 11 | `regime_stability` | `1.0 - sigmoid(5 * regime_velocity)` — inverse of transition strength | Derived from slot [8] |

### PEARLS

**P4: Regime Velocity Features**
`d(ADX)/dt` and `d(CUSUM)/dt` computed as bar-over-bar differences. The rate of change tells the model whether regime indicators are stable, entering, or leaving. These are the single most informative regime features missing from the current system. Computed in `isv_signal_update` by reading the states buffer (one sample's ADX/CUSUM) and comparing to an EMA.

**P5: Regime Disagreement**
When ADX says "strong trend" but CUSUM fires a change signal, that's regime ambiguity. High disagreement = the model should be conservative. The risk branch sees this through ISV and reduces exposure automatically.

### GEMS

**G4: Regime-Aware ISV**
Regime uncertainty becomes an ISV signal. When `regime_velocity` is high (transition in progress), the ISV encoder outputs lower branch gate weights and shorter gamma. The model automatically becomes conservative during transitions — through learning, not hardcoded rules.

**G5: Soft Regime Mixture**
The `branch_confidence_routing` kernel already uses ISV gate × Q-confidence. With regime signals in ISV, this becomes a **continuous regime mixture** — the model can express "60% trending, 40% mean-reverting" through the softmax branch gates. This replaces the old binary regime classification.

### NOVELS

**N6: Cross-Regime Temporal Decay**
The Mamba2 temporal scan uses fixed exponential decay for h_s2 history. With regime awareness, the decay becomes **regime-conditioned**: during transitions, recent bars get exponentially more weight (faster forgetting of old-regime patterns). During stable regimes, longer history is preserved.

Implementation: multiply Mamba2 decay by `regime_stability` (ISV slot [11]):
```
effective_decay = base_decay * regime_stability
```
When stability=1 (stable regime), full history. When stability=0.3 (transition), rapid forgetting. The model learns to adapt its memory horizon to regime dynamics.

**N7: Regime-Conditioned Gamma (stacks with drift-conditioned)**
During regime transitions, far-future Q-values estimated under the old regime are unreliable. The gamma modulation from ISV already shortens horizon during drift — adding regime_stability to the gamma computation makes it respond to BOTH training dynamics AND market regime:

```
gamma_mod = 0.5 + 0.5 * sigmoid(w_gamma @ [ISV_embedding including regime])
```

No separate mechanism needed — regime signals flow through the same ISV encoder to the same gamma output. The model learns when to be myopic (high drift + regime transition) vs patient (stable training + stable regime).

### Implementation Changes

| Component | Change |
|-----------|--------|
| `ISV_DIM` | 8 → 12 |
| `isv_signal_update` kernel | Add 4 regime signal computations. Needs `states_buf` pointer + sample index for ADX/CUSUM. |
| `isv_forward` kernel | w_fc1 grows from [16,8] to [16,12]. Input temporal averaging over 12 dims. |
| Weight tensor [68] | w_isv_fc1 grows from 16*8=128 to 16*12=192 elements |
| `isv_history` buffer | K*8=32 → K*12=48 floats |
| `isv_decay` buffer | 8 → 12 learned weights |
| Mamba2 temporal scan | Multiply decay by `regime_stability` from ISV slot [11] |
| Risk branch input | SH2+9 → SH2+13 (4 more ISV signals) |
| Weight tensor [64] | w_risk_fc grows from AH*(SH2+9) to AH*(SH2+13) |

### Regime Signal Computation (in `isv_signal_update`)

The kernel needs access to ADX/CUSUM from the current batch. Since `isv_signal_update` runs after `reduce_current_q_stats` which has already processed the batch, the states buffer is available. Pass `states_buf` pointer and read sample 0's features as representative:

```cuda
    /* [8] Regime velocity: |ΔADX| + |ΔCUSUM| */
    float adx  = states_ptr[0 * state_dim + 40];   /* sample 0 */
    float cusum = states_ptr[0 * state_dim + 41];
    float adx_delta = fabsf(adx - isv_signals[8]);  /* previous ADX stored here */
    float cusum_delta = fabsf(cusum - isv_signals[9]); /* previous CUSUM stored here */
    /* Temporarily store raw values for next step's delta */
    float prev_adx = isv_signals[8];
    float prev_cusum = isv_signals[9];
    isv_signals[8] = adx;  /* will be overwritten with velocity below */
    isv_signals[9] = cusum;

    float regime_vel = adx_delta + cusum_delta;
    isv_signals[8] = (1.0f - alpha) * prev_adx + alpha * regime_vel;  /* velocity EMA */

    /* [9] Regime disagreement: |normalized difference| */
    float norm_adx = adx / fmaxf(fabsf(adx) + fabsf(cusum), 1e-6f);
    float norm_cusum = cusum / fmaxf(fabsf(adx) + fabsf(cusum), 1e-6f);
    isv_signals[9] = fabsf(norm_adx - norm_cusum);

    /* [10] Regime transition EMA */
    isv_signals[10] = (1.0f - alpha) * isv_signals[10] + alpha * regime_vel;

    /* [11] Regime stability: 1 - sigmoid(5 * velocity) */
    isv_signals[11] = 1.0f - 1.0f / (1.0f + expf(-5.0f * regime_vel));
```

Note: This requires passing `states_buf` pointer and `state_dim` to the kernel. Updated signature adds 2 parameters.

---

## Phase 3: Dynamic Inference Brain — Feature Gating + Temporal Routing

Phases 1-2 give the model self-awareness (ISV) and market awareness (regime signals). Phase 3 uses the ISV embedding to dynamically modulate **what the model attends to** and **how much history it uses** — creating a model whose computational behavior adapts at inference without weight updates.

### The Complete Adaptive Pathway

At inference, without any gradient updates, the model dynamically adjusts 7 pathways:

```
              ISV Embedding [8]
                    │
    ┌───────┬───────┼───────┬───────┬──────────┬──────────┐
    │       │       │       │       │          │          │
    ▼       ▼       ▼       ▼       ▼          ▼          ▼
 Feature  Temporal Branch  Gamma   Risk    Mamba2     Branch
 Gating   Routing  Gating  Mod    Budget   Decay    Confidence
    │       │       │       │       │          │          │
 WHICH    HOW     WHICH   HOW    HOW      HOW        HOW
features  much   branches FAR    MUCH     MUCH      CERTAIN
 matter   history  matter  plan   risk     memory    each
          to use          ahead   to take  to keep   branch is
```

**Critical insight:** At inference, training dynamics signals (grad_norm, td_error, loss) freeze at their last training EMA values. But market signals (regime_velocity, regime_disagreement, Q-drift, ensemble_var) are LIVE — they're computed from the current market state. So the model adapts to unseen regimes through the live signals while retaining training-learned behavior through the frozen signals.

### PEARL: Feature-Level ISV Gating (P6)

The ISV embedding modulates WHICH features in h_s2 the branches should attend to. A single projection produces a per-feature gate:

```
feature_gate [SH2] = sigmoid(w_feature_gate @ isv_embedding + b_feature_gate)
h_s2_gated = h_s2 * feature_gate  // element-wise
```

Applied AFTER Mamba2 enrichment, BEFORE branch heads. During trends, momentum features (OFI signals) get amplified. During mean-reversion, depth imbalance features get amplified. The model learns the mapping.

This piggybacks on the existing Variable Selection Network (VSN, indices 26-33) which already does per-branch feature selection. ISV feature gating acts as a **pre-VSN filter** — it modulates which features reach the VSN at all.

```cuda
extern "C" __global__ void isv_feature_gate(
    float* __restrict__ h_s2,                    /* [B, SH2] in-place */
    const float* __restrict__ isv_embedding,     /* [8] shared */
    const float* __restrict__ w_feature_gate,    /* [SH2, 8] */
    const float* __restrict__ b_feature_gate,    /* [SH2] */
    int B, int SH2
) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= B) return;

    for (int j = 0; j < SH2; j++) {
        float gate_val = b_feature_gate[j];
        for (int k = 0; k < 8; k++)
            gate_val += w_feature_gate[(long long)j * 8 + k] * isv_embedding[k];
        float gate = 1.0f / (1.0f + expf(-gate_val));
        h_s2[(long long)i * SH2 + j] *= gate;
    }
}
```

**Init:** b_feature_gate = 2.0 (sigmoid(2)≈0.88 — near pass-through at start). The model only deviates when it helps. No feature is aggressively gated initially.

### GEM: Per-Feature Temporal Routing (G6)

Different features need different amounts of history. Momentum (OFI) is autocorrelated in trends but noise in ranges. Mean-reversion signals are the opposite. The ISV embedding controls how much Mamba2 temporal context each feature gets:

```
temporal_weight [SH2] = sigmoid(w_temporal_route @ isv_embedding + b_temporal_route)
h_enriched[j] = h_s2[j] + temporal_weight[j] * temporal_context[j]
```

This replaces the current fixed residual `h_enriched = h_s2 + temporal_context` in `mamba2_temporal_scan` with a per-feature weighted residual. Some features get full temporal enrichment, others get none — controlled by ISV.

Implementation: modify `mamba2_temporal_scan` to accept `temporal_weight [SH2]` buffer. The final line (59) changes:

```cuda
    /* Current: */
    out[j] = h_current[j] + ctx;
    
    /* With temporal routing: */
    out[j] = h_current[j] + temporal_weight[j] * ctx;
```

Where `temporal_weight` is pre-computed by a separate `isv_temporal_route` kernel:

```cuda
extern "C" __global__ void isv_temporal_route(
    const float* __restrict__ isv_embedding,     /* [8] shared */
    const float* __restrict__ w_temporal_route,  /* [SH2, 8] */
    const float* __restrict__ b_temporal_route,  /* [SH2] */
    float* __restrict__ temporal_weight,         /* [SH2] output */
    int SH2
) {
    if (threadIdx.x != 0) return;
    for (int j = 0; j < SH2; j++) {
        float val = b_temporal_route[j];
        for (int k = 0; k < 8; k++)
            val += w_temporal_route[(long long)j * 8 + k] * isv_embedding[k];
        temporal_weight[j] = 1.0f / (1.0f + expf(-val));
    }
}
```

**Init:** b_temporal_route = 2.0 (≈0.88 — mostly use temporal context at start, learn to suppress when appropriate).

### NOVEL: Dynamic Inference Adaptation to Unseen Data (N8)

The combination of all pathways creates a model that handles unseen market conditions through a principled cascade:

1. **Unseen regime appears** → regime_velocity spikes, regime_stability drops
2. **Mamba2 decay shortens** (N6) → model forgets old-regime patterns within 2-3 bars
3. **Feature gate shifts** (P6) → model amplifies features relevant to the new dynamics
4. **Temporal routing changes** (G6) → momentum features lose history weight, stable features keep it
5. **Branch gating shifts** → branches re-weight based on ISV's response to new signals
6. **Gamma shortens** → model focuses on near-term (doesn't extrapolate stale Q-values)
7. **Risk budget drops** → position size reduces until the model "understands" the new regime

All 7 steps happen within 2-3 bars of the regime change, without any weight updates. The model's INFERENCE BEHAVIOR adapts in real-time. This is the "dynamic brain" — not just a policy network but a **self-modifying inference system**.

### NOVEL: Inference Cold Start Graceful Degradation (N9)

At inference start, ISV history is empty (all zeros). The model gracefully handles this:
- Feature gate: sigmoid(b=2.0) ≈ 0.88 → near pass-through (no gating)
- Temporal routing: sigmoid(b=2.0) ≈ 0.88 → full temporal context
- Branch gate: softmax(0,0,0,0) = [0.25, 0.25, 0.25, 0.25] → uniform
- Gamma mod: 0.5 + 0.5*sigmoid(0) = 0.75 → slightly conservative horizon
- Risk budget: neutral ISV → moderate risk

Within 4-8 bars, ISV accumulates enough signal to start meaningful adaptation. The cold-start behavior is conservative-by-default — the model doesn't take large positions until it has enough self-awareness to justify them. This is the correct behavior for a trading system encountering new data.

### Phase 3 Weight Tensors

NUM_WEIGHT_TENSORS: 78 → 82 (4 new tensors).

| Index | Name | Shape | Purpose |
|-------|------|-------|---------|
| 78 | `w_feature_gate` | [SH2, 8] | ISV → per-feature amplitude gate |
| 79 | `b_feature_gate` | [SH2] | Feature gate bias (init=2.0) |
| 80 | `w_temporal_route` | [SH2, 8] | ISV → per-feature temporal depth |
| 81 | `b_temporal_route` | [SH2] | Temporal route bias (init=2.0) |

With SH2=256: 2048 + 256 + 2048 + 256 = **4608 new params**.

### Phase 3 Forward Pass

```
1. trunk forward (cuBLAS)                 → h_s2 [B, SH2]
2. isv_temporal_route()                   → temporal_weight [SH2]
3. mamba2_temporal_scan(temporal_weight)   → h_enriched [B, SH2] (weighted residual)
4. isv_forward()                          → embedding, gate, gamma_mod
5. isv_feature_gate()                     → h_s2_gated [B, SH2] (applied to h_enriched)
6. recursive_confidence_forward()         → predicted_error [B]
7. fill_gamma_buf()                       → gamma_buf [B]
8. risk_budget_forward()                  → risk_budget [B]
9. branch heads (cuBLAS)                  → Q values
10. branch_confidence_routing()           → gated Q
11. apply_risk_budget()                   → scaled Q
12. action selection
```

Steps 2, 5 are new in Phase 3. Step 3 is modified (adds temporal_weight parameter).

### Implementation Order (Phase 3)

15. **ISV feature gating** — write `isv_feature_gate` kernel. Add w_feature_gate [SH2,8] + b_feature_gate [SH2] at indices 78-79. Wire after Mamba2, before branch heads. NUM_WEIGHT_TENSORS: 78→82 (add all 4 at once).
16. **Per-feature temporal routing** — write `isv_temporal_route` kernel. Add w_temporal_route [SH2,8] + b_temporal_route [SH2] at indices 80-81. Modify `mamba2_temporal_scan` to accept + use temporal_weight buffer.
17. **Final verification (Phase 3)** — compute-sanitizer 0 errors, smoke test, full suite.

---

## Phase 4: Adaptive Minimum Hold Time

The model makes 619K trades across 5M bars — many are 1-2 bar coin flips that destroy Sharpe. A hard minimum hold time forces the model to commit to trades and follow trends. Making it adaptive via ISV + regime means the hold time adapts to market conditions.

### PEARL: Hard Minimum Hold (P7)

In `experience_env_step`, after action selection, if `bars_since_entry < min_hold_bars`, override the action to maintain current position. The model still sees the state and computes Q-values for learning, but can't exit until minimum hold expires. Base: 3 bars.

### GEM: Regime-Conditioned Hold Extension (G7)

Extend minimum hold based on regime signals from ISV:
- `adaptive_min_hold = base_hold + hold_extension`
- `hold_extension = floor(regime_stability * max_extension)`
- Trending regime (stability≈1.0): extend to 8-15 bars → ride the trend
- Transition regime (stability≈0.3): stick to base 3 bars → get out fast

### NOVEL: ISV-Driven Hold Confidence (N12)

The recursive confidence head (predicted_error) modulates hold time:
- Low predicted_error (model confident): extend hold by +3 bars → let conviction play out
- High predicted_error (model uncertain): minimum hold only → quick escape
- `hold_extension += floor((1.0 - predicted_error) * confidence_extension)`

This creates a system where the model learns: "when I'm confident AND the regime is trending, hold longer." Sharpe improves because the model captures the full trend instead of exiting after 1 bar.

### Implementation

In `experience_kernels.cu` `experience_env_step`, add after position update:

```cuda
    /* Adaptive minimum hold: prevent coin-flip exits */
    int bars_held = step - entry_step;
    float regime_stab = isv_signals[11];  /* regime_stability from ISV */
    float pred_err = predicted_error[env_id]; /* from recursive confidence */
    int base_hold = 3;
    int regime_ext = (int)(regime_stab * 12.0f);  /* 0-12 bars */
    int conf_ext = (int)((1.0f - pred_err) * 3.0f); /* 0-3 bars */
    int min_hold = base_hold + regime_ext + conf_ext; /* 3-18 bars */
    
    if (bars_held < min_hold && fabsf(position) > 0.001f) {
        /* Override: keep current position, don't allow exit */
        new_position = position;  /* maintain */
    }
```

Note: This only prevents EXITS. The model can still ENTER a new position freely. Once in, it must hold for min_hold bars.

### Expected Impact

With adaptive min hold of 3-18 bars:
- Trade count drops from 619K to ~100-200K (selective, longer trades)
- Average hold increases from 8 to 15-25 bars
- Each trade captures more of the trend → higher mean return per trade
- Sharpe improves by estimated +1-3 from reduced noise alone

---

## Risks and Mitigations

| Risk | Mitigation |
|------|-----------|
| ISV signals are global (same for all batch samples) — no per-sample variation | Per-sample predicted_error from recursive confidence head adds sample-level variation. Risk branch combines global ISV with per-sample h_s2. Branch confidence uses per-sample Q-value separation. |
| ISV encoder is tiny (582 params) — may underfit | 8→16→8 is sufficient for 8 input features. If needed, widen to 8→32→16. |
| One-step lag in ISV signals | Correct behavior — model reacts to recent dynamics, not current-batch. Same as how humans read yesterday's P&L to adjust today's risk. |
| Drift-conditioned gamma changes Bellman target — same gamma for all samples in batch | Single gamma_mod broadcast to all B samples is appropriate: ISV is a global training-dynamics signal, not market-dependent. Per-sample gamma variation comes from the base system (external adaptive_gamma). |
| Risk branch input change (SH2→SH2+9) breaks existing checkpoints | Zero-initialize new columns. Existing SH2 weights load normally. Risk behavior unchanged until ISV trains. |
| `w_risk_fc` resize ripples through risk_aligned_count, risk_grad_buf, backward stride | All must be updated atomically. Implementation plan step 6 handles this as a single unit. |
| `risk_budget_backward` stride change (SH2 → SH2+9) | Explicit kernel update with new stride and ISV+error gradient accumulation (see Backward Pass section). |
| Branch confidence routing × ISV gate product can approach 0 for one branch | Floor at 0.3 for confidence. ISV gate softmax ensures minimum ~0.05 per branch. Product floor: 0.05 × 0.3 = 0.015 — branch never fully muted. |
| Recursive confidence may learn to always predict 0.5 (uninformative) | MSE loss with actual TD-error provides real supervision signal. If TD-errors vary, the head WILL learn to predict them. Early epochs: TD-error is high and variable → strong learning signal. |
| `regime_branch_gate` call sites must all be replaced | enumerate: `apply_regime_gate()`, `reduce_current_q_stats()`, `submit_forward_ops_ddqn()`. Replace with `branch_confidence_routing`. Drop `regime_gate_kernel` field. |
| Target network uses ISV weights | ISV weights are online-only (see Target Network section). Target forward uses neutral defaults (base_gamma, uniform gates). |
| ISV_DIM expansion 8→12 invalidates Phase 1 buffers | Phase 2 tasks reallocate isv_signals [12], isv_history [K*12], isv_decay [12]. w_isv_fc1 grows. Graph recapture required. |
| Regime signals from sample 0 only (not batch mean) | Sample 0 is representative for ADX/CUSUM which are market-level (same for all samples in a time window). For per-sample variation, the regime flows through h_s2 → branches already. |
| Regime-conditioned Mamba2 decay changes temporal context | Bounded by regime_stability ∈ [0, 1]. Stable regime: no change. Transition: faster decay (recent bars weighted more). Gradual — never drops to 0. |
| Risk branch input grows again SH2+9→SH2+13 | Same atomic update pattern as Phase 1. All offset callers updated together. |
| Feature gate collapse (all 1s or all 0s) | Init b_feature_gate=2.0 → sigmoid≈0.88 (near pass-through). Model only deviates when it helps. |
| Temporal routing interference with Mamba2 | Init b_temporal_route=2.0 → ≈0.88 (mostly use temporal). Soft start, model adjusts gradually. |
| Too many ISV-conditioned pathways → single "ISV mode" | Each pathway has own learned projection (different W matrices). They CAN respond differently to same ISV state. |
| Inference cold start (no ISV history) | All gates initialize to near pass-through. Model starts conservative, adapts in 4-8 bars as ISV accumulates. |
| Phase 3 adds 4608 params → overfitting risk | 4608 / ~580K existing = 0.8%. Feature gate and temporal route are regularized by near-pass-through init. |
