fix(dqn): Flat opportunity cost breaks val Flat-equilibrium

Val-Flat-collapse fix #3 (task #94, 2026-04-24). Research agent
traced the dominant val failure mode: reward = 0.0 for Flat bars
produces structural Flat-beats-trading-actions in the argmax —
Flat accumulates Q ≈ 0 while CQL pulls trading-action Q's toward
pessimistic negative estimates. Result: val argmax picks Flat for
~all states (22 trades/epoch vs 30K+ in training via Boltzmann).

Prior comment defended the design: "No signal is correct — the
model should not be rewarded or penalized for correctly staying
flat when there's no edge." That reasoning was correct in isolation
but broken in context: the trading-action branch is penalized by
CQL's lower-bound estimator anyway, so Flat's neutral Q becomes the
de-facto upper bound on all action Q-values → Flat wins argmax
regardless of the policy's actual directional belief.

Add a per-bar opportunity cost for Flat bars, proportional to
realised volatility (ATR fraction):

  reward = -shaping_scale * holding_cost_rate * 0.5 * vol_proxy

This:
- Symmetric to the positioned-bar holding penalty (positioned
  pays `holding × |pos|`; Flat pays `holding × 0.5 × vol_proxy`).
- Vol-scaled: Flat during quiet markets approaches zero cost;
  Flat during volatile markets has a small negative signal
  representing missed-move expectation.
- Caps vol_proxy at 0.01 (1%) as numerical safety against
  single-bar vol spikes dominating reward.
- Recomputes the ATR locally (the segment-complete branch that
  normally does this is mutually exclusive with this Flat branch).
- Guarded on `features != NULL` and `bar_idx` range for smoke-test
  compatibility.

The `0.5` factor is a numerical-safety symmetric-to-holding-cost
bound per the carve-out in `feedback_isv_for_adaptive_bounds.md`,
not a tuned hyperparameter. The `holding_cost_rate` itself flows
from config (already adaptive via shaping_scale annealing).

Not part of the triad yet:
- plan_isv train/val parity (task #94 item #2): documented
  architectural gap requiring val-time plan-state generator; no
  trivial safe fix.
- Atom utilization (task #94 item #4): deep C51 representation
  work (atom-level entropy regularisation or target-variance
  engineering); deferred pending tie-break + opp-cost validation.

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

crates/ml/src/cuda_pipeline/experience_kernels.cu

@@ -2211,10 +2211,43 @@ extern "C" __global__ void experience_env_step(
     } else if (!segment_complete && fabsf(position) > 0.001f) {
         /* Fallback: original holding cost when micro_reward_scale=0 or no OFI */
         reward = -shaping_scale * holding_cost_rate * fabsf(position);
+    } else if (!segment_complete) {
+        /* ── Flat opportunity cost (val-Flat-collapse fix #3, 2026-04-24) ──
+         * Prior code kept reward = 0.0f for Flat bars with the note "no
+         * signal is correct". Empirically this produced the val-Flat-
+         * collapse: Flat Q accumulated to ~0 while trading actions were
+         * pulled toward negative Q by CQL's pessimistic lower-bound, so
+         * deterministic val argmax picked Flat regardless of state —
+         * producing 22 val trades per epoch vs 30K+ in training via
+         * Boltzmann. See `c51_loss_kernel.cu` CQL block.
+         *
+         * Flat is NOT free during volatile market regimes. A per-bar
+         * opportunity cost proportional to realised volatility (ATR)
+         * breaks the Flat-equilibrium: Flat during quiet markets
+         * approaches zero cost (vol tiny → cost tiny), Flat during
+         * volatile markets has a small negative signal representing the
+         * missed-move expectation. This pulls Flat Q below CQL-pessimistic
+         * trading Q only when genuinely idle market would have paid
+         * nothing anyway.
+         *
+         * Cost scale = holding_cost_rate × 0.5: symmetric to the holding
+         * penalty for positioned bars (positioned pays `holding × |pos|`,
+         * flat pays `holding × 0.5 × vol_proxy`). Volatility signal is
+         * the ATR fraction we already compute for segment-complete rewards.
+         * Recomputed locally here since the segment-complete branch is
+         * mutually exclusive with this one. Guarded on `features != NULL`
+         * and bar_idx in range so graph-safe & smoke-test compatible. */
+        if (features != NULL && bar_idx < total_bars && market_dim > 9) {
+            float atr_norm_flat = features[(long long)bar_idx * market_dim + 9];
+            float log_atr_flat = atr_norm_flat * 16.0f - 7.0f;
+            float atr_pct_flat = expf(log_atr_flat) / fmaxf(raw_close, 1.0f);
+            float vol_proxy_flat = fmaxf(atr_pct_flat, 0.0001f);
+            /* Cap vol_proxy at 0.01 (1%) so one extreme-vol bar cannot
+             * dominate reward — numerical safety, not tuning. */
+            if (vol_proxy_flat > 0.01f) vol_proxy_flat = 0.01f;
+            reward = -shaping_scale * holding_cost_rate * 0.5f * vol_proxy_flat;
+        }
     }
-    /* Flat (no position, no trade): reward stays 0.0f from initialization.
-     * No signal is correct — the model should not be rewarded or penalized
-     * for correctly staying flat when there's no edge. */
 
     /* ---- Drawdown penalty (from trade_physics.cuh) ──────────────────────
      * Phase 3: gated by shaping_scale. The drawdown penalty is *behavioral*