feat: spread-aware tx cost, confidence scaling, profit-taking bonus

Three new reward intelligence features, all zero-state GPU-native:

1. Spread-aware transaction costs: tx_cost scales by CUSUM volatility.
   Trading in choppy markets costs more — teaches the model to reduce
   frequency in volatile regimes. Real spread DOES widen with volatility.

2. Kelly-inspired confidence scaling: when realized_pnl > 0 (model has
   been right), amplify PnL weight 1.5x. When losing, amplify drawdown
   penalty 1.5x. Self-reinforcing: good decisions → stronger signal →
   better Q-values. Bad decisions → defensive mode → more exploration.

3. Profit-taking bonus: +0.1 reward when model reduces a position toward
   flat while cumulative episode PnL is positive. Explicitly rewards the
   ACT of taking profit, not just being in a winner. Teaches the model
   to lock in gains rather than riding them back to breakeven.

Total kernel additions: ~20 lines, ~10 FLOPs. Zero extra state beyond
what PORTFOLIO_STRIDE=12 already provides.

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

crates/ml/src/cuda_pipeline/experience_kernels.cu

@@ -470,12 +470,22 @@ extern "C" __global__ void experience_env_step(
     float target_exposure = exposure_idx_to_fraction(exposure_idx);
     float target_position = target_exposure * max_position;
 
-    /* ---- Position adjustment with flat proportional transaction cost ---- */
+    /* ---- Position adjustment with volatility-scaled transaction cost ---- */
+    /* Real spread widens in volatile markets. vol_scale (from CUSUM) is
+     * computed later, but we can read CUSUM here for the cost calculation.
+     * This teaches the model that trading in choppy markets is MORE expensive. */
+    float cusum_raw = 0.0f;
+    if (features != NULL && bar_idx < total_bars && market_dim > 41) {
+        cusum_raw = features[(long long)bar_idx * market_dim + 41];
+    }
+    float spread_scale = cusum_raw / 0.5f;
+    spread_scale = (spread_scale < 0.5f) ? 0.5f : ((spread_scale > 2.0f) ? 2.0f : spread_scale);
+
     float delta   = target_position - position;
     float tx_cost = 0.0f;
 
     if (delta != 0.0f) {
-        tx_cost   = fabsf(delta) * raw_close * tx_cost_multiplier;
+        tx_cost   = fabsf(delta) * raw_close * tx_cost_multiplier * spread_scale;
         cash     -= delta * raw_close;
         cash     -= tx_cost;
         position  = target_position;
@@ -484,6 +494,19 @@ extern "C" __global__ void experience_env_step(
     /* ---- Mark-to-market PnL for this timestep ---- */
     float raw_pnl = position * (raw_next - raw_close);
 
+    /* ---- Profit-taking bonus: reward closing profitable positions ---- */
+    /* Explicitly rewards the ACT of taking profit, not just holding winners.
+     * Triggered when: position reduced AND cumulative realized PnL is positive.
+     * prev_pos is the position BEFORE delta was applied (read at kernel start). */
+    float profit_take_bonus = 0.0f;
+    float prev_pos = ps[0]; /* still the pre-delta value (writeback at kernel end) */
+    if (fabsf(delta) > 0.001f && fabsf(position) < fabsf(prev_pos)) {
+        /* Position reduced toward flat */
+        if (ps[11] > 0.0f) { /* realized_pnl > 0: episode has been profitable */
+            profit_take_bonus = 0.1f;
+        }
+    }
+
     /* ==== Component 1: Differential Sharpe Ratio (Moody & Saffell 2001) ==== */
     /* Return: guard division by zero on prev_equity */
     float equity_denom = (prev_equity > 1.0f) ? prev_equity : 1.0f;
@@ -572,12 +595,23 @@ extern "C" __global__ void experience_env_step(
     float eff_w_dsr = w_dsr * (2.0f - trend_scale);
     float eff_w_dd  = w_dd  * vol_scale;
 
+    /* ==== Confidence scaling (Kelly-inspired) ==== */
+    /* When the model's recent decisions are profitable (realized_pnl > 0),
+     * amplify PnL weight — trust the model's edge. When losing, amplify
+     * drawdown/DSR — be defensive. Self-reinforcing learning loop:
+     * good decisions → stronger signal → better Q-values → better decisions. */
+    float realized = ps[11]; /* cumulative realized PnL this episode */
+    float confidence = (realized > 0.0f) ? 1.5f : 0.5f;
+    eff_w_pnl *= confidence;
+    eff_w_dd  *= (2.0f - confidence);  /* inverse: 0.5 when winning, 1.5 when losing */
+
     /* ==== Composite reward ==== */
     float reward = eff_w_dsr * dsr_t
                  + eff_w_pnl * asym_pnl
                  - eff_w_dd  * drawdown_penalty_t
                  - w_idle    * idle_penalty_t
                  - hold_time_penalty
+                 + profit_take_bonus
                  - tx_cost;
 
     /* ---- Done detection ---- */