# Reward v7.1: Six Critical Fixes Design

**Goal:** Fix 6 concerns identified in the initial v7 hyperopt run: i%3 Q-value degeneracy (exposure auxiliary loss), win rate below random (OFI epoch gate), CEA warmup too short, exit timing prev_sign bug, Kelly threshold with prior seeding, and reward noise magnitude scaling.

**Architecture:** Exposure auxiliary loss follows the existing IQN/CQL auxiliary gradient pattern — new lightweight CUDA kernel launched between C51 backward and Adam. Other 5 fixes are surgical changes to the experience kernel and Rust training loop config gating.

**Tech Stack:** CUDA (nvcc → cubin), Rust (cudarc), TOML config

---

## Fix 1: Exposure Auxiliary Loss with Scheduled Decay

### Problem
C51 backward gives identical `-1/9 × dL` gradient to all 8 non-selected exposure outputs. This causes them to converge into degenerate i%3 clusters. The pattern recurs even after CEA warmup breaks it temporarily.

### Root Cause
The dueling advantage mean subtraction `A'[i] = A[i] - mean(A)` creates gradient `dL/dA[selected] = 8/9 × dL` and `dL/dA[other] = -1/9 × dL`. With b0_size=9, the non-selected gradient is only 12.5% of the selected — too weak to break weight clusters from random initialization.

### Solution: Three Components

**1a. Experience kernel output** — In the CEA loop (already computing rewards for all 9 exposures), track `best_exposure_idx` and write to a new output buffer `out_best_exposure[N*L]`:

```cuda
int best_k = 0;
float best_alt = -1e30f;
for (int k = 0; k < b0_size; k++) {
    // ...existing CEA loop body...
    if (alt_reward > best_alt) { best_alt = alt_reward; best_k = k; }
}
if (segment_complete && segment_hold_time > 0.0f) {
    out_best_exposure[out_off] = best_k;
} else {
    out_best_exposure[out_off] = -1;  // no target (not a trade exit)
}
```

New kernel parameter: `int* out_best_exposure` added to `experience_env_step` signature.

**1b. New CUDA kernel** — `exposure_aux_grad_kernel` in `dqn_utility_kernels.cu`:

```cuda
extern "C" __global__ void exposure_aux_grad_kernel(
    const float* __restrict__ adv_logits,  // [batch, b0_size * num_atoms] exposure advantage logits
    const int*   __restrict__ targets,     // [batch] best exposure idx (-1 = skip)
    float*       __restrict__ grad_buf,    // [total_params] accumulate gradient
    int          adv_offset,               // offset into grad_buf for exposure adv weights
    int          b0_size,
    int          num_atoms,
    float        aux_weight,
    int          batch_size
) {
    int sample = blockIdx.x * blockDim.x + threadIdx.x;
    if (sample >= batch_size) return;

    int target = targets[sample];
    if (target < 0 || target >= b0_size) return;  // skip non-exit steps

    // Cross-entropy gradient for exposure branch:
    // For each atom z: grad[a][z] = softmax(logits[a][z]) - 1{a == target}
    const float* logits = adv_logits + (long long)sample * b0_size * num_atoms;

    for (int z = 0; z < num_atoms; z++) {
        // Compute softmax over b0_size for this atom
        float max_l = -1e30f;
        for (int a = 0; a < b0_size; a++) {
            float l = logits[a * num_atoms + z];
            if (l > max_l) max_l = l;
        }
        float sum_exp = 0.0f;
        for (int a = 0; a < b0_size; a++) {
            sum_exp += expf(logits[a * num_atoms + z] - max_l);
        }

        // Write gradients
        for (int a = 0; a < b0_size; a++) {
            float prob = expf(logits[a * num_atoms + z] - max_l) / sum_exp;
            float grad = prob - ((a == target) ? 1.0f : 0.0f);
            // Accumulate into grad_buf at the exposure advantage weight offset
            int param_idx = adv_offset + a * num_atoms + z;
            atomicAdd(&grad_buf[param_idx], aux_weight * grad / (float)batch_size);
        }
    }
}
```

**1c. Fused training pipeline** — Launch between C51 backward and Adam:
- Store `best_exposure` buffer in GPU replay buffer alongside actions/rewards
- In `submit_aux_ops()` or as a standalone post-backward step, launch `exposure_aux_grad_kernel`
- Weight schedule in training_loop.rs:
  ```rust
  let warmup = self.hyperparams.exposure_aux_warmup_epochs.max(1) as f32;
  let progress = (self.current_epoch as f32 / warmup).min(1.0);
  let aux_w = self.hyperparams.exposure_aux_weight as f32 * (1.0 - 0.9 * progress);
  // Decays from base (0.5) to 10% of base (0.05) over warmup_epochs
  ```

### Config
- `exposure_aux_weight: f64` — default 0.5, hyperopt [0.1, 1.0]
- `exposure_aux_warmup_epochs: usize` — default 5, hyperopt [3, 10]

---

## Fix 2: OFI Confidence Epoch Gate

### Problem
OFI multiplier amplifies wrong-direction trades in early training, reducing win rate below random (24% vs 30%).

### Solution
Gate OFI weight by epoch in training_loop.rs:

```rust
ofi_reward_weight: {
    let base = self.hyperparams.ofi_reward_weight as f32;
    if self.current_epoch < 5 { 0.0 } else { base }
},
```

No kernel changes. Model learns direction first (epochs 0-4), then OFI amplifies informed trades (epoch 5+).

---

## Fix 3: Extended CEA Warmup

### Problem
3-epoch hard cutoff warmup breaks symmetry but model reverts when CEA drops from 1.0 to 0.3 at epoch 4.

### Solution
Linear blend over 25% of training:

```rust
cea_weight: {
    let base = self.hyperparams.cea_weight as f32;
    let warmup_end = (self.hyperparams.epochs as f32 * 0.25).max(3.0);
    if (self.current_epoch as f32) < warmup_end {
        let progress = self.current_epoch as f32 / warmup_end;
        1.0 * (1.0 - progress) + base * progress
    } else {
        base
    }
},
```

For 20 epochs: warmup spans epochs 0-4, linearly blending 1.0→base. Smooth transition, no cliff.

---

## Fix 4: Exit Timing `prev_sign` Bug

### Problem
After a reversal (S100→L50), `prev_sign` holds the pre-trade direction. Exit timing quality uses `prev_sign` to determine which direction was closed, but for reversals, it references the OLD position, not the closing segment.

### Solution
`prev_sign` is computed at the top of the kernel from the pre-trade position — it IS the closing segment's direction. The variable name is confusing but the VALUE is correct for exits. For reversals, `prev_sign` is the direction being closed (the old position), which IS what we want.

However, the current code runs the timing block AFTER the reversal block has already updated `position` and `entry_price`. The `prev_sign` variable itself is still correct (computed from `pre_trade_position` earlier). But for clarity, save it explicitly:

```cuda
int closing_sign = prev_sign;  // capture before any trade logic

// ... later, in Layer 7 ...
if (exit_timing_weight > 0.0f && (exiting_trade || reversing_trade)) {
    float post_exit_move = (closing_sign > 0)
        ? (raw_next - raw_close) / fmaxf(raw_close, 1.0f)
        : (raw_close - raw_next) / fmaxf(raw_close, 1.0f);
    float timing_quality = -post_exit_move / vol_proxy;
    reward += exit_timing_weight * asymmetric_soft_clamp(timing_quality * 10.0f) * 0.1f;
}
```

This makes the intent explicit and guards against future refactors that might move `prev_sign` computation.

---

## Fix 5: Kelly Threshold with Bayesian Prior

### Problem
Kelly signal requires 5+ completed trades before computing. Early in training, episodes rarely accumulate enough trades for this threshold.

### Solution
Seed with a Bayesian prior (Jeffrey's prior: 2 pseudo-wins + 2 pseudo-losses with 1% average return each). This makes Kelly available from the first real trade while regularizing toward "no bet" with limited data:

```cuda
if (kelly_sizing_weight > 0.0f) {
    float prior_wins = 2.0f;
    float prior_losses = 2.0f;
    float prior_sum_wins = 0.01f;
    float prior_sum_losses = 0.01f;

    float eff_wins = win_count + prior_wins;
    float eff_losses = loss_count + prior_losses;
    float eff_total = eff_wins + eff_losses;

    // Always fires (eff_total >= 4.0 from priors)
    float win_rate_k = eff_wins / eff_total;
    float avg_win = (sum_wins + prior_sum_wins) / eff_wins;
    float avg_loss = (sum_losses + prior_sum_losses) / eff_losses;
    float payoff_ratio = avg_win / fmaxf(avg_loss, 0.0001f);
    float kelly_f = (payoff_ratio * win_rate_k - (1.0f - win_rate_k))
                   / fmaxf(payoff_ratio, 0.0001f);
    kelly_f = fmaxf(fminf(kelly_f, 1.0f), 0.0f);

    float actual_fraction = fabsf(position) / fmaxf(max_position, 1.0f);
    float kelly_closeness = 1.0f - fabsf(actual_fraction - kelly_f);
    kelly_closeness = fmaxf(kelly_closeness, 0.0f);

    if (segment_return > 0.0f) {
        reward += kelly_sizing_weight * kelly_closeness * 0.5f;
    }
}
```

With prior only: Kelly `f* = (1.0 × 0.5 - 0.5) / 1.0 = 0.0` (recommends no bet). As real trades accumulate, converges to empirical Kelly.

---

## Fix 6: Reward Noise Magnitude Scaling

### Problem
Fixed noise scale (±0.05) is 10-50% of small reward signals, drowning meaningful differences.

### Solution
Scale noise to a fixed PERCENTAGE of reward magnitude with a floor:

```cuda
if (reward_noise_scale > 0.0f && reward != 0.0f) {
    unsigned int hash = (unsigned int)(i * 31337 + bar_idx * 7919 + current_t * 1013);
    hash ^= (hash >> 16); hash *= 0x45d9f3bu; hash ^= (hash >> 16);
    float pseudo_noise = ((float)(hash & 0xFFFF) / 65535.0f - 0.5f) * 2.0f;
    float noise_magnitude = fmaxf(fabsf(reward) * reward_noise_scale, 0.01f);
    reward += noise_magnitude * pseudo_noise;
}
```

`reward_noise_scale` becomes a percentage (0.05 = 5% of reward magnitude). A +5.0 reward gets ±0.25 noise. A +0.1 reward gets ±0.01 noise (floor). Constant SNR across magnitudes.

---

## Files Modified

| File | Changes |
|------|---------|
| `crates/ml/src/cuda_pipeline/experience_kernels.cu` | Fixes 1a, 4, 5, 6 + new `out_best_exposure` arg |
| `crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu` | Fix 1b: new `exposure_aux_grad_kernel` |
| `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` | Fix 1c: load + launch aux kernel |
| `crates/ml/src/trainers/dqn/fused_training.rs` | Fix 1c: wire aux kernel in training step |
| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | Fixes 2, 3: OFI gate + CEA warmup schedule |
| `crates/ml/src/trainers/dqn/config.rs` | New config fields |
| `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` | New `out_best_exposure` buffer + arg wiring |
| `config/training/dqn-*.toml` (all 4) | New params |

## Testing

- `cargo check --workspace` — build verification
- `cargo test -p ml --lib` — all 891+ tests pass
- Smoke test: verify exposure Q-values do NOT show i%3 pattern by epoch 3
- Smoke test: verify win rate ≥ 30% by epoch 5 (above random baseline)
- Smoke test: verify Kelly signal fires from epoch 1 (prior seeding)
