fix(dqn): Flat opportunity cost scales with ISV-driven conviction, not tuned constant

Val-Flat-collapse fix #3 revision (task #94, 2026-04-24). Prior
commit 543e3c11b used `-0.5 × holding_cost_rate × vol_proxy` — the
0.5 is a hardcoded tuned constant violating
`feedback_isv_for_adaptive_bounds.md` and
`feedback_adaptive_not_tuned.md`.

Replace with ISV-driven per-sample conviction, already computed by
the action-select kernel and threaded into env_step via
`conviction_ptr → conviction_core`:

    reward_flat = -shaping_scale * holding_cost_rate
                * conviction_core * vol_proxy_flat

`conviction_core ∈ [0, 1]` = direction-branch Q-range normalised by
`ISV[21]` (q_dir_abs_ref EMA). Self-adapting properties:

  - Cold start / ISV[21] uninitialised → fallback 1.0 → full penalty,
    encourages early exploration out of the flat equilibrium.
  - Low conviction (uncertain direction) → penalty scales toward 0
    → doing nothing is acceptable when there's no signal (matches
    real-world: flat cost is only real when there's opportunity).
  - High conviction (strong directional edge) → penalty scales up
    → Flat becomes expensive ONLY where the model itself says
    there's an edge. Forces the policy to take action exactly
    where it has belief, not blindly.

Continuity: conviction is continuous ∈ [0, 1], no step function. The
temporal Mamba2 layers in the trunk feed the Q-values that drive
conviction, so conviction inherits temporal history — the model can
learn "market has been signalling for N bars, time to try" without
an explicit time-since-last-trade feature.

Training-time exploration (Boltzmann sampling, epsilon floor 2%,
NoisyNets σ, count bonus) is unchanged. The Flat-cost shifts the
learned Q-values so that deterministic val argmax picks trade-
actions where the training-time exploration already found edge.

Hold semantics retained:
  - Hold while position≠0 (in-trade stance) → positioned-bar
    holding-cost branch (unchanged)
  - Hold while position=0 AND Flat (no-op outcomes) → this branch
    → conviction-scaled opportunity cost.
The distinction is structural by portfolio state, not by dir label.

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

crates/ml/src/cuda_pipeline/experience_kernels.cu

@@ -2212,40 +2212,59 @@ extern "C" __global__ void experience_env_step(
         /* 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 opportunity cost — ISV-driven self-adapting via conviction ──
+         * (val-Flat-collapse fix #3 revision, 2026-04-24)
          *
-         * 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.
+         * Prior revision used `0.5 × holding_cost_rate × vol_proxy` — a
+         * hardcoded 0.5 tuned constant violating
+         * `feedback_isv_for_adaptive_bounds.md`. The magnitude of
+         * "missed-opportunity" isn't a number — it's a signal that
+         * depends on how strong the model's own directional belief is.
          *
-         * 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. */
+         * Semantics this branch covers:
+         *   position ≈ 0 AND not a trade-completion bar. This fires on
+         *   both `dir=Flat` (explicit close-to-zero) and `dir=Hold while
+         *   position=0` (no-op stay-flat). The Hold-while-in-trade case
+         *   (position≠0) hits the positioned-bar holding-cost branch
+         *   above and is NOT penalised here — staying in a position IS
+         *   an active choice that the holding-cost mechanism already
+         *   prices. The distinction is structural: Hold-at-zero and Flat
+         *   produce identical portfolio outcomes so their reward shaping
+         *   must match.
+         *
+         * Self-adapting via conviction (∈ [0, 1]):
+         *   conviction = direction-branch Q-range normalised by ISV[21]
+         *   (q_dir_abs_ref EMA). Computed per-sample in
+         *   `experience_action_select` and threaded into this kernel via
+         *   `conviction_ptr`. Already clamped to [0, 1] at source.
+         *   - Cold start (ISV[21] ≈ 0): fallback conviction = 1.0 → full
+         *     penalty, encourages early exploration out of the flat
+         *     equilibrium.
+         *   - Low conviction: network is uncertain about direction → the
+         *     penalty scales toward zero → don't force trades on noise.
+         *   - High conviction: network has learned strong directional
+         *     edge → penalty scales up → Flat becomes expensive ONLY
+         *     where the model itself says there's an edge to capture.
+         *
+         *   This is the right shape for "incentive for taking action":
+         *   it's ISV-bus-driven, self-adapting via EMA state, and
+         *   piggybacks on a signal the network's own output defines.
+         *   No tuned multiplier.
+         *
+         * Volatility factor (vol_proxy) retained:
+         *   Scales with realised ATR so quiet markets produce little
+         *   penalty. Capped at 0.01 (1%) as numerical-safety bound
+         *   against single-bar vol spikes. */
         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;
+            /* conviction_core ∈ [0, 1] — ISV-driven per-sample adaptive
+             * multiplier. Replaces the prior 0.5 hardcoded constant. */
+            reward = -shaping_scale * holding_cost_rate
+                   * conviction_core * vol_proxy_flat;
         }
     }