// Vectorized backtest environment step kernel.
// One thread per walk-forward window. Each thread steps sequentially.
//
// Portfolio state layout per window [8 f32]:
//   [0] value        - current portfolio value
//   [1] position     - current position size (-1.0 to +1.0)
//   [2] cash         - cash balance
//   [3] entry_price  - entry price of current position (0 if flat)
//   [4] max_equity   - peak equity for drawdown tracking
//   [5] hold_time    - bars held in current position (0 = flat)
//   [6] cum_return   - cumulative log return
//   [7] step_count   - number of completed steps
//
// The 8-slot layout matches `SL_PORTFOLIO_BASE_DIM` — this buffer is read
// directly by `backtest_state_gather` (experience_kernels.cu) which indexes
// with stride 8. Changing this stride would corrupt window boundaries in
// the gather kernel.
//
// Kelly win/loss statistics used by the shared apply_kelly_cap helper
// live in a SEPARATE kelly_stats buffer of stride 4 per window:
//   [0] win_count, [1] loss_count, [2] sum_wins, [3] sum_losses
// This keeps the gather semantics clean (the network doesn't see Kelly
// stats) while letting the env kernel apply a Kelly-constrained
// position cap and record completed-trade outcomes via the shared
// record_kelly_trade_outcome helper in trade_physics.cuh.
//
// Trade physics (action decoding, tx costs, capital floor, Kelly cap,
// Kelly stat updates) are ALL shared with the training kernel via
// trade_physics.cuh — no duplicated step logic between the two envs.

#define PORTFOLIO_STATE_SIZE 8
#define KELLY_STATS_SIZE 4
#include "trade_physics.cuh"

// Capital floor breach: full episode restart for this window.
// cum_return and step_count are intentionally zeroed — the breached episode's
// metrics are captured in done_flags and step_returns before this reset.
__device__ void handle_capital_floor_breach(
    float* portfolio_state, int ps,
    float new_capital, float step_ret,
    float* step_rewards, float* step_returns, int* actions_history, int* done_flags,
    int w, int max_len, int current_step,
    int b0_size, int b1_size, int b2_size, int b3_size
) {
    step_rewards[w] = step_ret;
    step_returns[w * max_len + current_step] = step_ret;
    // Flat = direction=3, magnitude=0, order=any(0), urgency=any(0)
    // Flat action = 3 * b1*b2*b3 + 0 * b2*b3 + 0 * b3 + 0
    actions_history[w * max_len + current_step] = 3 * b1_size * b2_size * b3_size;
    // Full episode restart — metrics captured before this reset
    portfolio_state[ps + 0] = new_capital;
    portfolio_state[ps + 1] = 0.0f;
    portfolio_state[ps + 2] = new_capital;
    portfolio_state[ps + 3] = 0.0f;
    portfolio_state[ps + 4] = new_capital;
    portfolio_state[ps + 5] = 0.0f;
    portfolio_state[ps + 6] = 0.0f;
    portfolio_state[ps + 7] = 0.0f;
    done_flags[w] = 1;
}

extern "C" __global__ void backtest_env_step(
    // Market data (read-only, uploaded once)
    const float* __restrict__ prices,       // [n_windows * max_len * 4] (OHLC)
    const int*   __restrict__ window_lens,  // [n_windows]
    // Pre-uploaded feature tensor [n_windows * max_len * market_dim] — read-only.
    // Needed for regime-adaptive trailing stop (ADX at idx 40, CUSUM at idx 41).
    // Pass NULL (and market_dim <= 41) to fall back to fixed-width stop.
    const float* __restrict__ features,
    int          market_dim,

    // Actions from model for current step
    const int*   __restrict__ actions,      // [n_windows] factored action index

    // Portfolio state (read-write, persistent across steps)
    float*       portfolio_state,           // [n_windows * PORTFOLIO_STATE_SIZE]

    // Step outputs
    float*       step_rewards,              // [n_windows]
    float*       step_returns,              // [n_windows * max_len] (accumulated)
    int*         done_flags,                // [n_windows]
    int*         actions_history,           // [n_windows * max_len] (accumulated, for metrics)

    // Config — host scalars stay float
    int          n_windows,
    int          max_len,
    float        max_position,
    float        tx_cost_bps,
    float        spread_cost,
    int          current_step,
    int          b0_size,
    int          b1_size,
    int          b2_size,
    int          b3_size,
    float        holding_cost_rate,     // continuous inventory penalty per bar per unit position
    float        churn_threshold,      // bars threshold for churn penalty
    float        churn_penalty_scale,  // scale factor for churn penalty
    float        contract_multiplier,  // e.g. 50.0 for ES, 20.0 for NQ
    float        margin_pct,           // e.g. 0.06 (6% initial margin)
    float        opp_cost_scale,       // opportunity cost scale (0.0 = disabled during backtest)
    /* ── Kelly position cap (shared with training via trade_physics.cuh) ── */
    float*       kelly_stats,          // [n_windows * KELLY_STATS_SIZE=4]: win_count, loss_count, sum_wins, sum_losses
    const float* __restrict__ isv_signals,  // [13] pinned; [12]=health. NULL → 0.5 fallback.
    /* ── Phase 3 unified-env control scalars (signature alignment) ──
     * exploration_scale ∈ [0,1]: gates training-only stochastic perturbations
     *   (saboteur, plan_params, counterfactuals). NO-OP in backtest — val
     *   has no saboteur, no plan, no CF. Kept for ABI alignment with the
     *   future unified_env_step kernel.
     * shaping_scale ∈ [0,1]: gates additive reward-shaping bundles.
     *   NO-OP in backtest — step_returns are pure per-bar P&L (the
     *   inventory/churn/opp_cost penalties were removed in commit
     *   1ffdf38dd). Kept for ABI alignment.
     * Both pointers may be NULL — kernel ignores them (NO-OP today). */
    const float* __restrict__ exploration_scale_ptr,
    const float* __restrict__ shaping_scale_ptr,
    /* Per-sample conviction [n_windows] for adaptive Kelly warmup floor.
     * Written by experience_action_select (out_conviction). NULL → fallback
     * 1.0 (trust policy face-value). Threaded into unified_env_step_core. */
    const float* __restrict__ conviction_ptr,
    /* Per-sample magnitude-branch conviction [n_windows] — same shape as
     * conviction_ptr, parallel signal. Written by experience_action_select
     * (out_magnitude_conviction). Threaded into unified_env_step_core where
     * the cap formula composes max(dir_conv, mag_conv). NULL → fallback 0.0
     * (helper composes against direction conviction alone). */
    const float* __restrict__ magnitude_conviction_ptr,
    /* SP21 T2.2 Phase 1 (2026-05-10) — per-trade tape outputs.
     * Single thread per window writes at `[w * max_trades + count[w]]`
     * on close events; count[w] increments per trade. NULL-tolerant
     * for callers that don't need the tape (training-side launchers,
     * pre-Phase-1 callers preserve bit-identical behavior).
     * Layout `[n_windows * max_trades_per_window]` for the SoA buffers,
     * `[n_windows]` for the count.
     * Consumer: `enrichment::run_enrichments` reads per-trade Vec<EvalTrade>
     * from these buffers post-window for E1/E2/E3/E4/E5/E6/E7/E8. */
    float*        __restrict__ per_trade_pnl_out,           /* [n_windows * max_trades] or NULL */
    unsigned int* __restrict__ per_trade_holding_bars_out,  /* [n_windows * max_trades] or NULL */
    unsigned int* __restrict__ per_trade_bar_index_out,     /* [n_windows * max_trades] or NULL */
    unsigned int* __restrict__ per_trade_dir_mag_out,       /* [n_windows * max_trades] or NULL */
    unsigned int* __restrict__ per_trade_count,             /* [n_windows] or NULL */
    int           max_trades_per_window,
    /* SP21 T2.2 Phase 1.5 (2026-05-11) — entry_q tracking.
     * `entry_q_state` is per-window persistent storage: kernel writes the
     * predicted Q-value of the chosen action at trade open / reversal,
     * reads it back at trade close to emit on the per-trade tape (E1
     * Q-correction). `q_values_per_window` is the Q-VALUE buffer for the
     * CURRENT step ([n_windows * num_actions]); single-step kernel reads
     * Q[w, action_taken]. NULL-tolerant on both sides:
     *   - entry_q_state NULL → no entry_q snapshot, no per-trade emission
     *   - q_values_per_window NULL → no entry_q update on opens (state
     *     stays whatever it was; cold-start = 0.0 from buffer alloc)
     * The val pipeline (chunked path) supplies real buffers; non-DQN /
     * single-step `evaluate()` callers pass NULL since their forward_fn
     * closure exposes only action indices (not Q-values). Per-trade
     * `predicted_q` for trades closed by NULL callers is left at
     * buffer-init 0.0 — these tapes are not consumed by enrichment. */
    float*        __restrict__ entry_q_state,                 /* [n_windows] read-write or NULL */
    const float*  __restrict__ q_values_per_window,           /* [n_windows * num_actions] or NULL */
    int           num_actions,                                /* = b0 * b1 * b2 * b3 */
    float*        __restrict__ per_trade_predicted_q_out      /* [n_windows * max_trades] or NULL */
) {
    /* Both scales are no-ops in backtest today — ignored. The signature
     * alignment is so the future unified_env_kernel can be a drop-in
     * replacement for both backtest_env_step and experience_env_step
     * with the only behavioral difference being these two scalar values. */
    (void)exploration_scale_ptr;
    (void)shaping_scale_ptr;

    __shared__ float shmem_pf[256 * PORTFOLIO_STATE_SIZE];

    int w = blockIdx.x * blockDim.x + threadIdx.x;
    int local_tid = threadIdx.x;

    // Cooperative tile load
    int ps = w * PORTFOLIO_STATE_SIZE;
    if (w < n_windows) {
        for (int i = 0; i < PORTFOLIO_STATE_SIZE; i++) {
            shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + i] = portfolio_state[ps + i];
        }
    }
    __syncthreads();

    if (w >= n_windows) return;
    if (done_flags[w]) return;

    int wlen = window_lens[w];
    if (current_step >= wlen) {
        done_flags[w] = 1;
        return;
    }

    // Read current prices — native f32
    int price_base = (w * max_len + current_step) * 4;
    float open  = prices[price_base + 0];
    float high  = prices[price_base + 1];
    float low   = prices[price_base + 2];
    float close = prices[price_base + 3];

    // Read portfolio state from shared memory tile — native f32
    float value      = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 0];
    float position   = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 1];
    float cash       = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 2];
    float entry_price = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 3];
    float max_equity = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 4];
    float hold_time  = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 5];
    float cum_return = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 6];

    // Kelly stats from separate buffer (stride 4 per window)
    int ks = w * KELLY_STATS_SIZE;
    float win_count  = kelly_stats[ks + 0];
    float loss_count = kelly_stats[ks + 1];
    float sum_wins   = kelly_stats[ks + 2];
    float sum_losses = kelly_stats[ks + 3];

    // ── Capital floor circuit breaker (shared: trade_physics.cuh) ────────
    if (check_capital_floor(value, max_equity)) {
        // Force flat: close any open position at current price (notional model)
        if (fabsf(position) > 0.001f && close > 0.0f) {
            float exit_cost = compute_tx_cost(position, close,
                                              tx_cost_bps, spread_cost,
                                              max_position, 0, -1.0f);
            cash = cash + position * close;   // sell position at market (notional)
            cash = cash - exit_cost;
            position = 0.0f;
        }
        float liq_value = cash;
        float liq_ret = (value > 0.01f)
            ? (liq_value - value) / value
            : 0.0f;
        handle_capital_floor_breach(portfolio_state, ps, liq_value, liq_ret,
                                    step_rewards, step_returns, actions_history, done_flags,
                                    w, max_len, current_step, b0_size, b1_size, b2_size, b3_size);
        return;
    }

    /* Suppress unused HLOC: kernel only uses close (notional MTM model). */
    (void)open; (void)high; (void)low;

    /* ── UNIFIED CORE STEP (shared: trade_physics.cuh::unified_env_step_core) ──
     * One source of truth for the canonical env step. The training kernel
     * (experience_env_step) calls the same helper for the same computation,
     * guaranteeing step_return parity at the source level. */
    int   action_val = actions[w];
    int   prev_position_sign = 0;
    int   actual_dir = -1, actual_mag = 0;
    int   trail_triggered = 0;
    float prev_position = position;
    float new_value     = 0.0f;
    float step_ret_core = 0.0f;
    float health_k      = (isv_signals != NULL) ? isv_signals[12] : 0.5f;

    /* Regime-adaptive trailing-stop scales — same helper as training so the
     * stop fires on identical ADX/CUSUM thresholds. features buffer is
     * [n_windows, max_len, market_dim] so the flat bar index is w*max_len+step. */
    float trail_vol_scale   = 1.0f;
    float trail_trend_scale = 1.0f;
    compute_regime_trail_scales(
        features, (long long)w * max_len + current_step, market_dim,
        &trail_vol_scale, &trail_trend_scale
    );

    /* Per-sample conviction for Kelly warmup floor — matches training.
     * NULL conviction_ptr defaults to 1.0 (trust policy face-value, same as
     * pre-mag-conviction behaviour). NULL magnitude_conviction_ptr defaults
     * to 0.0 — the helper's max(dir_conv, mag_conv) then falls back to the
     * direction signal alone, preserving prior single-conviction behaviour. */
    float conviction_core = (conviction_ptr != NULL)
        ? fminf(1.0f, fmaxf(0.0f, conviction_ptr[w]))
        : 1.0f;
    float mag_conviction_core = (magnitude_conviction_ptr != NULL)
        ? fminf(1.0f, fmaxf(0.0f, magnitude_conviction_ptr[w]))
        : 0.0f;

    /* SP21 T2.2 Phase 1 (2026-05-10): snapshot entry_price + hold_time
     * BEFORE the call. unified_env_step_core mutates these in place
     * (zeros entry_price on close, resets hold_time to 0 / increments
     * by 1 on hold continuation). The post-call close-detection
     * predicate below mirrors `record_kelly_trade_outcome` (in
     * `trade_physics.cuh`) — single source of truth for "what is a
     * trade close" — and uses these snapshots to recompute the
     * realized return for the per-trade tape emission.
     * Snapshot is bit-stable: same kernel thread, same memory, no race. */
    const float pre_entry_price = entry_price;
    const int   pre_hold_bars   = (int)hold_time;
    /* SP21 T2.2 Phase 1.5 (2026-05-11): snapshot entry_q from per-window
     * persistent state. Set on previous open/reversal; read here at
     * potential close. NULL-tolerant — falls back to 0.0 when the caller
     * doesn't supply the buffer (training-side launchers, smoke tests). */
    const float pre_entry_q     = (entry_q_state != NULL) ? entry_q_state[w] : 0.0f;

    unified_env_step_core(
        action_val, close,
        b1_size, b2_size, b3_size,
        max_position, max_position, tx_cost_bps, spread_cost,
        contract_multiplier, margin_pct,
        health_k,
        trail_vol_scale, trail_trend_scale,
        conviction_core,             /* direction-branch Kelly warmup floor */
        mag_conviction_core,         /* magnitude-branch Kelly warmup floor */
        &position, &cash, &entry_price, &hold_time,
        &max_equity,
        value,                       /* prev_equity (snapshot) */
        &win_count, &loss_count, &sum_wins, &sum_losses,
        &step_ret_core,              /* pure P&L (caller may add shaping) */
        &new_value,
        &prev_position_sign,
        &actual_dir,
        &actual_mag,
        &trail_triggered,
        isv_signals              /* SP5 Layer B: Pearl 6 (kelly_f_smooth) + Pearl 8 (trail_dist_per_dir) */
    );
    (void)prev_position; (void)prev_position_sign; (void)trail_triggered;

    /* SP21 T2.2 Phase 1 (2026-05-10) — per-trade tape emission.
     * Mirrors `record_kelly_trade_outcome`'s close predicate
     * (`is_exit || is_reversal` with positive entry_price + valid
     * pre-trade equity guards). NULL-tolerant — when buffers aren't
     * wired (training-side launchers, pre-Phase-1 callers), the emit
     * is a no-op. Single thread per window writes its own slot ⇒
     * race-free counter increment without atomicAdd, per
     * `feedback_no_atomicadd`. */
    if (per_trade_count != NULL) {
        const int prev_pos_active = (fabsf(prev_position) > 0.001f);
        const int is_exit         = prev_pos_active && (fabsf(position) < 0.001f);
        const int is_reversal     = (prev_position * position < 0.0f);
        if ((is_exit || is_reversal) && pre_entry_price > 0.0f && value > 1.0f) {
            const float realized     = prev_position * (close - pre_entry_price);
            const float trade_return = realized / value;
            const unsigned int slot  = per_trade_count[w];
            if ((int)slot < max_trades_per_window) {
                const unsigned int idx =
                    (unsigned int)w * (unsigned int)max_trades_per_window + slot;
                per_trade_pnl_out[idx]          = trade_return;
                per_trade_holding_bars_out[idx] = (unsigned int)pre_hold_bars;
                per_trade_bar_index_out[idx]    = (unsigned int)current_step;
                /* dir_mag packed: high 16 = direction (Short=0,Hold=1,Long=2,Flat=3),
                 * low 16 = magnitude (Quarter=0, Half=1, Full=2). */
                per_trade_dir_mag_out[idx] =
                    ((unsigned int)actual_dir << 16) |
                    ((unsigned int)actual_mag & 0xFFFFu);
                /* SP21 T2.2 Phase 1.5 (2026-05-11): predicted_q at trade
                 * open. NULL-tolerant — when caller doesn't wire the
                 * buffer, the field stays at buffer-init 0.0. */
                if (per_trade_predicted_q_out != NULL) {
                    per_trade_predicted_q_out[idx] = pre_entry_q;
                }
                per_trade_count[w] = slot + 1u;
            }
        }
    }

    // ── Post-trade floor check: catch intra-step breaches ────────────
    // The pre-trade check (top of kernel) catches breaches from the
    // PREVIOUS step. This catches breaches from THIS step's price move
    // or transaction costs. Without it, max_dd overshoots the 75% floor
    // because the metrics kernel records the breached value before the
    // next step's pre-trade check fires.
    if (check_capital_floor(new_value, max_equity)) {
        // Emergency liquidation: close position at current price (notional model)
        if (fabsf(position) > 0.001f) {
            float f_exit_cost = compute_tx_cost(position, close,
                                                tx_cost_bps, spread_cost,
                                                max_position, 0, -1.0f);
            cash = cash + position * close;   // sell position at market (notional)
            cash = cash - f_exit_cost;
            position = 0.0f;
            new_value = cash;           // flat: equity = cash
            entry_price = 0.0f;
        }
        float step_ret = (value > 0.01f)
            ? (new_value - value) / value
            : 0.0f;
        handle_capital_floor_breach(portfolio_state, ps, new_value, step_ret,
                                    step_rewards, step_returns, actions_history, done_flags,
                                    w, max_len, current_step, b0_size, b1_size, b2_size, b3_size);
        return;
    }

    /* step_return is PURE per-bar P&L computed by unified_env_step_core
     * (Phase 3 unified-env policy). Shaping args kept for ABI stability
     * but ignored. */
    float step_ret = step_ret_core;
    (void)holding_cost_rate;
    (void)churn_threshold;
    (void)churn_penalty_scale;
    (void)opp_cost_scale;

    float new_cum_return = cum_return + step_ret;
    float new_max = max_equity;  /* already updated inside unified_env_step_core */

    /* SP21 T2.2 Phase 1.5 (2026-05-11) — entry_q capture on fresh open or
     * reversal. unified_env_step_core has executed; position now reflects
     * the post-step state. A reversal closed the prior trade and opened
     * a new one in the opposite direction — the per-trade tape block
     * above already emitted the close (with the old pre_entry_q); now
     * we capture the NEW position's entry_q for the next close emission.
     * Both buffers must be wired (real q_values + writable state) and
     * the action must be in-range. */
    if (entry_q_state != NULL && q_values_per_window != NULL && num_actions > 0) {
        const int prev_pos_active2 = (fabsf(prev_position) > 0.001f);
        const int now_pos_active   = (fabsf(position)      > 0.001f);
        const int is_open          = (!prev_pos_active2) && now_pos_active;
        const int is_reverse_open  = (prev_position * position < 0.0f);
        if ((is_open || is_reverse_open)
            && action_val >= 0 && action_val < num_actions) {
            entry_q_state[w] = q_values_per_window[
                (long long)w * (long long)num_actions + (long long)action_val];
        }
    }

    // Write portfolio state
    portfolio_state[ps + 0] = new_value;
    portfolio_state[ps + 1] = position;
    portfolio_state[ps + 2] = cash;
    portfolio_state[ps + 3] = entry_price;
    portfolio_state[ps + 4] = new_max;
    portfolio_state[ps + 5] = hold_time;
    portfolio_state[ps + 6] = new_cum_return;
    portfolio_state[ps + 7] = portfolio_state[ps + 7] + 1.0f;

    // Kelly stats — separate buffer, updated by record_kelly_trade_outcome
    kelly_stats[ks + 0] = win_count;
    kelly_stats[ks + 1] = loss_count;
    kelly_stats[ks + 2] = sum_wins;
    kelly_stats[ks + 3] = sum_losses;

    // Outputs
    step_rewards[w] = step_ret;
    step_returns[w * max_len + current_step] = step_ret;
    /* Record ACTUAL action taken (post-enforcement, post-trailing-stop).
     * actual_dir + actual_mag already computed by unified_env_step_core;
     * we just stitch back the caller's original order/urgency. */
    {
        int orig_order   = decode_order_4b(action_val, b2_size, b3_size);
        int orig_urgency = decode_urgency_4b(action_val, b3_size);
        int actual_action = actual_dir * b1_size * b2_size * b3_size
                          + actual_mag * b2_size * b3_size
                          + orig_order * b3_size
                          + orig_urgency;
        actions_history[w * max_len + current_step] = actual_action;
    }
}

/* ══════════════════════════════════════════════════════════════════════
 * BATCHED ENV STEP — process chunk_len steps in a single launch.
 *
 * Eliminates chunk_len individual kernel launches per chunk. Each thread
 * handles one window and loops over all steps in the chunk sequentially.
 * Actions come from the chunked_actions buffer (already computed by the
 * model forward + action_select for the entire chunk).
 *
 * Launch: grid=(ceil(n_windows/256)), block=(256).
 * ══════════════════════════════════════════════════════════════════════ */
extern "C" __global__ void backtest_env_step_batch(
    const float* __restrict__ prices,
    const int*   __restrict__ window_lens,
    const float* __restrict__ features,    /* [n_windows, max_len, market_dim] or NULL */
    int          market_dim,               /* <=41 disables regime-adaptive trail */
    const int*   __restrict__ chunked_actions,  /* [chunk_len * n_windows] */
    float*       portfolio_state,
    float*       step_rewards,
    float*       step_returns,
    int*         done_flags,
    int*         actions_history,
    int          n_windows,
    int          max_len,
    float        max_position,
    float        tx_cost_bps,
    float        spread_cost,
    int          start_step,          /* first step index in this chunk */
    int          chunk_len,           /* number of steps to process */
    int          b0_size,
    int          b1_size,
    int          b2_size,
    int          b3_size,
    float        holding_cost_rate,
    float        churn_threshold,
    float        churn_penalty_scale,
    float        contract_multiplier,
    float        margin_pct,
    float        opp_cost_scale,       /* opportunity cost scale (0.0 = disabled during backtest) */
    /* ── Kelly position cap (shared with training via trade_physics.cuh) ── */
    float*       kelly_stats,          /* [n_windows * KELLY_STATS_SIZE=4] win_count, loss_count, sum_wins, sum_losses */
    const float* __restrict__ isv_signals,  /* [13] pinned; [12]=health. NULL → 0.5 fallback. */
    /* Phase 3 unified-env scalars — see single-step variant for rationale.
     * NO-OP in backtest; kept for ABI alignment with future unified kernel. */
    const float* __restrict__ exploration_scale_ptr,
    const float* __restrict__ shaping_scale_ptr,
    /* Per-sample conviction [chunk_len * n_windows] — layout matches
     * chunked_actions ([step_offset * n_windows + window]). NULL → fallback
     * 1.0. Threaded into unified_env_step_core for adaptive Kelly warmup. */
    const float* __restrict__ conviction_ptr,
    /* Per-sample magnitude-branch conviction — same shape and layout as
     * conviction_ptr. NULL → fallback 0.0; the helper's max(dir_conv,
     * mag_conv) composer falls back to the direction signal alone. */
    const float* __restrict__ magnitude_conviction_ptr,
    /* SP21 T2.2 Phase 1 (2026-05-10) — per-trade tape outputs.
     * See single-step variant header for full rationale + layout.
     * NULL-tolerant; identical semantics. */
    float*        __restrict__ per_trade_pnl_out,           /* [n_windows * max_trades] or NULL */
    unsigned int* __restrict__ per_trade_holding_bars_out,  /* [n_windows * max_trades] or NULL */
    unsigned int* __restrict__ per_trade_bar_index_out,     /* [n_windows * max_trades] or NULL */
    unsigned int* __restrict__ per_trade_dir_mag_out,       /* [n_windows * max_trades] or NULL */
    unsigned int* __restrict__ per_trade_count,             /* [n_windows] or NULL */
    int           max_trades_per_window,
    /* SP21 T2.2 Phase 1.5 (2026-05-11) — entry_q tracking. See single-step
     * variant header for full rationale. Layout difference: in the batch
     * kernel `q_values_per_window` is `[chunk_len * n_windows * num_actions]`
     * — per-step Q-values produced once per chunk by the cuBLAS forward,
     * indexed `q[s, w, a] = q_values_per_window[s * n_windows * num_actions
     * + w * num_actions + a]`. `entry_q_state` is per-window persistent
     * state (single value per window), unchanged shape vs single-step. */
    float*        __restrict__ entry_q_state,                /* [n_windows] read-write or NULL */
    const float*  __restrict__ q_values_per_window,          /* [chunk_len * n_windows * num_actions] or NULL */
    int           num_actions,                               /* = b0 * b1 * b2 * b3 */
    float*        __restrict__ per_trade_predicted_q_out     /* [n_windows * max_trades] or NULL */
) {
    /* Phase 3 unified-env: shaping params retained for ABI but no longer
     * applied to step_returns. See single-step variant header for rationale. */
    (void)holding_cost_rate;
    (void)churn_threshold;
    (void)churn_penalty_scale;
    (void)opp_cost_scale;
    (void)exploration_scale_ptr;
    (void)shaping_scale_ptr;

    int w = blockIdx.x * blockDim.x + threadIdx.x;
    if (w >= n_windows) return;

    int wlen = window_lens[w];
    int ps = w * PORTFOLIO_STATE_SIZE;
    int ks = w * KELLY_STATS_SIZE;

    /* Load portfolio into registers — persistent across the step loop */
    float value      = portfolio_state[ps + 0];
    float position   = portfolio_state[ps + 1];
    float cash       = portfolio_state[ps + 2];
    float entry_price = portfolio_state[ps + 3];
    float max_equity = portfolio_state[ps + 4];
    float hold_time  = portfolio_state[ps + 5];
    float cum_return = portfolio_state[ps + 6];

    /* Kelly stats from dedicated buffer — kept in registers, updated on trade
     * exits below via record_kelly_trade_outcome (shared: trade_physics.cuh). */
    float win_count  = kelly_stats[ks + 0];
    float loss_count = kelly_stats[ks + 1];
    float sum_wins   = kelly_stats[ks + 2];
    float sum_losses = kelly_stats[ks + 3];

    int steps_processed = 0;

    for (int s = 0; s < chunk_len; s++) {
        int current_step = start_step + s;
        if (done_flags[w] || current_step >= wlen) break;

        /* Read prices for this step */
        int price_base = (w * max_len + current_step) * 4;
        float open  = prices[price_base + 0];
        float high  = prices[price_base + 1];
        float low   = prices[price_base + 2];
        float close = prices[price_base + 3];

        /* ── Capital floor circuit breaker (pre-trade) ────────────────── */
        /* Bug 1 fix: use close > 0.0f (matches original) */
        if (check_capital_floor(value, max_equity)) {
            if (fabsf(position) > 0.001f && close > 0.0f) {
                float exit_cost = compute_tx_cost(position, close,
                    tx_cost_bps, spread_cost, max_position, 0, -1.0f);
                cash = cash + position * close;
                cash = cash - exit_cost;
                position = 0.0f;
            }
            /* Bug 2 fix: compute liq_ret properly (matches original) */
            float liq_value = cash;
            float liq_ret = (value > 0.01f)
                ? (liq_value - value) / value
                : 0.0f;
            handle_capital_floor_breach(portfolio_state, ps, liq_value, liq_ret,
                step_rewards, step_returns, actions_history, done_flags,
                w, max_len, current_step, b0_size, b1_size, b2_size, b3_size);
            /* Bug 9 fix: return instead of break to avoid tail write */
            return;
        }

        /* Suppress unused HLOC: notional MTM uses close only. */
        (void)open; (void)high; (void)low;

        /* Read action from chunked buffer: [step_offset, window] layout */
        int action_val = chunked_actions[s * n_windows + w];

        /* ── UNIFIED CORE STEP (shared: trade_physics.cuh) ──
         * One source of truth — same helper called by backtest_env_step
         * (single-step variant) and, in a follow-up commit, by
         * experience_env_step (training). Guarantees step_return parity. */
        int   prev_position_sign = 0;
        int   actual_dir         = -1;
        int   actual_mag         = 0;
        int   trail_triggered    = 0;
        float prev_position      = position;
        float new_value          = 0.0f;
        float step_ret_core      = 0.0f;
        float health_k_batch     = (isv_signals != NULL) ? isv_signals[12] : 0.5f;

        /* Regime-adaptive trailing stop — shared helper keeps formulas
         * identical to the single-step kernel and the training kernel. */
        float trail_vol_scale_b   = 1.0f;
        float trail_trend_scale_b = 1.0f;
        compute_regime_trail_scales(
            features, (long long)w * max_len + current_step, market_dim,
            &trail_vol_scale_b, &trail_trend_scale_b
        );

        /* Per-sample conviction for adaptive Kelly warmup floor. Layout
         * matches chunked_actions: [s * n_windows + w]. NULL magnitude_
         * conviction_ptr defaults to 0.0 — the helper's max(dir, mag) composer
         * falls back to the direction signal alone, preserving prior behaviour. */
        float conviction_core = (conviction_ptr != NULL)
            ? fminf(1.0f, fmaxf(0.0f, conviction_ptr[s * n_windows + w]))
            : 1.0f;
        float mag_conviction_core = (magnitude_conviction_ptr != NULL)
            ? fminf(1.0f, fmaxf(0.0f, magnitude_conviction_ptr[s * n_windows + w]))
            : 0.0f;

        /* SP21 T2.2 Phase 1 (2026-05-10): snapshot entry_price + hold_time
         * BEFORE the unified step call — see single-step variant for the
         * full rationale. The post-call close-detection predicate below
         * mirrors `record_kelly_trade_outcome`. */
        const float pre_entry_price_b = entry_price;
        const int   pre_hold_bars_b   = (int)hold_time;
        /* SP21 T2.2 Phase 1.5 (2026-05-11): snapshot entry_q from per-
         * window persistent state for potential close emission. */
        const float pre_entry_q_b     = (entry_q_state != NULL) ? entry_q_state[w] : 0.0f;

        unified_env_step_core(
            action_val, close,
            b1_size, b2_size, b3_size,
            max_position, max_position, tx_cost_bps, spread_cost,
            contract_multiplier, margin_pct,
            health_k_batch,
            trail_vol_scale_b, trail_trend_scale_b,
            conviction_core,                /* direction-branch Kelly warmup floor */
            mag_conviction_core,            /* magnitude-branch Kelly warmup floor */
            &position, &cash, &entry_price, &hold_time,
            &max_equity,
            value,                          /* prev_equity (snapshot for step_return) */
            &win_count, &loss_count, &sum_wins, &sum_losses,
            &step_ret_core,                 /* PURE per-bar P&L */
            &new_value,
            &prev_position_sign,
            &actual_dir,
            &actual_mag,
            &trail_triggered,
            isv_signals              /* SP5 Layer B: Pearl 6 (kelly_f_smooth) + Pearl 8 (trail_dist_per_dir) */
        );
        (void)prev_position; (void)prev_position_sign; (void)trail_triggered;

        /* SP21 T2.2 Phase 1 (2026-05-10) — per-trade tape emission.
         * Same predicate as single-step variant; fires inside the
         * batched step loop so each step's close emits within the
         * same kernel call. NULL-tolerant. */
        if (per_trade_count != NULL) {
            const int prev_pos_active_b = (fabsf(prev_position) > 0.001f);
            const int is_exit_b         = prev_pos_active_b && (fabsf(position) < 0.001f);
            const int is_reversal_b     = (prev_position * position < 0.0f);
            if ((is_exit_b || is_reversal_b)
                && pre_entry_price_b > 0.0f
                && value > 1.0f) {
                const float realized_b     = prev_position * (close - pre_entry_price_b);
                const float trade_return_b = realized_b / value;
                const unsigned int slot_b  = per_trade_count[w];
                if ((int)slot_b < max_trades_per_window) {
                    const unsigned int idx_b =
                        (unsigned int)w * (unsigned int)max_trades_per_window + slot_b;
                    per_trade_pnl_out[idx_b]          = trade_return_b;
                    per_trade_holding_bars_out[idx_b] = (unsigned int)pre_hold_bars_b;
                    per_trade_bar_index_out[idx_b]    = (unsigned int)current_step;
                    per_trade_dir_mag_out[idx_b] =
                        ((unsigned int)actual_dir << 16) |
                        ((unsigned int)actual_mag & 0xFFFFu);
                    /* SP21 T2.2 Phase 1.5 (2026-05-11): predicted_q at trade
                     * open. NULL-tolerant. */
                    if (per_trade_predicted_q_out != NULL) {
                        per_trade_predicted_q_out[idx_b] = pre_entry_q_b;
                    }
                    per_trade_count[w] = slot_b + 1u;
                }
            }
        }

        /* ── Post-trade floor check: catch intra-step breaches ────────── */
        if (check_capital_floor(new_value, max_equity)) {
            if (fabsf(position) > 0.001f) {
                float f_exit_cost = compute_tx_cost(position, close,
                    tx_cost_bps, spread_cost, max_position, 0, -1.0f);
                cash = cash + position * close;
                cash = cash - f_exit_cost;
                position = 0.0f;
                new_value = cash;
                entry_price = 0.0f;
            }
            float step_ret = (value > 0.01f)
                ? (new_value - value) / value : 0.0f;
            handle_capital_floor_breach(portfolio_state, ps, new_value, step_ret,
                step_rewards, step_returns, actions_history, done_flags,
                w, max_len, current_step, b0_size, b1_size, b2_size, b3_size);
            return; /* Exit kernel — episode over */
        }

        /* step_return is PURE per-bar P&L from unified_env_step_core. */
        float step_ret = step_ret_core;

        /* SP21 T2.2 Phase 1.5 (2026-05-11) — entry_q capture on fresh open
         * or reversal. Mirrors the single-step kernel; the only difference
         * is the chunked Q-values index `[s, w, a]`. Placed AFTER the
         * post-trade breach check so we don't capture entry_q for a
         * trade that gets force-closed by the same step. */
        if (entry_q_state != NULL && q_values_per_window != NULL && num_actions > 0) {
            const int prev_pos_active2 = (fabsf(prev_position) > 0.001f);
            const int now_pos_active   = (fabsf(position)      > 0.001f);
            const int is_open          = (!prev_pos_active2) && now_pos_active;
            const int is_reverse_open  = (prev_position * position < 0.0f);
            if ((is_open || is_reverse_open)
                && action_val >= 0 && action_val < num_actions) {
                entry_q_state[w] = q_values_per_window[
                    (long long)s * (long long)n_windows * (long long)num_actions
                    + (long long)w * (long long)num_actions
                    + (long long)action_val];
            }
        }

        cum_return = cum_return + step_ret;
        value = new_value;   /* max_equity already updated inside helper */

        steps_processed++;

        /* Outputs */
        step_rewards[w] = step_ret;
        step_returns[w * max_len + current_step] = step_ret;

        /* Record actual action (post-enforcement): helper computed actual_dir
         * and actual_mag; stitch back the caller's order + urgency. */
        {
            int orig_order   = decode_order_4b(action_val, b2_size, b3_size);
            int orig_urgency = decode_urgency_4b(action_val, b3_size);
            int actual_action = actual_dir * b1_size * b2_size * b3_size
                              + actual_mag * b2_size * b3_size
                              + orig_order * b3_size
                              + orig_urgency;
            actions_history[w * max_len + current_step] = actual_action;
        }
    } /* end step loop */

    /* Write portfolio state back — but NOT if a floor breach already wrote it */
    /* Bug 8 fix: use steps_processed instead of chunk_len */
    if (!done_flags[w]) {
        portfolio_state[ps + 0] = value;
        portfolio_state[ps + 1] = position;
        portfolio_state[ps + 2] = cash;
        portfolio_state[ps + 3] = entry_price;
        portfolio_state[ps + 4] = max_equity;
        portfolio_state[ps + 5] = hold_time;
        portfolio_state[ps + 6] = cum_return;
        portfolio_state[ps + 7] = portfolio_state[ps + 7] + (float)steps_processed;
    }
    /* Kelly stats — always written back (accumulate across episodes, even
     * if a capital-floor breach reset portfolio_state). */
    kelly_stats[ks + 0] = win_count;
    kelly_stats[ks + 1] = loss_count;
    kelly_stats[ks + 2] = sum_wins;
    kelly_stats[ks + 3] = sum_losses;
}
