// Per-window metrics reduction kernel — native F32.
#include "state_layout.cuh"
// One block per window. Threads cooperate to reduce step_returns.
//
// SP15 Phase 1.1.b (2026-05-06): sharpe split out into the dedicated
// `sharpe_per_bar_kernel` (launched separately by gpu_backtest_evaluator).
// The host applies annualization to the new launcher's per-bar output and
// stores it into `WindowMetrics.sharpe`. Output stride here drops from
// 14 → 13 floats per window; old slot 0 (sharpe) is removed and every
// remaining metric shifts down by one.
//
// 2026-05-09 — per-regime WR instrumentation (audit follow-up). Output
// stride grows 13 → 19 floats per window: 6 new slots [13..19) carrying
// per-regime win-rate + trade-count for the {Trending, Ranging, Volatile}
// bucketing introduced by `gpu_walk_forward.rs::classify_regime_from_features`.
// Regime is read from feature[40] (ADX-norm, post-`RollingPercentileRank`)
// per state_layout.cuh's market block; bucket thresholds are empirically
// grounded structural defaults (NOT tuned constants) — Trending: ADX>0.4,
// Ranging: ADX<0.2, Volatile otherwise. Each trade is attributed to the
// regime at its OPENING bar; boundary-stitched trades that span chunks
// inherit the regime from the chunk where they first appeared (block
// boundary buffers carry the open-bar regime alongside the existing
// prefix/suffix pair).
//
// Output per window [19 f32 values written to metrics_out]:
//   [0] total_pnl (multiplicative compounding)
//   [1] max_drawdown (exact sequential scan)
//   [2] sortino_ratio
//   [3] win_rate
//   [4] total_trades
//   [5] var_95
//   [6] cvar_95
//   [7] calmar_ratio
//   [8] omega_ratio
//   [9] buy_count
//  [10] sell_count
//  [11] hold_count
//  [12] unique_action_count
//  [13] win_rate_trending  (regime T: ADX > 0.4)
//  [14] trade_count_trending
//  [15] win_rate_ranging   (regime R: ADX < 0.2)
//  [16] trade_count_ranging
//  [17] win_rate_volatile  (regime V: otherwise)
//  [18] trade_count_volatile

// Per-regime ADX thresholds. Empirically-grounded structural defaults
// (post-`RollingPercentileRank` normalisation, where ADX-norm ∈ [0,1]):
//   - ADX > 0.4 ⇒ Trending  (raw ADX > 25 in the upper 60th percentile)
//   - ADX < 0.2 ⇒ Ranging   (raw ADX < 10 in the lower 20th percentile)
//   - else      ⇒ Volatile  (mid-ADX, often choppy)
// Per `feedback_isv_for_adaptive_bounds` these are conceptually adaptive
// (could derive from observed distribution). They live as kernel-local
// constants only because per-regime WR is observability-only (does not
// feed any controller); if a controller ever consumes a per-regime WR
// signal, lift the thresholds to ISV slots.
#define REGIME_TRENDING_THRESHOLD 0.4f
#define REGIME_RANGING_THRESHOLD  0.2f
// Feature index for ADX-norm in the 42-dim market feature block. Mirrors
// `gpu_walk_forward.rs::classify_regime_from_features` (uses index 40)
// and `state_layout.cuh`'s SL_MARKET block — slot 40 is the post-rolling-
// percentile-ranked ADX, ∈ [0,1].
#define REGIME_ADX_FEATURE_IDX    40

// Regime bucket labels. -1 sentinel = "no data" (chunk-empty / boundary).
#define REGIME_BUCKET_T  0
#define REGIME_BUCKET_R  1
#define REGIME_BUCKET_V  2
#define REGIME_BUCKET_NONE -1

// Classify a bar's regime from its feature vector at `bar_idx`. Returns
// REGIME_BUCKET_T / R / V (never NONE — that sentinel is reserved for
// "no observation" in boundary buffers).
__device__ __forceinline__ int classify_bar_regime(
    const float* __restrict__ features,
    int bar_idx,
    int feature_dim
) {
    float adx = features[bar_idx * feature_dim + REGIME_ADX_FEATURE_IDX];
    if (adx > REGIME_TRENDING_THRESHOLD) return REGIME_BUCKET_T;
    if (adx < REGIME_RANGING_THRESHOLD)  return REGIME_BUCKET_R;
    return REGIME_BUCKET_V;
}

extern "C" __global__ void compute_backtest_metrics(
    const float* __restrict__ step_returns,   // [n_windows * max_len]
    const float* __restrict__ portfolio_state, // [n_windows * 8]
    const int*   __restrict__ window_lens,    // [n_windows]
    const int*   __restrict__ actions_history, // [n_windows * max_len]
    const float* __restrict__ features,       // [n_windows * max_len * feature_dim]
    int          feature_dim,
    float* metrics_out,                       // [n_windows * 19]
    int          n_windows,
    int          max_len,
    float        annualization_factor,        // f32 scalar (sortino × factor; calmar × factor²)
    int          num_actions,
    int          order_actions,
    int          urgency_actions
) {
    int w = blockIdx.x;
    if (w >= n_windows) return;

    int wlen = window_lens[w];
    int tid = threadIdx.x;
    int stride = blockDim.x;
    int base = w * max_len;
    // Per-window features base offset. Each row at `bar_idx` lives at
    // `features[feat_base + bar_idx * feature_dim + k]`.
    int feat_base = w * max_len * feature_dim;

    // Shared memory layout (5 reduction arrays + 6 per-regime arrays +
    // sort scratch). SP15 Phase 1.1.b removed s_sq_sum (was sharpe-only)
    // — variance/std are now computed by the dedicated
    // `sharpe_per_bar_kernel`. The remaining metrics use s_sum (mean for
    // sortino + calmar daily_mean), s_down_sq (sortino's down_std), and
    // the action-distribution counters (buys/sells/holds).
    //
    // 2026-05-09 — per-regime WR adds 6 reduction arrays for internal
    // (chunk-local) per-regime trade/win counters. Boundary-stitched
    // trades are attributed by thread 0 using the open-regime carried
    // through `s_bnd[7..9*stride]`.
    extern __shared__ float shmem[];
    float* s_sum     = shmem;
    float* s_down_sq = shmem + stride;
    float* s_buys    = shmem + 2*stride;
    float* s_sells   = shmem + 3*stride;
    float* s_holds   = shmem + 4*stride;
    float* s_trades_T = shmem + 5*stride;
    float* s_wins_T   = shmem + 6*stride;
    float* s_trades_R = shmem + 7*stride;
    float* s_wins_R   = shmem + 8*stride;
    float* s_trades_V = shmem + 9*stride;
    float* s_wins_V   = shmem + 10*stride;
    float* s_sorted  = shmem + 11*stride;    // [4096] for sort + boundary data

    // Per-thread local accumulators — native F32
    float local_sum = 0.0f;
    float local_down = 0.0f;
    float local_cum = 1.0f;
    int local_buys = 0, local_sells = 0, local_holds = 0;
    int local_action_mask = 0;

    // Per-regime trade counters — populated for INTERNAL complete trades
    // (i.e., trades whose open AND close both fall within this thread's
    // chunk). Boundary-stitched trades are handled by thread 0 below using
    // the open-bar regime carried in `bnd_open_regime` shared-memory slot.
    int local_trades_T = 0, local_wins_T = 0;
    int local_trades_R = 0, local_wins_R = 0;
    int local_trades_V = 0, local_wins_V = 0;

    // Boundary-aware trade tracking.
    // Sentinel = -2 (was -1). After the trade-key change to signed_dir
    // {-1=Short, 0=Hold/Flat, +1=Long}, -1 is a legitimate direction, so
    // "no data in this chunk" must use a value outside the signed_dir range.
    int bnd_first_action = -2;
    int bnd_last_action = -2;
    float bnd_prefix_return = 0.0f;
    float bnd_suffix_return = 0.0f;
    int bnd_complete_trades = 0;
    int bnd_complete_wins = 0;
    int bnd_num_changes = 0;
    int bnd_cur_action = -1;
    float bnd_cur_return = 0.0f;
    // Regime at the bar where THIS chunk's first trade opens (i.e., the
    // chunk's first valid action). Carried across the boundary stitch so
    // thread 0's trade-counting loop attributes stitched trades to the
    // correct regime bucket. REGIME_BUCKET_NONE if the chunk has no
    // valid action (entire chunk is `act < 0` sentinel).
    int bnd_open_regime = REGIME_BUCKET_NONE;
    // Regime at the open bar of the trade currently being accumulated
    // INSIDE this thread's chunk. Tracked locally so internal-trade
    // closes can attribute themselves to the correct bucket; reset on
    // each transition.
    int bnd_cur_open_regime = REGIME_BUCKET_NONE;

    // Consecutive chunks for correct drawdown tracking
    int chunk_size = (wlen + stride - 1) / stride;
    int chunk_start = tid * chunk_size;
    int chunk_end = chunk_start + chunk_size;
    if (chunk_end > wlen) chunk_end = wlen;

    /* Trade boundary key: DIRECTION only, not (direction × magnitude).
     *
     * Previously the boundary key was `exp_idx = act / (order×urgency)` which
     * decodes to (dir × mag). That counted Long-Half → Long-Full as a NEW
     * trade even though the position sign didn't change. Each "new" trade
     * absorbed the bar-of-change tx_cost as its first step_return, biasing
     * win-rate downward asymmetrically vs. training's win-rate (which counts
     * trades only on position-sign cycles, see experience_kernels.cu:1592).
     *
     * 4-branch action encoding: action = dir*27 + mag*9 + ord*3 + urg.
     * Direction values: Short(0), Hold(1), Long(2), Flat(3).
     * dir_idx = action / (mag×order×urgency) = action / 27 = exp_idx / 3.
     *
     * For trade-boundary purposes we collapse Hold(1) and Flat(3) — both
     * mean "no exposure" — so a Long → Hold → Long is also one trade cycle.
     * This matches training's "position sign change" semantics. */
    for (int i = chunk_start; i < chunk_end; i++) {
        float r = step_returns[base + i];
        local_sum += r;
        if (r < 0.0f) local_down += r * r;

        local_cum *= (1.0f + r);

        int act = actions_history[base + i];
        /* Skip unwritten slots (act < 0): after `done_flags[w]=1` (capital
         * floor breach), backtest_env_step early-returns without writing
         * actions_history for the remaining slots in [done_step, max_len).
         * `reset_evaluation_state` initialises those to -1 sentinel. Without
         * this guard, C integer division (`-1 / 9 = 0`) decoded unwritten
         * slots as Short Quarter (action 0), inflating sell_count and
         * collapsing active_frac/dir_entropy to a measurement artifact
         * rather than real model behaviour. step_returns at these slots are
         * still zero (memset_zeros'd, never written), so summing them with
         * `r=0` is a no-op and remains correct. */
        if (act < 0) continue;
        int exp_idx = act / (order_actions * urgency_actions);
        int raw_dir = exp_idx / 3;  /* 0=Short, 1=Hold, 2=Long, 3=Flat */
        /* Collapse Hold(1) and Flat(3) → no-exposure (-1).
         * Keep Short(0) and Long(2) as distinct sign states. */
        int signed_dir = (raw_dir == DIR_SHORT) ? -1 :
                         (raw_dir == DIR_LONG)  ?  1 :
                                                   0;  /* Hold/Flat = no exposure */

        /* Action-distribution counters: now reflect ACTUAL directions, not
         * legacy 9-action exposure-band thresholds. Counts: long, short,
         * flat-or-hold (treated as "neutral"). */
        if (signed_dir > 0)      local_buys++;
        else if (signed_dir < 0) local_sells++;
        else                     local_holds++;

        if (exp_idx >= 0 && exp_idx < num_actions)
            local_action_mask |= (1 << exp_idx);

        // Classify regime at the current bar (for trade-open attribution).
        int cur_regime = classify_bar_regime(features, feat_base / feature_dim + i, feature_dim);

        if (i == chunk_start) {
            bnd_first_action = signed_dir;
            bnd_cur_action = signed_dir;
            // Open-bar regime for this chunk's prefix trade. Persists
            // across the boundary stitch so thread 0 can attribute the
            // stitched trade to the correct bucket.
            bnd_open_regime = cur_regime;
            bnd_cur_open_regime = cur_regime;
        }

        if (signed_dir != bnd_cur_action) {
            bnd_num_changes++;
            if (bnd_num_changes == 1) {
                bnd_prefix_return = bnd_cur_return;
            } else {
                // Internal complete trade close. Attribute to the regime
                // captured at the trade's OPEN bar (held in
                // `bnd_cur_open_regime` since the last transition).
                bnd_complete_trades++;
                bool is_win = (bnd_cur_return > 0.0f);
                if (is_win) bnd_complete_wins++;
                if (bnd_cur_open_regime == REGIME_BUCKET_T) {
                    local_trades_T++;
                    if (is_win) local_wins_T++;
                } else if (bnd_cur_open_regime == REGIME_BUCKET_R) {
                    local_trades_R++;
                    if (is_win) local_wins_R++;
                } else if (bnd_cur_open_regime == REGIME_BUCKET_V) {
                    local_trades_V++;
                    if (is_win) local_wins_V++;
                }
            }
            bnd_cur_return = 0.0f;
            bnd_cur_action = signed_dir;
            // New trade just opened at bar `i` — capture its regime.
            bnd_cur_open_regime = cur_regime;
        }

        bnd_cur_return += r;
        bnd_last_action = signed_dir;
    }

    bnd_suffix_return = bnd_cur_return;
    if (bnd_num_changes == 0) {
        bnd_prefix_return = bnd_cur_return;
    }
    // Regime of the trade currently "open" at chunk-end. Equals
    // `bnd_open_regime` if no transition occurred (single ongoing trade),
    // else `bnd_cur_open_regime` (regime of the last transition's open).
    int bnd_last_open_regime = bnd_cur_open_regime;

    // Store to shared memory for parallel reduction
    s_sum[tid] = local_sum;
    s_down_sq[tid] = local_down;
    s_buys[tid] = (float)local_buys;
    s_sells[tid] = (float)local_sells;
    s_holds[tid] = (float)local_holds;
    s_trades_T[tid] = (float)local_trades_T;
    s_wins_T[tid]   = (float)local_wins_T;
    s_trades_R[tid] = (float)local_trades_R;
    s_wins_R[tid]   = (float)local_wins_R;
    s_trades_V[tid] = (float)local_trades_V;
    s_wins_V[tid]   = (float)local_wins_V;
    __syncthreads();

    // Block-level parallel reduction — native f32 += f32
    for (int s = stride / 2; s > 0; s >>= 1) {
        if (tid < s) {
            s_sum[tid] += s_sum[tid + s];
            s_down_sq[tid] += s_down_sq[tid + s];
            s_buys[tid] += s_buys[tid + s];
            s_sells[tid] += s_sells[tid + s];
            s_holds[tid] += s_holds[tid + s];
            s_trades_T[tid] += s_trades_T[tid + s];
            s_wins_T[tid]   += s_wins_T[tid + s];
            s_trades_R[tid] += s_trades_R[tid + s];
            s_wins_R[tid]   += s_wins_R[tid + s];
            s_trades_V[tid] += s_trades_V[tid + s];
            s_wins_V[tid]   += s_wins_V[tid + s];
        }
        __syncthreads();
    }

    // Equity product reduction (multiplicative)
    s_sorted[tid] = local_cum;
    __syncthreads();
    for (int s = stride / 2; s > 0; s >>= 1) {
        if (tid < s) {
            s_sorted[tid] *= s_sorted[tid + s];
        }
        __syncthreads();
    }

    // Export boundary trade data to shared memory. Slots 7-8 added for
    // per-regime stitching: `bnd_open_regime` carries the regime at this
    // chunk's first action (used when a stitched trade ENDS at this
    // chunk's first action), `bnd_last_open_regime` carries the regime
    // at this chunk's last trade-open bar (used when a stitched trade
    // CONTINUES from this chunk's suffix into a subsequent chunk).
    {
        float* s_bnd = &s_sorted[stride];
        s_bnd[0 * stride + tid] = (float)bnd_first_action;
        s_bnd[1 * stride + tid] = (float)bnd_last_action;
        s_bnd[2 * stride + tid] = bnd_prefix_return;
        s_bnd[3 * stride + tid] = bnd_suffix_return;
        s_bnd[4 * stride + tid] = (float)bnd_complete_trades;
        s_bnd[5 * stride + tid] = (float)bnd_complete_wins;
        s_bnd[6 * stride + tid] = (float)bnd_num_changes;
        s_bnd[7 * stride + tid] = (float)bnd_open_regime;
        s_bnd[8 * stride + tid] = (float)bnd_last_open_regime;
    }
    __syncthreads();

    // Thread 0: final metrics from fully reduced values
    if (tid == 0) {
        float n = (float)wlen;
        float mean = s_sum[0] / n;
        float down_std = sqrtf(fmaxf(s_down_sq[0] / n, 1e-10f));

        // Exact max drawdown via sequential scan
        float exact_equity = 1.0f;
        float exact_peak = 1.0f;
        float exact_max_dd = 0.0f;
        for (int i = 0; i < wlen; i++) {
            float r = step_returns[base + i];
            exact_equity *= (1.0f + r);
            exact_peak = fmaxf(exact_peak, exact_equity);
            float dd = (exact_peak - exact_equity) / fmaxf(exact_peak, 1e-8f);
            exact_max_dd = fmaxf(exact_max_dd, dd);
        }

        // Boundary stitching: exact trade count + win rate.
        //
        // 2026-05-09 — also accumulate per-regime trade/win counts.
        // INTERNAL trade closes (within a chunk's body) were attributed
        // per-thread; here we sum those across threads. STITCHED trade
        // closes (boundary-spanning) are attributed using the open-bar
        // regime carried via `bnd_open_regime` / `bnd_last_open_regime`
        // in `s_bnd[7..]` and `s_bnd[8..]`.
        float* s_bnd = &s_sorted[stride];
        int total_trades = 0;
        int total_wins = 0;
        // Per-regime accumulators. INTERNAL counts come from the
        // pre-reduced `s_trades_T/R/V` + `s_wins_T/R/V` arrays
        // (slots 0 after the parallel reduction above).
        int trades_T = (int)s_trades_T[0];
        int wins_T   = (int)s_wins_T[0];
        int trades_R = (int)s_trades_R[0];
        int wins_R   = (int)s_wins_R[0];
        int trades_V = (int)s_trades_V[0];
        int wins_V   = (int)s_wins_V[0];

        for (int t = 0; t < stride; t++) {
            total_trades += (int)s_bnd[4 * stride + t];
            total_wins   += (int)s_bnd[5 * stride + t];
        }

        /* Boundary stitching across CUDA blocks.
         * fa/la are now signed_dir values {-1, 0, +1}; sentinel "no data"
         * is -2. Replaced `< 0` checks with `< -1` (only -2 means empty). */
        float open_return = 0.0f;
        int open_action = -2;
        // Open-bar regime of the running stitched trade. Sentinel = NONE
        // (carries through if the stitch spans regime-less blocks; in
        // practice every block with `fa != -2` has a valid regime).
        int open_regime = REGIME_BUCKET_NONE;

        for (int t = 0; t < stride; t++) {
            int fa = (int)s_bnd[0 * stride + t];
            int la = (int)s_bnd[1 * stride + t];
            float pr = s_bnd[2 * stride + t];
            float sr = s_bnd[3 * stride + t];
            int nc = (int)s_bnd[6 * stride + t];
            int blk_open_regime = (int)s_bnd[7 * stride + t];
            int blk_last_open_regime = (int)s_bnd[8 * stride + t];

            if (fa < -1) continue;  /* -2 = no data in this block */

            if (open_action < -1) {
                // First active block — capture initial open regime.
                open_action = fa;
                open_return = 0.0f;
                open_regime = blk_open_regime;
            } else if (fa != open_action) {
                // Stitched trade closes here. Attribute to the regime
                // captured at its open (carried in `open_regime`).
                total_trades++;
                bool is_win = (open_return > 0.0f);
                if (is_win) total_wins++;
                if (open_regime == REGIME_BUCKET_T) {
                    trades_T++;
                    if (is_win) wins_T++;
                } else if (open_regime == REGIME_BUCKET_R) {
                    trades_R++;
                    if (is_win) wins_R++;
                } else if (open_regime == REGIME_BUCKET_V) {
                    trades_V++;
                    if (is_win) wins_V++;
                }
                // New stitched trade opens at this block's first action;
                // its regime is the block's open-bar regime.
                open_action = fa;
                open_return = 0.0f;
                open_regime = blk_open_regime;
            }

            if (nc == 0) {
                open_return += pr;
            } else {
                // Block has internal transitions. Prefix portion closes
                // the running stitched trade; the suffix opens a NEW
                // stitched trade at the block's last transition bar
                // (regime = `blk_last_open_regime`).
                open_return += pr;
                total_trades++;
                bool is_win = (open_return > 0.0f);
                if (is_win) total_wins++;
                if (open_regime == REGIME_BUCKET_T) {
                    trades_T++;
                    if (is_win) wins_T++;
                } else if (open_regime == REGIME_BUCKET_R) {
                    trades_R++;
                    if (is_win) wins_R++;
                } else if (open_regime == REGIME_BUCKET_V) {
                    trades_V++;
                    if (is_win) wins_V++;
                }
                open_action = la;
                open_return = sr;
                open_regime = blk_last_open_regime;
            }
        }

        if (open_action >= -1) {
            // Final trailing trade — attribute to its tracked open regime.
            total_trades++;
            bool is_win = (open_return > 0.0f);
            if (is_win) total_wins++;
            if (open_regime == REGIME_BUCKET_T) {
                trades_T++;
                if (is_win) wins_T++;
            } else if (open_regime == REGIME_BUCKET_R) {
                trades_R++;
                if (is_win) wins_R++;
            } else if (open_regime == REGIME_BUCKET_V) {
                trades_V++;
                if (is_win) wins_V++;
            }
        }

        int out_base = w * 19;
        metrics_out[out_base + 0] = s_sorted[0] - 1.0f;
        metrics_out[out_base + 1] = exact_max_dd;
        metrics_out[out_base + 2] = (mean / down_std) * annualization_factor;

        float total_trades_f = (float)total_trades;
        metrics_out[out_base + 3] = (total_trades > 0)
            ? (float)total_wins / total_trades_f : 0.0f;
        metrics_out[out_base + 4] = total_trades_f;

        metrics_out[out_base + 9]  = s_buys[0];
        metrics_out[out_base + 10] = s_sells[0];
        metrics_out[out_base + 11] = s_holds[0];

        // Per-regime WR + trade-count outputs (slots 13-18).
        // WR is `wins / trades` per bucket; emits 0.0 when the bucket
        // has zero trades (matches the aggregate WR's empty-trades
        // convention at slot 3 above).
        float trades_T_f = (float)trades_T;
        float trades_R_f = (float)trades_R;
        float trades_V_f = (float)trades_V;
        metrics_out[out_base + 13] = (trades_T > 0) ? (float)wins_T / trades_T_f : 0.0f;
        metrics_out[out_base + 14] = trades_T_f;
        metrics_out[out_base + 15] = (trades_R > 0) ? (float)wins_R / trades_R_f : 0.0f;
        metrics_out[out_base + 16] = trades_R_f;
        metrics_out[out_base + 17] = (trades_V > 0) ? (float)wins_V / trades_V_f : 0.0f;
        metrics_out[out_base + 18] = trades_V_f;
    }

    // Unique action count via bitmask OR-reduction
    {
        int* s_action_mask = (int*)s_sorted;
        s_action_mask[tid] = local_action_mask;
        __syncthreads();
        for (int s = stride / 2; s > 0; s >>= 1) {
            if (tid < s) s_action_mask[tid] |= s_action_mask[tid + s];
            __syncthreads();
        }
        if (tid == 0) {
            int out_base = w * 19;
            metrics_out[out_base + 12] = (float)__popc(s_action_mask[0]);
        }
    }

    // VaR, CVaR, Calmar, Omega via bitonic sort
    int sort_len = wlen < 4096 ? wlen : 4096;
    int sample_step = (wlen + sort_len - 1) / sort_len;
    for (int i = tid; i < sort_len; i += stride) {
        int src = i * sample_step;
        s_sorted[i] = (src < wlen) ? step_returns[base + src] : 0.0f;
    }

    int padded_len = 1;
    while (padded_len < sort_len) padded_len <<= 1;

    float f32_inf = 1e30f;
    for (int i = sort_len + tid; i < padded_len; i += stride) {
        s_sorted[i] = f32_inf;
    }
    __syncthreads();

    // Bitonic sort — ascending, native f32 comparisons
    for (int k = 2; k <= padded_len; k <<= 1) {
        for (int j = k >> 1; j > 0; j >>= 1) {
            for (int i = tid; i < padded_len; i += stride) {
                int ixj = i ^ j;
                if (ixj > i) {
                    bool ascending = ((i & k) == 0);
                    if ((ascending  && s_sorted[i] > s_sorted[ixj]) ||
                        (!ascending && s_sorted[i] < s_sorted[ixj])) {
                        float tmp = s_sorted[i];
                        s_sorted[i] = s_sorted[ixj];
                        s_sorted[ixj] = tmp;
                    }
                }
            }
            __syncthreads();
        }
    }

    // Thread 0: extended metrics from sorted array
    if (tid == 0) {
        int out_base = w * 19;

        int var_idx = (int)(0.05f * (float)sort_len);
        if (var_idx < 0) var_idx = 0;
        if (var_idx >= sort_len) var_idx = sort_len - 1;
        float var_95 = s_sorted[var_idx];

        int cvar_count = var_idx + 1;
        float cvar_sum = 0.0f;
        for (int i = 0; i < cvar_count; i++) {
            cvar_sum += s_sorted[i];
        }
        float cvar_95 = cvar_sum / (float)cvar_count;

        float daily_mean = s_sum[0] / (float)wlen;
        // max_dd at slot [out_base + 1] after Phase 1.1.b offset shift
        // (sharpe removed → every metric below it shifted down by 1).
        float max_dd = metrics_out[out_base + 1];
        float calmar = (max_dd > 0.001f)
            ? (daily_mean * annualization_factor * annualization_factor) / max_dd
            : 0.0f;
        calmar = fmaxf(-100.0f, fminf(100.0f, calmar));

        float gain_sum = 0.0f, loss_sum = 0.0f;
        for (int i = 0; i < sort_len; i++) {
            if (s_sorted[i] > 0.0f) gain_sum += s_sorted[i];
            else                     loss_sum -= s_sorted[i];
        }
        float omega = (loss_sum > 0.001f)
            ? gain_sum / loss_sum : 0.0f;
        omega = fmaxf(0.0f, fminf(100.0f, omega));

        metrics_out[out_base + 5] = var_95;
        metrics_out[out_base + 6] = cvar_95;
        metrics_out[out_base + 7] = calmar;
        metrics_out[out_base + 8] = omega;
    }
}
