# Action Diversity + IQN 4-Branch Fix 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:** Fix IQN 3→4 branch kernel mismatch (corrupting 75% of gradient budget), add order-type counterfactual relabeling, per-branch entropy regularization, and position histogram for all 4 branches — raising order diversity from 1/3 to 3/3.

**Architecture:** Five compounding causes of order branch collapse, fixed bottom-up: (1) IQN kernel uses pre-4-branch 3-branch defines with BRANCH_0_SIZE=5, corrupting 2/3 of gradient signal; (2) no counterfactual experience for order types; (3) no per-branch entropy boost for order; (4) position histogram only tracks dir+mag; (5) Flat-state trains order Q-values on null transitions. All changes are in CUDA kernels + build.rs + Rust config — no CUDA graph structure changes.

**Tech Stack:** CUDA 12.4 (nvcc cubin), Rust 1.85, cudarc

---

## File Map

| File | Action | Responsibility |
|------|--------|---------------|
| `crates/ml/build.rs` | Modify | Add `-D` defines for IQN branch sizes |
| `crates/ml/src/cuda_pipeline/iqn_dual_head_kernel.cu` | Modify | 3→4 branch: defines, offsets, forward, backward, decode |
| `crates/ml/src/cuda_pipeline/gpu_iqn_head.rs` | Modify | Update buffer sizes for 4 branches, fix `total_branch_actions` in save buffers |
| `crates/ml/src/cuda_pipeline/experience_kernels.cu` | Modify | 3-cycle counterfactual + order/urgency histogram |
| `crates/ml/src/cuda_pipeline/c51_grad_kernel.cu` | Modify | Per-branch entropy boost for order (d==2) |
| `crates/ml/src/cuda_pipeline/mse_grad_kernel.cu` | Modify | Per-branch entropy boost for order (d==2) |

---

### Task 1: IQN kernel — runtime branch sizes, eliminate compile-time defines

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/iqn_dual_head_kernel.cu:60-99,188-199`

No build.rs changes needed — branch sizes become runtime kernel parameters like `hidden_dim` and `embed_dim` already are. This prevents the entire class of stale-define bugs (the compile-time BRANCH_0_SIZE=5 was never updated when 4-branch landed).

- [ ] **Step 1: Remove compile-time BRANCH_*_SIZE defines, add runtime parameters**

In `iqn_dual_head_kernel.cu`, delete the compile-time branch defines (lines 62-72) and the `IQN_W_B*_SIZE` / `IQN_TOTAL_PARAMS` macros (lines 83-95). Replace with a comment:

```c
/* Branch sizes are RUNTIME parameters — passed to every kernel.
 * Eliminates stale compile-time define bugs (BRANCH_0_SIZE was 5 when 4-branch uses 3).
 * Weight layout: W_embed[H,D], b_embed[H], then per-branch W_bd[Bd,H], b_bd[Bd]. */
```

- [ ] **Step 2: Update offset function to use runtime branch sizes**

Replace `iqn_compute_offsets` (lines 188-199) to accept runtime branch sizes:

```c
/** Compute runtime weight offsets from hidden_dim, embed_dim, and 4 branch sizes.
 *  Returns offsets[10]: W_EMBED, B_EMBED, W_B0, B_B0, W_B1, B_B1, W_B2, B_B2, W_B3, B_B3 */
__device__ __forceinline__ void iqn_compute_offsets(
    int hidden_dim, int embed_dim,
    int b0, int b1, int b2, int b3,
    int offsets[10]
) {
    offsets[0] = 0;                                          /* W_EMBED */
    offsets[1] = offsets[0] + hidden_dim * embed_dim;        /* B_EMBED */
    offsets[2] = offsets[1] + hidden_dim;                    /* W_B0 */
    offsets[3] = offsets[2] + b0 * hidden_dim;               /* B_B0 */
    offsets[4] = offsets[3] + b0;                            /* W_B1 */
    offsets[5] = offsets[4] + b1 * hidden_dim;               /* B_B1 */
    offsets[6] = offsets[5] + b1;                            /* W_B2 */
    offsets[7] = offsets[6] + b2 * hidden_dim;               /* B_B2 */
    offsets[8] = offsets[7] + b2;                            /* W_B3 */
    offsets[9] = offsets[8] + b3 * hidden_dim;               /* B_B3 */
}
```

- [ ] **Step 3: Add branch size parameters to all 9 kernel signatures**

Every IQN kernel signature gets 4 new `int` parameters: `int b0_size, int b1_size, int b2_size, int b3_size`. Replace all internal uses of `BRANCH_0_SIZE` etc. with these runtime params. Also compute `int tba = b0_size + b1_size + b2_size + b3_size;` at the top of each kernel (replaces `TOTAL_BRANCH_ACTIONS`).

For `iqn_forward_loss_kernel`, `iqn_backward_kernel`, `iqn_forward_kernel`:
```c
/* Add after existing params (hidden_dim, embed_dim): */
    int b0_size, int b1_size, int b2_size, int b3_size
```

For `iqn_decode_actions_kernel`, `iqn_grad_norm_kernel`, `iqn_adam_kernel`, `iqn_ema_kernel`, `iqn_trunk_forward_kernel`, `iqn_sample_taus_kernel`: add the same 4 params even if unused (consistent interface, like `shared_h1` is already passed to some kernels unused).

- [ ] **Step 4: Update Rust kernel launch calls in `gpu_iqn_head.rs`**

Every `.launch_builder()` call for IQN kernels needs 4 additional `.arg()` calls:
```rust
.arg(&(self.config.branch_0_size as i32))
.arg(&(self.config.branch_1_size as i32))
.arg(&(self.config.branch_2_size as i32))
.arg(&(self.config.branch_3_size as i32))
```

- [ ] **Step 5: Verify kernel compiles and all tests pass**

```bash
SQLX_OFFLINE=true cargo build -p ml --lib 2>&1 | grep -E "iqn_dual_head|error|FAILED"
SQLX_OFFLINE=true cargo check -p ml --lib
```
Expected: `Compiled iqn_dual_head_kernel.cu -> iqn_dual_head_kernel.cubin` with no errors.

- [ ] **Step 6: Commit**

```bash
git add crates/ml/src/cuda_pipeline/iqn_dual_head_kernel.cu crates/ml/src/cuda_pipeline/gpu_iqn_head.rs
git commit -m "fix(iqn): runtime branch sizes — eliminate stale compile-time defines"
```

---

### Task 2: IQN kernel — 4-branch forward, backward, and decode

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/iqn_dual_head_kernel.cu:215-475,488-664,766-988`

- [ ] **Step 1: Update decode kernel to 4-branch**

Replace `iqn_decode_actions_kernel` (lines 962-988). Output becomes `[B, 4]` (was `[B, 3]`):

```c
/* Output: branch_actions[B*4] -- i32 [dir, order, urgency, magnitude] per sample */
extern "C" __global__
void iqn_decode_actions_kernel(
    const int* __restrict__ flat_actions,
    int* __restrict__ branch_actions,
    int batch_size,
    int shared_h1,
    int hidden_dim,
    int embed_dim
)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= batch_size) return;

    int flat = flat_actions[i];
    int max_action = BRANCH_0_SIZE * BRANCH_1_SIZE * BRANCH_2_SIZE * BRANCH_3_SIZE;
    if (flat < 0 || flat >= max_action) flat = 0;

    /* 4-branch decode: action = dir*b1*b2*b3 + mag*b2*b3 + order*b3 + urgency */
    int stride_0 = BRANCH_1_SIZE * BRANCH_2_SIZE * BRANCH_3_SIZE;
    int stride_1 = BRANCH_2_SIZE * BRANCH_3_SIZE;
    int stride_2 = BRANCH_3_SIZE;

    int dir     = flat / stride_0;
    int rem     = flat % stride_0;
    int mag     = rem / stride_1;
    rem         = rem % stride_1;
    int order   = rem / stride_2;
    int urgency = rem % stride_2;

    branch_actions[i * 4 + 0] = dir;
    branch_actions[i * 4 + 1] = mag;
    branch_actions[i * 4 + 2] = order;
    branch_actions[i * 4 + 3] = urgency;
}
```

- [ ] **Step 2: Update forward+loss kernel for 4 branches**

In `iqn_forward_loss_kernel`, add branch 3 to the online forward (after branch 2 block at line ~380), target forward (after line ~437), and sum of taken Q-values. The pattern is identical to branches 0-2 — just add:

```c
/* Branch 3 (magnitude) — online */
float q3_taken = 0.0f;
for (int a = 0; a < BRANCH_3_SIZE; a++) {
    float partial = 0.0f;
    const float* w_row = w_b3 + a * hidden_dim;
    for (int h = tid; h < hidden_dim; h += IQN_BLOCK_SIZE)
        partial += w_row[h] * comb_dist[h / IQN_BLOCK_SIZE];
    float acc = b_b3[a] + iqn_block_sum(partial, shmem_reduce);
    if (tid == 0) {
        save_q_online[q_save_offset + BRANCH_0_SIZE + BRANCH_1_SIZE + BRANCH_2_SIZE + a] = acc;
        if (a == a3) q3_taken = acc;
    }
}
if (tid == 0) shmem_reduce[0] = q3_taken;
__syncthreads();
q3_taken = shmem_reduce[0];
```

Update the sum line: `online_q_a[t] = q0_taken + q1_taken + q2_taken + q3_taken;`

Similarly for the target branch and `target_q_a` save.

Also update action decode at the top of the kernel:
```c
int a0 = actions[sample * 4 + 0]; // was sample * 3
int a1 = actions[sample * 4 + 1];
int a2 = actions[sample * 4 + 2];
int a3 = actions[sample * 4 + 3];
```

And add weight pointers:
```c
const float* w_b3 = online_params + off[8];
const float* b_b3 = online_params + off[9];
const float* tw_b3 = target_params + off[8];
const float* tb_b3 = target_params + off[9];
```

- [ ] **Step 3: Update backward kernel for 4 branches**

In `iqn_backward_kernel`, add branch 3 gradient computation after the branch 2 block (after line ~624). Pattern identical:

```c
/* Branch 3 (magnitude) */
for (int h = tid; h < hidden_dim; h += IQN_BLOCK_SIZE) {
    dL_dcomb[h / IQN_BLOCK_SIZE] += w_b3[a3 * hidden_dim + h] * dL_dq;
    atomicAdd(&grad_buf[off[8] + a3 * hidden_dim + h],
              inv_batch * dL_dq * comb_dist[h / IQN_BLOCK_SIZE]);
}
if (tid == 0)
    atomicAdd(&grad_buf[off[9] + a3], inv_batch * dL_dq);
```

Update action reads: `int a3 = actions[sample * 4 + 3];` and add `const float* w_b3 = online_params + off[8];`

Update offset array: `int off[10]; iqn_compute_offsets(hidden_dim, embed_dim, off);`

- [ ] **Step 4: Update forward-only (inference) kernel for 4 branches**

The `iqn_forward_kernel` (line 766) also needs branch 3. Same pattern as forward+loss but without loss computation. Add branch 3 Q-value computation and update `q_total += q3`.

- [ ] **Step 5: Verify kernel compiles**

```bash
SQLX_OFFLINE=true cargo build -p ml --lib 2>&1 | grep -E "iqn_dual_head|error"
```

- [ ] **Step 6: Commit**

```bash
git add crates/ml/src/cuda_pipeline/iqn_dual_head_kernel.cu
git commit -m "fix(iqn): 4-branch forward, backward, decode — was 3-branch corrupting 75% gradient"
```

---

### Task 3: IQN Rust side — update buffer sizes and decode for 4 branches

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_iqn_head.rs:113-127,233-293,339-345`

- [ ] **Step 1: Update `branch_actions` allocation from `B*3` to `B*4`**

In `GpuIqnHead::new()`, change:
```rust
// Was: alloc_zeros::<i32>(b * 3)
let branch_actions = stream.alloc_zeros::<i32>(b * 4).map_err(|e| {
    MLError::ModelError(format!("IQN alloc branch_actions: {e}"))
})?;
```

- [ ] **Step 2: Update `save_q_online` and `save_q_target` buffer sizes**

These use `tba = config.total_branch_actions()` which already returns `b0+b1+b2+b3=12`. Verify `tba` is used (not hardcoded). Check lines 292-293:
```rust
let save_q_online = alloc_f32(&stream, b * n * tba, "iqn_save_q_online")?;
let save_q_target = alloc_f32(&stream, b * n * tba, "iqn_save_q_target")?;
```
These should already be correct since `total_branch_actions()` includes b3.

- [ ] **Step 3: Update IQR buffer size**

Line 306: `let iqr_buf = alloc_f32(&stream, tba, "iqn_iqr")?;` — already uses `tba=12`. OK.

- [ ] **Step 4: Verify all IQN-related tests pass**

```bash
SQLX_OFFLINE=true cargo check -p ml --lib
SQLX_OFFLINE=true cargo test -p ml --lib -- gradient_budget --nocapture
```

- [ ] **Step 5: Commit**

```bash
git add crates/ml/src/cuda_pipeline/gpu_iqn_head.rs
git commit -m "fix(iqn): update Rust buffer sizes for 4-branch decode (B*4 actions)"
```

---

### Task 4: Order-type counterfactual relabeling (3-cycle)

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

- [ ] **Step 1: Replace 2-cycle counterfactual with 3-cycle**

Replace the `if ((current_t & 1) == 0)` block (lines 1799-1846) with a 3-cycle that adds order-type counterfactuals:

```c
int cf_cycle = current_t % 3;

if (cf_cycle == 0) {
    /* ── Directional mirror (existing) ── */
    int cf_dir = (b0_size - 1) - dir_idx;
    cf_action = cf_dir * b1_size * b2_size * b3_size
              + mag_idx * b2_size * b3_size
              + orig_order * b3_size + orig_urgency;
    cf_reward = -reward;
} else if (cf_cycle == 1) {
    /* ── Hindsight magnitude relabeling (existing) ── */
    const float mag_fracs[3] = { 0.50f, 0.707f, 1.00f };
    float actual_mag_frac = mag_fracs[mag_idx < 3 ? mag_idx : 0];

    if (fabsf(reward) > 0.001f) {
        int alt_mag;
        int cycle = (current_t / 3) % 3;
        if (cycle == 0)      alt_mag = (mag_idx > 0) ? mag_idx - 1 : 2;
        else if (cycle == 1) alt_mag = (mag_idx < 2) ? mag_idx + 1 : 0;
        else                 alt_mag = 2 - mag_idx;

        float alt_mag_frac = mag_fracs[alt_mag];
        cf_action = dir_idx * b1_size * b2_size * b3_size
                  + alt_mag * b2_size * b3_size
                  + orig_order * b3_size + orig_urgency;
        cf_reward = reward * (alt_mag_frac / actual_mag_frac);
        cf_reward = fminf(fmaxf(cf_reward, -10.0f), 10.0f);
    } else {
        int cf_dir = (b0_size - 1) - dir_idx;
        cf_action = cf_dir * b1_size * b2_size * b3_size
                  + mag_idx * b2_size * b3_size
                  + orig_order * b3_size + orig_urgency;
        cf_reward = -reward;
    }
} else {
    /* ── NEW: Order-type counterfactual ──
     * Cycle through alternative order types. Compute execution cost
     * difference and amplify it to be comparable to directional signal.
     * Only meaningful when position is non-zero (order type affects cost). */
    int alt_order = (orig_order + 1 + ((current_t / 3) & 1)) % b2_size;

    if (fabsf(position) > 0.001f) {
        /* Compute cost delta between taken and alternative order type */
        float taken_cost = compute_tx_cost(
            position - pre_trade_position, raw_close, tx_cost_multiplier,
            spread_cost, max_position, orig_order, spread_scale);
        float alt_cost = compute_tx_cost(
            position - pre_trade_position, raw_close, tx_cost_multiplier,
            spread_cost, max_position, alt_order, spread_scale);
        float cost_delta = (taken_cost - alt_cost)
                         / fmaxf(prev_equity, 1.0f);

        cf_action = dir_idx * b1_size * b2_size * b3_size
                  + mag_idx * b2_size * b3_size
                  + alt_order * b3_size + orig_urgency;
        /* Amplify 50×: raw cost delta ~0.0001-0.0005, directional ~0.01-0.1 */
        cf_reward = reward + 50.0f * cost_delta;
        cf_reward = fminf(fmaxf(cf_reward, -10.0f), 10.0f);
    } else {
        /* Flat position: order type irrelevant, fall back to directional mirror */
        int cf_dir = (b0_size - 1) - dir_idx;
        cf_action = cf_dir * b1_size * b2_size * b3_size
                  + mag_idx * b2_size * b3_size
                  + orig_order * b3_size + orig_urgency;
        cf_reward = -reward;
    }
}
```

Note: `position`, `pre_trade_position`, `raw_close`, `tx_cost_multiplier`, `spread_cost`, `max_position`, `spread_scale` are all already in scope from the experience_env_step kernel's local variables. `compute_tx_cost` is declared in `trade_physics.cuh` which is already included.

- [ ] **Step 2: Verify kernel compiles**

```bash
SQLX_OFFLINE=true cargo build -p ml --lib 2>&1 | grep -E "experience_kernels|error"
```

- [ ] **Step 3: Commit**

```bash
git add crates/ml/src/cuda_pipeline/experience_kernels.cu
git commit -m "feat: order-type counterfactual relabeling (3-cycle: dir/mag/order)"
```

---

### Task 5: Per-branch entropy regularization for order

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/c51_grad_kernel.cu:53-57`
- Modify: `crates/ml/src/cuda_pipeline/mse_grad_kernel.cu:53-58`

- [ ] **Step 1: Add order entropy boost to C51 grad kernel**

Replace lines 53-57 in `c51_grad_kernel.cu`:

```c
/* Entropy regularization — per-branch coefficients.
 * Magnitude (d==1): 5× boost for atom distribution collapse prevention.
 * Order (d==2): 3× boost to prevent order type collapse.
 * Direction (d==0), urgency (d==3): baseline 1×. */
if (entropy_coeff > 0.0f) {
    float ent_scale;
    if (d == 1) ent_scale = 5.0f;
    else if (d == 2) ent_scale = 3.0f;
    else ent_scale = 1.0f;
    float lp_clamped = fmaxf(lp, -10.0f);
    d_combined += ent_scale * entropy_coeff * (1.0f + lp_clamped);
}
```

- [ ] **Step 2: Add order entropy boost to MSE grad kernel**

Replace lines 53-58 in `mse_grad_kernel.cu`:

```c
/* Per-branch entropy boost: prevent distribution collapse.
 * Magnitude (d==1): 0.005 = 5× base (0.001). Order (d==2): 0.003 = 3× base. */
if (d == 1 || d == 2) {
    float ent_weight = (d == 1) ? 0.005f : 0.003f;
    float lp_approx = fmaxf(logf(fmaxf(p_j, 1e-8f)), -10.0f);
    d_combined += ent_weight * (1.0f + lp_approx);
}
```

- [ ] **Step 3: Verify kernel compiles**

```bash
SQLX_OFFLINE=true cargo build -p ml --lib 2>&1 | grep -E "c51_grad\|mse_grad|error"
```

- [ ] **Step 4: Commit**

```bash
git add crates/ml/src/cuda_pipeline/c51_grad_kernel.cu crates/ml/src/cuda_pipeline/mse_grad_kernel.cu
git commit -m "feat: per-branch entropy regularization — order gets 3× boost"
```

---

### Task 6: Position histogram for all 4 branches

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

- [ ] **Step 1: Expand histogram from 6 bins to 12 bins (3×4 branches)**

Replace the histogram tracking block (lines 1664-1707):

```c
/* ── #19 Position entropy: track all 4 branches (3 bins each = 12 total) ── */
if (position_histogram != NULL && position_entropy_weight > 0.0f) {
    /* Histogram layout: [N, 12] — 3 bins per branch: dir[0:3], mag[3:6], order[6:9], urgency[9:12] */
    int order_idx = decode_order_4b(action_idx, b2_size, b3_size);
    int urgency_idx = decode_urgency_4b(action_idx, b3_size);
    if (dir_idx >= 0 && dir_idx < 3)
        position_histogram[(long long)i * 12 + dir_idx] += 1.0f;
    if (mag_idx >= 0 && mag_idx < 3)
        position_histogram[(long long)i * 12 + 3 + mag_idx] += 1.0f;
    if (order_idx >= 0 && order_idx < 3)
        position_histogram[(long long)i * 12 + 6 + order_idx] += 1.0f;
    if (urgency_idx >= 0 && urgency_idx < 3)
        position_histogram[(long long)i * 12 + 9 + urgency_idx] += 1.0f;
}
```

- [ ] **Step 2: Update entropy bonus computation to cover all 4 branches**

Replace the done-block entropy bonus (lines 1679-1707):

```c
/* #19 Position entropy bonus at episode end — all 4 branches */
if (done && position_histogram != NULL && position_entropy_weight > 0.0f) {
    float* hist = position_histogram + (long long)i * 12;
    float combined_entropy = 0.0f;
    /* Compute entropy over each branch's 3 bins */
    for (int branch = 0; branch < 4; branch++) {
        float* branch_hist = hist + branch * 3;
        float total = 0.0f;
        for (int b = 0; b < 3; b++) total += branch_hist[b];
        float ent = 0.0f;
        if (total > 1.0f) {
            for (int b = 0; b < 3; b++) {
                float p = branch_hist[b] / total;
                if (p > 1e-6f) ent -= p * logf(p);
            }
            ent /= 1.0986f; /* Normalize by log(3) → [0, 1] */
        }
        combined_entropy += ent;
    }
    /* Average across all 4 branches (was 2) */
    combined_entropy *= 0.25f;
    reward += position_entropy_weight * combined_entropy;
}
```

- [ ] **Step 3: Update position_histogram allocation in Rust**

Search for `position_histogram` allocation in `gpu_experience_collector.rs` and change from `alloc_episodes * 6` to `alloc_episodes * 12`:

```rust
// Was: alloc_episodes * 6
let position_histogram = stream.alloc_zeros::<f32>(alloc_episodes * 12)
    .map_err(|e| MLError::ModelError(format!("alloc position_histogram: {e}")))?;
```

- [ ] **Step 4: Verify compilation and run tests**

```bash
SQLX_OFFLINE=true cargo check -p ml --lib
SQLX_OFFLINE=true cargo test -p ml-dqn --lib
SQLX_OFFLINE=true cargo test -p ml --lib -- gradient_budget config monitoring
```

- [ ] **Step 5: Commit**

```bash
git add crates/ml/src/cuda_pipeline/experience_kernels.cu crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
git commit -m "feat: position histogram for all 4 branches — order+urgency get entropy bonus"
```

---

### Task 7: Full integration test — smoke test + compute-sanitizer

**Files:** (no changes — verification only)

- [ ] **Step 1: Run full smoke test**

```bash
FOXHUNT_TEST_DATA=test_data/futures-baseline SQLX_OFFLINE=true cargo test -p ml --lib -- smoke_tests::feature_coverage --ignored --test-threads=1 --nocapture 2>&1 | grep -E "Action diversity|order|entropy|LOW ORDER"
```

Expected: `Action diversity` should show order > 1/3. `LOW ORDER DIVERSITY` warnings should disappear or reduce.

- [ ] **Step 2: Run compute-sanitizer**

```bash
SQLX_OFFLINE=true cargo test -p ml --lib --no-run 2>&1 | tail -1
FOXHUNT_TEST_DATA=test_data/futures-baseline compute-sanitizer --tool memcheck target/debug/deps/ml-*.test "smoke_tests::feature_coverage" --ignored --test-threads=1
```

Expected: `ERROR SUMMARY: 0 errors`

- [ ] **Step 3: Run full test suite**

```bash
SQLX_OFFLINE=true cargo test -p ml-dqn --lib
SQLX_OFFLINE=true cargo test -p ml --lib
FOXHUNT_TEST_DATA=test_data/futures-baseline SQLX_OFFLINE=true cargo test -p ml --lib -- smoke_tests --ignored --test-threads=1
```

Expected: All pass (286 + 903 + 19).

- [ ] **Step 4: Final commit with all changes**

```bash
git add -A
git commit -m "fix: IQN 4-branch + order counterfactual + per-branch entropy — action diversity restored"
```
