/**
 * Shared trade physics for DQN environment step kernels.
 *
 * This header defines the SINGLE SOURCE OF TRUTH for:
 * - Factored action decoding (exposure index extraction)
 * - Exposure-to-position mapping (index → fraction → contracts)
 * - Transaction cost computation (Almgren-Chriss impact model)
 * - Capital floor circuit breaker (75% of peak equity)
 * - Hold time tracking (bars in current position)
 *
 * Both backtest_env_kernel.cu and experience_kernels.cu MUST include
 * this header to ensure train/eval consistency. Any change here
 * automatically propagates to both.
 *
 * All functions are __device__ __forceinline__ for zero call overhead.
 */

#ifndef TRADE_PHYSICS_CUH
#define TRADE_PHYSICS_CUH

/* ── SP5 Layer B: ISV slot constants consumed in unified_env_step_core ── */

/* Pearl 6 (cross-fold-persistent Kelly): smoothed Kelly fraction.
 * ISV[KELLY_F_SMOOTH_INDEX] = ISV[280].
 * Mirrors sp5_isv_slots.rs::KELLY_F_SMOOTH_INDEX. */
#define SP5_KELLY_F_SMOOTH_INDEX 280

/* Pearl 8 (per-direction trailing stop distance): base trail distance per dir.
 * ISV[TRAIL_DIST_PER_DIR_BASE + dir] = ISV[270..274).
 * dir: 0=Short, 1=Hold, 2=Long, 3=Flat.
 * Mirrors sp5_isv_slots.rs::TRAIL_DIST_PER_DIR_BASE. */
#define SP5_TRAIL_DIST_PER_DIR_BASE 270

/* SP9 Fix 37 (2026-05-03): Kelly cold-start warmup floor — adaptive,
 * fully ISV-driven, OR'd statistical/behavioral/temporal confidence.
 * ISV[KELLY_WARMUP_FLOOR_INDEX] = ISV[330].
 * Mirrors sp5_isv_slots.rs::KELLY_WARMUP_FLOOR_INDEX. Consumed in
 * `unified_env_step_core` as an additional max() floor on `health_safety`,
 * threaded through `apply_kelly_cap`'s `health_floor` parameter. Zero
 * sentinel at fold boundary (Pearl A bootstrap) — kernel writes the
 * first observation directly so the slot is well-defined from epoch 1
 * onward. The consumer's `fmaxf(...)` gracefully degrades: while the
 * slot is still at sentinel zero (pre-first-observation, fold 0 step 0)
 * the floor contribution is zero and the existing `health_safety_sp5`
 * dominates.
 *
 * Per `pearl_controller_anchors_isv_driven.md` + `pearl_cold_start_exit_signal_or.md`. */
#define SP9_KELLY_WARMUP_FLOOR_INDEX 330

/* ── Factored action → exposure index ─────────────────────────────────── */

__device__ __forceinline__ int decode_exposure_index(
    int action_val, int b0_size, int b1_size, int b2_size
) {
    int exposure_idx;
    if (b1_size > 0 && b2_size > 0) {
        int denom = b1_size * b2_size;
        exposure_idx = action_val / denom;
    } else {
        exposure_idx = action_val;
    }
    if (exposure_idx < 0) exposure_idx = 0;
    if (exposure_idx >= b0_size) exposure_idx = b0_size - 1;
    return exposure_idx;
}

/* ── Exposure index → target position ─────────────────────────────────── */

__device__ __forceinline__ float compute_target_position(
    int exposure_idx, int b0_size, float max_position
) {
    float step_size = (b0_size > 1) ? 2.0f / (float)(b0_size - 1) : 0.0f;
    float fraction = -1.0f + (float)exposure_idx * step_size;
    return fraction * max_position;
}

/* ── 4-branch hierarchical position mapping (Phase 2) ────────────────
 * direction(4) × magnitude(3) replaces flat exposure(9).
 * Direction: 0=Short(-1), 1=Hold(current), 2=Long(+1), 3=Flat(0)
 * Magnitude: 0=Quarter(0.25), 1=Half(0.50), 2=Full(1.00)
 * Position = direction × magnitude × max_position
 *
 * Hold (dir_idx=1) returns NaN sentinel — caller MUST check dir_idx==1
 * and skip trade execution, preserving current position unchanged.
 * The NaN sentinel ensures a missed check produces visible corruption
 * rather than silent wrong trades.
 *
 * Grid: S25  S50  S100
 *       (hold — current position unchanged)
 *       L25  L50  L100
 *       0    0    0     (Flat × any = 0)
 */
#include "state_layout.cuh"

__device__ __forceinline__ float compute_target_position_4branch(
    int dir_idx, int mag_idx, float max_position
) {
    /* Hold sentinel: caller must handle dir_idx==1 before calling.
     * Returns NaN as safety net — any arithmetic with NaN propagates visibly. */
    if (dir_idx == DIR_HOLD) return __int_as_float(0x7FC00000); /* NaN sentinel */
    float direction = (dir_idx == DIR_SHORT) ? -1.0f : (dir_idx == DIR_LONG) ? 1.0f : 0.0f;
    float magnitude = (mag_idx == MAG_QUARTER) ? 0.25f : (mag_idx == MAG_HALF) ? 0.5f : 1.0f;
    return direction * magnitude * max_position;
}

/* ── 4-branch action decoding ────────────────────────────────────────
 * Action = dir * (b1*b2*b3) + mag * (b2*b3) + order * b3 + urgency
 * With b0=4, b1=b2=b3=3: action = dir*27 + mag*9 + order*3 + urgency */
__device__ __forceinline__ int decode_direction_4b(int action, int b1, int b2, int b3) {
    return action / (b1 * b2 * b3);
}
__device__ __forceinline__ int decode_magnitude_4b(int action, int b1, int b2, int b3) {
    return (action / (b2 * b3)) % b1;
}
__device__ __forceinline__ int decode_order_4b(int action, int b2, int b3) {
    return (action / b3) % b2;
}
__device__ __forceinline__ int decode_urgency_4b(int action, int b3) {
    return action % b3;
}

/* ── Margin-aware position cap with drawdown scaling ──────────────────── */
/* Real brokers enforce margin requirements AND risk managers cut exposure
 * as drawdown increases. Without drawdown scaling, the margin cap is static
 * and multi-bar cascading losses can push DD far past the circuit breaker.
 *
 * Drawdown scaling:
 *   dd_frac = (peak - equity) / peak     [0 = no DD, 0.25 = floor]
 *   dd_scale = max(0, 1 - dd_frac / 0.25)  [1.0 at 0% DD, 0.0 at 25% DD]
 *   effective_margin = margin_per_contract / dd_scale  [tighter as DD grows]
 *
 * At 0% DD: full margin (100% of normal position)
 * At 12.5% DD: 50% of normal position
 * At 20% DD: 20% of normal position
 * At 25% DD: floor triggers (check_capital_floor handles this)
 *
 * This ensures max single-bar loss decreases proportionally with DD,
 * preventing the 38-51% overshoot seen with static margin. */

__device__ __forceinline__ float apply_margin_cap(
    float target_position,
    float equity,
    float margin_per_contract,
    float peak_equity
) {
    if (margin_per_contract <= 0.0f || equity <= 0.0f) return target_position;

    /* Drawdown-based position scaling */
    float dd_scale = 1.0f;
    if (peak_equity > 1.0f) {
        float dd_frac = (peak_equity - equity) / peak_equity;
        dd_frac = fmaxf(0.0f, fminf(dd_frac, 1.0f));
        /* Linear ramp: 1.0 at 0% DD → 0.0 at 25% DD (floor threshold) */
        dd_scale = fmaxf(0.05f, 1.0f - dd_frac / 0.25f);  /* floor at 5% to avoid zero */
    }

    float max_contracts = (equity / margin_per_contract) * dd_scale;
    if (max_contracts < 0.01f) max_contracts = 0.0f;
    float abs_target = fabsf(target_position);
    if (abs_target > max_contracts) {
        float sign = (target_position > 0.0f) ? 1.0f : -1.0f;
        return sign * max_contracts;
    }
    return target_position;
}

/* ── Order type index from factored action ────────────────────────────── */

__device__ __forceinline__ int decode_order_type(
    int action_val, int b1_size, int b2_size
) {
    return (b1_size > 0 && b2_size > 0) ? (action_val / b2_size) % b1_size : 0;
}

/* ── Transaction cost (Almgren-Chriss impact model) ──────────────────── */

__device__ __forceinline__ float compute_tx_cost(
    float delta,                  /* position change (signed) */
    float close,                  /* current price */
    float tx_cost_bps,            /* base cost multiplier */
    float spread_cost,            /* bid-ask spread cost per unit */
    float max_position,           /* for impact scaling */
    int   order_type_idx,         /* 0=Market, 1=IoC, 2=LimitMaker */
    float spread_scale_override   /* <= 0 = compute from sqrt(delta/max_pos); > 0 = use directly (CUSUM) */
) {
    float abs_delta = fabsf(delta);
    float size_ratio = abs_delta / fmaxf(max_position, 0.01f);

    /* Spread scaling: CUSUM-based from market features when available,
     * otherwise sqrt(delta/max_pos) static model. */
    float spread_scale;
    if (spread_scale_override > 0.0f) {
        spread_scale = spread_scale_override;
    } else {
        spread_scale = sqrtf(size_ratio);
        if (spread_scale < 1.0f) spread_scale = 1.0f;
    }

    /* Square-root market impact (Almgren & Chriss, 2000).
     * impact_scale = 1 + sqrt(|delta|/max_position)
     * Both training and backtest use the same impact model. */
    float impact_scale = 1.0f + sqrtf(size_ratio);

    float order_premium = (order_type_idx == 0) ? 0.0f
                        : (order_type_idx == 1) ? 0.0002f    /* IoC: +2 bps */
                        : -0.0005f;                           /* LimitMaker: -5 bps rebate */
    return abs_delta * close * (tx_cost_bps * 0.0001f * spread_scale * impact_scale + order_premium)
         + abs_delta * spread_cost * 0.5f;
}

/* ── Unified trade execution (notional cash model) ───────────────────── */
/* Both training and backtest kernels MUST use this for P&L consistency.
 *
 * The notional model:
 *   cash -= delta * price    (buying costs cash, selling adds cash)
 *   cash -= tx_cost          (transaction costs always deducted)
 *   position = target        (position updated to new level)
 *
 * This correctly handles partial fills (L50->L100: delta=0.5, cash -= 0.5*price)
 * without needing entry_price tracking for P&L.
 *
 * Returns: transaction cost charged (for logging/stats). */

__device__ __forceinline__ float execute_trade(
    float* position,
    float* cash,
    float target_position,
    float price,
    float tx_cost_bps,
    float spread_cost,
    float max_position,
    int   order_type_idx,
    float spread_scale_override
) {
    float delta = target_position - *position;
    if (fabsf(delta) <= 0.001f || price <= 0.0f) return 0.0f;

    float cost = compute_tx_cost(delta, price, tx_cost_bps, spread_cost,
                                 max_position, order_type_idx, spread_scale_override);
    *cash -= delta * price;    /* notional: buy costs, sell adds */
    *cash -= cost;             /* transaction cost */
    *position = target_position;
    return cost;
}

/* ── Drawdown computation ────────────────────────────────────────────── */
/* Returns (peak - equity) / peak, clamped to [0, 1].                    */

__device__ __forceinline__ float compute_drawdown(
    float equity, float peak_equity
) {
    if (peak_equity <= 1.0f) return 0.0f;
    float dd = (peak_equity - equity) / peak_equity;
    return fmaxf(dd, 0.0f);
}

/* ── Kelly-optimal position size cap ──────────────────────────────────── */
/* Returns the maximum absolute position the agent is allowed to hold,
 * given the empirical win/loss statistics accumulated so far in this run.
 *
 * Replaces the old `kelly_sizing_weight` behavioral reward (which penalized
 * deviation from Kelly). Physics constraint instead of reward shaping:
 * the environment refuses to let the agent over-lever, rather than
 * scoring the agent for matching a formula.
 *
 * Priors (prior_wins=2, prior_losses=2, prior_sum_wins=0.01,
 * prior_sum_losses=0.01) prevent the cap from being degenerate on cold-start
 * runs where no trades have completed yet — without priors Kelly would
 * be undefined for the first ~10 trades.
 *
 * The safety_multiplier is composed by the caller from two orthogonal
 * adaptive signals:
 *   safety = max(0.5 + 0.5 × health,  conviction)
 *                 \---------------/  \---------/
 *                 training-stability  per-sample
 *                 floor               policy confidence
 * Bounded to [0.5, 1.0] by the health-safety floor. Same temporal-coupling
 * pattern as the distillation SAXPY alpha and the Q-target label smoothing
 * eps. Applied uniformly in both training and backtest envs so there's no
 * train/val mismatch.
 *
 * The `conviction` parameter (∈ [0, 1], caller-normalised) drives the
 * cold-start warmup floor — no longer a static 0.5f. Derived by the
 * caller from the policy's per-sample direction Q-spread normalised by
 * the `q_dir_abs_ref` ISV EMA (isv_signals[21]). High conviction → high
 * floor (trust the policy to size up even before Kelly stats accumulate).
 * Low conviction → low floor (safety dominates). `warmup_floor` only
 * matters during the cold-start window; once `maturity → 1` (trade count
 * ≥ 10) the blend flows to pure kelly_f and the floor contribution
 * vanishes. Conviction is [0, 1] clamped here — a structural bound
 * (fraction of max_position), not a tuned numeric knob; the raw signal
 * IS the floor.
 */
__device__ __forceinline__ float kelly_position_cap(
    float win_count,
    float loss_count,
    float sum_wins,
    float sum_losses,
    float max_position,
    float safety_multiplier,
    float conviction,
    float health_floor
) {
    const float prior_wins = 2.0f;
    const float prior_losses = 2.0f;
    const float prior_sum_wins = 0.01f;
    const float prior_sum_losses = 0.01f;

    float eff_wins = win_count + prior_wins;
    float eff_losses = loss_count + prior_losses;
    float eff_total = eff_wins + eff_losses;
    float win_rate = eff_wins / eff_total;
    float avg_win = (sum_wins + prior_sum_wins) / eff_wins;
    float avg_loss = (sum_losses + prior_sum_losses) / eff_losses;
    float payoff = avg_win / fmaxf(avg_loss, 0.0001f);
    float kelly_f = (payoff * win_rate - (1.0f - win_rate)) / fmaxf(payoff, 0.0001f);
    kelly_f = fmaxf(fminf(kelly_f, 1.0f), 0.0f);

    /* Cold-start warmup: with balanced priors (equal wins/losses, equal
     * sums), the formula above yields kelly_f ≈ 0 until real trades
     * accumulate. That would clamp the agent to zero position forever
     * and starve Q-learning before any signal can form.
     *
     * Blend towards a warmup floor weighted by stats maturity:
     * maturity = min(1, total_trades / 10). Early on the cap is dominated
     * by the floor; as real trades accumulate it transitions to true
     * Kelly. At 10+ completed trades the floor contribution is zero —
     * pure data-driven Kelly.
     *
     * Adaptive warmup floor: derived from policy-observed conviction
     * (normalised direction Q-spread, computed at action_select and
     * passed through by the caller). High conviction → trust the policy
     * to size up even at cold start. Low conviction → safety dominates
     * (floor collapses toward zero). Combined with the health-coupled
     * safety multiplier (0.5 + 0.5×health ∈ [0.5, 1.0]), cold-start cap
     * adapts to both network health AND per-sample Q-spread.
     *
     * Clamped to [0, 1] — structural bound (a fraction of max_position).
     * No hard-coded numeric knob. Conviction flows adaptively from the
     * network's own Q-spread signal, evolving temporally with training.
     * Once maturity → 1 the floor term drops out entirely, so conviction
     * only matters during the cold-start window.
     *
     * 10-trade maturity is short enough that learning isn't suppressed
     * for long; the priors (4 virtual trades) already ensure the formula
     * 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 — 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))
    );
    /* 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;
}

/* Apply Kelly cap: clamp a signed target_position to [-cap, +cap].
 * Convenience for call sites that want a single clamp call.
 * `conviction` ∈ [0, 1] threads through to kelly_position_cap's adaptive
 * warmup_floor — see docstring there. */
__device__ __forceinline__ float apply_kelly_cap(
    float target_position,
    float win_count,
    float loss_count,
    float sum_wins,
    float sum_losses,
    float max_position,
    float safety_multiplier,
    float conviction,
    float health_floor
) {
    float cap = kelly_position_cap(win_count, loss_count, sum_wins, sum_losses,
                                   max_position, safety_multiplier, conviction,
                                   health_floor);
    return fmaxf(-cap, fminf(cap, target_position));
}

/* ── Record a completed-trade outcome into Kelly statistics ───────────
 *
 * Shared by both training (experience_kernels.cu) and validation
 * (backtest_env_kernel.cu) env steps so the Kelly stats feeding
 * apply_kelly_cap are computed from identical trade-exit semantics in
 * both environments. Eliminates the stats-update duplication that would
 * otherwise drift — exactly the mismatch we're trying to fix with the
 * shared trade-physics module.
 *
 * A "trade completion" is: positioned → flat, OR direction reversal.
 * Both represent a realized P&L event on the prior position. Partial
 * scaling (L50 → L100) is NOT counted — no prior position was closed.
 *
 * The realized trade return is computed as
 *   (prev_position × (close − entry_price)) / pre_trade_equity
 * normalized by PRE-trade equity so winners and losers are measured on
 * the same scale. Break-even trades (return exactly 0) are ignored.
 *
 * Stats are passed by pointer so the caller's register-resident floats
 * are updated in place. Caller owns writing back to global/shared memory. */
__device__ __forceinline__ void record_kelly_trade_outcome(
    float prev_position,   /* position before the trade executed this step */
    float curr_position,   /* position after execute_trade this step */
    float entry_price,     /* entry price recorded for the closed trade */
    float close,           /* mark-to-market close price this step */
    float pre_trade_equity,/* portfolio value BEFORE this step's P&L */
    float* win_count,      /* in/out: winning trades counter */
    float* loss_count,     /* in/out: losing trades counter */
    float* sum_wins,       /* in/out: cumulative winning returns */
    float* sum_losses      /* in/out: cumulative |losing returns| */
) {
    int prev_pos_active = (fabsf(prev_position) > 0.001f);
    int is_exit     = prev_pos_active && (fabsf(curr_position) < 0.001f);
    int is_reversal = (prev_position * curr_position < 0.0f);
    if ((is_exit || is_reversal) && entry_price > 0.0f && pre_trade_equity > 1.0f) {
        float realized = prev_position * (close - entry_price);
        float trade_return = realized / pre_trade_equity;
        if (trade_return > 0.0f) {
            *win_count  += 1.0f;
            *sum_wins   += trade_return;
        } else if (trade_return < 0.0f) {
            *loss_count += 1.0f;
            *sum_losses += fabsf(trade_return);
        }
    }
}

/* ── SP12 v3 (2026-05-04): per-trade event-driven reward composition ──────
 * Three small device functions extracted from `experience_kernels.cu` so the
 * math is testable in isolation. Production `experience_env_step` calls these
 * inside the `segment_complete` branch (around line ~2800). The Rust unit
 * tests in `tests/sp12_reward_math_tests.rs` exercise them via a single-thread
 * wrapper kernel (`sp12_reward_math_test_kernel.cu`) — synthetic inputs,
 * analytical expected outputs, no CPU reference impl per
 * `feedback_no_cpu_test_fallbacks`.
 *
 * Constants live in `state_layout.cuh` (REWARD_NEG_CAP, REWARD_POS_CAP,
 * MIN_HOLD_TARGET, MIN_HOLD_PENALTY_MAX, MIN_HOLD_TEMPERATURE_*). The device
 * functions take the bounds/parameters as arguments rather than reading the
 * macros directly so the test wrapper can drive the same code path with
 * arbitrary values without #define gymnastics.
 * ──────────────────────────────────────────────────────────────────────── */

/* SP12 v3 Change 1: asymmetric bounded reward cap (prospect-theory loss
 * aversion). With production constants `neg_cap=-10, pos_cap=+5` the 2:1
 * ratio matches the Kahneman/Tversky meta-analysis estimate of human loss
 * aversion (~2.0-2.25). Restores selectivity that SP11's symmetric ±10 cap
 * (commit 35db31089) erased while preserving the bilateral bound that
 * fixed slot-63 PopArt EMA inflation. Per
 * `pearl_audit_unboundedness_for_implicit_asymmetry`. */
__device__ __forceinline__ float compute_asymmetric_capped_pnl(
    float base_reward,
    float neg_cap,   /* production: REWARD_NEG_CAP = -10.0f */
    float pos_cap    /* production: REWARD_POS_CAP =  +5.0f */
) {
    return fmaxf(neg_cap, fminf(base_reward, pos_cap));
}

/* SP12 v3 Change 2: min-hold soft penalty for voluntary trade exits.
 * Returns the penalty MAGNITUDE (positive). Caller subtracts it from the
 * reward (production multiplies by `shaping_scale` at the call site).
 *
 * Soft saturation: `factor = deficit / (deficit + T)` — bounded [0, 1],
 * smooth (no cliff), monotone increasing in deficit, monotone decreasing
 * in T. When `hold_time >= min_hold_target` the deficit is non-positive
 * and the function returns 0 — gating handled here so the call site
 * stays a single expression.
 *
 * Numerical guard: production uses `fmaxf(deficit + T, 1e-6f)` to defend
 * against simultaneous (deficit ≤ 0, T ≤ 0) malformed configs. The early
 * `hold_time >= target` exit makes the guard unreachable for sensible
 * inputs but the bound mirrors production exactly. */
__device__ __forceinline__ float compute_min_hold_penalty(
    float hold_time,
    float min_hold_target,
    float min_hold_temperature,
    float min_hold_penalty_max
) {
    if (hold_time >= min_hold_target) return 0.0f;
    const float deficit     = min_hold_target - hold_time;
    const float soft_factor = deficit / fmaxf(deficit + min_hold_temperature, 1e-6f);
    return min_hold_penalty_max * soft_factor;
}

/* SP12 v3 Change 3: lump-sum carrying-cost at trade close.
 * Returns the (negative) reward magnitude. Production wires the result into
 * `r_opp_cost` (rc[4]) only on `segment_complete` bars; per-bar callers
 * MUST gate before invocation (the device function itself does NOT check
 * exiting_trade — that gating is a call-site invariant of
 * `experience_env_step`'s segment_complete branch).
 *
 * Per `pearl_event_driven_reward_density_alignment`: preserves the
 * carrying-cost economic concept (k × |position| × hold_time) while moving
 * its delivery from per-bar density-positive shaping to a single event-
 * driven payment at trade close. */
__device__ __forceinline__ float compute_lump_sum_opp_cost(
    float shaping_scale,
    float holding_cost_rate,
    float position,
    float hold_time
) {
    /* Parens preserve the production multiplication order
     * (`-(scale * rate * |pos| * hold)`) so the refactor is bit-equivalent
     * under f32 rounding to the inline computation it replaces. */
    return -(shaping_scale * holding_cost_rate * fabsf(position) * hold_time);
}

/* ── Drawdown penalty: smooth ramp from threshold to floor ───────────── */
/* Returns a penalty in [-5*w_dd, 0]. Applied every step so the model    */
/* learns to reduce position size DURING drawdown.                       */

__device__ __forceinline__ float compute_drawdown_penalty(
    float equity, float peak_equity,
    float dd_threshold, float w_dd
) {
    if (w_dd <= 0.0f) return 0.0f;
    float dd = compute_drawdown(equity, peak_equity);
    if (dd <= dd_threshold) return 0.0f;
    float floor_dd = 0.25f; // capital floor = 25% DD
    float dd_excess = fminf((dd - dd_threshold) / (floor_dd - dd_threshold), 1.0f);
    return -5.0f * dd_excess * w_dd;
}

/* ── Capital floor circuit breaker (PDT $25K rule) ───────────────────── */
/* Triggers when equity drops 25% from peak (floor = 75% of peak).       */
/* With $35K initial: floor = $26.25K — above $25K PDT minimum.          */
/* Matches training kernel (experience_env_step) capital_floor logic.     */

__device__ __forceinline__ int check_capital_floor(
    float value, float max_equity
) {
    float capital_floor = max_equity * 0.75f;
    return (value < capital_floor && max_equity > 1.0f) ? 1 : 0;
}

/* ── Hold time tracking ──────────────────────────────────────────────── */

__device__ __forceinline__ float update_hold_time(
    float prev_position, float new_position, float hold_time
) {
    if (fabsf(new_position) > 0.001f) {
        if (fabsf(prev_position) < 0.001f) {
            return 1.0f;  /* Entry from flat */
        }
        int prev_sign = (prev_position > 0.001f) ? 1 : -1;
        int new_sign  = (new_position > 0.001f) ? 1 : -1;
        if (new_sign == prev_sign) {
            return hold_time + 1.0f;  /* Same direction — increment */
        }
        return 1.0f;  /* Reversal — new trade */
    }
    return 0.0f;  /* Flat */
}

/* ── Trailing stop: exit when profit retreats from peak ──────────────── */
/* Only triggers when:
 *   - Position held >= 1 bar (must have at least one bar of data)
 *   - Peak equity > 1.0 (valid equity tracking)
 *   - Peak return > trail_distance (profit must exist before trailing)
 *   - Unrealized return < trail_floor (profit retreated beyond threshold)
 *
 * vol_scale and trend_scale allow regime-adaptive widening (ADX + CUSUM).
 * Pass 1.0f for both in contexts without feature access (e.g. backtest).
 * Returns 1 if trailing stop triggered, 0 otherwise. */

__device__ __forceinline__ int check_trailing_stop(
    float hold_time,
    float peak_equity,
    float prev_equity,
    float current_trade_return,
    float base_trail_distance,  /* base threshold, e.g. 0.005 = 0.5% */
    float vol_scale,            /* CUSUM-based volatility scaling (1.0 = no scale) */
    float trend_scale           /* ADX-based trend scaling (1.0 = no scale) */
) {
    if (hold_time < 1.0f || peak_equity <= 1.0f) return 0;
    float trail_distance = base_trail_distance * vol_scale * trend_scale;
    float peak_return = (peak_equity - prev_equity) / fmaxf(prev_equity, 1.0f);
    if (peak_return > trail_distance) {
        float trail_floor = peak_return - trail_distance;
        if (current_trade_return < trail_floor) return 1;
    }
    return 0;
}

/* ── Regime-adaptive trailing-stop scales from ADX/CUSUM features ──────
 * Both training and val call this with their own features[] + bar_idx so
 * the computation is bit-identical. Feature layout: index 40 = ADX
 * (normalized, threshold 0.25 for trending), index 41 = CUSUM direction
 * (in [-1,+1], |value| > 0.5 indicates regime change).
 *
 * Caller passes NULL features to fall back to 1.0/1.0 (fixed-width stop),
 * which also covers kernels built before features were plumbed. */
__device__ __forceinline__ void compute_regime_trail_scales(
    const float* __restrict__ features,
    long long    bar_flat_idx,     /* already resolved to features[bar_flat_idx * market_dim + ...] */
    int          market_dim,
    float*       vol_scale_out,
    float*       trend_scale_out
) {
    if (features == NULL || market_dim <= 41) {
        *vol_scale_out   = 1.0f;
        *trend_scale_out = 1.0f;
        return;
    }
    float adx   = features[bar_flat_idx * (long long)market_dim + 40];
    float cusum = features[bar_flat_idx * (long long)market_dim + 41];
    /* Defensive clamp — feature pipeline guarantees [0,1]/[-1,1] but guard
     * against out-of-range from data gaps to keep the stop bounded. */
    adx = fminf(fmaxf(adx, 0.0f), 1.0f);
    float cusum_abs = fminf(fmaxf(fabsf(cusum), 0.0f), 1.0f);
    /* Formulas mirror the pre-unification reward v5 (commit 263997ad3),
     * translated to the normalized feature convention:
     *   trend_scale widens when ADX > 0.25 (trending), cap 2.5
     *   vol_scale   widens when |CUSUM| > 0.5 (regime-change), cap 2.5 */
    float trend_scale = 1.0f + fmaxf(adx       - 0.25f, 0.0f) * 2.0f;
    float vol_scale   = 1.0f + fmaxf(cusum_abs - 0.5f,  0.0f) * 2.0f;
    *trend_scale_out = fminf(trend_scale, 2.5f);
    *vol_scale_out   = fminf(vol_scale,   2.5f);
}

/* ══════════════════════════════════════════════════════════════════════
 * UNIFIED ENV STEP CORE (Phase 3 — single source of truth)
 *
 * Captures the canonical environment step that BOTH experience_env_step
 * (training) and backtest_env_step (validation) execute identically:
 *
 *   1. Action decode (4-branch: dir, mag, order)
 *   2. Hold action passthrough
 *   3. Target position via compute_target_position_4branch
 *   4. Margin cap (apply_margin_cap)
 *   5. Kelly cap with health-coupled safety AND per-branch policy conviction
 *      (apply_kelly_cap). Two conviction signals (direction + magnitude),
 *      each ∈ [0, 1] — per-sample, derived from each branch's Q-spread
 *      normalised by its ISV |Q|-scale reference (ISV[21] for direction,
 *      ISV[16] for magnitude). Composed with `max()` then composed with
 *      health_safety with `max()`. Drives the cold-start warmup floor so
 *      Kelly cap adapts to whichever branch expresses high confidence,
 *      instead of a static 0.5 floor; once maturity → 1 it's a no-op.
 *   6. Trailing stop check — base 0.5%, regime-adaptive via vol_scale and
 *      trend_scale (caller resolves from ADX/CUSUM features; identical
 *      formulas in both training and val to preserve train/eval parity)
 *   7. execute_trade with sqrt-impact tx_cost (spread_scale = -1.0)
 *   8. record_kelly_trade_outcome
 *   9. entry_price update on entry/reversal/flat
 *  10. hold_time tick via update_hold_time
 *  11. step_return = (new_value − old_value) / old_value (PURE per-bar P&L)
 *
 * Pass-by-pointer for in/out state. Capital-floor check is done by caller
 * BEFORE invoking this helper (different output buffer formats between
 * training/val make in-helper handling awkward).
 *
 * Saboteur and shaping are NOT applied here — those are training-only
 * concerns the caller layers on top by gating their own additional code
 * with exploration_scale / shaping_scale before/after this call.
 *
 * Both kernels MUST call this for their step_return computation. Any
 * change here propagates to both atomically — drift becomes structurally
 * impossible.
 * ══════════════════════════════════════════════════════════════════════ */
__device__ __forceinline__ void unified_env_step_core(
    /* ── Action + price (caller-provided) ── */
    int   action_val,
    float close,
    /* ── Branch sizes (4-branch action layout) ── */
    int   b1_size, int b2_size, int b3_size,
    /* ── Static config ──
     * TWO max-position params so training's Kelly-from-atoms variance sizing
     * can shrink the target without also tightening Kelly cap and tx_cost
     * physics (which must reflect the real broker's max contracts). Val
     * callers pass the same value for both. */
    float max_position_target,   /* compute_target_position scale — training shrinks by Var[Q] */
    float max_position_physics,  /* apply_kelly_cap + execute_trade scale — broker's real max */
    float tx_cost_multiplier,    /* training: includes saboteur/curriculum scaling; val: tx_cost_bps as-is */
    float spread_cost,           /* training: scaled by saboteur/curriculum; val: as-is */
    float contract_multiplier,
    float margin_pct,
    float health_for_kelly,      /* clamped to [0,1] before passing — usually isv_signals[12] */
    /* ── Regime-adaptive trailing stop scales (caller resolves from features) ──
     * Trail distance = 0.005f * vol_scale * trend_scale. Pass 1.0f/1.0f for
     * fixed-width behaviour. Caller formulas (kept symmetric train↔val):
     *   vol_scale   = min(2.5, 1 + max(|CUSUM| - 0.5, 0) * 2.0)
     *   trend_scale = min(2.5, 1 + max(ADX   - 0.25, 0) * 2.0)
     * with ADX, CUSUM read from features[bar * market_dim + 40/41]. */
    float vol_scale,
    float trend_scale,
    /* ── Policy conviction ∈ [0, 1] for adaptive Kelly warmup floor ──
     * Per-sample direction Q-spread normalised by q_dir_abs_ref (ISV[21]).
     * Drives `kelly_position_cap` warmup_floor — high conviction relaxes
     * the cold-start cap so the policy can size up; low conviction keeps
     * safety dominant. Only affects the cap when maturity < 1 (<10 trades);
     * once Kelly stats graduate the floor term drops out. Clamped [0,1]
     * inside kelly_position_cap; caller may pass out-of-range and the
     * helper bounds it. */
    float conviction,
    /* ── Magnitude-branch conviction ∈ [0, 1] for adaptive Kelly warmup floor ──
     * Per-sample magnitude Q-spread normalised by q_mag_abs_ref (ISV[16]).
     * Composed with direction conviction inside this helper as
     *   policy_conviction = max(direction_conviction, magnitude_conviction)
     * so a confident magnitude preference unblocks the Kelly cap when the
     * direction head is near-uniform (typical at smoke-horizon: direction
     * Q-values cluster while the magnitude head sharply prefers Full). Same
     * shape/clamp/semantics as `conviction` — caller may pass out-of-range
     * and the helper bounds it. NULL/zero from the caller is benign — the
     * max() composer falls back to the direction signal alone. */
    float magnitude_conviction,
    /* ── Per-thread mutable state (in/out) ── */
    float* position,
    float* cash,
    float* entry_price,
    float* hold_time,
    float* max_equity,
    float  prev_equity,          /* read-only snapshot for step_return denominator + trailing stop */
    /* ── Kelly stats (in/out) ── */
    float* win_count, float* loss_count,
    float* sum_wins,  float* sum_losses,
    /* ── Outputs ── */
    float* step_return_out,      /* pure per-bar P&L = (new_value - prev_equity) / prev_equity */
    float* new_value_out,        /* post-trade equity = cash + position * close */
    int*   prev_position_sign_out,  /* sign of position BEFORE the trade — for caller's reversal logic */
    int*   actual_dir_out,       /* post-enforcement direction (Hold / Long / Short / Flat) — for caller's actions_history */
    int*   actual_mag_out,       /* post-enforcement magnitude bucket (0/1/2 for Quarter/Half/Full; 0 if Hold/Flat) */
    int*   trail_triggered_out,  /* 1 if trailing stop forced position to 0 */
    /* ── SP5 Layer B: ISV signal bus for Pearl 6 + Pearl 8 ──
     * NULL-tolerant: when NULL, falls back to static defaults
     * (0.5f for kelly_f_smooth, 0.005f for trail distance). */
    const float* __restrict__ isv_signals_ptr  /* [SP5_SLOT_END] pinned ISV signals; NULL = static defaults */
) {
    /* ── Step 1: Decode action (4-branch) ── */
    int dir_idx        = decode_direction_4b(action_val, b1_size, b2_size, b3_size);
    int mag_idx        = decode_magnitude_4b(action_val, b1_size, b2_size, b3_size);
    int order_type_idx = decode_order_4b(action_val, b2_size, b3_size);

    int is_hold_action = (dir_idx == DIR_HOLD);

    /* ── Step 2-3: Target position (uses target scale, which training shrinks) ── */
    float target_position;
    if (is_hold_action) {
        target_position = *position;
    } else {
        target_position = compute_target_position_4branch(dir_idx, mag_idx, max_position_target);
    }

    /* ── Step 4: Margin cap ── */
    if (!is_hold_action) {
        float margin_per_contract = close * contract_multiplier * margin_pct;
        target_position = apply_margin_cap(
            target_position, prev_equity, margin_per_contract, *max_equity
        );
    }

    /* ── Step 5: Kelly cap (safety composed from three orthogonal signals
     * + conviction-adaptive warmup floor; uses physics scale).
     *
     * policy_conviction = max(direction_conviction, magnitude_conviction)
     * safety_multiplier = max(health_safety, policy_conviction)
     *   health_safety           = 0.5 + 0.5 × health  [training-stability floor]
     *   direction_conviction    ∈ [0, 1]              [direction-branch confidence]
     *   magnitude_conviction    ∈ [0, 1]              [magnitude-branch confidence]
     *
     * Health measures training stability (is the optimiser making coherent
     * progress?). Per-branch conviction measures the policy's confidence on
     * each axis of the factored action: direction conviction reflects WHERE
     * to trade, magnitude conviction reflects HOW MUCH. These are orthogonal —
     * the policy can be near-uniform on one axis while sharply preferring a
     * single bin on the other. The Kelly cap composes both per-branch signals
     * via max(): if EITHER axis expresses high confidence, the policy is
     * committed enough to its choice that the cap should relax. Health then
     * composes via max() too as a training-stability floor. Bounded to
     * [0.5, 1.0] by the health floor.
     *
     * All signals flow per-sample from the action_select kernel's per-branch
     * Q-spreads, normalised by ISV per-branch |Q|-scale EMAs (ISV[21] for
     * direction, ISV[16] for magnitude). Adaptive / temporal — no static knobs.
     *
     * Per project_magnitude_eval_collapse_kelly_capped.md: at the smoke-test
     * horizon, typical health=0.49 gives health_safety=0.745. The direction
     * branch is often near-uniform (q_dir_gap / q_dir_abs_ref ≈ 0.04 when
     * Hold competes with Long), so direction_conviction alone leaves the cap
     * formula at 0.5625 × max_pos — pinned to the Half bucket [0.375, 0.75).
     * The magnitude branch simultaneously expresses sharp preferences
     * (q_f / q_h ≈ 1.9 → mag_conviction ≈ 0.5). Threading mag_conviction
     * unblocks Full realisation for confident magnitude picks even when the
     * direction head hesitates, without requiring health graduation
     * (structurally unreachable in 20-epoch smokes). */
    if (!is_hold_action) {
        float h = fminf(1.0f, fmaxf(0.0f, health_for_kelly));
        float health_safety = 0.5f + 0.5f * h;
        float dir_conv_clamped = fminf(1.0f, fmaxf(0.0f, conviction));
        float mag_conv_clamped = fminf(1.0f, fmaxf(0.0f, magnitude_conviction));
        float policy_conviction = fmaxf(dir_conv_clamped, mag_conv_clamped);

        /* SP5 Layer B (Pearl 6): cross-fold-persistent smoothed Kelly fraction
         * from ISV[KELLY_F_SMOOTH_INDEX=280]. When ISV is available and the
         * slot has been populated (value > 0), incorporate it as an additional
         * floor on `health_safety` — both are fractions ∈ [0, 1] measuring
         * long-run sizing confidence. The max() is safe: when Kelly says
         * "go bigger" than health alone would permit, we trust the data. When
         * health says "be conservative", health dominates. Bootstrap-safe
         * because ISV[280] starts at 0 and Pearl 6 uses EWMA α=0.01 — it
         * stays near zero until real Kelly data accumulate.
         * EPS_CLAMP_FLOOR = 0.001f: prevents a very small cross-fold Kelly
         * from being amplified by the max() in a numerically unstable way
         * (same cold-start pattern as SP4 EPS_CLAMP_FLOOR). */
        float kelly_f_smooth = (isv_signals_ptr != NULL)
            ? fmaxf(0.001f, fminf(1.0f, isv_signals_ptr[SP5_KELLY_F_SMOOTH_INDEX]))
            : 0.5f;
        float health_safety_sp5 = fmaxf(health_safety, kelly_f_smooth);
        /* SP9 Fix 37: ISV-driven adaptive warmup_floor max'd into the
         * health_safety chain. Per `pearl_cold_start_exit_signal_or.md`:
         * the SP9 producer writes a value that decays as OR'd statistical /
         * behavioral / temporal confidence accumulates; the consumer here
         * lifts the cap above that floor regardless of measured Kelly_f.
         * Zero sentinel (fold boundary, pre-first-observation) → no-op via
         * fmaxf — the prior path's health_safety_sp5 dominates. */
        float kelly_warmup_floor_sp9 = (isv_signals_ptr != NULL)
            ? fmaxf(0.0f, fminf(1.0f, isv_signals_ptr[SP9_KELLY_WARMUP_FLOOR_INDEX]))
            : 0.0f;
        float health_safety_sp9 = fmaxf(health_safety_sp5, kelly_warmup_floor_sp9);
        float safety = fmaxf(health_safety_sp9, policy_conviction);
        target_position = apply_kelly_cap(
            target_position,
            *win_count, *loss_count,
            *sum_wins,  *sum_losses,
            max_position_physics, safety,
            policy_conviction,
            health_safety_sp9
        );
    }

    /* ── Step 6: Trailing stop ── */
    int trail_triggered = 0;
    if (!is_hold_action && fabsf(*position) > 0.001f) {
        float trade_ret = 0.0f;
        if (*entry_price > 0.0f && prev_equity > 1.0f) {
            float unrealized = (*position) * (close - *entry_price);
            trade_ret = unrealized / prev_equity;
        }
        /* SP5 Layer B (Pearl 8): per-direction trailing stop base distance from
         * ISV[TRAIL_DIST_PER_DIR_BASE + dir_idx] = ISV[270 + dir_idx].
         * dir_idx ∈ {0=Short,1=Hold,2=Long,3=Flat}; Hold/Flat are excluded by
         * the enclosing !is_hold_action guard so only Short (0) and Long (2)
         * reach here in practice. Bootstrap floor: when ISV is NULL or the
         * slot reads as 0 (pre-first-observation), fall back to the SP4
         * static value of 0.005f (0.5% trail). */
        float base_trail_dist = (isv_signals_ptr != NULL)
            ? fmaxf(0.001f, isv_signals_ptr[SP5_TRAIL_DIST_PER_DIR_BASE + dir_idx])
            : 0.005f;
        if (check_trailing_stop(
                *hold_time, *max_equity, prev_equity,
                trade_ret, base_trail_dist, vol_scale, trend_scale
            )) {
            target_position = 0.0f;
            trail_triggered = 1;
        }
    }
    *trail_triggered_out = trail_triggered;

    /* ── Step 7: Execute trade — spread_scale = -1.0 (sqrt impact, no override) ── */
    float prev_position = *position;
    int prev_sign = (prev_position > 0.001f) ? 1
                  : (prev_position < -0.001f) ? -1 : 0;
    *prev_position_sign_out = prev_sign;

    float tx_cost = execute_trade(
        position, cash,
        target_position, close,
        tx_cost_multiplier, spread_cost, max_position_physics,
        order_type_idx, -1.0f
    );
    (void)tx_cost;

    /* ── Step 8: Kelly stats update on trade exit/reversal ── */
    record_kelly_trade_outcome(
        prev_position, *position,
        *entry_price, close, prev_equity,
        win_count, loss_count, sum_wins, sum_losses
    );

    /* ── Step 9: entry_price update ── */
    if (fabsf(prev_position) < 0.001f && fabsf(*position) > 0.001f) {
        *entry_price = close;            /* new entry from flat */
    } else if (prev_position * (*position) < 0.0f) {
        *entry_price = close;            /* reversal */
    } else if (fabsf(*position) < 0.001f) {
        *entry_price = 0.0f;             /* went flat */
    }
    /* same-sign scaling: keep original entry_price */

    /* ── Step 10: hold_time tick ── */
    *hold_time = update_hold_time(prev_position, *position, *hold_time);

    /* ── Step 11: Mark-to-market + step_return ── */
    float new_value = (*cash) + (*position) * close;
    *new_value_out = new_value;
    *max_equity = fmaxf(*max_equity, new_value);

    *step_return_out = (prev_equity > 0.0f)
        ? (new_value - prev_equity) / prev_equity
        : 0.0f;

    /* ── Action re-encoding for actions_history (post-enforcement direction)
     * Caller stitches order_type/urgency back if they need the full action,
     * since the post-enforcement direction is the only thing this helper
     * can determine. */
    int actual_dir;
    if (is_hold_action && fabsf(*position - prev_position) < 0.001f) {
        actual_dir = 1;  /* Hold (model chose Hold and position unchanged) */
    } else {
        float pos_sign = *position;
        actual_dir = (pos_sign > 0.001f) ? 2
                   : (pos_sign < -0.001f) ? 0
                   : 3;  /* Flat */
    }
    float abs_pos = fabsf(*position) / fmaxf(max_position_physics, 0.01f);
    int actual_mag;
    if (abs_pos < 0.375f)      actual_mag = 0;  /* Quarter */
    else if (abs_pos < 0.75f)  actual_mag = 1;  /* Half */
    else                       actual_mag = 2;  /* Full */
    if (actual_dir == 1 || actual_dir == 3) actual_mag = 0;  /* Hold/Flat: mag irrelevant */

    *actual_dir_out = actual_dir;
    *actual_mag_out = actual_mag;
}

#endif /* TRADE_PHYSICS_CUH */
