diff --git a/crates/ml/src/cuda_pipeline/trade_physics.cuh b/crates/ml/src/cuda_pipeline/trade_physics.cuh
index acf5c8fd0..baddef62f 100644
--- a/crates/ml/src/cuda_pipeline/trade_physics.cuh
+++ b/crates/ml/src/cuda_pipeline/trade_physics.cuh
@@ -317,23 +317,42 @@ __device__ __forceinline__ float kelly_position_cap(
      * doesn't swing wildly on small samples. */
     float total_trades = win_count + loss_count;
     float maturity = fminf(1.0f, total_trades * 0.1f);
-    /* warmup_floor must never collapse to zero: at cold start (maturity=0) the
-     * effective Kelly cap reduces to `(1 - maturity) * warmup_floor = warmup_floor`,
-     * so a zero floor pins target_position to 0 and the policy can never open a
-     * position to accumulate the trades that would graduate Kelly stats — the
-     * same bootstrap deadlock pattern that gated the IQN trunk SAXPY. The
-     * adaptive floor is therefore the max of (a) per-sample policy conviction
-     * and (b) the caller's training-stability floor (`health_floor` = 0.5 +
-     * 0.5 × health, ∈ [0.5, 1.0], computed from ISV[LEARNING_HEALTH]). When
-     * the policy is confident, conviction wins; when training is healthy but
-     * Q-spread is tight, the health floor wins. Both are adaptive ISV-driven
-     * signals; no tuned constants. The floor only matters during cold start —
-     * once `maturity → 1` (≥10 completed trades) the term drops out entirely. */
+    /* warmup_floor must never collapse to zero — and `effective_kelly` must
+     * never collapse to zero either, because that pins target_position to 0
+     * and the policy can never open a position to accumulate the trades that
+     * would update Kelly stats. Same bootstrap-deadlock pattern as the IQN
+     * trunk SAXPY.
+     *
+     * The original blend `maturity * kelly_f + (1 - maturity) * warmup_floor`
+     * protected the cold-start branch (maturity → 0 keeps warmup_floor live)
+     * but left the warm branch unprotected: once maturity = 1, the blend
+     * collapses to `kelly_f`, and `kelly_f = 0` is the natural state of the
+     * Kelly formula with balanced priors and small actual returns
+     *   `(payoff*win_rate − (1 − win_rate)) / payoff`
+     * trivially clamps to 0 when payoff*win_rate ≤ (1 − win_rate). Validation
+     * environments with pure per-bar P&L (no saboteur/shaping like training)
+     * settle into this regime within ~10 trades, after which non-Hold targets
+     * get clamped to 0 deterministically — the val-Flat-collapse pathology.
+     *
+     * Fix: take `max(kelly_f, warmup_floor)`. The conviction-and-health-driven
+     * floor remains a permanent minimum cap; `kelly_f` only takes over when
+     * the policy has demonstrated enough edge to *exceed* the floor. This
+     * preserves the design intent (Kelly drives sizing once stats mature)
+     * while preventing the bootstrap deadlock in environments where balanced
+     * trades naturally yield kelly_f = 0.
+     *
+     * The adaptive floor is the max of (a) per-sample policy conviction and
+     * (b) the caller's training-stability floor (`health_floor` = 0.5 +
+     * 0.5 × health, ∈ [0.5, 1.0], computed from ISV[LEARNING_HEALTH]). Both
+     * are adaptive ISV-driven signals; no tuned constants. */
     float warmup_floor = fmaxf(
         fminf(1.0f, fmaxf(0.0f, conviction)),
         fminf(1.0f, fmaxf(0.0f, health_floor))
     );
-    float effective_kelly = maturity * kelly_f + (1.0f - maturity) * warmup_floor;
+    /* maturity reserved for downstream signals; the cap itself uses the
+     * permanent-floor max formulation. */
+    (void)maturity;
+    float effective_kelly = fmaxf(kelly_f, warmup_floor);
 
     return effective_kelly * max_position * safety_multiplier;
 }
diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md
index a1599885a..6a05e6185 100644
--- a/docs/dqn-wire-up-audit.md
+++ b/docs/dqn-wire-up-audit.md
@@ -277,7 +277,7 @@ P5T5 Phase E (2026-04-26): per-fold reset for `MetricBandsRegistry` regression-d
 | `backtest_env_kernel.cu` (`backtest_env_step`) | `gpu_backtest_evaluator.rs` | Wired | Already listed above | — |
 | `common_device_functions.cuh` | prepended to all kernels at compile time by `build.rs` | Wired | Shared device helpers (BF16 wrappers, warp reduce) | — |
 | `state_layout.cuh` | included by kernels that reference state buffer layout | Wired | Named index constants for state buffer | — |
-| `trade_physics.cuh` | included by `backtest_env_kernel.cu` and others at compile time | Wired | Kelly / physics helpers shared across kernels | — |
+| `trade_physics.cuh` | included by `backtest_env_kernel.cu` and others at compile time | Wired | Kelly / physics helpers shared across kernels. `kelly_position_cap` warm-branch deadlock fixed: `effective_kelly = max(kelly_f, warmup_floor)` so the cap can never collapse to zero when balanced trades naturally yield kelly_f=0 at maturity=1; preserves design intent (Kelly drives sizing once stats mature with real edge) while preventing the bootstrap deadlock that produced 100% Long/Short→Flat conversion in val. | — |
 
 ## Model Modules — Supervised-Only (OUT-of-DQN-scope per Part E)