// crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu
//
// SP11 (Fix 39, 2026-05-04): per-component reward-magnitude-ratio canary
// producer with z-score normalisation (B1b smoke-recovery 2026-05-04).
//
// B1b fix-up (2026-05-04) split the popart axis off slot 63 (which is
// PopArt's normalisation input = total reward magnitude EMA; pre-SP11
// invariant preserved) onto a dedicated slot 360 =
// `POPART_COMPONENT_MAG_EMA_INDEX` produced by `popart_component_ema_kernel`.
// B1b smoke-recovery (2026-05-04) added per-component variance EMAs at
// slots 361..367 to drive z-score normalisation. This kernel reads
// non-contiguous magnitude sources after the fix-up:
//
//   axis 0 (popart)   = isv[popart_specific_slot]      ← slot 360 (B1b fix-up)
//   axis 1 (cf)       = isv[cf_others_base_slot + 0]   ← slot 64
//   axis 2 (trail)    = isv[cf_others_base_slot + 1]   ← slot 65
//   axis 3 (micro)    = isv[cf_others_base_slot + 2]   ← slot 66
//   axis 4 (opp_cost) = isv[cf_others_base_slot + 3]   ← slot 67
//   axis 5 (bonus)    = isv[cf_others_base_slot + 4]   ← slot 68
//
// SP22 H6 Phase 3 (2026-05-13): aux_align (component idx 6 in the
// 7-stride reward_components_per_sample buffer) is INTENTIONALLY
// excluded from this kernel's 6-axis ratio computation. The aux_align
// EMA lives at ISV[REWARD_AUX_ALIGN_EMA_INDEX=536] (non-contiguous
// with cf_others_base_slot at 64..68), and the SP11 controller's
// `w_aux_align` adaptive weight has its own anchor on that slot
// directly (see reward_subsystem_controller_kernel.cu Phase 3 extension).
// Keeping aux_align out of this kernel's ratio preserves the
// pre-Phase-3 6-axis mag/var z-score machinery bit-for-bit; only the
// SP11 controller's emit list gains the 7th weight.
//
// And per-component variance sources (B1b smoke-recovery):
//
//   axis 0 var = isv[popart_var_slot]                  ← slot 361
//   axis 1 var = isv[cf_others_var_base_slot + 0]      ← slot 362
//   ...
//   axis 5 var = isv[cf_others_var_base_slot + 4]      ← slot 366
//
// Why z-score (B1b smoke-recovery, spec §4 amendment lines 564-619):
// the linear ratio formula `winner_weight = mag_ratio` amplified
// whichever reward component had the largest intrinsic magnitude.
// Popart (trade P&L) is O(100) per fire; cf/trail/micro/opp_cost/bonus
// are O(0.1-2). Result: `w_pop` saturated toward MAX_WEIGHT in the
// controller, curiosity_b exploded, sharpe collapsed (B1b smoke
// `smoke-test-6wd2c` on commit `61b2fa962`). Z-score normalisation
// makes ratios scale-invariant — every component contributes "shares
// of own scale" rather than "shares of absolute magnitude". The
// controller's redistribution logic (`winner_weight = ratio`,
// unchanged) works correctly when it sees ratios ≈ 1/6 each at
// stability rather than ratios skewed by intrinsic O(100) vs O(1)
// scale differences.
//
// Why two slot-pair indices: the popart axis lives at non-contiguous
// magnitude (360) and variance (361) slots, while the other 5 axes
// have contiguous magnitude (64..68) and variance (362..366) slots.
// Pre-fix-up the kernel read 6 contiguous mag slots at slot 63..69 —
// see the popart_component_ema_kernel.cu header for the slot-63
// overload diagnosis. The cf/trail/micro/opp_cost/bonus axes still
// share their producer (`reward_component_ema_kernel`); only popart
// is sourced from a different magnitude+variance slot pair.
//
// Writes 7 floats to `scratch_out`:
//   scratch_out[0..6) = ratios (each = z[c] / Σz, Σ=1 modulo EPS_DIV floor),
//                       where z[c] = mag[c] / max(sqrt(var[c]), EPS_DIV).
//   scratch_out[6]    = mag[0] (popart raw magnitude) — mirrored to
//                       PNL_REWARD_MAGNITUDE_EMA_INDEX = 359 via the
//                       chained Pearls A+D launch (n_slots=1).
//
// Mirror semantic preserved post-z-score: scratch_out[6] stays as raw
// `mag[0]` (popart magnitude EMA, slot 360 source) because slot 359
// = `PNL_REWARD_MAGNITUDE_EMA_INDEX` is consumed by curiosity bounding
// (spec §3.4.2) and saboteur engagement thresholding (spec §3.3.1) —
// both need raw P&L scale, NOT z-score. Z-scoring is purely the
// normalisation step inside the ratio computation; the mirror feeds
// downstream consumers that need physical scale information.
//
// Pearls A+D applied chain-style by `apply_pearls_ad_kernel` (n_slots=6
// for the contiguous ratio block at ISV[REWARD_COMPONENT_MAG_RATIO_BASE
// =352..358), then n_slots=1 for the singleton mirror at ISV[359]) by
// the launcher `launch_sp11_mag_ratio_compute` in `gpu_dqn_trainer.rs`.
// The kernel emits all 7 scratch values in one shot.
//
// Single block, 6 threads (one per component); thread 0 also computes
// the sum and writes the popart mirror.

#include <cuda_runtime.h>

extern "C" __global__ void reward_component_mag_ratio_compute_kernel(
    /* ISV signal bus (read-only). */
    const float* __restrict__ isv,
    /* Producer scratch buffer. Caller indexes the base into this array
     * via the `SCRATCH_SP11_MAG_RATIO_BASE` constant via pointer
     * arithmetic in the launcher. */
    float* __restrict__ scratch_out,
    /* ISV slot index of popart-component-specific magnitude EMA
     * = POPART_COMPONENT_MAG_EMA_INDEX = 360 (B1b fix-up, 2026-05-04). */
    int popart_specific_slot,
    /* ISV slot index of the first non-popart per-component magnitude EMA
     * = REWARD_POPART_EMA_INDEX + 1 = 64 (cf at 64, trail 65, micro 66,
     * opp_cost 67, bonus 68 — produced by SP4 reward_component_ema_kernel). */
    int cf_others_base_slot,
    /* ISV slot index of popart-component-specific variance EMA
     * = REWARD_COMPONENT_VAR_EMA_BASE = 361 (B1b smoke-recovery, 2026-05-04). */
    int popart_var_slot,
    /* ISV slot index of the first non-popart per-component variance EMA
     * = REWARD_COMPONENT_VAR_EMA_BASE + 1 = 362 (cf var at 362, trail 363,
     * micro 364, opp_cost 365, bonus 366 — produced by SP4 reward_component_ema
     * extended in B1b smoke-recovery). */
    int cf_others_var_base_slot
) {
    if (blockIdx.x != 0) return;
    if (threadIdx.x >= 6) return;

    /* Block-load all 6 component magnitudes AND variances into shared
     * mem. The popart axis (component 0) is sourced from a separate
     * ISV slot than the rest — see the header for the B1b-fix-up
     * rationale. The variance pair is identically split (slot 361 vs
     * slots 362..366). */
    __shared__ float mag[6];
    __shared__ float var[6];
    __shared__ float sum_z;

    if (threadIdx.x == 0) {
        mag[0] = isv[popart_specific_slot];
        var[0] = isv[popart_var_slot];
    } else {
        /* Components 1..5 → ISV[cf_others_base_slot + (c - 1)] for mag,
         * ISV[cf_others_var_base_slot + (c - 1)] for var. */
        mag[threadIdx.x] = isv[cf_others_base_slot     + (int)threadIdx.x - 1];
        var[threadIdx.x] = isv[cf_others_var_base_slot + (int)threadIdx.x - 1];
    }
    __syncthreads();

    /* Z-score per component. EPS_DIV (1e-6) is the Invariant-1
     * numerical-stability anchor matching SP4/SP5/SP9/SP10/SP11
     * conventions. Guards the impossible "all-zero variance" case
     * during the first observations after a fold reset (Pearl A
     * bootstraps each component variance EMA from its first
     * observation, but timing of Pearls A+D fires across the 6
     * sources is per-slot, so a literal 0/0 is conceivable in the
     * cold-start window — and on a constant input where var = 0
     * exactly). When var[c] = 0 the floor sets σ = EPS_DIV ⇒
     * z[c] = mag[c] × 1e6 (huge but well-defined); since all 6
     * z[c] scale linearly with the same EPS_DIV factor when all
     * variances are zero, the ratios remain proportional and well-
     * defined (each component still contributes its share of the
     * total z-score envelope). */
    const float EPS_DIV = 1.0e-6f;
    float z_local = mag[threadIdx.x] / fmaxf(sqrtf(var[threadIdx.x]), EPS_DIV);

    /* Block-publish per-thread z so thread 0 can compute the sum. We
     * reuse the `mag` shared array as a temporary z-score store —
     * `mag` is no longer needed past this point. */
    mag[threadIdx.x] = z_local;
    __syncthreads();

    if (threadIdx.x == 0) {
        float s = 0.0f;
        #pragma unroll
        for (int c = 0; c < 6; c++) s += mag[c];
        /* EPS_DIV (1e-6) floor on the sum — same Invariant-1 anchor.
         * Guards the all-zero-z case (which can't happen given σ-floor
         * above, but keeps the formula structurally robust). */
        sum_z = fmaxf(s, EPS_DIV);
        /* Side-output mirror: popart-component RAW magnitude (not the
         * z-score) → ISV[359] via chained Pearls A+D singleton launch.
         * Slot 359 = `PNL_REWARD_MAGNITUDE_EMA_INDEX` is consumed by
         * curiosity bounding + saboteur engagement thresholds — both
         * need physical P&L scale, not z-score. Read directly from
         * `var[0]`'s old neighbour `isv[popart_specific_slot]` — but
         * we no longer have it as a shared (mag[0] is now z[0]). Re-
         * read from ISV; thread 0 is the only writer to scratch_out[6]. */
        scratch_out[6] = isv[popart_specific_slot];
    }
    __syncthreads();

    /* Each of threads 0..5 writes its own ratio. */
    scratch_out[threadIdx.x] = mag[threadIdx.x] / sum_z;

    if (threadIdx.x == 0) {
        __threadfence_system();
    }
}
