// state_layout.cuh — single source of truth for GPU state vector layout.
// Included by experience_kernels.cu and any kernel that assembles states.
// Must match crates/ml-core/src/state_layout.rs exactly.

#pragma once

// ── Dimensions ──
// D.8 Plan 2 Task 6C: TLOB attention (16-dim) inserted after OFI.
// Layout: Market(42) + OFI(32) + TLOB(16) + MTF(16) + Portfolio(8) + PlanISV(7) + Padding(7) = 128.
// SL_STATE_DIM 112→128 (matches SL_STATE_DIM_PADDED; padding reclaimed to 7 for 8-alignment).
#define SL_STATE_DIM          128
#define SL_MARKET_DIM          42
#define SL_OFI_DIM             32
#define SL_TLOB_DIM            16
#define SL_MTF_DIM             16
#define SL_PORTFOLIO_BASE_DIM   8
#define SL_PORTFOLIO_PLAN_DIM   7
#define SL_PADDING_DIM          7

// ── Offsets ──
#define SL_MARKET_START        0
#define SL_OFI_START           (SL_MARKET_START + SL_MARKET_DIM)
#define SL_TLOB_START          (SL_OFI_START + SL_OFI_DIM)
#define SL_MTF_START           (SL_TLOB_START + SL_TLOB_DIM)
#define SL_PORTFOLIO_START     (SL_MTF_START + SL_MTF_DIM)
#define SL_PLAN_ISV_START      (SL_PORTFOLIO_START + SL_PORTFOLIO_BASE_DIM)
#define SL_PADDING_START       (SL_PLAN_ISV_START + SL_PORTFOLIO_PLAN_DIM)

// ── cuBLAS alignment ──
#define SL_STATE_DIM_PADDED   128

// ────────────────────────────────────────────────────────────────────────────
// Portfolio state slot indices (ps[0..PORTFOLIO_STRIDE=41)).
// Invariant 8: every slot has a named constant. Raw ps[N] access outside the
// definition site in experience_kernels.cu is a lint violation.
//
// Slots 3-6 and 22 also carry named aliases defined locally in
// experience_kernels.cu (DSR_A_SLOT, DSR_B_SLOT, DSR_TRADE_COUNT_SLOT,
// PREV_CLOSE_SLOT, PREV_MID_SLOT). The PS_* names below are canonical for
// cross-file reference; the local aliases remain for in-file use.
// ────────────────────────────────────────────────────────────────────────────
#define PS_POSITION                  0   // current contract position (signed)
#define PS_CASH                      1   // cash balance
#define PS_PORTFOLIO_VALUE           2   // mark-to-market total (cash + pos*price)
#define PS_DSR_A                     3   // EMA of realised trade returns (DSR)
#define PS_DSR_B                     4   // EMA of squared trade returns (DSR)
#define PS_DSR_TRADE_COUNT           5   // completed trades counter (DSR warmup)
#define PS_PREV_CLOSE                6   // prev bar raw_close for DSR per-bar returns
#define PS_PEAK_EQUITY               7   // high-water mark (init to initial_capital)
#define PS_FLAT_COUNTER              8   // consecutive flat steps
#define PS_PREV_EQUITY               9   // equity at previous step
#define PS_HOLD_TIME                10   // consecutive steps with position
#define PS_REALIZED_PNL             11   // cumulative realised PnL
#define PS_ENTRY_PRICE              12   // price when trade was entered
#define PS_TRADE_START_PNL          13   // realized_pnl snapshot at trade entry
#define PS_KELLY_WIN_COUNT          14   // Kelly: number of profitable trade exits
#define PS_KELLY_LOSS_COUNT         15   // Kelly: number of losing trade exits
#define PS_KELLY_SUM_WINS           16   // Kelly: cumulative profit from winners
#define PS_KELLY_SUM_LOSSES         17   // Kelly: cumulative |loss| from losers
#define PS_KELLY_SUM_RETURNS        18   // Kelly: cumulative net returns (for mu)
#define PS_KELLY_SUM_SQ_RETURNS     19   // Kelly: cumulative squared returns (sigma^2)
#define PS_INTRA_TRADE_MAX_DD       20   // worst unrealised drawdown during trade (v8)
#define PS_INTRA_TRADE_MAX_PNL      21   // best unrealised P&L during trade (hindsight)
#define PS_PREV_MID                 22   // prev bar MBP-10 mid-price
#define PS_PLAN_TARGET_BARS         23   // plan: max hold bars (>0.5 = plan active)
#define PS_PLAN_PROFIT_TARGET       24   // plan: raw profit threshold
#define PS_PLAN_STOP_LOSS           25   // plan: raw stop threshold
#define PS_PLAN_SCALE_AGGRESSION    26   // plan: position ramp speed
#define PS_PLAN_CONVICTION          27   // plan: conviction at entry [0,1]
#define PS_PLAN_ASYMMETRY           28   // plan: profit/stop asymmetry
#define PS_PLAN_ENTRY_REGIME        29   // plan: regime_stability at entry
// Slots 30-37: prev-bar OFI scratch for delta computation (see experience_kernels.cu)
#define PS_OFI_PREV_BASE            30   // first OFI prev-bar slot (stride 8: [30..37])
#define PS_PEAK_PNL_BAR                   38   // Plan 3 C.4: hold_time at which MAX_PNL was hit (temporal timing bonus)
#define PS_INTRA_TRADE_MIN_PNL            39   // Plan 3 D.4a: worst unrealised PnL during trade (symmetric to MAX_PNL at 21)
#define PS_REGIME_SHIFT_BAR               40   // Plan 3 D.4b: hold_time at first regime-shift detection (0 if no shift)
#define PS_PRE_ENTRY_CONVICTION_EMA       41   // Plan 3 D.4c: EMA of conviction_core during Flat bars
#define PS_PRE_ENTRY_CONVICTION_VAR_EMA   42   // Plan 3 D.4c: EMA of (conviction_core − mean)² during Flat bars
#define PS_STRIDE                         43   // total portfolio state stride (== PORTFOLIO_STRIDE), grown 41→43 by Plan 3 D.4c conviction consistency

// ────────────────────────────────────────────────────────────────────────────
// Plan-ISV slot indices (plan_isv[0..PLAN_ISV_DIM=7)).
// Invariant 8: every slot has a named constant.
// These map to the SL_PLAN_ISV_START block of the state vector.
// ────────────────────────────────────────────────────────────────────────────
#define PLAN_ISV_PROGRESS            0   // hold_time / target_bars (clamped to 2.0)
#define PLAN_ISV_PNL_VS_TARGET       1   // unrealized / (profit_target × equity)
#define PLAN_ISV_PNL_VS_STOP         2   // -unrealized / (stop_loss × equity)
#define PLAN_ISV_ENTRY_CONVICTION    3   // conviction at trade entry [0, 1]
#define PLAN_ISV_CONVICTION_DRIFT    4   // current conviction − entry conviction
#define PLAN_ISV_REGIME_SHIFT        5   // |isv[11] − entry_regime_stability|
#define PLAN_ISV_REMAINING_FRACTION  6   // max(0, min(1, (target_bars − hold_time) / target_bars)) when active else 0
#define PLAN_ISV_DIM                 7

// ────────────────────────────────────────────────────────────────────────────
// Plan MLP output slot indices (plan_params[0..PLAN_PARAM_DIM=6)).
// Invariant 8: every slot has a named constant.
// Used for pp[N] and plan_params_ptr[i * 6 + N] access patterns.
// ────────────────────────────────────────────────────────────────────────────
#define PLAN_PARAM_TARGET_BARS       0   // max hold bars
#define PLAN_PARAM_PROFIT_TARGET     1   // raw profit threshold %
#define PLAN_PARAM_STOP_LOSS         2   // raw stop loss threshold %
#define PLAN_PARAM_SCALE_AGGRESSION  3   // position ramp speed
#define PLAN_PARAM_CONVICTION        4   // position size fraction [0, 1]
#define PLAN_PARAM_ASYMMETRY         5   // profit/stop ratio
#define PLAN_PARAM_DIM               6

// ────────────────────────────────────────────────────────────────────────────
// Branch selector indices (4-branch factored action).
// Invariant 8: every branch has a named constant.
// These are used to index per-branch arrays (norms, targets, sizes, etc.).
// ────────────────────────────────────────────────────────────────────────────
#define BRANCH_DIR                   0   // direction branch (4 actions: Short/Hold/Long/Flat)
#define BRANCH_MAG                   1   // magnitude branch (3 actions: Quarter/Half/Full)
#define BRANCH_ORD                   2   // order-type branch (3 actions)
#define BRANCH_URG                   3   // urgency branch (3 actions)
#define NUM_BRANCHES                 4

// ────────────────────────────────────────────────────────────────────────────
// Direction action sub-indices (dir_idx in [0..NUM_DIRECTIONS)).
// Invariant 8: every direction value has a named constant.
// ────────────────────────────────────────────────────────────────────────────
#define DIR_SHORT                    0   // open/maintain short position
#define DIR_HOLD                     1   // keep current position (no-op)
#define DIR_LONG                     2   // open/maintain long position
#define DIR_FLAT                     3   // close all positions to zero
#define NUM_DIRECTIONS               4

// ────────────────────────────────────────────────────────────────────────────
// Magnitude action sub-indices (mag_idx in [0..NUM_MAGNITUDES)).
// Invariant 8: every magnitude value has a named constant.
// ────────────────────────────────────────────────────────────────────────────
#define MAG_QUARTER                  0   // 0.25× max_position
#define MAG_HALF                     1   // 0.50× max_position
#define MAG_FULL                     2   // 1.00× max_position
#define NUM_MAGNITUDES               3

// ────────────────────────────────────────────────────────────────────────────
// Order-type and urgency branch sizes. Production default in
// `ml-dqn::dqn::DQNConfig` is 3 each (`num_order_types = 3`,
// `num_urgency_levels = 3`); the runtime config has never been tuned away
// from these and the kernels that consume packed factored actions assume
// the same value tile-wide. SP13 P0a.T4 (2026-05-04) lifts these into
// `state_layout.cuh` so kernels that need to decode the packed
// `action_idx = dir*M*O*U + mag*O*U + ord*U + urg` (e.g.
// `hold_rate_observer`) can compute the divisor without taking new
// launcher args. If runtime config ever changes, every consumer of the
// packed encoding must migrate atomically per
// `feedback_no_partial_refactor.md`.
// ────────────────────────────────────────────────────────────────────────────
#define NUM_ORD                      3
#define NUM_URG                      3

// ────────────────────────────────────────────────────────────────────────────
// ISV bus slot indices referenced by kernels that receive the ISV pointer.
// Kept in state_layout.cuh because both experience_kernels.cu and
// backtest_plan_kernel.cu include this header and need these indices.
// The authoritative named constants live in gpu_dqn_trainer.rs; the values
// here must match exactly. Any renumber requires updating both files in the
// same commit (per feedback_no_partial_refactor.md).
// ────────────────────────────────────────────────────────────────────────────
#define ISV_PLAN_THRESHOLD_IDX   49   // == PLAN_THRESHOLD_INDEX — plan activation threshold (shifted +3 by D.2 per-branch gamma, Plan 2 Task 3)
#define ISV_SHARPE_EMA_IDX       22   // == SHARPE_EMA_INDEX — EMA of training Sharpe ratio (Plan 3 D.4b adaptive threshold)
#define ISV_STATE_KL_AMP_IDX     79   // == STATE_KL_AMPLIFICATION_INDEX — Flat-trap escape amplifier ∈ [1, 2] (Plan 3 C.3)
#define ISV_CQL_ALPHA_IDX        48   // == CQL_ALPHA_INDEX — CQL pessimism base coefficient (Plan 1 Task 12; producer upgraded to GPU `cql_alpha_seed_update` by Plan 3 Task 9 C.5)
#define ISV_SEED_STEPS_TARGET_IDX 82  // == SEED_STEPS_TARGET_INDEX — config replay_seed_steps target (Plan 3 Task 8 B.3)
#define ISV_SEED_STEPS_DONE_IDX   83  // == SEED_STEPS_DONE_INDEX — cumulative seed-phase steps completed (Plan 3 Task 8 B.3)
#define ISV_SEED_FRAC_EMA_IDX     84  // == SEED_FRAC_EMA_INDEX — adaptive EMA of (1 - done/target) ∈ [0, 1] (Plan 3 Task 8 B.3; consumed by Task 9 CQL ramp)
#define ISV_EVAL_THOMPSON_TEMP_IDX 339 // == EVAL_THOMPSON_TEMP_INDEX — eval Thompson selector temperature (SP10 / Fix 38 2026-05-03; ISV-driven temperature blend on direction-branch Thompson sample)

// ────────────────────────────────────────────────────────────────────────────
// === SP13 SLOT INDICES === (2026-05-04, redefine success for predictive skill)
// Mirror of crates/ml/src/cuda_pipeline/sp13_isv_slots.rs. Slots [372..383)
// at the top of the SP5 ISV linear span; [367..372) is an intentional 5-slot
// boundary gap separating the SP11 controller block from this SP13 block.
// Slots 380..383 are the v3 P0a.T3 (2026-05-04) Hold-pricing controller
// (replaces v2's atomic Hold elimination — preserves the 4-way action space,
// prices Hold instead of removing it).
// Spec: docs/superpowers/specs/2026-05-04-sp13-redefine-success-for-predictive-skill.md
// ────────────────────────────────────────────────────────────────────────────
#define ISV_TARGET_DIR_ACC_IDX        372  // directional accuracy target (static, default 0.55)
#define ISV_AUX_DIR_ACC_SHORT_EMA_IDX 373  // aux dir-acc fast EMA (per-fold, sentinel 0.5)
#define ISV_AUX_DIR_ACC_LONG_EMA_IDX  374  // aux dir-acc slow EMA (per-fold, sentinel 0.5) — stagnation detector
#define ISV_AUX_DIR_PREDICTION_IDX    375  // aux head per-bar prediction (per-bar overwrite, [-1, +1])
#define ISV_DIR_SKILL_BONUS_ALPHA_IDX 376  // dir-skill bonus α (per-fold, default 1.0) — magnitude on correct calls
#define ISV_DIR_SKILL_BONUS_BETA_IDX  377  // dir-skill bonus β (per-fold, default 1.0) — penalty on wrong calls
#define ISV_LUCK_WIN_DISCOUNT_IDX     378  // lucky-win discount factor (per-fold, default 0.3)
#define ISV_SKILL_BONUS_CAP_RATIO_IDX 379  // skill-bonus cap ratio (per-fold, default 0.3) — bonus ≤ ratio × |α|
#define ISV_HOLD_COST_IDX             380  // SP13 v3 per-bar Hold cost (controller output, sentinel = HOLD_COST_BASE=0.001)
#define ISV_HOLD_RATE_TARGET_IDX      381  // SP13 v3 target Hold-pick rate (static, default 0.20)
#define ISV_HOLD_RATE_OBSERVED_EMA_IDX 382 // SP13 v3 observed Hold-pick rate EMA (per-fold, sentinel 0.0)

#define ISV_TOTAL_DIM 383

// ────────────────────────────────────────────────────────────────────────────
// Feature-group ranges — Plan 4 Task 1A (E.1 VSN prerequisite).
// Referenced by: VSN feature selection (E.1), attention diagnostics (E.5),
// encoder-decoder separation (E.4). Ranges are half-open `[BEGIN, END)`.
// `SL_PLAN_ISV_GROUP_END` aliases `SL_PADDING_START` because the 7-element
// plan-ISV block is the last group before the 8-alignment padding.
// Invariant 8: every range has a named constant; raw `SL_*_START` arithmetic
// for group bounds outside this header is a lint violation.
// ────────────────────────────────────────────────────────────────────────────
#define SL_MARKET_GROUP_BEGIN    SL_MARKET_START
#define SL_MARKET_GROUP_END      SL_OFI_START
#define SL_OFI_GROUP_BEGIN       SL_OFI_START
#define SL_OFI_GROUP_END         SL_TLOB_START
#define SL_TLOB_GROUP_BEGIN      SL_TLOB_START
#define SL_TLOB_GROUP_END        SL_MTF_START
#define SL_MTF_GROUP_BEGIN       SL_MTF_START
#define SL_MTF_GROUP_END         SL_PORTFOLIO_START
#define SL_PORTFOLIO_GROUP_BEGIN SL_PORTFOLIO_START
#define SL_PORTFOLIO_GROUP_END   SL_PLAN_ISV_START
#define SL_PLAN_ISV_GROUP_BEGIN  SL_PLAN_ISV_START
#define SL_PLAN_ISV_GROUP_END    SL_PADDING_START

#define SL_NUM_FEATURE_GROUPS    6
#define SL_MAX_FEATURE_GROUP_DIM 42  // == SL_MARKET_DIM (largest group)

// ────────────────────────────────────────────────────────────────────────────
// SP12 v3 (2026-05-04): per-trade event-driven reward composition constants.
// Spec: docs/superpowers/specs/2026-05-04-sp12-quality-over-quantity-reward-shaping.md
//
// REWARD_POS_CAP / REWARD_NEG_CAP — asymmetric bounded reward cap (loss aversion).
//   2:1 ratio matches Kahneman/Tversky prospect-theory meta-analysis pegging
//   human loss aversion at ~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. Site:
//   experience_kernels.cu line ~2758 (segment_complete vol-normalized P&L).
//   Per pearl_audit_unboundedness_for_implicit_asymmetry.
//
// MIN_HOLD_TARGET — patience requirement on voluntary trade exits, in bars.
//   30 bars = MFT-leaning but not pure MFT (full MFT ≈ 100 bars,
//   HFT-tolerant ≈ 10 bars). Matches the user-profile target band
//   "between HFT and MFT".
//
// MIN_HOLD_PENALTY_MAX — max penalty magnitude when voluntary exit at hold=0.
//   3.0 = 60% of REWARD_POS_CAP. Meaningful gradient pressure but not
//   paralyzing — a profitable scalp can still net positive with the penalty.
//
// MIN_HOLD_TEMPERATURE_{START,END,DECAY} — annealing schedule for the
//   soft-saturation temperature in the deficit/(deficit + T) formula. Computed
//   in Rust per epoch via T(e) = T_end + (T_start - T_end) × exp(-e/decay)
//   and passed to the kernel as a scalar (not a constant lookup). Curriculum:
//   early training forgiving (T=50), late training sharp (T=5).
//
// Phase 1 = Invariant-1 numerical anchors (static constants). Phase 2 lifts
// the target / penalty_max / temperature schedule to ISV-driven controllers
// per pearl_kelly_cap_signal_driven_floors and feedback_isv_for_adaptive_bounds
// IF Phase 1 validation reveals adaptive need. Spec §implementation phases.
// ────────────────────────────────────────────────────────────────────────────
#define REWARD_POS_CAP             5.0f
#define REWARD_NEG_CAP            -10.0f
#define MIN_HOLD_TARGET            30.0f
#define MIN_HOLD_PENALTY_MAX        3.0f
#define MIN_HOLD_TEMPERATURE_START 50.0f
#define MIN_HOLD_TEMPERATURE_END    5.0f
#define MIN_HOLD_TEMPERATURE_DECAY 20.0f

// ── Compile-time checks ──
static_assert(SL_PADDING_START + SL_PADDING_DIM == SL_STATE_DIM,
    "State layout dimensions must sum to SL_STATE_DIM");
static_assert(SL_STATE_DIM % 8 == 0,
    "SL_STATE_DIM must be 8-aligned for tensor core cuBLAS");
static_assert(SL_STATE_DIM <= SL_STATE_DIM_PADDED,
    "SL_STATE_DIM must fit within SL_STATE_DIM_PADDED");
// Feature groups must contiguously cover the non-padding state range.
static_assert(SL_MARKET_GROUP_BEGIN == 0,
    "feature groups must start at offset 0");
static_assert(SL_PLAN_ISV_GROUP_END == SL_PADDING_START,
    "feature groups must end at SL_PADDING_START — padding is not a group");
static_assert(SL_MARKET_DIM <= SL_MAX_FEATURE_GROUP_DIM,
    "Market group dim must not exceed SL_MAX_FEATURE_GROUP_DIM");
static_assert(SL_OFI_DIM    <= SL_MAX_FEATURE_GROUP_DIM,
    "OFI group dim must not exceed SL_MAX_FEATURE_GROUP_DIM");
static_assert(SL_TLOB_DIM   <= SL_MAX_FEATURE_GROUP_DIM,
    "TLOB group dim must not exceed SL_MAX_FEATURE_GROUP_DIM");
static_assert(SL_MTF_DIM    <= SL_MAX_FEATURE_GROUP_DIM,
    "MTF group dim must not exceed SL_MAX_FEATURE_GROUP_DIM");
static_assert(SL_PORTFOLIO_BASE_DIM <= SL_MAX_FEATURE_GROUP_DIM,
    "Portfolio group dim must not exceed SL_MAX_FEATURE_GROUP_DIM");
static_assert(SL_PORTFOLIO_PLAN_DIM <= SL_MAX_FEATURE_GROUP_DIM,
    "Plan-ISV group dim must not exceed SL_MAX_FEATURE_GROUP_DIM");

// ── Shared state assembly function ──
// Called by both experience_state_gather (training) and backtest_state_gather (validation).
// Each caller prepares its own data and calls this to produce identical layouts.
//
// All pointers are per-thread (one state row). Caller is responsible for
// computing the correct source pointers for this thread's episode/window.

__device__ __forceinline__ void assemble_state(
    float* __restrict__ out,
    const float* __restrict__ market,
    const float* __restrict__ ofi,
    const float  mtf[SL_MTF_DIM],
    const float  portfolio[SL_PORTFOLIO_BASE_DIM],
    const float  plan_isv[SL_PORTFOLIO_PLAN_DIM]
) {
    // Market features [0..42)
    for (int k = 0; k < SL_MARKET_DIM; k++)
        out[SL_MARKET_START + k] = market[k];

    // OFI [42..74)
    for (int k = 0; k < SL_OFI_DIM; k++)
        out[SL_OFI_START + k] = ofi[k];

    // TLOB [74..90) — zeros on assembly; GpuTlob::forward overwrites post-gather (D.8 Plan 2 Task 6C)
    for (int k = 0; k < SL_TLOB_DIM; k++)
        out[SL_TLOB_START + k] = 0.0f;

    // MTF [90..106)
    for (int k = 0; k < SL_MTF_DIM; k++)
        out[SL_MTF_START + k] = mtf[k];

    // Portfolio base [106..114)
    for (int k = 0; k < SL_PORTFOLIO_BASE_DIM; k++)
        out[SL_PORTFOLIO_START + k] = portfolio[k];

    // Plan/ISV [114..121)
    for (int k = 0; k < SL_PORTFOLIO_PLAN_DIM; k++)
        out[SL_PLAN_ISV_START + k] = plan_isv[k];

    // Padding [121..128) — 7 zeros for 8-alignment (D.8 Plan 2 Task 6C)
    for (int k = 0; k < SL_PADDING_DIM; k++)
        out[SL_PADDING_START + k] = 0.0f;
}
