/**
 * backtest_plan_state_isv — val plan_isv parity with training.
 *
 * Val-Flat-collapse fix #5 (task #94, 2026-04-24). Research-agent audit
 * traced val trade count of 21 per 214,654 bars (0.0098% trade rate)
 * to a structural train/val distribution shift: state positions [86..92)
 * (`plan_isv[6]`) carry real plan signals during training but were
 * zero-filled in val. The deterministic val policy defaulted to
 * Flat/Hold on OOD states regardless of directional edge.
 *
 * The training plan machinery:
 *   1. Plan MLP runs on trunk hidden state (h_s2) → plan_params[B, 6]:
 *      {target_bars, profit_target, stop_loss, scale_aggression,
 *       conviction, asymmetry}.
 *   2. On Flat→Positioned transitions in env_step, plan_params are
 *      COPIED into persistent portfolio slots ps[23..29] (the "stored
 *      plan" the trade is operating under). This lets the plan persist
 *      across bars within the trade.
 *   3. State assembly reads ps[23..29] and computes plan_isv[0..6] as
 *      real-time ratios: plan progress (hold_time/target), PnL vs
 *      profit target, PnL vs stop, entry conviction, conviction drift,
 *      regime shift, remaining_fraction (D.6 Plan 2 Task 6A).
 *
 * Val replicates this with a SEPARATE plan_state[N, 7] buffer (val's
 * portfolio state is 8 slots and can't hold plan fields without a
 * structural change; a parallel buffer is cleaner). Layout:
 *   plan_state[w*7 + 0] = target_bars  (== stored ps[23]; plan_active iff > 0.5)
 *   plan_state[w*7 + 1] = profit_target        (stored ps[24])
 *   plan_state[w*7 + 2] = stop_loss            (stored ps[25])
 *   plan_state[w*7 + 3] = scale_aggression     (stored ps[26])
 *   plan_state[w*7 + 4] = conviction_at_entry  (stored ps[27])
 *   plan_state[w*7 + 5] = asymmetry            (stored ps[28])
 *   plan_state[w*7 + 6] = entry_regime_stability (stored ps[29])
 *
 * Portfolio layout (val's 8-slot buffer):
 *   [0] value (equity), [1] position, [2] cash, [3] entry_price,
 *   [4] max_equity, [5] hold_time, [6] cum_return, [7] step_count
 *
 * This kernel runs ONCE PER VAL STEP, AFTER env_step and BEFORE the
 * NEXT step's state gather. It performs three operations per window:
 *
 *   (a) Plan ACTIVATION on Flat→Positioned: if plan_state[0] < 0.5
 *       (no active plan) and position just became non-zero, copy the
 *       fresh plan_params into plan_state + store entry regime.
 *   (b) Plan DEACTIVATION on return to Flat: if position ≈ 0, zero
 *       out plan_state so the next entry activates a fresh plan.
 *   (c) plan_isv[0..6] computation from plan_state (stored, stable
 *       across the trade) + plan_params (live, for conviction drift)
 *       + portfolio (for hold_time and unrealized P&L).
 *
 * The seven plan_isv features mirror training formula
 * (experience_kernels.cu):
 *   plan_isv[PLAN_ISV_PROGRESS]           = hold_time / target_bars  (plan progress, clamped to 2.0)
 *   plan_isv[PLAN_ISV_PNL_VS_TARGET]      = unrealized_pnl / (profit_target × equity)
 *   plan_isv[PLAN_ISV_PNL_VS_STOP]        = -unrealized_pnl / (stop_loss × equity)
 *   plan_isv[PLAN_ISV_ENTRY_CONVICTION]   = conviction_at_entry
 *   plan_isv[PLAN_ISV_CONVICTION_DRIFT]   = plan_params[conviction] - conviction_at_entry
 *   plan_isv[PLAN_ISV_REGIME_SHIFT]       = |isv[11] - entry_regime_stability|
 *   plan_isv[PLAN_ISV_REMAINING_FRACTION] = max(0, min(1, (target_bars - hold_time) / target_bars))  (D.6)
 *
 * Graph safety: fixed grid/block (N windows, 256-thread blocks). No
 * atomics, no dynamic alloc, no host syncs. Safe to capture inside
 * the val backtest's main loop.
 *
 * Diagnostic: writes per-window {has_plan, hold_time, conviction_drift}
 * to a small scratch buffer so the Rust side can log plan activity
 * statistics at epoch boundaries without a big DtoH.
 */

#include "state_layout.cuh"

extern "C" __global__ void backtest_plan_state_isv(
    const float* __restrict__ plan_params,   /* [N, 6] live output from plan MLP */
    float*       __restrict__ plan_state,    /* [N, 7] persistent plan state (in-place) */
    const float* __restrict__ portfolio,     /* [N, 8] val portfolio state */
    const float* __restrict__ isv_signals,   /* [ISV_TOTAL_DIM] for slot [11] */
    const float* __restrict__ prices,        /* [N * max_len * 4] raw OHLC; close at col 3. NULL → skip P&L ratios */
    int          current_step,               /* current bar index within window */
    int          max_len,                    /* bars per window */
    float*       __restrict__ plan_isv_out,  /* [N, SL_PORTFOLIO_PLAN_DIM=7] plan_isv output (consumed by next state gather) */
    int n_windows
) {
    int w = blockIdx.x * blockDim.x + threadIdx.x;
    if (w >= n_windows) return;

    float position    = portfolio[w * 8 + 1];
    float cash        = portfolio[w * 8 + 2];
    float entry_price = portfolio[w * 8 + 3];
    float hold_time   = portfolio[w * 8 + 5];

    /* Fix 30 Stale-B (2026-05-02): pre-fix this kernel read raw_close from
     * `features[bar*feat_dim + 0]`, but post Bug-1 (commit 5a5dd0fed)
     * `features[..+0]` is z-normed log-return, NOT raw_close. The
     * resulting `equity = cash + position*raw_close` and
     * `unrealized = position*(raw_close - entry_price)` formulas mixed
     * z-normed-log-return as a dollar price, corrupting val plan_isv
     * slots [PNL_VS_TARGET] and [PNL_VS_STOP]. Post-fix raw_close comes
     * from the upload-once `prices` buffer (layout [n*max_len*4] OHLC,
     * close column = index 3) — the same source the env_step kernel
     * reads from for portfolio mark-to-market (`backtest_env_kernel.cu:64`).
     * The other plan_isv slots (progress, conviction, drift, regime,
     * remaining) don't depend on raw_close and were correct pre-fix. */
    float raw_close = 0.0f;
    bool have_close = false;
    if (prices != nullptr && current_step >= 0 && current_step < max_len) {
        long long price_off = ((long long)w * max_len + current_step) * 4 + 3;
        raw_close = prices[price_off];
        have_close = true;
    }

    int in_position = (fabsf(position) > 0.001f);

    /* ── Plan deactivation: return to Flat zeros the stored plan. ── */
    if (!in_position) {
        #pragma unroll
        for (int k = 0; k < 7; k++) plan_state[w * 7 + k] = 0.0f;
    }

    /* ── Plan activation: Flat→Positioned transition.
     * Detection: plan_state[0] < 0.5f (no active plan) AND currently
     * in a position. This approximates training's
     * (was_flat && now_positioned) check using the state we track.
     * Caveats: if a trade flips direction in a single step, we'll see
     * in_position=true with a stale plan_state — the plan for the
     * prior direction. In practice trade flips are rare and the plan
     * gets re-populated on the NEXT Flat→Position transition after
     * the flip closes. ── */
    float plan_thr = (isv_signals != nullptr) ? isv_signals[ISV_PLAN_THRESHOLD_IDX] : 0.5f;
    float active_tgt = plan_state[w * 7 + 0];
    if (in_position && active_tgt < plan_thr) {
        const float* pp = plan_params + w * 6;
        plan_state[w * 7 + 0] = pp[PLAN_PARAM_TARGET_BARS];  /* target_bars */
        plan_state[w * 7 + 1] = pp[PLAN_PARAM_PROFIT_TARGET];  /* profit_target */
        plan_state[w * 7 + 2] = pp[PLAN_PARAM_STOP_LOSS];  /* stop_loss */
        plan_state[w * 7 + 3] = pp[PLAN_PARAM_SCALE_AGGRESSION];  /* scale_aggression */
        plan_state[w * 7 + 4] = pp[PLAN_PARAM_CONVICTION];  /* conviction (at entry) */
        plan_state[w * 7 + 5] = pp[PLAN_PARAM_ASYMMETRY];  /* asymmetry */
        plan_state[w * 7 + 6] = (isv_signals != nullptr)
            ? isv_signals[11] : 1.0f;   /* regime_stability at entry */
    }

    /* ── Compute plan_isv from stored plan_state + live signals. ── */
    float plan_tgt_bars = plan_state[w * 7 + 0];
    float plan_profit   = plan_state[w * 7 + 1];
    float plan_stop     = plan_state[w * 7 + 2];
    float plan_conv_ent = plan_state[w * 7 + 4];
    float entry_stab    = plan_state[w * 7 + 6];

    /* Unrealized P&L: position × (raw_close − entry_price).
     * Equity is current mark-to-market: cash + position×raw_close.
     * When raw_close unavailable (smoke-test callers), fall back to
     * unrealized ≈ 0 (still produces plan_isv[PLAN_ISV_PROGRESS] and [3..5], skips [1][2]). */
    float equity = cash + position * (have_close ? raw_close : entry_price);
    float unrealized = (have_close && entry_price > 0.0f)
        ? position * (raw_close - entry_price) : 0.0f;

    float pisv[SL_PORTFOLIO_PLAN_DIM];

    /* [0] plan progress [0, 2]: hold_time/target_bars */
    pisv[PLAN_ISV_PROGRESS] = (plan_tgt_bars > plan_thr)
        ? fminf(hold_time / plan_tgt_bars, 2.0f) : 0.0f;

    /* [1] PnL vs profit target [symmetric bound -2.0..2.0]
     * unrealized can be negative → fminf alone leaves unbounded lower tail.
     * Mirror of SP11 symmetric reward cap (35db31089). */
    pisv[PLAN_ISV_PNL_VS_TARGET] = (plan_profit > 1e-6f)
        ? fmaxf(-2.0f, fminf(unrealized / (plan_profit * equity + 1e-6f), 2.0f)) : 0.0f;

    /* [2] PnL vs stop (negative PnL vs stop → loss approaching) [symmetric bound -2.0..2.0] */
    pisv[PLAN_ISV_PNL_VS_STOP] = (plan_stop > 1e-6f)
        ? fmaxf(-2.0f, fminf(-unrealized / (plan_stop * equity + 1e-6f), 2.0f)) : 0.0f;

    /* [3] Entry conviction (stable across the trade) */
    pisv[PLAN_ISV_ENTRY_CONVICTION] = plan_conv_ent;

    /* [4] Conviction drift: current MLP conviction - entry conviction */
    pisv[PLAN_ISV_CONVICTION_DRIFT] = (plan_tgt_bars > plan_thr)
        ? (plan_params[w * 6 + PLAN_PARAM_CONVICTION] - plan_conv_ent) : 0.0f;

    /* [5] Regime shift magnitude since entry */
    pisv[PLAN_ISV_REGIME_SHIFT] = (isv_signals != nullptr && plan_tgt_bars > plan_thr)
        ? fabsf(isv_signals[11] - entry_stab) : 0.0f;

    /* [6] D.6 remaining_fraction (Plan 2 Task 6A): temporal pressure.
     * Value 1.0 = plan just activated, 0.0 = at or past plan expiry. */
    pisv[PLAN_ISV_REMAINING_FRACTION] = (plan_tgt_bars > plan_thr)
        ? fmaxf(0.0f, fminf(1.0f, (plan_tgt_bars - hold_time) / plan_tgt_bars))
        : 0.0f;

    #pragma unroll
    for (int k = 0; k < SL_PORTFOLIO_PLAN_DIM; k++) plan_isv_out[w * SL_PORTFOLIO_PLAN_DIM + k] = pisv[k];
}

/* ================================================================== */
/* Diagnostic reduction: count active plans + mean plan_isv magnitudes */
/* ================================================================== */

/**
 * backtest_plan_diag_reduce — single-block reduction emitting 8 diagnostic
 * floats summarising plan activity across all val windows:
 *
 *   diag_out[0] = fraction of windows with active plan (plan_state[0] > 0.5)
 *   diag_out[1] = mean |plan_isv[PLAN_ISV_PROGRESS]| (progress) across windows
 *   diag_out[2] = mean |plan_isv[PLAN_ISV_PNL_VS_TARGET]| (pnl vs target)
 *   diag_out[3] = mean |plan_isv[PLAN_ISV_PNL_VS_STOP]| (pnl vs stop)
 *   diag_out[4] = mean  plan_isv[PLAN_ISV_ENTRY_CONVICTION]  (entry conviction)
 *   diag_out[5] = mean |plan_isv[PLAN_ISV_CONVICTION_DRIFT]| (conviction drift)
 *   diag_out[6] = mean |plan_isv[PLAN_ISV_REGIME_SHIFT]| (regime shift)
 *   diag_out[7] = n_active windows (raw count, useful for confidence)
 *
 * Launched once per epoch (cold path) — rust side DtoH-reads these 8
 * floats for HEALTH_DIAG emission. One block × 256 threads is enough
 * for typical n_windows <= 1024; larger N still fits because each
 * thread strides over multiple windows.
 */
extern "C" __global__ void backtest_plan_diag_reduce(
    const float* __restrict__ plan_state,   /* [N, 7] */
    const float* __restrict__ plan_isv,     /* [N, SL_PORTFOLIO_PLAN_DIM=7] */
    float*       __restrict__ diag_out,     /* [8] */
    int n_windows
) {
    __shared__ float sh_active[256];
    __shared__ float sh_sum0[256];
    __shared__ float sh_sum1[256];
    __shared__ float sh_sum2[256];
    __shared__ float sh_sum3[256];
    __shared__ float sh_sum4[256];
    __shared__ float sh_sum5[256];

    int tid = threadIdx.x;
    float active = 0.0f;
    float s0 = 0.0f, s1 = 0.0f, s2 = 0.0f, s3 = 0.0f, s4 = 0.0f, s5 = 0.0f;

    for (int w = tid; w < n_windows; w += blockDim.x) {
        float has_plan = (plan_state[w * 7 + 0] > 0.5f) ? 1.0f : 0.0f;
        active += has_plan;
        s0 += fabsf(plan_isv[w * SL_PORTFOLIO_PLAN_DIM + 0]);
        s1 += fabsf(plan_isv[w * SL_PORTFOLIO_PLAN_DIM + 1]);
        s2 += fabsf(plan_isv[w * SL_PORTFOLIO_PLAN_DIM + 2]);
        s3 += plan_isv[w * SL_PORTFOLIO_PLAN_DIM + 3];
        s4 += fabsf(plan_isv[w * SL_PORTFOLIO_PLAN_DIM + 4]);
        s5 += fabsf(plan_isv[w * SL_PORTFOLIO_PLAN_DIM + 5]);
    }
    sh_active[tid] = active;
    sh_sum0[tid] = s0; sh_sum1[tid] = s1; sh_sum2[tid] = s2;
    sh_sum3[tid] = s3; sh_sum4[tid] = s4; sh_sum5[tid] = s5;
    __syncthreads();

    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) {
            sh_active[tid] += sh_active[tid + s];
            sh_sum0[tid]   += sh_sum0[tid + s];
            sh_sum1[tid]   += sh_sum1[tid + s];
            sh_sum2[tid]   += sh_sum2[tid + s];
            sh_sum3[tid]   += sh_sum3[tid + s];
            sh_sum4[tid]   += sh_sum4[tid + s];
            sh_sum5[tid]   += sh_sum5[tid + s];
        }
        __syncthreads();
    }

    if (tid == 0) {
        float inv_n = 1.0f / fmaxf((float)n_windows, 1.0f);
        diag_out[0] = sh_active[0] * inv_n;       /* active fraction */
        diag_out[1] = sh_sum0[0]   * inv_n;
        diag_out[2] = sh_sum1[0]   * inv_n;
        diag_out[3] = sh_sum2[0]   * inv_n;
        diag_out[4] = sh_sum3[0]   * inv_n;
        diag_out[5] = sh_sum4[0]   * inv_n;
        diag_out[6] = sh_sum5[0]   * inv_n;
        diag_out[7] = sh_active[0];               /* raw active count */
    }
}
