/* health_diag_kernel — kernel family that fills `HealthDiagSnapshot` (the
 * single mapped-pinned struct read by the per-epoch HEALTH_DIAG emit
 * path; see `health_diag.rs`) directly from device-resident producers.
 *
 * Phase 2A — `health_diag_isv_mirror`. This is the simplest member of the
 * family: pure scalar copies from the ISV signal bus into the
 * corresponding snapshot fields. No reductions, no masking, no
 * per-sample loops. Single-block, single-thread launch (the work is
 * ~22 unaligned global loads + stores; spawning a wider block would
 * waste every thread except thread 0). Mirrors the footprint of
 * `aux_heads_loss_ema_update` in `aux_heads_loss_ema_kernel.cu`.
 *
 * The snapshot layout is `#[repr(C)]` with every field 4-byte aligned
 * (every f32, u32, or array thereof — see `health_diag.rs`'s
 * `snapshot_size_is_stable` test which pins the size to 147 × 4 bytes).
 * The kernel addresses the snapshot through a `float*` base pointer plus
 * compile-time constant word indices. **The Rust caller and this
 * kernel both maintain the same field-index table — adding/reordering
 * fields requires updating both sites and bumping the size assertion**
 * (host-side: `snapshot_size_is_stable`; kernel-side: the `STATIC_ASSERT`
 * of `WORD_TOTAL` below).
 *
 * Architectural rules upheld:
 *   - No HtoD / DtoH (writes through mapped-pinned device pointer).
 *   - No `atomicAdd` (single thread does all writes).
 *   - No DtoD memcpy (kernel-load + kernel-store via pointers, not
 *     `cudaMemcpy*`).
 *   - Mapped-pinned coherence covers visibility on the host after the
 *     stream sync the producer-family launchers already perform via the
 *     captured-graph fence; this kernel does not call
 *     `__threadfence_system()` itself — the *finalise* kernel (the last
 *     in the family, Phase 2E) emits the single fence covering all
 *     writes from earlier members. This keeps the system fence count at
 *     one per epoch, regardless of how the family fans out.
 */

/* Word indices into the snapshot. MUST match the Rust struct
 * `HealthDiagSnapshot` field order in `health_diag.rs` (every field 4
 * bytes; arrays count as their length). Updated only when the struct
 * itself changes — review the struct + this header + the
 * `snapshot_size_is_stable` test together. */

/* Health aggregator (8 words, [0..8)). */
#define WORD_HEALTH                  0
#define WORD_COMP_Q_GAP              1
#define WORD_COMP_Q_VAR              2
#define WORD_COMP_ATOMS              3
#define WORD_COMP_GRAD_STABLE        4
#define WORD_COMP_ENS_AGREE          5
#define WORD_COMP_GRAD_COS           6
#define WORD_COMP_SPECTRAL           7

/* Effective hyperparams (8 words, [8..16)). */
#define WORD_EFF_CQL_ALPHA           8
#define WORD_EFF_IQN_BUDGET          9
#define WORD_EFF_CQL_BUDGET         10
#define WORD_EFF_C51_BUDGET         11
#define WORD_EFF_TAU                12
#define WORD_EFF_SARSA_TAU          13
#define WORD_EFF_GAMMA              14
#define WORD_EFF_CF_RATIO           15

/* Mechanism state (7 words, [16..23)). */
#define WORD_NOVEL_DISTILL_ON       16  /* u32 */
#define WORD_NOVEL_BARRIER          17
#define WORD_NOVEL_PLASTICITY_READY 18  /* u32 */
#define WORD_NOVEL_IB               19
#define WORD_NOVEL_ENSEMBLE_COLL    20
#define WORD_NOVEL_CONTRARIAN_ON    21  /* u32 */
#define WORD_NOVEL_META_Q_PRED      22

/* Diag (2 words, [23..25)). */
#define WORD_DIAG_SHARPE_EMA        23
#define WORD_DIAG_ACTION_ENTROPY    24

/* Gems (1 word, [25..26)). */
#define WORD_GEM_G12_PREDICTIVE     25

/* Magnitude block (10 words, [26..36)). */
#define WORD_MAG_Q_FULL             26
#define WORD_MAG_Q_HALF             27
#define WORD_MAG_Q_QUARTER          28
#define WORD_MAG_VAR_SCALE          29
#define WORD_MAG_KELLY_F            30
#define WORD_MAG_AVG_WIN_RATIO      31
#define WORD_MAG_GRAD_RATIO_MAG_DIR 32
#define WORD_MAG_DIST_Q             33
#define WORD_MAG_DIST_H             34
#define WORD_MAG_DIST_F             35

/* Grad split bwd (4 words, [36..40)). */
#define WORD_GRAD_SPLIT_BWD_IQN     36
#define WORD_GRAD_SPLIT_BWD_CQL     37
#define WORD_GRAD_SPLIT_BWD_C51     38
#define WORD_GRAD_SPLIT_BWD_ENS     39

/* Grad split aux (5 words, [40..45)). */
#define WORD_GRAD_SPLIT_AUX_DISTILL 40
#define WORD_GRAD_SPLIT_AUX_REC     41
#define WORD_GRAD_SPLIT_AUX_PRED    42
#define WORD_GRAD_SPLIT_AUX_CQL_SX  43
#define WORD_GRAD_SPLIT_AUX_C51_BS  44

/* Grad trunk (9 words, [45..54)). */
#define WORD_GRAD_TRUNK_BASE        45

/* Grad abs (2 words, [54..56)). */
#define WORD_GRAD_ABS_DIR           54
#define WORD_GRAD_ABS_MAG           55

/* Trail rates + holds (6 words, [56..62)). */
#define WORD_TRAIL_FIRE_Q           56
#define WORD_TRAIL_FIRE_H           57
#define WORD_TRAIL_FIRE_F           58
#define WORD_TRAIL_HOLD_Q           59
#define WORD_TRAIL_HOLD_H           60
#define WORD_TRAIL_HOLD_F           61

/* Per-magnitude winrate + variance (6 words, [62..68)). */
#define WORD_MAGSTAT_WR_Q           62
#define WORD_MAGSTAT_WR_H           63
#define WORD_MAGSTAT_WR_F           64
#define WORD_MAGSTAT_VAR_Q          65
#define WORD_MAGSTAT_VAR_H          66
#define WORD_MAGSTAT_VAR_F          67

/* Noisy/VSN/drift (6 words, [68..74)). */
#define WORD_NOISY_VSN_MAG          68
#define WORD_NOISY_VSN_DIR          69
#define WORD_NOISY_SIGMA_MAG        70
#define WORD_NOISY_SIGMA_DIR        71
#define WORD_NOISY_DRIFT_MAG        72
#define WORD_NOISY_DRIFT_DIR        73

/* Eval dist (3 words, [74..77)). */
#define WORD_EVAL_DIST_Q            74
#define WORD_EVAL_DIST_H            75
#define WORD_EVAL_DIST_F            76

/* Intent dist (3 words, [77..80)). Pure-policy magnitude bucket BEFORE
 * Kelly cap clipping; sibling of eval_dist for separating "policy isn't
 * learning Full" from "Kelly cap is suppressing Full". Per memory pearl
 * `project_magnitude_eval_collapse_kelly_capped.md`. */
#define WORD_INTENT_DIST_Q          77
#define WORD_INTENT_DIST_H          78
#define WORD_INTENT_DIST_F          79

/* Reward contributions (4 words, [80..84)). */
#define WORD_REWARD_CONTRIB_POPART  80
#define WORD_REWARD_CONTRIB_CF      81
#define WORD_REWARD_CONTRIB_TRAIL   82
#define WORD_REWARD_CONTRIB_MICRO   83

/* Controller fire bools + frac (7 words, [84..91)). */
#define WORD_CTRL_FIRE_ANTI_LR      84  /* u32 */
#define WORD_CTRL_FIRE_TAU          85  /* u32 */
#define WORD_CTRL_FIRE_GAMMA        86  /* u32 */
#define WORD_CTRL_FIRE_CLIP         87  /* u32 */
#define WORD_CTRL_FIRE_CQL          88  /* u32 */
#define WORD_CTRL_FIRE_COST         89  /* u32 */
#define WORD_CTRL_FIRE_FRAC         90

/* Explore (3 words, [91..94)). */
#define WORD_EXPLORE_ENT_MAG        91
#define WORD_EXPLORE_ENT_DIR        92
#define WORD_EXPLORE_SIGMA_MEAN     93

/* Validation block (16 words, [94..110)). */
#define WORD_VAL_BASE               94

/* Validation per-direction action distribution (4 words, [110..114)). */
#define WORD_VAL_DIR_DIST_BASE     110

/* Validation per-direction picked distribution (4 words, [114..118)). */
#define WORD_VAL_PICKED_DIR_DIST_BASE 114

/* Reward split (6 words, [118..124)). */
#define WORD_REWARD_SPLIT_POPART   118
#define WORD_REWARD_SPLIT_CF       119
#define WORD_REWARD_SPLIT_TRAIL    120
#define WORD_REWARD_SPLIT_MICRO    121
#define WORD_REWARD_SPLIT_OPP_COST 122
#define WORD_REWARD_SPLIT_BONUS    123

/* Aux block (3 words, [124..127)). SP13 B1.0 (2026-05-05): retired
 * `WORD_AUX_LABEL_SCALE` together with ISV[117] when the next-bar MSE
 * flipped to scale-free; downstream offsets shifted down by 1. */
#define WORD_AUX_NEXT_BAR_MSE      124
#define WORD_AUX_REGIME_CE         125
#define WORD_AUX_WEIGHT            126

/* Aux MoE (10 words, [127..137) — 8 utils + ent + λ_eff). */
#define WORD_AUX_MOE_UTIL_BASE     127
#define WORD_AUX_MOE_ENT           135
#define WORD_AUX_MOE_LAMBDA_EFF    136

/* Action-count histogram (12 words, [137..149), all u32). */
#define WORD_ACTION_COUNTS_BASE    137

#define WORD_TOTAL                 149

/* Pin the layout: matches `snapshot_size_is_stable` on the Rust side. */
static_assert(WORD_TOTAL == 149,
              "HealthDiagSnapshot layout drift — update kernel WORD_* and the "
              "Rust struct in lockstep, then bump snapshot_size_is_stable.");

/* ──────────────────────────────────────────────────────────────────────────
 * health_diag_isv_mirror — Phase 2A.
 *
 * Copies a small set of ISV signal-bus slots into the corresponding snapshot
 * fields. Source contracts (kept in sync with the constants in
 * `gpu_dqn_trainer.rs`):
 *
 *   ISV[63] REWARD_POPART_EMA_INDEX        → snap.reward_split_popart   (word 115)
 *   ISV[64] REWARD_CF_EMA_INDEX            → snap.reward_split_cf       (word 116)
 *   ISV[65] REWARD_TRAIL_EMA_INDEX         → snap.reward_split_trail    (word 117)
 *   ISV[66] REWARD_MICRO_EMA_INDEX         → snap.reward_split_micro    (word 118)
 *   ISV[67] REWARD_OPP_COST_EMA_INDEX      → snap.reward_split_opp_cost (word 119)
 *   ISV[68] REWARD_BONUS_EMA_INDEX         → snap.reward_split_bonus    (word 120)
 *   ISV[87] VSN_MAG_EMA_INDEX              → snap.noisy_vsn_mag         (word 68)
 *   ISV[88] VSN_DIR_EMA_INDEX              → snap.noisy_vsn_dir         (word 69)
 *   ISV[92] TARGET_DRIFT_MAG_EMA_INDEX     → snap.noisy_drift_mag       (word 72)
 *   ISV[93] TARGET_DRIFT_DIR_EMA_INDEX     → snap.noisy_drift_dir       (word 73)
 *   ISV[113] AUX_NEXT_BAR_MSE_EMA_INDEX    → snap.aux_next_bar_mse      (word 124)
 *   ISV[114] AUX_REGIME_CE_EMA_INDEX       → snap.aux_regime_ce         (word 125)
 *   ISV[118..126) MOE_EXPERT_UTIL_EMA_BASE → snap.aux_moe_util[0..8]    (words 127..135)
 *   ISV[126] MOE_GATE_ENTROPY_EMA_INDEX    → snap.aux_moe_ent           (word 135)
 *   ISV[128] MOE_LAMBDA_EFF_INDEX          → snap.aux_moe_lambda_eff    (word 136)
 *
 * SP13 B1.0 (2026-05-05): the former ISV[117] AUX_LABEL_SCALE_EMA_INDEX →
 * snap.aux_label_scale mirror was dropped together with that slot's
 * retirement. All snap word offsets at and beyond the former
 * `WORD_AUX_LABEL_SCALE=127` shifted down by 1.
 *
 * Indices passed in as kernel args (rather than `#define`d) so we don't
 * couple the kernel to a specific numbering — the Rust caller passes
 * `REWARD_POPART_EMA_INDEX`, `MOE_EXPERT_UTIL_EMA_BASE`, etc. from the
 * single source of truth in `gpu_dqn_trainer.rs`. If those indices
 * shift the kernel still works without recompilation.
 *
 * Cost: ~22 global loads + 22 global stores from a single thread.
 * Comfortably under any kernel-launch overhead. Producer-only — the
 * snapshot fields are downstream-consumed by the format-string emit
 * after Phase 4.
 */
extern "C" __global__ void health_diag_isv_mirror(
    const float* __restrict__ isv,         /* ISV bus [ISV_TOTAL_DIM] */
    float*       __restrict__ snap,        /* mapped-pinned HealthDiagSnapshot,
                                              addressable as float* (every
                                              field 4-byte aligned). */
    int reward_popart_idx,                 /* ISV[63] */
    int reward_cf_idx,                     /* ISV[64] */
    int reward_trail_idx,                  /* ISV[65] */
    int reward_micro_idx,                  /* ISV[66] */
    int reward_opp_cost_idx,               /* ISV[67] */
    int reward_bonus_idx,                  /* ISV[68] */
    int vsn_mag_idx,                       /* ISV[87] */
    int vsn_dir_idx,                       /* ISV[88] */
    int target_drift_mag_idx,              /* ISV[92] */
    int target_drift_dir_idx,              /* ISV[93] */
    int aux_next_bar_mse_idx,              /* ISV[113] */
    int aux_regime_ce_idx,                 /* ISV[114] */
    /* SP13 B1.0 (2026-05-05): ISV[117] AUX_LABEL_SCALE retired —
     * argument dropped from kernel signature. */
    int moe_expert_util_base_idx,          /* ISV[118], 8 contiguous slots */
    int moe_gate_entropy_idx,              /* ISV[126] */
    int moe_lambda_eff_idx                 /* ISV[128] */
) {
    /* Single-block, single-thread contract. Mirrors aux_heads_loss_ema's
     * grid guard: the work is ~22 scalar EMA reads + writes. */
    if (blockIdx.x != 0) return;
    if (threadIdx.x != 0) return;

    /* Reward-split block (6 floats — already EMA'd per-component
     * |reward| means; producer is reward_component_ema_kernel). */
    snap[WORD_REWARD_SPLIT_POPART]    = isv[reward_popart_idx];
    snap[WORD_REWARD_SPLIT_CF]        = isv[reward_cf_idx];
    snap[WORD_REWARD_SPLIT_TRAIL]     = isv[reward_trail_idx];
    snap[WORD_REWARD_SPLIT_MICRO]     = isv[reward_micro_idx];
    snap[WORD_REWARD_SPLIT_OPP_COST]  = isv[reward_opp_cost_idx];
    snap[WORD_REWARD_SPLIT_BONUS]     = isv[reward_bonus_idx];

    /* VSN attention focus EMAs (Plan 4 Task 5 Mode A producer:
     * attention_focus_ema_kernel). */
    snap[WORD_NOISY_VSN_MAG]          = isv[vsn_mag_idx];
    snap[WORD_NOISY_VSN_DIR]          = isv[vsn_dir_idx];

    /* Target drift EMAs (Plan 4 follow-up producer: target_drift_kernel). */
    snap[WORD_NOISY_DRIFT_MAG]        = isv[target_drift_mag_idx];
    snap[WORD_NOISY_DRIFT_DIR]        = isv[target_drift_dir_idx];

    /* Aux-head loss EMAs (Plan 4 Task 6 Commit A producer:
     * aux_heads_loss_ema_kernel). SP13 B1.0 (2026-05-05): the third
     * mirror (ISV[117] → snap.aux_label_scale) was retired together with
     * the slot when the next-bar MSE flipped to scale-free. */
    snap[WORD_AUX_NEXT_BAR_MSE]       = isv[aux_next_bar_mse_idx];
    snap[WORD_AUX_REGIME_CE]          = isv[aux_regime_ce_idx];

    /* MoE per-expert utilisation EMA (8 contiguous slots). Producer:
     * moe_expert_util_ema_update in moe_kernels.cu. */
    #pragma unroll
    for (int k = 0; k < 8; ++k) {
        snap[WORD_AUX_MOE_UTIL_BASE + k] = isv[moe_expert_util_base_idx + k];
    }

    /* MoE gate entropy EMA + λ_eff. Producers:
     * moe_expert_util_ema_update (entropy slot 126) and
     * moe_lambda_eff_kernel (slot 128). */
    snap[WORD_AUX_MOE_ENT]            = isv[moe_gate_entropy_idx];
    snap[WORD_AUX_MOE_LAMBDA_EFF]     = isv[moe_lambda_eff_idx];

    /* No __threadfence_system() here — the finalise kernel (Phase 2E)
     * issues a single fence covering all writes from the family. */
}
