// 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)

// ────────────────────────────────────────────────────────────────────────────
// === SP14-C.4b SLOT INDICES === (2026-05-05, aux trunk adaptive horizon)
// Mirror of crates/ml/src/cuda_pipeline/sp14_isv_slots.rs slot 451.
// Pearl-A-bootstrapped Welford EMA of observed winning trade hold times (bars).
// Sentinel = 0.0 (SENTINEL_AVG_WIN_HOLD_TIME_BARS). First winning trade close
// replaces sentinel directly (first-observation bootstrap). Consumer here
// (Class A P0-C): replaces MIN_HOLD_TARGET hardcoded constant in
// `compute_min_hold_penalty` call site — see experience_kernels.cu.
// ────────────────────────────────────────────────────────────────────────────
#define AVG_WIN_HOLD_TIME_BARS_INDEX 451 // == AVG_WIN_HOLD_TIME_BARS_INDEX — EMA of winning-trade hold duration (bars); Pearl-A bootstrap sentinel 0.0

// ────────────────────────────────────────────────────────────────────────────
// === 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)

// ────────────────────────────────────────────────────────────────────────────
// === SP16 T3 MHT SLOTS (2026-05-08, Wiener-optimal adaptive α) ===
// Mirror of crates/ml/src/cuda_pipeline/sp14_isv_slots.rs slots [468..474).
// 6 Welford accumulators driving the surviving MIN_HOLD_TEMPERATURE
// producer (T1) per `pearl_wiener_optimal_adaptive_alpha`. FoldReset
// sentinel 0.0; cold-start fallback (sample_count < 3 → α = 1.0
// REPLACE).
// ────────────────────────────────────────────────────────────────────────────
#define ISV_MHT_TARGET_MEAN_IDX     468  // Welford mean of target_temp (T1 producer)
#define ISV_MHT_TARGET_M2_IDX       469  // Welford ΣΔ² (target_temp)
#define ISV_MHT_DIFF_MEAN_IDX       470  // Welford mean of consecutive diff (target_temp)
#define ISV_MHT_DIFF_M2_IDX         471  // Welford ΣΔ² (diff)
#define ISV_MHT_PREV_TARGET_IDX     472  // previous-epoch target_temp
#define ISV_MHT_SAMPLE_COUNT_IDX    473  // Welford counter (T1 producer)
// Wiener-α floor for non-stationary control loops (2026-05-08). Wiener-
// optimal α = diff_var / (diff_var + sample_var + ε) is MSE-optimal for
// stationary signals. The SP16 T3 producer (min_hold_temperature) tracks
// a hold-rate overrun driven by a co-adapting policy — non-stationary by
// construction. Once Welford accumulates
// 3+ samples and the policy converges toward target, diff_var collapses
// and α_wiener → 0, leaving the controller unable to react when the
// policy subsequently drifts. The 0.4 floor preserves ~2.5-step
// catch-up bandwidth. This is a structural-control constraint (Wiener
// stationarity assumption violated), NOT a tuned constant; floor only
// applies once sample_count ≥ WELFORD_MIN_SAMPLE_COUNT_F (cold-start
// branch keeps α = 1.0 REPLACE per Pearl A). See
// `pearl_wiener_alpha_floor_for_nonstationary.md`. Empirical anchor:
// train-multi-seed-b5gmp sha 641aa0dfd Fold 1 collapse (HD4 α=0.248,
// blended_scale lag → Hold% 52.6% → Sharpe 55.55 vs 88.65 baseline).
#define WELFORD_ALPHA_MIN_F         0.4f
#define WELFORD_ALPHA_MAX_F         0.95f
#define WELFORD_MIN_SAMPLE_COUNT_F  3.0f
#define WELFORD_EPS_F               1.0e-6f

// NOTE: this macro is kept at 383 for the historic SP13-era boundary marker
// and is referenced only in commentary inside .cu/.cuh files. The actual
// ISV bus dimension is the Rust `ISV_TOTAL_DIM` constant in
// `gpu_dqn_trainer.rs` (currently 458, post Class A P1-Producer
// 2026-05-08). Kernels take `const float*` pointers and use raw indices —
// no .cu/.cuh file dimensions an array with this macro. The Rust constant
// is the single source of truth for the bus size.
#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 — cold-start fallback for the patience requirement on
//   voluntary trade exits, in bars. Used when ISV[AVG_WIN_HOLD_TIME_BARS_
//   INDEX=451] is still at sentinel (0.0 = no winning trades observed yet).
//   Once the first winning trade closes and the EMA bootstraps, the
//   adaptive ISV value takes over in experience_kernels.cu. Value 30 bars
//   = MFT-leaning fallback ("between HFT and MFT" user-profile target band).
//   Class A P0-C (2026-05-08): per feedback_isv_for_adaptive_bounds.
//
// 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} — DELETED by Class A audit-fix
//   Batch 4-B (2026-05-08). The epoch-driven anneal schedule
//   `T_end + (T_start - T_end) × exp(-epoch / decay)` was replaced with
//   the signal-driven ISV slot `MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460`
//   (sp14_isv_slots.rs). Driving signal: dir_acc skill from
//   ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]. Per
//   `feedback_no_legacy_aliases.md` — once the ISV-driven version is
//   wired, the legacy schedule constants are deleted (no half-migrated
//   state). The single epoch-driven entry point
//   `min_hold_temperature_for_epoch` in training_loop.rs is also
//   deleted; the training loop now reads ISV[460] (with cold-start
//   fallback to MIN_HOLD_TEMPERATURE_START_FALLBACK=50.0) when seeding
//   the kernel-passed scalar.
//
// 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
// SP21 T2.4 (2026-05-10): MIN_HOLD_TARGET and MIN_HOLD_PENALTY_MAX
// #defines DELETED. The SP12 v3 min-hold-penalty production callers
// were removed in SP20 Phase 2 Task 2.2 (per `experience_kernels.cu:2118-2139`
// — "the SP12 v3 min-hold-penalty args were DELETED atomically with the
// SP12 v3 reward block replacement"). The 4-quadrant SP20 reward at
// trade close has no patience-via-soft-penalty notion — patience emerges
// from advantage-centered alpha. The `compute_min_hold_penalty` device
// function in `trade_physics.cuh` is RETAINED (still called by
// `sp12_reward_math_test_kernel.cu` with literal values), but no
// production caller. Per `feedback_no_legacy_aliases`.
// Cold-start fallback for the kernel-passed `min_hold_temperature` scalar
// when ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460] is at sentinel.
// Bit-identical to pre-Batch-4B epoch-0 behavior under the deleted
// MIN_HOLD_TEMPERATURE_START=50.0 anchor.
#define MIN_HOLD_TEMPERATURE_FALLBACK 50.0f

// Class A P0-A (2026-05-08) — adaptive REWARD_POS/NEG_CAP ISV slots.
// REWARD_POS_CAP / REWARD_NEG_CAP above remain as Phase-1 cold-start
// fallback values; consumers read the ISV slots first and fall back to
// the macros when the slot is at sentinel (bit-identical pre-P0-A
// behavior until the first valid observation). Slot indices mirror
// `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` (locked by
// `sp14_p0a_reward_cap_slot_layout_locked` test). Producer kernel:
// `reward_cap_update_kernel.cu`; sentinels match the macros above for
// bit-identical bootstrap. Per Class A audit, the highest-suspected-
// impact fix for the months-long WR-stuck-at-46-48% plateau across
// 11 superprojects — controller signals (sharpe EMA, var_q, q_gap)
// derive from the CAPPED reward buffer, so the controller cannot
// select for trades it cannot see. Per
// `pearl_audit_unboundedness_for_implicit_asymmetry`, the 2:1
// asymmetry stays — moved from hardcoded scalar to producer-time
// multiplier (single source of truth).
#define REWARD_POS_CAP_ADAPTIVE_INDEX 452
#define REWARD_NEG_CAP_ADAPTIVE_INDEX 453
#define SENTINEL_REWARD_POS_CAP        5.0f
#define SENTINEL_REWARD_NEG_CAP      -10.0f
// Defensive consumer-side bounds (mirror Rust constants
// REWARD_POS_CAP_MIN/MAX in sp14_isv_slots.rs). Values outside this
// window fall back to the macro on the consumer side as a defense
// against malformed prior state. Producer kernel clamps to the same
// range, so a healthy ISV slot will always pass the guard.
#define REWARD_POS_CAP_MIN_BOUND       1.0f
#define REWARD_POS_CAP_MAX_BOUND      50.0f

// Class A P1-Producer (2026-05-08) — adaptive Bayesian Kelly priors.
// Replaces hardcoded `prior_wins=2.0f / prior_losses=2.0f /
// prior_sum_wins=0.01f / prior_sum_losses=0.01f` in kelly_cap_update_kernel
// and trade_physics.cuh::kelly_position_cap with ISV-driven slow-EMA values
// fed by the `kelly_bayesian_priors_update_kernel` producer at per-fold
// boundary (per-epoch boundary in current scheduling — same cadence as
// kelly_cap_update). Sentinels match the pre-P1-Producer hardcoded values
// so cold-start (pre-first-observation) is bit-identical to the prior
// behavior. Mirrors `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs`
// constants of the same names (locked by
// `sp14_p1_kelly_prior_slot_layout_locked` test).
#define KELLY_PRIOR_WINS_INDEX            454
#define KELLY_PRIOR_LOSSES_INDEX          455
#define KELLY_PRIOR_SUM_WINS_INDEX        456
#define KELLY_PRIOR_SUM_LOSSES_INDEX      457
#define SENTINEL_KELLY_PRIOR_WINS         2.0f
#define SENTINEL_KELLY_PRIOR_LOSSES       2.0f
#define SENTINEL_KELLY_PRIOR_SUM_WINS     0.01f
#define SENTINEL_KELLY_PRIOR_SUM_LOSSES   0.01f
// Cold-start fallback for the consumer-side ISV-NULL guard (mirrors the
// existing `kelly_warmup_floor_sp9` and `kelly_f_smooth` patterns in
// trade_physics.cuh). When `isv_signals_ptr == NULL` (validation envs that
// don't pass the bus through) OR the slot value is at sentinel, the
// consumer falls back to the static defaults below. Bit-identical
// pre-P1-Producer behavior.
#define KELLY_PRIOR_WINS_DEFAULT          2.0f
#define KELLY_PRIOR_LOSSES_DEFAULT        2.0f
#define KELLY_PRIOR_SUM_WINS_DEFAULT      0.01f
#define KELLY_PRIOR_SUM_LOSSES_DEFAULT    0.01f

// Class A audit-fix Batch 4-A (2026-05-08) — adaptive DD saturation floor.
// Replaces the hardcoded `0.25f` in `trade_physics.cuh::apply_margin_cap`
// with an ISV-driven adaptive bound derived from the realized per-env
// DD_MAX distribution. Producer kernel: `dd_saturation_floor_update_kernel`
// (per-epoch boundary cadence). Sentinel matches the pre-fix hardcoded
// value (0.25) for bit-identical cold-start. Consumer at the call site
// uses the standard sentinel-detect + range-guard pattern (mirrors the
// SP12-cap and Kelly-prior consumer patterns) to fall back to
// DD_SATURATION_FLOOR_DEFAULT when the slot is at sentinel or
// out-of-range. Mirrors `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs`
// constants of the same names (locked by
// `sp14_audit_4a_dd_saturation_floor_slot_layout_locked` test).
#define DD_SATURATION_FLOOR_ADAPTIVE_INDEX 458
#define SENTINEL_DD_SATURATION_FLOOR      0.25f
// Defensive consumer-side bounds (mirror Rust constants
// DD_SATURATION_FLOOR_MIN/MAX in sp14_isv_slots.rs). Values outside this
// window fall back to the macro on the consumer side as a defense
// against malformed prior state. Producer kernel clamps to the same
// range, so a healthy ISV slot will always pass the guard.
#define DD_SATURATION_FLOOR_MIN_BOUND     0.10f
#define DD_SATURATION_FLOOR_MAX_BOUND     0.50f
// Cold-start fallback for the consumer-side ISV-NULL guard (mirrors the
// SP12-cap fall-back pattern in `sp15_apply_sp12_cap`). When
// `isv_signals_ptr == NULL` OR the slot value is at sentinel within EPS
// OR out of [MIN_BOUND, MAX_BOUND], the consumer falls back to this
// static default. Bit-identical pre-Batch-4A behavior.
#define DD_SATURATION_FLOOR_DEFAULT       0.25f

// Class A audit-fix Batch 4-B (2026-05-08) — adaptive plan_threshold floor.
// Replaces the hardcoded `0.1f` lower bound in
// `plan_threshold_update_kernel.cu:44` (`fmaxf(0.1f, 0.5f * ema)`). The
// floor is INLINE-produced by the same kernel — on each invocation the
// kernel writes both the new plan_threshold (slot 49) and a slow-EMA
// floor (slot 459) and uses the floor as the lower bound on the next
// invocation. Sentinel matches the pre-fix value (0.1) for bit-identical
// cold-start. Mirrors `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs`
// constants of the same names (locked by
// `sp14_audit_4b_plan_threshold_floor_slot_layout_locked` test).
#define PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX 459
#define SENTINEL_PLAN_THRESHOLD_FLOOR     0.1f
#define PLAN_THRESHOLD_FLOOR_MIN_BOUND    0.05f
#define PLAN_THRESHOLD_FLOOR_MAX_BOUND    0.50f
#define PLAN_THRESHOLD_FLOOR_DEFAULT      0.1f
#define PLAN_THRESHOLD_FLOOR_EMA_ALPHA    0.005f

// Class A audit-fix Batch 4-B (2026-05-08) — adaptive MIN_HOLD_TEMPERATURE.
// Replaces the deleted epoch-driven schedule
// `T_end + (T_start - T_end) × exp(-epoch / decay)` with an ISV-driven
// signal-driven temperature derived from the model's directional
// accuracy skill (ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]). Producer kernel:
// `min_hold_temperature_update_kernel` (per-epoch boundary cadence).
// Sentinel matches the pre-fix MIN_HOLD_TEMPERATURE_START=50 for
// bit-identical cold-start (first epoch under both regimes had T=50).
// Mirrors `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` constants of
// the same names (locked by
// `sp14_audit_4b_min_hold_temperature_slot_layout_locked` test).
#define MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX 460
#define SENTINEL_MIN_HOLD_TEMPERATURE     50.0f
#define MIN_HOLD_TEMPERATURE_MIN_BOUND     5.0f
#define MIN_HOLD_TEMPERATURE_MAX_BOUND    50.0f
#define MIN_HOLD_TEMPERATURE_DEFAULT      50.0f
// SP16 T3 (2026-05-08): the hardcoded `MIN_HOLD_TEMPERATURE_EMA_ALPHA = 0.05f`
// constant was deleted. α is now Wiener-optimal, derived from the Welford
// accumulators in slots [468..474) per `pearl_wiener_optimal_adaptive_alpha`.
// See the SP16 T3 SLOTS section above for the slot indices.
#define DIR_ACC_RANDOM_BASELINE            0.5f

// ── SP18 (2026-05-08) — combined Hold-attractor fix: D-leg + B-leg ────
// 22 ISV slots [483..505) split across two atomic legs per the v2
// combined plan. Mirrors `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs`
// constants of the same names (locked by `sp18_combined_slot_layout_locked`
// test). Like SP16/SP17 mirrors, the C-side header tracks slot indices
// + sentinels + bounds for kernels that read the bus directly via raw
// indices; the Rust constant is the single source of truth (see
// ISV_TOTAL_DIM macro note above).
//
// D-leg ("structural Hold opportunity-cost reward") replaces the SP13
// P0a + SP16 P2 + SP16 T3 reactive per-bar Hold-cost chain (DD7=c) with
// a structural reward at the 3 Hold-cost subtraction sites in
// experience_kernels.cu via the new device function
// `compute_sp18_hold_opportunity_cost` (added to trade_physics.cuh in
// Phase 2). HOLD_REWARD_POS_CAP / HOLD_REWARD_NEG_CAP are adaptive bounds
// producer-tracked from p99(|step_ret|) over Long/Short trade closes ×
// HOLD_REWARD_CAP_SAFETY_FACTOR=1.5 (DD3=b) with NEG = -2 × POS
// (DD5=b mirrored loss-aversion).
//
// B-leg ("TD(λ) Q(s') bootstrap") replaces the q_next = rewards
// self-bootstrap at gpu_experience_collector.rs:4143 with a real
// per-branch DDQN `Q_target(s', argmax_a Q_online(s', a))` computed via
// target-net forward pass on next_states at experience-collection time
// (Phase 4-5). All B-leg slots are observability-only (HEALTH_DIAG
// diagnostics) plus the one-shot POPART_RESET_FLAG that gates the
// per-fold PopArt EMA reset at the SP18 deployment boundary (B-DD11).
//
// Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md

// D-leg slots [483..493) — 10 slots
#define HOLD_REWARD_POS_CAP_INDEX             483
#define HOLD_REWARD_NEG_CAP_INDEX             484
#define HOLD_REWARD_DECOMP_DIAG_INDEX         485
#define HOLD_OPP_COST_FIRE_RATE_EMA_INDEX     486
#define HRC_TARGET_MEAN_INDEX                 487
#define HRC_TARGET_M2_INDEX                   488
#define HRC_DIFF_MEAN_INDEX                   489
#define HRC_DIFF_M2_INDEX                     490
#define HRC_PREV_TARGET_INDEX                 491
#define HRC_SAMPLE_COUNT_INDEX                492

// D-leg sentinels — Pearl-A first-observation bootstrap. POS=5.0 /
// NEG=-10.0 match the SP14 P0-A REWARD_POS/NEG_CAP_ADAPTIVE pattern
// (bit-identical cold-start to the position-side adaptive cap).
#define SENTINEL_HOLD_REWARD_POS_CAP            5.0f
#define SENTINEL_HOLD_REWARD_NEG_CAP          -10.0f
#define SENTINEL_HOLD_REWARD_DIAG               0.0f
#define SENTINEL_HOLD_OPP_FIRE_RATE             0.0f

// D-leg defensive consumer-side bounds (mirror Rust constants
// HOLD_REWARD_POS_CAP_MIN/MAX_BOUND in sp14_isv_slots.rs). Values
// outside this window fall back to sentinels on the consumer side as a
// defense against malformed prior state. Producer kernel clamps to the
// same range, so a healthy ISV slot will always pass the guard.
#define HOLD_REWARD_POS_CAP_MIN_BOUND           0.5f
#define HOLD_REWARD_POS_CAP_MAX_BOUND          50.0f

// D-leg producer constant. Safety factor 1.5× p99 mirrors the SP14 P0-A
// REWARD_POS_CAP producer pattern.
#define HOLD_REWARD_CAP_SAFETY_FACTOR           1.5f

// B-leg slots [493..505) — 12 slots
#define TD_ERROR_MAG_EMA_INDEX                493
#define Q_NEXT_TARGET_P99_INDEX               494
#define Q_NEXT_MINUS_REWARD_P99_INDEX         495
#define V_SHARE_TREND_DIAG_INDEX              496
#define POPART_RESET_FLAG_INDEX               497
#define TDB_TARGET_MEAN_INDEX                 498
#define TDB_TARGET_M2_INDEX                   499
#define TDB_DIFF_MEAN_INDEX                   500
#define TDB_DIFF_M2_INDEX                     501
#define TDB_PREV_TARGET_INDEX                 502
#define TDB_SAMPLE_COUNT_INDEX                503
#define SP18_B_LEG_RESERVED_INDEX             504

// B-leg sentinels — Pearl-A first-observation bootstrap. Diagnostic
// slots bootstrap to 0.0 (direct-replace on first valid observation).
// POPART_RESET_FLAG sentinel = 1.0 (one-shot — gates the per-fold
// PopArt slot 63 EMA reset on the first SP18 epoch per B-DD11; host
// writes 1.0 at deploy, kernel zeroes after consuming).
#define SENTINEL_TD_ERROR_MAG_EMA              0.0f
#define SENTINEL_Q_NEXT_TARGET_P99             0.0f
#define SENTINEL_Q_NEXT_MINUS_REWARD_P99       0.0f
#define SENTINEL_V_SHARE_TREND_DIAG            0.0f
#define SENTINEL_POPART_RESET_FLAG             1.0f

// ── SP18 v2 (2026-05-09) — ISV-adaptive shrink-and-perturb at fold transition ──
// 2 ISV slots [505..507) drive `reset_for_fold`'s `shrink_and_perturb`
// call replacing prior hardcoded `α=0.8 / σ=0.01` constants per
// `feedback_isv_for_adaptive_bounds`. Mirrors
// `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` constants of the same
// names (locked by `sp18_shrink_perturb_slot_layout_locked` test). The
// kernel `weight_drift_diag_kernel.cu` writes both slots in the same
// launch as the existing drift diagnostic — single producer, atomic
// per-epoch. Pearl-A first-observation bootstrap sentinels match prior
// hardcoded values for bit-identical cold-start (first fold, no
// `best_params_snapshot` → host skips kernel → slots stay at
// sentinels). Bounds Category-1 dimensional floors per
// `feedback_isv_for_adaptive_bounds`. Bilateral clamp per
// `pearl_symmetric_clamp_audit`.
#define SHRINK_ALPHA_ADAPTIVE_INDEX           505
#define SHRINK_SIGMA_ADAPTIVE_INDEX           506

#define SENTINEL_SHRINK_ALPHA                  0.8f
#define SENTINEL_SHRINK_SIGMA                  0.01f

#define SHRINK_ALPHA_MIN                       0.5f
#define SHRINK_ALPHA_MAX                       0.95f
#define SHRINK_SIGMA_MIN                       1.0e-4f
#define SHRINK_SIGMA_MAX                       0.1f

// Producer reference scale (Category-1 dimensional anchor — typical
// mean RMS of healthy stably-trained DQN weights). At
// `||best||_RMS = 0.5`, σ = 0.01 = SENTINEL — equivalence with prior
// hardcoded value for typical healthy weights.
#define SHRINK_SIGMA_RMS_REFERENCE             0.5f

// ── SP19 (2026-05-09) — producer-side multi-horizon reward blend weights ──
// 3 ISV slots [507..510) RESERVED for future controller-driven multi-horizon
// reward blending. Path (B) of the SP19 spec applies the blend at fxcache
// write time (`precompute_features.rs` / `data_loading.rs`) using a hardcoded
// equal-thirds default; these slots are slot reservations, NOT yet read by
// any kernel consumer. Future Path (A) refactor would bump TARGET_DIM to
// carry per-horizon log-returns through the kernel pipeline.
//
// Mirrors `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` constants of the
// same names (locked by `sp19_reward_horizon_slot_layout_locked` test).
// Sentinels match the producer-time hardcoded `SP19_HORIZON_BLEND_WEIGHTS`
// (1/3 each) for bit-identical cold-start with the on-disk fxcache.
#define REWARD_HORIZON_WEIGHT_1BAR_INDEX      507
#define REWARD_HORIZON_WEIGHT_5BAR_INDEX      508
#define REWARD_HORIZON_WEIGHT_30BAR_INDEX     509

#define SENTINEL_REWARD_HORIZON_WEIGHT_DEFAULT (1.0f / 3.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],
    const float  aux_dir_prob
) {
    // 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) — SP22 H6 Phase 2 (2026-05-12): slot 0 = aux
    // directional probability from previous step's aux softmax, RECENTERED
    // to `2*p_up - 1 ∈ [-1, +1]` (matches every other state slot's
    // "no signal = 0" baseline per `pearl_first_observation_bootstrap`).
    // Slots 1..7 zero for 8-alignment. Caller passes 0.0 sentinel at
    // cold-start / FoldReset and for eval-time NULL fallback (A3).
    out[SL_PADDING_START + 0] = aux_dir_prob;
    for (int k = 1; k < SL_PADDING_DIM; k++)
        out[SL_PADDING_START + k] = 0.0f;
}

// ────────────────────────────────────────────────────────────────────────────
// SP22 H6 vNext Phase C-2 (2026-05-14) — K=3 trade-outcome state assembly
// ────────────────────────────────────────────────────────────────────────────
//
// Three-slot variant of `assemble_state` that writes the K=3 trade-outcome
// head's softmax probabilities into the padding region instead of the K=2
// head's single recentered p_up scalar. Slot layout:
//
//   state[SL_PADDING_START + 0 = 121] = p_Profit   ∈ [0, 1]
//   state[SL_PADDING_START + 1 = 122] = p_Stop     ∈ [0, 1]
//   state[SL_PADDING_START + 2 = 123] = p_Timeout  ∈ [0, 1]
//   state[SL_PADDING_START + 3..7]    = 0.0  (remaining padding for 8-align)
//
// Encoding rationale (vs K=2 head's recentered `2*p - 1`):
//   - Softmax outputs are inherently in [0, 1] and sum to 1 per row
//   - Three-slot probabilistic representation is more informative than the
//     K=2 head's single "p_up − p_down" scalar
//   - Cold-start sentinel = 0.0 for all three slots = "no Profit/Stop/Timeout
//     prediction yet" (e.g., before the first rollout step or under
//     producer-unwired conditions)
//
// Reuses every other slot in the state layout (market, OFI, MTF, portfolio,
// plan_isv) exactly as `assemble_state` does — only the padding region's
// content differs. Mirrors the standard `assemble_state` signature except
// for the trailing 3-element outcome probs vector replacing the scalar.
__device__ __forceinline__ void assemble_state_outcome_k3(
    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],
    const float  aux_outcome_probs[3]
) {
    // 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
    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) — K=3 trade-outcome head's 3-slot softmax probs.
    // Replaces the K=2 head's single-slot `2*p_up - 1` encoding at slot 121.
    // Slots [124..128) zero for 8-alignment.
    out[SL_PADDING_START + 0] = aux_outcome_probs[0];  // p_Profit
    out[SL_PADDING_START + 1] = aux_outcome_probs[1];  // p_Stop
    out[SL_PADDING_START + 2] = aux_outcome_probs[2];  // p_Timeout
    for (int k = 3; k < SL_PADDING_DIM; k++)
        out[SL_PADDING_START + k] = 0.0f;
}
