diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu
index 9c0daedce..c534e60b7 100644
--- a/crates/ml/src/cuda_pipeline/experience_kernels.cu
+++ b/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*