fix(dqn): Kelly cap warm-branch deadlock — effective_kelly never collapses to zero

VAL DIAGNOSTIC PROOF (train-4fpzx step=400):
  pick=Long Full (target=0.66) prev_pos=0.0 → pos_post=0.0 actual_dir=Flat
  trail=0, margin can't clip to 0 → only Kelly cap zeroed the target.

ROOT CAUSE:
  effective_kelly = maturity * kelly_f + (1 - maturity) * warmup_floor

  Cold-start fix (commit 2c97e0436) protected `warmup_floor` so it never
  collapses to zero at maturity=0. But the warm branch was left exposed:
  once maturity → 1 (>=10 completed trades), the blend collapses to
  `kelly_f` alone, and `kelly_f = 0` is the natural state of
    (payoff*win_rate - (1-win_rate)) / payoff
  with balanced priors and small actual returns. Val environments with
  pure per-bar P&L (no saboteur/shaping perturbations like training)
  settle into this regime within ~10 trades — after which every non-Hold
  target gets clamped to 0 deterministically. Identical bootstrap-deadlock
  pattern to the IQN trunk SAXPY.

EVIDENCE:
  val_picked_dir_dist [short=0.19 hold=0.20 long=0.39 flat=0.21]  (kernel pick)
  val_dir_dist        [short=0.0001 hold=0.20 long=0.0000 flat=0.80]  (post-physics)
  100% of Long picks and ~99.95% of Short picks become actual_dir=Flat.
  Hold passes through 1:1 (Hold skips margin/Kelly/trail in env_step).

  VALDIAG step=400: act=77 (Long Full, target=0.66) prev=0.0 pos_post=0.0
  -> confirms target zeroed before execute_trade; trail=0 rules out trail;
  margin cap can't produce 0 with equity=$35K vs margin/contract=$17.9K.

FIX:
  effective_kelly = max(kelly_f, warmup_floor)

  The conviction-and-health-driven warmup_floor (in [0.5, 1.0]) becomes a
  permanent minimum cap. `kelly_f` only takes over when the policy has
  demonstrated enough edge to *exceed* the floor. Preserves design intent
  (Kelly drives sizing once stats mature with real edge) while preventing
  the bootstrap deadlock in environments where balanced trades naturally
  yield kelly_f = 0.

  All adaptive ISV-driven signals; no tuned constants.

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

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
+    /* warmup_floor must never collapse to zero — and `effective_kelly` must
-     * effective Kelly cap reduces to `(1 - maturity) * warmup_floor = warmup_floor`,
+     * never collapse to zero either, because that pins target_position to 0
-     * so a zero floor pins target_position to 0 and the policy can never open a
+     * and the policy can never open a position to accumulate the trades that
-     * position to accumulate the trades that would graduate Kelly stats — the
+     * would update Kelly stats. Same bootstrap-deadlock pattern as the IQN
-     * same bootstrap deadlock pattern that gated the IQN trunk SAXPY. The
+     * trunk SAXPY.
-     * adaptive floor is therefore the max of (a) per-sample policy conviction
+     *
-     * and (b) the caller's training-stability floor (`health_floor` = 0.5 +
+     * The original blend `maturity * kelly_f + (1 - maturity) * warmup_floor`
-     * 0.5 × health, ∈ [0.5, 1.0], computed from ISV[LEARNING_HEALTH]). When
+     * protected the cold-start branch (maturity → 0 keeps warmup_floor live)
-     * the policy is confident, conviction wins; when training is healthy but
+     * but left the warm branch unprotected: once maturity = 1, the blend
-     * Q-spread is tight, the health floor wins. Both are adaptive ISV-driven
+     * collapses to `kelly_f`, and `kelly_f = 0` is the natural state of the
-     * signals; no tuned constants. The floor only matters during cold start —
+     * Kelly formula with balanced priors and small actual returns
-     * once `maturity → 1` (≥10 completed trades) the term drops out entirely. */
+     *   `(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;
 }

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)