foxhunt/crates/ml/tests/sp5_producer_unit_tests.rs

#![allow(unsafe_code)] // CUDA kernel launch + mapped-pinned memory.

//! SP5 Layer A producer kernel unit tests.
//!
//! Tests cover the SP5 Layer A CUDA kernels across Tasks A1–A5:
//!
//!   Task A1 (Pearl 1 atom-span + Q-stats source):
//!     1. `sp5_q_branch_stats_uniform_input_matches_analytical_ground_truth`
//!     2. `pearl_1_atom_update_v_half_tracks_max_abs_dev_with_bootstrap_headroom`
//!   Task A2 (Pearl 3 NoisyNet σ):
//!     3. `pearl_3_sigma_bootstrap_at_target_entropy`
//!     4. `pearl_3_sigma_fraction_increases_when_entropy_below_target`
//!   Task A3 (Pearl 2 loss budget):
//!     5. `pearl_2_budget_flat_regime_iqn_floor_c51_takes_remainder`
//!     6. `pearl_2_budget_sharp_regime_iqn_dominates`
//!   Task A4 (Pearl 4 per-group Adam β/β/ε + auxiliary cosine-sim):
//!     7. `pearl_4_adam_hparams_high_stability_yields_long_memory`
//!     8. `pearl_4_adam_hparams_low_stability_yields_short_memory`
//!     9. `grad_cosine_sim_per_group_dot_norm_and_writeback`
//!   Task A5 (Pearl 5 per-branch IQN τ schedule):
//!    10. `pearl_5_iqn_tau_symmetric_q_yields_uniform_default_schedule`
//!    11. `pearl_5_iqn_tau_left_skew_shifts_quantiles_negatively`
//!
//!   Task A6 (Pearl 6 cross-fold-persistent Kelly):
//!    12. `pearl_6_kelly_within_fold_ewma_blend`
//!    13. `pearl_6_kelly_cross_fold_sample_count_persists_via_max`
//!
//!   Task A7 (Pearl 8 per-direction trail distance):
//!    14. `pearl_8_trail_per_direction_factors`
//!    15. `pearl_8_trail_atr_floor_at_quiet_market`
//!
//!   Task A8 (Pearl 1-ext per-branch num_atoms):
//!    16. `pearl_1_ext_num_atoms_threshold_cascade`
//!    17. `pearl_1_ext_num_atoms_exact_threshold_falls_to_next_branch`
//!
//!   Layer D Task D1 (PnL aggregation):
//!    18. `pnl_aggregation_kernel_correctness`
//!
//!   Layer D Task D2 (Health composition):
//!    19. `health_composition_kernel_correctness`
//!
//!   Layer D Task D3 (Training metrics EMA):
//!    20. `training_metrics_ema_kernel_correctness`
//!
//! All tests use analytically-known synthetic inputs with no CPU reference
//! implementation, per `feedback_no_cpu_test_fallbacks.md`. All mapped-pinned
//! memory (no DtoH), all `#[ignore]`-gated for GPU.
//!
//! Run with:
//!
//!     cargo test -p ml --tests sp5_producer_unit_tests -- --ignored --nocapture

use std::sync::Arc;

use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg};
use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedI32Buffer};

// ── Cubins ────────────────────────────────────────────────────────────────────

const SP5_Q_BRANCH_STATS_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/q_branch_stats_kernel.cubin"));

const SP5_PEARL_1_ATOM_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/pearl_1_atom_kernel.cubin"));

const SP5_PEARL_3_SIGMA_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/pearl_3_sigma_kernel.cubin"));

const SP5_PEARL_2_BUDGET_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/pearl_2_budget_kernel.cubin"));

const SP5_PEARL_4_ADAM_HPARAMS_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/pearl_4_adam_hparams_kernel.cubin"));

const SP5_GRAD_COSINE_SIM_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/grad_cosine_sim_kernel.cubin"));

const SP5_Q_SKEW_KURTOSIS_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/q_skew_kurtosis_kernel.cubin"));

const SP5_PEARL_5_IQN_TAU_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/pearl_5_iqn_tau_kernel.cubin"));

const SP5_PEARL_6_KELLY_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/pearl_6_kelly_kernel.cubin"));

const SP5_PEARL_8_TRAIL_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/pearl_8_trail_kernel.cubin"));

const SP5_PEARL_1_EXT_NUM_ATOMS_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/pearl_1_ext_num_atoms_kernel.cubin"));

const SP5_PNL_AGGREGATION_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/pnl_aggregation_kernel.cubin"));

const SP5_HEALTH_COMPOSITION_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/health_composition_kernel.cubin"));

const SP5_TRAINING_METRICS_EMA_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/training_metrics_ema_kernel.cubin"));

const SP7_LOSS_BALANCE_CONTROLLER_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/loss_balance_controller_kernel.cubin"));

// ── Helpers ───────────────────────────────────────────────────────────────────

fn load_pearl_3_sigma_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
    let module = stream
        .context()
        .load_cubin(SP5_PEARL_3_SIGMA_CUBIN.to_vec())
        .expect("load pearl_3_sigma_kernel cubin");
    module
        .load_function("pearl_3_sigma_update")
        .expect("load pearl_3_sigma_update function")
}

fn load_pearl_2_budget_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
    let module = stream
        .context()
        .load_cubin(SP5_PEARL_2_BUDGET_CUBIN.to_vec())
        .expect("load pearl_2_budget_kernel cubin");
    module
        .load_function("pearl_2_budget_update")
        .expect("load pearl_2_budget_update function")
}

fn load_pearl_4_adam_hparams_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
    let module = stream
        .context()
        .load_cubin(SP5_PEARL_4_ADAM_HPARAMS_CUBIN.to_vec())
        .expect("load pearl_4_adam_hparams_kernel cubin");
    module
        .load_function("pearl_4_adam_hparams_update")
        .expect("load pearl_4_adam_hparams_update function")
}

fn load_grad_cosine_sim_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
    let module = stream
        .context()
        .load_cubin(SP5_GRAD_COSINE_SIM_CUBIN.to_vec())
        .expect("load grad_cosine_sim_kernel cubin");
    module
        .load_function("grad_cosine_sim_update")
        .expect("load grad_cosine_sim_update function")
}

fn make_test_stream() -> Arc<CudaStream> {
    let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
    ctx.default_stream()
}

fn load_q_branch_stats_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
    let module = stream
        .context()
        .load_cubin(SP5_Q_BRANCH_STATS_CUBIN.to_vec())
        .expect("load q_branch_stats_kernel cubin");
    module
        .load_function("q_branch_stats_update")
        .expect("load q_branch_stats_update function")
}

fn load_pearl_1_atom_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
    let module = stream
        .context()
        .load_cubin(SP5_PEARL_1_ATOM_CUBIN.to_vec())
        .expect("load pearl_1_atom_kernel cubin");
    module
        .load_function("pearl_1_atom_update")
        .expect("load pearl_1_atom_update function")
}

// ── Test 1: q_branch_stats_update ─────────────────────────────────────────────

/// SP5 Task A1 unit test: `q_branch_stats_update` correctly computes per-branch
/// statistics on a synthetic uniform batch.
///
/// Synthetic input: B=4 batch, q_out_buf [4, 13] with all values = constant c,
/// where c = 2.0.
///
/// Analytical ground truth for constant input (c = 2.0):
///   - mean           = 2.0            (trivially)
///   - max_abs_dev    = 0.0            (all deviations from mean are 0)
///   - var            = 0.0            (no spread)
///   - entropy        ≈ log(n_actions) (uniform distribution after softmax)
///
/// For Direction branch (n_actions=4): entropy = log(4) ≈ 1.3862944
/// For Magnitude/Order/Urgency branches (n_actions=3): entropy = log(3) ≈ 1.0986123
///
/// The uniform softmax arises because all Q-values are equal (after subtracting
/// the max, all differences are 0, yielding exp(0)/exp(0) = 1/n for each action).
/// This is an analytically-known result that does NOT require any CPU computation
/// per `feedback_no_cpu_test_fallbacks.md`.
#[test]
#[ignore = "requires GPU"]
fn sp5_q_branch_stats_uniform_input_matches_analytical_ground_truth() {
    const BATCH_SIZE: usize = 4;
    const TOTAL_ACTIONS: usize = 13;
    const Q_CONST: f32 = 2.0_f32;
    const SCRATCH_IDX_BASE: usize = 71;  // SCRATCH_Q_STATS

    let stream = make_test_stream();
    let kernel = load_q_branch_stats_kernel(&stream);

    // q_out_buf [B, total_actions=13] all = Q_CONST.
    let q_buf = unsafe { MappedF32Buffer::new(BATCH_SIZE * TOTAL_ACTIONS) }
        .expect("alloc q_out_buf");
    q_buf.write_from_slice(&vec![Q_CONST; BATCH_SIZE * TOTAL_ACTIONS]);

    // action_counts = {4, 3, 3, 3}
    let counts_buf = unsafe { MappedI32Buffer::new(4) }
        .expect("alloc action_counts_buf");
    counts_buf.write_from_slice(&[4_i32, 3, 3, 3]);

    // action_offsets = {0, 4, 7, 10}
    let offsets_buf = unsafe { MappedI32Buffer::new(4) }
        .expect("alloc action_offsets_buf");
    offsets_buf.write_from_slice(&[0_i32, 4, 7, 10]);

    // scratch [71..87) for 16 outputs (4 per branch × 4 branches).
    let scratch_buf = unsafe { MappedF32Buffer::new(SCRATCH_IDX_BASE + 16) }
        .expect("alloc scratch_buf");

    let q_ptr            = q_buf.dev_ptr;
    let b_i32            = BATCH_SIZE as i32;
    let counts_dev       = counts_buf.dev_ptr;
    let offsets_dev      = offsets_buf.dev_ptr;
    let scratch_dev      = scratch_buf.dev_ptr;
    let scratch_base_i32 = SCRATCH_IDX_BASE as i32;

    unsafe {
        stream
            .launch_builder(&kernel)
            .arg(&q_ptr)
            .arg(&b_i32)
            .arg(&counts_dev)
            .arg(&offsets_dev)
            .arg(&scratch_dev)
            .arg(&scratch_base_i32)
            .launch(LaunchConfig {
                grid_dim:         (1, 1, 1),
                block_dim:        (4, 1, 1),
                shared_mem_bytes: 0,
            })
            .expect("launch q_branch_stats_update");
    }
    stream.synchronize().expect("sync after q_branch_stats_update");

    let scratch = scratch_buf.read_all();

    // Analytical values:
    let action_counts = [4_usize, 3, 3, 3];
    for branch in 0..4_usize {
        let base = SCRATCH_IDX_BASE + branch * 4;
        let mean        = scratch[base + 0];
        let max_abs_dev = scratch[base + 1];
        let var         = scratch[base + 2];
        let entropy     = scratch[base + 3];

        let n = action_counts[branch] as f32;
        let expected_entropy = n.ln();  // log(n) for uniform distribution

        println!(
            "branch={branch} n={n:.0} mean={mean:.5} max_abs_dev={max_abs_dev:.5} \
             var={var:.5} entropy={entropy:.5} (expected entropy={expected_entropy:.5})"
        );

        // Mean = Q_CONST exactly (constant input).
        assert!(
            (mean - Q_CONST).abs() < 1e-5,
            "branch={branch}: mean={mean} != Q_CONST={Q_CONST}"
        );
        // max_abs_dev = 0 for constant input.
        assert!(
            max_abs_dev.abs() < 1e-5,
            "branch={branch}: max_abs_dev={max_abs_dev} should be 0 for constant input"
        );
        // var = 0 for constant input.
        assert!(
            var.abs() < 1e-5,
            "branch={branch}: var={var} should be 0 for constant input"
        );
        // entropy = log(n_actions) for uniform distribution, within 1% relative.
        let rel_err = ((entropy - expected_entropy) / expected_entropy).abs();
        assert!(
            rel_err < 0.01,
            "branch={branch}: entropy={entropy:.5} expected={expected_entropy:.5} rel_err={rel_err:.5}"
        );
    }
}

// ── Test 2: pearl_1_atom_update ────────────────────────────────────────────────

/// SP5 Task A1 unit test: `pearl_1_atom_update` correctly computes the
/// atom-span headroom controller on known synthetic inputs.
///
/// Synthetic input:
///   - mean = 1.0 (all branches), max_abs_dev = 0.5 (all branches)
///   - atoms_clip_count = [0, 0, 0, 0] (Layer B not wired yet)
///   - headroom from ISV = 0.0 (Pearl A sentinel → bootstrap to 6.0)
///
/// Analytical ground truth:
///   - clip_rate = 0 / (B × 13) = 0.0
///   - new_headroom = 6.0 + 0.01 × (0.0 - 0.01) = 6.0 - 0.0001 = 5.9999
///     (clamped to [2.0, 20.0] → 5.9999)
///   - v_half = max_abs_dev × new_headroom = 0.5 × 5.9999 ≈ 2.99995
///   - v_center = mean = 1.0
///
/// This verifies:
///   1. Pearl A sentinel detection (headroom ≤ 0 → bootstrap 6.0).
///   2. Headroom controller arithmetic (one step from bootstrap).
///   3. v_half = max_abs_dev × new_headroom.
///   4. Output slot layout (v_center/v_half/headroom/clip_rate in SCRATCH_ATOM_OUT groups).
#[test]
#[ignore = "requires GPU"]
fn sp5_pearl_1_atom_headroom_bootstrap_and_controller_step() {
    const BATCH_SIZE: usize = 4;
    const SCRATCH_Q_STATS: usize = 71;
    const SCRATCH_ATOM_OUT: usize = 87;

    let stream = make_test_stream();
    let kernel = load_pearl_1_atom_kernel(&stream);

    // q_branch_stats scratch: 16 floats at [71..87).
    // Layout: branch b at scratch[71 + b*4 + {0,1,2,3}] = {mean, max_abs_dev, var, entropy}
    let scratch_buf = unsafe { MappedF32Buffer::new(SCRATCH_ATOM_OUT + 16) }
        .expect("alloc scratch_buf");
    let scratch_data: Vec<f32> = (0..scratch_buf.len)
        .map(|i| {
            let rel = i as i32 - SCRATCH_Q_STATS as i32;
            if rel >= 0 && rel < 16 {
                // branch b at offset b*4: mean=1.0, max_abs_dev=0.5, var=0.1, entropy=1.0
                let field = (rel % 4) as usize;
                match field {
                    0 => 1.0_f32,   // mean
                    1 => 0.5_f32,   // max_abs_dev
                    2 => 0.1_f32,   // var
                    3 => 1.0_f32,   // entropy
                    _ => unreachable!(),
                }
            } else {
                0.0_f32
            }
        })
        .collect();
    scratch_buf.write_from_slice(&scratch_data);

    // atoms_clip_count = [0, 0, 0, 0] (Layer B not wired).
    let clip_count_buf = unsafe { MappedI32Buffer::new(4) }
        .expect("alloc atoms_clip_count_buf");
    // Already zero from alloc; write explicitly for clarity.
    clip_count_buf.write_from_slice(&[0_i32; 4]);

    // action_counts = {4, 3, 3, 3} — per-branch Q-value counts. The clip_rate
    // denominator is batch_size × action_counts[b], not a global sum.
    let action_counts_buf = unsafe { MappedI32Buffer::new(4) }
        .expect("alloc action_counts_buf");
    action_counts_buf.write_from_slice(&[4_i32, 3, 3, 3]);

    // ISV signals: all zero (headroom sentinel → bootstrap to 6.0).
    // Need at least ATOM_HEADROOM_BASE+4 = 182+4 = 186 slots.
    const ISV_SIZE: usize = 200;
    let isv_buf = unsafe { MappedF32Buffer::new(ISV_SIZE) }
        .expect("alloc isv_signals");
    // All zeros — headroom ≤ 0 triggers Pearl A bootstrap in kernel.

    let scratch_dev       = scratch_buf.dev_ptr;
    let scratch_stats_i32 = SCRATCH_Q_STATS as i32;
    let clip_dev          = clip_count_buf.dev_ptr;
    let action_counts_dev = action_counts_buf.dev_ptr;
    let b_i32             = BATCH_SIZE as i32;
    // SCRATCH_ATOM_OUT layout: [87..91)=v_center, [91..95)=v_half, [95..99)=headroom, [99..103)=clip_rate
    let v_center_base_i32  = SCRATCH_ATOM_OUT as i32;
    let v_half_base_i32    = (SCRATCH_ATOM_OUT + 4) as i32;
    let headroom_base_i32  = (SCRATCH_ATOM_OUT + 8) as i32;
    let clip_rate_base_i32 = (SCRATCH_ATOM_OUT + 12) as i32;
    let isv_dev            = isv_buf.dev_ptr;
    let isv_headroom_i32   = 182_i32;  // ATOM_HEADROOM_BASE

    unsafe {
        stream
            .launch_builder(&kernel)
            .arg(&scratch_dev)
            .arg(&scratch_stats_i32)
            .arg(&clip_dev)
            .arg(&action_counts_dev)
            .arg(&b_i32)
            .arg(&scratch_dev)
            .arg(&v_center_base_i32)
            .arg(&v_half_base_i32)
            .arg(&headroom_base_i32)
            .arg(&clip_rate_base_i32)
            .arg(&isv_dev)
            .arg(&isv_headroom_i32)
            .launch(LaunchConfig {
                grid_dim:         (1, 1, 1),
                block_dim:        (4, 1, 1),
                shared_mem_bytes: 0,
            })
            .expect("launch pearl_1_atom_update");
    }
    stream.synchronize().expect("sync after pearl_1_atom_update");

    let scratch = scratch_buf.read_all();

    // Analytical expected values.
    let expected_clip_rate = 0.0_f32;
    // new_headroom = 6.0 + 0.01 × (0.0 - 0.01) = 5.9999
    let expected_headroom  = 6.0_f32 + 0.01 * (expected_clip_rate - 0.01);
    let expected_v_half    = 0.5_f32 * expected_headroom;
    let expected_v_center  = 1.0_f32;

    for b in 0..4_usize {
        let v_center  = scratch[SCRATCH_ATOM_OUT + b];
        let v_half    = scratch[SCRATCH_ATOM_OUT + 4 + b];
        let headroom  = scratch[SCRATCH_ATOM_OUT + 8 + b];
        let clip_rate = scratch[SCRATCH_ATOM_OUT + 12 + b];

        println!(
            "branch={b} v_center={v_center:.5} v_half={v_half:.5} \
             headroom={headroom:.5} clip_rate={clip_rate:.5}"
        );

        // v_center = mean = 1.0
        assert!(
            (v_center - expected_v_center).abs() < 1e-4,
            "branch={b}: v_center={v_center} expected={expected_v_center}"
        );
        // v_half = 0.5 × 5.9999 ≈ 2.99995 (1e-4 tolerance for fp32 rounding)
        assert!(
            (v_half - expected_v_half).abs() < 1e-3,
            "branch={b}: v_half={v_half} expected={expected_v_half:.5}"
        );
        // headroom = 5.9999 (bootstrap 6.0 minus one step)
        assert!(
            (headroom - expected_headroom).abs() < 1e-4,
            "branch={b}: headroom={headroom} expected={expected_headroom:.5}"
        );
        // clip_rate = 0 (no clips yet)
        assert!(
            clip_rate.abs() < 1e-6,
            "branch={b}: clip_rate={clip_rate} expected=0.0"
        );
    }
}

// ── Test 3: pearl_3_sigma_update bootstrap at target entropy ──────────────────

/// SP5 Task A2 unit test: `pearl_3_sigma_update` bootstraps SF=0.1 when ISV
/// sigma_fraction=0 (Pearl A sentinel), computes new_sigma = SF * v_half.
///
/// Synthetic input:
///   - ISV[ATOM_V_HALF_BASE+b] = 2.0 (all branches)
///   - ISV[BRANCH_ENTROPY_BASE+b] = log(n_actions[b]) * 0.7 (target entropy exactly)
///   - ISV[SIGMA_FRACTION_BASE+b] = 0.0 (Pearl A sentinel → bootstrap to 0.1)
///   - action_counts = {4, 3, 3, 3}
///
/// Analytical ground truth (target == actual → no SF update step, only bootstrap):
///   - sigma_fraction bootstraps to 0.1 (sentinel 0 detection)
///   - new_sf = 0.1 + 0.005 * (target - actual) = 0.1 + 0.005 * 0 = 0.1 (exact)
///   - new_sigma = 0.1 * max(2.0, 1.0) = 0.1 * 2.0 = 0.2 for each branch
///
/// This verifies Pearl A sentinel detection, bootstrap value, and SF * v_half formula.
/// No CPU reference oracle per `feedback_no_cpu_test_fallbacks.md`.
#[test]
#[ignore = "requires GPU"]
fn pearl_3_sigma_bootstrap_at_target_entropy() {
    const SCRATCH_SIGMA: usize = 103;
    const SCRATCH_SF:    usize = 107;

    // Need enough ISV slots: ATOM_V_HALF_BASE=178, SIGMA_FRACTION_BASE=214,
    // BRANCH_ENTROPY_BASE=218. Max = 218+4 = 222.
    const ISV_SIZE: usize = 230;

    let stream = make_test_stream();
    let kernel = load_pearl_3_sigma_kernel(&stream);

    let action_counts = [4_i32, 3, 3, 3];
    let action_counts_buf = unsafe { MappedI32Buffer::new(4) }
        .expect("alloc action_counts_buf");
    action_counts_buf.write_from_slice(&action_counts);

    // Build ISV: v_half=2.0 at [178..182), SF=0.0 at [214..218) (sentinel),
    // entropy=log(n)*0.7 at [218..222).
    let isv_buf = unsafe { MappedF32Buffer::new(ISV_SIZE) }
        .expect("alloc isv_buf");
    let mut isv_data = vec![0.0_f32; ISV_SIZE];
    for b in 0..4_usize {
        isv_data[178 + b] = 2.0_f32;  // ATOM_V_HALF_BASE = 178
        // SIGMA_FRACTION_BASE = 214: leave at 0.0 (sentinel)
        // BRANCH_ENTROPY_BASE = 218: set to target_entropy = log(n) * 0.7
        let n = action_counts[b] as f32;
        isv_data[218 + b] = n.ln() * 0.7_f32;
    }
    isv_buf.write_from_slice(&isv_data);

    let scratch_buf = unsafe { MappedF32Buffer::new(SCRATCH_SF + 4) }
        .expect("alloc scratch_buf");

    let isv_dev         = isv_buf.dev_ptr;
    let v_half_base_i32          = 178_i32;  // ATOM_V_HALF_BASE
    let branch_entropy_base_i32  = 218_i32;  // BRANCH_ENTROPY_BASE
    let sigma_fraction_base_i32  = 214_i32;  // SIGMA_FRACTION_BASE
    let counts_dev               = action_counts_buf.dev_ptr;
    let scratch_dev              = scratch_buf.dev_ptr;
    let sigma_idx_i32            = SCRATCH_SIGMA as i32;
    let sf_idx_i32               = SCRATCH_SF as i32;

    unsafe {
        stream
            .launch_builder(&kernel)
            .arg(&isv_dev)
            .arg(&v_half_base_i32)
            .arg(&branch_entropy_base_i32)
            .arg(&sigma_fraction_base_i32)
            .arg(&counts_dev)
            .arg(&scratch_dev)
            .arg(&sigma_idx_i32)
            .arg(&sf_idx_i32)
            .launch(LaunchConfig {
                grid_dim:         (1, 1, 1),
                block_dim:        (4, 1, 1),
                shared_mem_bytes: 0,
            })
            .expect("launch pearl_3_sigma_update (bootstrap test)");
    }
    stream.synchronize().expect("sync after pearl_3_sigma_update (bootstrap test)");

    let scratch = scratch_buf.read_all();

    for b in 0..4_usize {
        let sigma = scratch[SCRATCH_SIGMA + b];
        let sf    = scratch[SCRATCH_SF + b];

        println!("branch={b} sigma={sigma:.6} sf={sf:.6}");

        // SF bootstraps to 0.1 (sentinel 0 → 0.1), no update because target=actual.
        let rel_err_sf = (sf - 0.1_f32).abs() / 0.1_f32;
        assert!(
            rel_err_sf < 0.01,
            "branch={b}: sf={sf:.6} expected=0.1 rel_err={rel_err_sf:.5}"
        );

        // sigma = sf * v_half = 0.1 * 2.0 = 0.2
        let expected_sigma = 0.1_f32 * 2.0_f32;
        let rel_err_sigma = (sigma - expected_sigma).abs() / expected_sigma;
        assert!(
            rel_err_sigma < 0.01,
            "branch={b}: sigma={sigma:.6} expected={expected_sigma:.6} rel_err={rel_err_sigma:.5}"
        );
    }
}

// ── Test 4: pearl_3_sigma_update fraction increases when entropy below target ──

/// SP5 Task A2 unit test: when actual entropy is BELOW target, the entropy-deficit
/// controller increases SF.
///
/// Synthetic input for branch 0 (Direction, n_actions=4):
///   - ISV[ATOM_V_HALF_BASE+0] = 2.0
///   - ISV[BRANCH_ENTROPY_BASE+0] = 0.0  (extreme case: zero entropy)
///   - ISV[SIGMA_FRACTION_BASE+0] = 0.1  (already bootstrapped)
///   - target_entropy = log(4) * 0.7 = 1.3862944 * 0.7 = 0.97040609
///
/// Analytical ground truth for branch 0:
///   new_sf = 0.1 + 0.005 * (0.97040609 - 0.0) = 0.1 + 0.00485203 = 0.10485203
///   new_sigma = 0.10485203 * 2.0 = 0.20970406
///
/// Same inputs for branches 1/2/3 (n=3, entropy=0, sf=0.1):
///   target_entropy = log(3) * 0.7 = 1.0986123 * 0.7 = 0.76902860
///   new_sf = 0.1 + 0.005 * 0.76902860 = 0.10384514
///   new_sigma = 0.20769028
///
/// Checks SF strictly increases when entropy is below target. 1% rel-err tolerance.
#[test]
#[ignore = "requires GPU"]
fn pearl_3_sigma_fraction_increases_when_entropy_below_target() {
    const SCRATCH_SIGMA: usize = 103;
    const SCRATCH_SF:    usize = 107;
    const ISV_SIZE: usize = 230;

    let stream = make_test_stream();
    let kernel = load_pearl_3_sigma_kernel(&stream);

    let action_counts = [4_i32, 3, 3, 3];
    let action_counts_buf = unsafe { MappedI32Buffer::new(4) }
        .expect("alloc action_counts_buf");
    action_counts_buf.write_from_slice(&action_counts);

    // ISV: v_half=2.0, SF=0.1 (already bootstrapped), entropy=0 (below target).
    let isv_buf = unsafe { MappedF32Buffer::new(ISV_SIZE) }
        .expect("alloc isv_buf");
    let mut isv_data = vec![0.0_f32; ISV_SIZE];
    for b in 0..4_usize {
        isv_data[178 + b] = 2.0_f32;  // ATOM_V_HALF_BASE = 178
        isv_data[214 + b] = 0.1_f32;  // SIGMA_FRACTION_BASE = 214 (already bootstrapped)
        // BRANCH_ENTROPY_BASE = 218: leave at 0.0 (zero entropy)
    }
    isv_buf.write_from_slice(&isv_data);

    let scratch_buf = unsafe { MappedF32Buffer::new(SCRATCH_SF + 4) }
        .expect("alloc scratch_buf");

    let isv_dev                  = isv_buf.dev_ptr;
    let v_half_base_i32          = 178_i32;
    let branch_entropy_base_i32  = 218_i32;
    let sigma_fraction_base_i32  = 214_i32;
    let counts_dev               = action_counts_buf.dev_ptr;
    let scratch_dev              = scratch_buf.dev_ptr;
    let sigma_idx_i32            = SCRATCH_SIGMA as i32;
    let sf_idx_i32               = SCRATCH_SF as i32;

    unsafe {
        stream
            .launch_builder(&kernel)
            .arg(&isv_dev)
            .arg(&v_half_base_i32)
            .arg(&branch_entropy_base_i32)
            .arg(&sigma_fraction_base_i32)
            .arg(&counts_dev)
            .arg(&scratch_dev)
            .arg(&sigma_idx_i32)
            .arg(&sf_idx_i32)
            .launch(LaunchConfig {
                grid_dim:         (1, 1, 1),
                block_dim:        (4, 1, 1),
                shared_mem_bytes: 0,
            })
            .expect("launch pearl_3_sigma_update (entropy-below-target test)");
    }
    stream.synchronize().expect("sync after pearl_3_sigma_update (entropy-below-target test)");

    let scratch = scratch_buf.read_all();

    // Branch 0: n=4, target=log(4)*0.7
    {
        let b = 0_usize;
        let sigma = scratch[SCRATCH_SIGMA + b];
        let sf    = scratch[SCRATCH_SF + b];
        let target_entropy = (4.0_f32).ln() * 0.7_f32;
        let expected_sf    = 0.1_f32 + 0.005_f32 * target_entropy;
        let expected_sigma = expected_sf * 2.0_f32;
        println!("branch={b} n=4 sf={sf:.7} expected_sf={expected_sf:.7} \
                  sigma={sigma:.7} expected_sigma={expected_sigma:.7}");
        let rel_sf = (sf - expected_sf).abs() / expected_sf;
        assert!(rel_sf < 0.01, "branch={b}: sf={sf:.7} expected={expected_sf:.7} rel={rel_sf:.5}");
        let rel_sig = (sigma - expected_sigma).abs() / expected_sigma;
        assert!(rel_sig < 0.01, "branch={b}: sigma={sigma:.7} expected={expected_sigma:.7} rel={rel_sig:.5}");
        // SF strictly increased from 0.1
        assert!(sf > 0.1_f32, "branch={b}: SF should have increased from 0.1, got {sf}");
    }

    // Branches 1-3: n=3, target=log(3)*0.7
    for b in 1..4_usize {
        let sigma = scratch[SCRATCH_SIGMA + b];
        let sf    = scratch[SCRATCH_SF + b];
        let target_entropy = (3.0_f32).ln() * 0.7_f32;
        let expected_sf    = 0.1_f32 + 0.005_f32 * target_entropy;
        let expected_sigma = expected_sf * 2.0_f32;
        println!("branch={b} n=3 sf={sf:.7} expected_sf={expected_sf:.7} \
                  sigma={sigma:.7} expected_sigma={expected_sigma:.7}");
        let rel_sf = (sf - expected_sf).abs() / expected_sf;
        assert!(rel_sf < 0.01, "branch={b}: sf={sf:.7} expected={expected_sf:.7} rel={rel_sf:.5}");
        let rel_sig = (sigma - expected_sigma).abs() / expected_sigma;
        assert!(rel_sig < 0.01, "branch={b}: sigma={sigma:.7} expected={expected_sigma:.7} rel={rel_sig:.5}");
        assert!(sf > 0.1_f32, "branch={b}: SF should have increased from 0.1, got {sf}");
    }
}

// ── Tests 5 + 6: pearl_2_budget_update ────────────────────────────────────────

/// SP5 Task A3 unit test: `pearl_2_budget_update` on flat-Q regime.
///
/// Synthetic ISV: σ=1.0 (all branches), var_q=2.0 (all branches).
/// flatness = clamp(2.0 / (1.0² + 1e-8), 0, 1) = clamp(~2.0, 0, 1) = 1.0
///
/// Analytical expected (flatness = 1.0):
///   BASE_IQN       = 0.11
///   BASE_C51_SHARE = 0.84 / (1 - 0.11) = 0.84 / 0.89 ≈ 0.94382...
///   budget_iqn = 0.11 + (1 - 1.0) * 0.89 = 0.11
///   budget_c51 = 1.0 * 0.89 * 0.94382 ≈ 0.84
///   budget_cql = 0
///   budget_ens = max(0, 1 - 0.11 - 0.84 - 0) = 0.05
///
/// All 4 branches receive identical inputs, so identical outputs.
/// No CPU reference oracle per `feedback_no_cpu_test_fallbacks.md`.
#[test]
#[ignore = "requires GPU"]
fn pearl_2_budget_flat_regime_iqn_floor_c51_takes_remainder() {
    // Scratch layout mirrors SCRATCH_PEARL_2_* constants:
    //   [111..115) = budget_c51[4]
    //   [115..119) = budget_iqn[4]
    //   [119..123) = budget_cql[4]
    //   [123..127) = budget_ens[4]
    //   [127..131) = flatness[4]
    const SCRATCH_C51:      usize = 111;
    const SCRATCH_IQN:      usize = 115;
    const SCRATCH_CQL:      usize = 119;
    const SCRATCH_ENS:      usize = 123;
    const SCRATCH_FLATNESS: usize = 127;
    const SCRATCH_SIZE:     usize = 131;

    // ISV: Q_VAR_PER_BRANCH_BASE=222, NOISY_SIGMA_BASE=210.
    // Need ISV size >= 222 + 4 = 226.
    const ISV_SIZE: usize = 230;
    // Q_VAR_PER_BRANCH_BASE = 222, NOISY_SIGMA_BASE = 210.
    const Q_VAR_BASE: usize = 222;
    const SIGMA_BASE: usize = 210;

    let stream = make_test_stream();
    let kernel = load_pearl_2_budget_kernel(&stream);

    let isv_buf = unsafe { MappedF32Buffer::new(ISV_SIZE) }
        .expect("alloc isv_buf");
    let mut isv_data = vec![0.0_f32; ISV_SIZE];
    for b in 0..4_usize {
        isv_data[Q_VAR_BASE + b] = 2.0_f32;  // var_q = 2.0 (above σ² → clamp to 1.0 flatness)
        isv_data[SIGMA_BASE + b] = 1.0_f32;  // σ = 1.0
    }
    isv_buf.write_from_slice(&isv_data);

    let scratch_buf = unsafe { MappedF32Buffer::new(SCRATCH_SIZE) }
        .expect("alloc scratch_buf");

    let isv_dev           = isv_buf.dev_ptr;
    let q_var_base_i32    = Q_VAR_BASE as i32;
    let sigma_base_i32    = SIGMA_BASE as i32;
    let scratch_dev       = scratch_buf.dev_ptr;
    let c51_base_i32      = SCRATCH_C51 as i32;
    let iqn_base_i32      = SCRATCH_IQN as i32;
    let cql_base_i32      = SCRATCH_CQL as i32;
    let ens_base_i32      = SCRATCH_ENS as i32;
    let flatness_base_i32 = SCRATCH_FLATNESS as i32;

    unsafe {
        stream
            .launch_builder(&kernel)
            .arg(&isv_dev)
            .arg(&q_var_base_i32)
            .arg(&sigma_base_i32)
            .arg(&scratch_dev)
            .arg(&c51_base_i32)
            .arg(&iqn_base_i32)
            .arg(&cql_base_i32)
            .arg(&ens_base_i32)
            .arg(&flatness_base_i32)
            .launch(LaunchConfig {
                grid_dim:         (1, 1, 1),
                block_dim:        (4, 1, 1),
                shared_mem_bytes: 0,
            })
            .expect("launch pearl_2_budget_update (flat regime)");
    }
    stream.synchronize().expect("sync after pearl_2_budget_update (flat regime)");

    let scratch = scratch_buf.read_all();

    // Analytical values for flatness = 1.0 (saturated ceiling):
    //   BASE_IQN = 0.11, BASE_C51_SHARE = 0.84 / 0.89
    const BASE_IQN: f32       = 0.11_f32;
    const BASE_C51_SHARE: f32 = 0.84_f32 / (1.0_f32 - BASE_IQN);
    let expected_flatness = 1.0_f32;
    let expected_iqn      = BASE_IQN + (1.0_f32 - expected_flatness) * (1.0_f32 - BASE_IQN);
    let expected_c51      = expected_flatness * (1.0_f32 - BASE_IQN) * BASE_C51_SHARE;
    let expected_cql      = 0.0_f32;
    let expected_ens      = (1.0_f32 - expected_iqn - expected_c51 - expected_cql).max(0.0_f32);

    for b in 0..4_usize {
        let c51      = scratch[SCRATCH_C51 + b];
        let iqn      = scratch[SCRATCH_IQN + b];
        let cql      = scratch[SCRATCH_CQL + b];
        let ens      = scratch[SCRATCH_ENS + b];
        let flatness = scratch[SCRATCH_FLATNESS + b];

        println!(
            "branch={b} flatness={flatness:.6} iqn={iqn:.6} c51={c51:.6} \
             cql={cql:.6} ens={ens:.6}"
        );

        let rel_flatness = (flatness - expected_flatness).abs();
        assert!(
            rel_flatness < 1e-5,
            "branch={b}: flatness={flatness:.6} expected={expected_flatness:.6}"
        );
        let rel_iqn = (iqn - expected_iqn).abs() / expected_iqn.max(1e-9);
        assert!(
            rel_iqn < 0.01,
            "branch={b}: iqn={iqn:.6} expected={expected_iqn:.6} rel={rel_iqn:.5}"
        );
        let rel_c51 = (c51 - expected_c51).abs() / expected_c51.max(1e-9);
        assert!(
            rel_c51 < 0.01,
            "branch={b}: c51={c51:.6} expected={expected_c51:.6} rel={rel_c51:.5}"
        );
        assert!(
            cql.abs() < 1e-7,
            "branch={b}: cql={cql:.6} expected=0"
        );
        let rel_ens = (ens - expected_ens).abs() / expected_ens.max(1e-9);
        assert!(
            rel_ens < 0.01,
            "branch={b}: ens={ens:.6} expected={expected_ens:.6} rel={rel_ens:.5}"
        );
        // Normalization invariant: budgets sum to 1.0 (ens is computed as the
        // residual `max(0, 1 - iqn - c51 - cql)`, so the structural claim is
        // sum ≈ 1.0). A coefficient typo (e.g. wrong BASE_IQN) would still
        // pass the per-component relative checks but break this invariant.
        let sum = c51 + iqn + cql + ens;
        assert!(
            (sum - 1.0_f32).abs() < 1e-4,
            "branch={b}: budget sum={sum:.6} should be ≈ 1.0 (c51={c51:.6} iqn={iqn:.6} cql={cql:.6} ens={ens:.6})"
        );
    }
}

/// SP5 Task A3 unit test: `pearl_2_budget_update` on sharp-Q regime.
///
/// Synthetic ISV: σ=2.0 (all branches), var_q=0.04 (all branches).
/// flatness = clamp(0.04 / (4.0 + 1e-8), 0, 1) = 0.04/4.0 = 0.01
///
/// Analytical expected (flatness = 0.01):
///   BASE_IQN       = 0.11
///   BASE_C51_SHARE = 0.84 / 0.89 ≈ 0.94382
///   budget_iqn = 0.11 + (1 - 0.01) * 0.89 = 0.11 + 0.8811 = 0.9911
///   budget_c51 = 0.01 * 0.89 * 0.94382 ≈ 0.0084
///   budget_cql = 0
///   budget_ens = max(0, 1 - 0.9911 - 0.0084 - 0) ≈ 0.0005 (rounds to ~0)
///
/// IQN dominates in the sharp-Q regime, C51 share is minimal.
/// No CPU reference oracle per `feedback_no_cpu_test_fallbacks.md`.
#[test]
#[ignore = "requires GPU"]
fn pearl_2_budget_sharp_regime_iqn_dominates() {
    const SCRATCH_C51:      usize = 111;
    const SCRATCH_IQN:      usize = 115;
    const SCRATCH_CQL:      usize = 119;
    const SCRATCH_ENS:      usize = 123;
    const SCRATCH_FLATNESS: usize = 127;
    const SCRATCH_SIZE:     usize = 131;
    const ISV_SIZE:         usize = 230;
    const Q_VAR_BASE:       usize = 222;
    const SIGMA_BASE:       usize = 210;

    let stream = make_test_stream();
    let kernel = load_pearl_2_budget_kernel(&stream);

    let isv_buf = unsafe { MappedF32Buffer::new(ISV_SIZE) }
        .expect("alloc isv_buf");
    let mut isv_data = vec![0.0_f32; ISV_SIZE];
    for b in 0..4_usize {
        isv_data[Q_VAR_BASE + b] = 0.04_f32;  // var_q small
        isv_data[SIGMA_BASE + b] = 2.0_f32;   // σ large → flatness small
    }
    isv_buf.write_from_slice(&isv_data);

    let scratch_buf = unsafe { MappedF32Buffer::new(SCRATCH_SIZE) }
        .expect("alloc scratch_buf");

    let isv_dev           = isv_buf.dev_ptr;
    let q_var_base_i32    = Q_VAR_BASE as i32;
    let sigma_base_i32    = SIGMA_BASE as i32;
    let scratch_dev       = scratch_buf.dev_ptr;
    let c51_base_i32      = SCRATCH_C51 as i32;
    let iqn_base_i32      = SCRATCH_IQN as i32;
    let cql_base_i32      = SCRATCH_CQL as i32;
    let ens_base_i32      = SCRATCH_ENS as i32;
    let flatness_base_i32 = SCRATCH_FLATNESS as i32;

    unsafe {
        stream
            .launch_builder(&kernel)
            .arg(&isv_dev)
            .arg(&q_var_base_i32)
            .arg(&sigma_base_i32)
            .arg(&scratch_dev)
            .arg(&c51_base_i32)
            .arg(&iqn_base_i32)
            .arg(&cql_base_i32)
            .arg(&ens_base_i32)
            .arg(&flatness_base_i32)
            .launch(LaunchConfig {
                grid_dim:         (1, 1, 1),
                block_dim:        (4, 1, 1),
                shared_mem_bytes: 0,
            })
            .expect("launch pearl_2_budget_update (sharp regime)");
    }
    stream.synchronize().expect("sync after pearl_2_budget_update (sharp regime)");

    let scratch = scratch_buf.read_all();

    // Analytical values for flatness = 0.04 / (4.0 + 1e-8) ≈ 0.01:
    let expected_sigma_sq = 2.0_f32 * 2.0_f32 + 1e-8_f32;
    let expected_flatness = (0.04_f32 / expected_sigma_sq).clamp(0.0, 1.0);
    const BASE_IQN: f32       = 0.11_f32;
    const BASE_C51_SHARE: f32 = 0.84_f32 / (1.0_f32 - BASE_IQN);
    let expected_iqn = BASE_IQN + (1.0_f32 - expected_flatness) * (1.0_f32 - BASE_IQN);
    let expected_c51 = expected_flatness * (1.0_f32 - BASE_IQN) * BASE_C51_SHARE;
    let expected_cql = 0.0_f32;
    let expected_ens = (1.0_f32 - expected_iqn - expected_c51 - expected_cql).max(0.0_f32);

    for b in 0..4_usize {
        let c51      = scratch[SCRATCH_C51 + b];
        let iqn      = scratch[SCRATCH_IQN + b];
        let cql      = scratch[SCRATCH_CQL + b];
        let ens      = scratch[SCRATCH_ENS + b];
        let flatness = scratch[SCRATCH_FLATNESS + b];

        println!(
            "branch={b} flatness={flatness:.6} iqn={iqn:.6} c51={c51:.6} \
             cql={cql:.6} ens={ens:.6}"
        );

        let rel_flatness = (flatness - expected_flatness).abs() / expected_flatness.max(1e-9);
        assert!(
            rel_flatness < 0.01,
            "branch={b}: flatness={flatness:.6} expected={expected_flatness:.6} rel={rel_flatness:.5}"
        );
        let rel_iqn = (iqn - expected_iqn).abs() / expected_iqn.max(1e-9);
        assert!(
            rel_iqn < 0.01,
            "branch={b}: iqn={iqn:.6} expected={expected_iqn:.6} rel={rel_iqn:.5}"
        );
        // IQN must dominate: iqn >> c51
        assert!(
            iqn > 0.9_f32,
            "branch={b}: iqn={iqn:.6} should exceed 0.9 in sharp regime"
        );
        let rel_c51 = (c51 - expected_c51).abs() / expected_c51.max(1e-9);
        assert!(
            rel_c51 < 0.01,
            "branch={b}: c51={c51:.6} expected={expected_c51:.6} rel={rel_c51:.5}"
        );
        assert!(
            cql.abs() < 1e-7,
            "branch={b}: cql={cql:.6} expected=0"
        );
        // ens is near-zero in sharp regime; just verify it is non-negative
        assert!(
            ens >= 0.0_f32,
            "branch={b}: ens={ens:.6} must be non-negative"
        );
        let _ = expected_ens; // expected is ~0.0005, just verify non-negative above
        // Normalization invariant: budgets sum to 1.0 (ens is computed as
        // residual). Catches coefficient typos that would still pass the
        // per-component dominance + relative-tolerance checks.
        let sum = c51 + iqn + cql + ens;
        assert!(
            (sum - 1.0_f32).abs() < 1e-4,
            "branch={b}: budget sum={sum:.6} should be ≈ 1.0 (c51={c51:.6} iqn={iqn:.6} cql={cql:.6} ens={ens:.6})"
        );
    }
}

// ── Tests 7 + 8: pearl_4_adam_hparams_update ──────────────────────────────────

/// SP5 Task A4 unit test: `pearl_4_adam_hparams_update` with high stability input.
///
/// Synthetic input: cosine_sim=[1.0; 8] (maximum stability), grad_norm=[1.0; 8].
///
/// Analytical expected:
///   stability = fmaxf(0, fminf(1, cosine_sim)) = 1.0
///   β1 = 0.85 + 0.10 × 1.0 = 0.95   (long-memory ceiling)
///   β2 = 0.99 + 0.0095 × 1.0 = 0.9995 (long-memory ceiling)
///   ε  = clamp(1.0 × 1e-7, 1e-10, 1e-6) = 1e-7 → clamped to 1e-7 (within [1e-10, 1e-6])
///
/// Structural envelopes (Invariant 1 anchors): β1∈[0.85,0.95], β2∈[0.99,0.9995].
/// No CPU reference oracle per `feedback_no_cpu_test_fallbacks.md`.
#[test]
#[ignore = "requires GPU"]
fn pearl_4_adam_hparams_high_stability_yields_long_memory() {
    // Scratch layout (Pearl 4 outputs):
    //   [0..8)  = β1[8] (beta1_idx_base=0)
    //   [8..16) = β2[8] (beta2_idx_base=8)
    //   [16..24) = ε[8] (eps_idx_base=16)
    const BETA1_BASE:    usize = 0;
    const BETA2_BASE:    usize = 8;
    const EPS_BASE:      usize = 16;
    const SCRATCH_SIZE:  usize = 24;
    const N_GROUPS:      usize = 8;

    let stream = make_test_stream();
    let kernel = load_pearl_4_adam_hparams_kernel(&stream);

    // cosine_sim = 1.0 for all 8 groups → maximum stability.
    let cosine_buf = unsafe { MappedF32Buffer::new(N_GROUPS) }
        .expect("alloc cosine_buf");
    cosine_buf.write_from_slice(&[1.0_f32; N_GROUPS]);

    // grad_norm = 1.0 for all 8 groups → ε = 1e-7 (within envelope).
    let l2_norm_buf = unsafe { MappedF32Buffer::new(N_GROUPS) }
        .expect("alloc l2_norm_buf");
    l2_norm_buf.write_from_slice(&[1.0_f32; N_GROUPS]);

    let scratch_buf = unsafe { MappedF32Buffer::new(SCRATCH_SIZE) }
        .expect("alloc scratch_buf");

    let cosine_dev      = cosine_buf.dev_ptr;
    let l2_dev          = l2_norm_buf.dev_ptr;
    let scratch_dev     = scratch_buf.dev_ptr;
    let beta1_base_i32  = BETA1_BASE as i32;
    let beta2_base_i32  = BETA2_BASE as i32;
    let eps_base_i32    = EPS_BASE   as i32;

    unsafe {
        stream
            .launch_builder(&kernel)
            .arg(&cosine_dev)
            .arg(&l2_dev)
            .arg(&scratch_dev)
            .arg(&beta1_base_i32)
            .arg(&beta2_base_i32)
            .arg(&eps_base_i32)
            .launch(LaunchConfig {
                grid_dim:         (1, 1, 1),
                block_dim:        (8, 1, 1),
                shared_mem_bytes: 0,
            })
            .expect("launch pearl_4_adam_hparams_update (high stability)");
    }
    stream.synchronize().expect("sync after pearl_4_adam_hparams_update (high stability)");

    let scratch = scratch_buf.read_all();

    // Analytical values for stability = 1.0:
    let expected_beta1: f32 = 0.85 + 0.10 * 1.0;   // = 0.95
    let expected_beta2: f32 = 0.99 + 0.0095 * 1.0;  // = 0.9995
    // ε = clamp(1.0 × 1e-7, 1e-10, 1e-6) = 1e-7
    let expected_eps: f32   = 1.0_f32 * 1e-7_f32;   // within [1e-10, 1e-6]

    for g in 0..N_GROUPS {
        let beta1 = scratch[BETA1_BASE + g];
        let beta2 = scratch[BETA2_BASE + g];
        let eps   = scratch[EPS_BASE   + g];

        println!("group={g} beta1={beta1:.6} beta2={beta6:.7} eps={eps:.2e}",
            beta6 = beta2);

        // β1 ceiling: 0.95 ± 1 ULP.
        let rel_b1 = (beta1 - expected_beta1).abs() / expected_beta1.max(1e-9);
        assert!(
            rel_b1 < 1e-5,
            "group={g}: beta1={beta1:.6} expected={expected_beta1:.6} rel={rel_b1:.2e}"
        );
        // β2 ceiling: 0.9995 ± 1 ULP.
        let rel_b2 = (beta2 - expected_beta2).abs() / expected_beta2.max(1e-9);
        assert!(
            rel_b2 < 1e-5,
            "group={g}: beta2={beta2:.7} expected={expected_beta2:.7} rel={rel_b2:.2e}"
        );
        // ε must be within structural envelope.
        assert!(
            eps >= 1e-10_f32 && eps <= 1e-6_f32,
            "group={g}: eps={eps:.2e} outside structural envelope [1e-10, 1e-6]"
        );
        let rel_eps = (eps - expected_eps).abs() / expected_eps.max(1e-15);
        assert!(
            rel_eps < 1e-4,
            "group={g}: eps={eps:.2e} expected={expected_eps:.2e} rel={rel_eps:.2e}"
        );
        // β1 must be at structural ceiling for maximum stability.
        assert!(
            (beta1 - 0.95_f32).abs() < 1e-5,
            "group={g}: beta1={beta1:.6} should be at ceiling 0.95 for stability=1.0"
        );
        // β2 must be at structural ceiling for maximum stability.
        assert!(
            (beta2 - 0.9995_f32).abs() < 1e-5,
            "group={g}: beta2={beta2:.7} should be at ceiling 0.9995 for stability=1.0"
        );
    }
}

/// SP5 Task A4 unit test: `pearl_4_adam_hparams_update` with low stability input.
///
/// Synthetic input: cosine_sim=[0.0; 8] (minimum stability), grad_norm=[0.01; 8].
///
/// Analytical expected:
///   stability = fmaxf(0, fminf(1, 0.0)) = 0.0
///   β1 = 0.85 + 0.10 × 0.0 = 0.85   (low-stability floor)
///   β2 = 0.99 + 0.0095 × 0.0 = 0.99  (low-stability floor)
///   ε  = clamp(0.01 × 1e-7, 1e-10, 1e-6) = 1e-9 → clamped to 1e-10
///
/// No CPU reference oracle per `feedback_no_cpu_test_fallbacks.md`.
#[test]
#[ignore = "requires GPU"]
fn pearl_4_adam_hparams_low_stability_yields_short_memory() {
    const BETA1_BASE:    usize = 0;
    const BETA2_BASE:    usize = 8;
    const EPS_BASE:      usize = 16;
    const SCRATCH_SIZE:  usize = 24;
    const N_GROUPS:      usize = 8;

    let stream = make_test_stream();
    let kernel = load_pearl_4_adam_hparams_kernel(&stream);

    // cosine_sim = 0.0 for all 8 groups → minimum stability.
    let cosine_buf = unsafe { MappedF32Buffer::new(N_GROUPS) }
        .expect("alloc cosine_buf");
    cosine_buf.write_from_slice(&[0.0_f32; N_GROUPS]);

    // grad_norm = 0.01 → ε = 0.01 × 1e-7 = 1e-9 → clamped to 1e-10.
    let l2_norm_buf = unsafe { MappedF32Buffer::new(N_GROUPS) }
        .expect("alloc l2_norm_buf");
    l2_norm_buf.write_from_slice(&[0.01_f32; N_GROUPS]);

    let scratch_buf = unsafe { MappedF32Buffer::new(SCRATCH_SIZE) }
        .expect("alloc scratch_buf");

    let cosine_dev      = cosine_buf.dev_ptr;
    let l2_dev          = l2_norm_buf.dev_ptr;
    let scratch_dev     = scratch_buf.dev_ptr;
    let beta1_base_i32  = BETA1_BASE as i32;
    let beta2_base_i32  = BETA2_BASE as i32;
    let eps_base_i32    = EPS_BASE   as i32;

    unsafe {
        stream
            .launch_builder(&kernel)
            .arg(&cosine_dev)
            .arg(&l2_dev)
            .arg(&scratch_dev)
            .arg(&beta1_base_i32)
            .arg(&beta2_base_i32)
            .arg(&eps_base_i32)
            .launch(LaunchConfig {
                grid_dim:         (1, 1, 1),
                block_dim:        (8, 1, 1),
                shared_mem_bytes: 0,
            })
            .expect("launch pearl_4_adam_hparams_update (low stability)");
    }
    stream.synchronize().expect("sync after pearl_4_adam_hparams_update (low stability)");

    let scratch = scratch_buf.read_all();

    // Analytical values for stability = 0.0:
    let expected_beta1: f32 = 0.85;   // low-stability floor
    let expected_beta2: f32 = 0.99;   // low-stability floor
    // ε = clamp(0.01 × 1e-7, 1e-10, 1e-6) = clamp(1e-9, 1e-10, 1e-6) = 1e-9
    // Actually 1e-9 is within [1e-10, 1e-6], so no clamping needed.
    let expected_eps: f32   = (0.01_f32 * 1e-7_f32).clamp(1e-10, 1e-6);  // = 1e-9

    for g in 0..N_GROUPS {
        let beta1 = scratch[BETA1_BASE + g];
        let beta2 = scratch[BETA2_BASE + g];
        let eps   = scratch[EPS_BASE   + g];

        println!("group={g} beta1={beta1:.6} beta2={beta2:.7} eps={eps:.2e}");

        // β1 floor: 0.85 ± 1 ULP.
        let rel_b1 = (beta1 - expected_beta1).abs() / expected_beta1.max(1e-9);
        assert!(
            rel_b1 < 1e-5,
            "group={g}: beta1={beta1:.6} expected={expected_beta1:.6} rel={rel_b1:.2e}"
        );
        // β2 floor: 0.99 ± 1 ULP.
        let rel_b2 = (beta2 - expected_beta2).abs() / expected_beta2.max(1e-9);
        assert!(
            rel_b2 < 1e-5,
            "group={g}: beta2={beta2:.7} expected={expected_beta2:.7} rel={rel_b2:.2e}"
        );
        // ε must be within structural envelope.
        assert!(
            eps >= 1e-10_f32 && eps <= 1e-6_f32,
            "group={g}: eps={eps:.2e} outside structural envelope [1e-10, 1e-6]"
        );
        let rel_eps = (eps - expected_eps).abs() / expected_eps.max(1e-15);
        assert!(
            rel_eps < 1e-4,
            "group={g}: eps={eps:.2e} expected={expected_eps:.2e} rel={rel_eps:.2e}"
        );
        // β1 must be at structural floor for zero stability.
        assert!(
            (beta1 - 0.85_f32).abs() < 1e-5,
            "group={g}: beta1={beta1:.6} should be at floor 0.85 for stability=0.0"
        );
        // β2 must be at structural floor for zero stability.
        assert!(
            (beta2 - 0.99_f32).abs() < 1e-5,
            "group={g}: beta2={beta2:.7} should be at floor 0.99 for stability=0.0"
        );
    }
}

// ── Test 9: grad_cosine_sim_update ────────────────────────────────────────────

/// SP5 Task A4 unit test: `grad_cosine_sim_update` correctly computes per-group
/// cosine similarity, L2 norm, and writes back grad_curr → grad_prev.
///
/// Synthetic input: 8 groups × 4 params = 32 total params.
///   group 0: grad_curr = [1, 0, 0, 0], grad_prev = [1, 0, 0, 0]    → cos = +1.0
///   group 1: grad_curr = [1, 0, 0, 0], grad_prev = [0, 1, 0, 0]    → cos =  0.0 (orthogonal)
///   group 2: grad_curr = [1, 0, 0, 0], grad_prev = [-1, 0, 0, 0]   → cos = -1.0 (antiparallel; raw, NOT clamped here)
///   group 3: grad_curr = [3, 4, 0, 0], grad_prev = [3, 4, 0, 0]    → cos = +1.0; |g_curr|=5
///   group 4: grad_curr = [0, 0, 0, 0], grad_prev = [1, 0, 0, 0]    → cos =  0/(0×1+EPS) = 0.0; |g_curr|=0 (cold-start curr)
///   group 5: grad_curr = [1, 0, 0, 0], grad_prev = [0, 0, 0, 0]    → cos =  0/(1×0+EPS) = 0.0; |g_curr|=1 (cold-start prev)
///   group 6..7: identical to group 0 for repetition.
///
/// Then verify writeback: after the kernel runs, grad_prev_buf should equal
/// grad_curr_buf for all 32 elements (writeback is the load-bearing state
/// machinery for cross-step cosine evolution).
///
/// NO CPU reference oracle — all expected values are analytically known unit-vector cosines.
#[test]
#[ignore = "requires GPU"]
fn grad_cosine_sim_per_group_dot_norm_and_writeback() {
    const N_GROUPS: usize = 8;
    const PARAMS_PER_GROUP: usize = 4;
    const TOTAL_PARAMS: usize = N_GROUPS * PARAMS_PER_GROUP;
    const SCRATCH_BASE_COSINE: usize = 0;
    const SCRATCH_BASE_L2: usize = 8;

    let stream = make_test_stream();
    let kernel = load_grad_cosine_sim_kernel(&stream);

    let grad_curr_buf = unsafe { MappedF32Buffer::new(TOTAL_PARAMS) }
        .expect("alloc grad_curr_buf");
    let grad_prev_buf = unsafe { MappedF32Buffer::new(TOTAL_PARAMS) }
        .expect("alloc grad_prev_buf");
    let offsets_buf = unsafe { MappedI32Buffer::new(N_GROUPS + 1) }
        .expect("alloc group_param_offsets_buf");
    let scratch_buf = unsafe { MappedF32Buffer::new(16) }  // 8 cosine + 8 L2
        .expect("alloc scratch_buf");

    // Group offsets: [0, 4, 8, 12, 16, 20, 24, 28, 32]
    let offsets: Vec<i32> = (0..=N_GROUPS as i32).map(|i| i * PARAMS_PER_GROUP as i32).collect();
    offsets_buf.write_from_slice(&offsets);

    // Per-group synthetic vectors (analytical cos_sim and L2 norms).
    let mut grad_curr = vec![0.0_f32; TOTAL_PARAMS];
    let mut grad_prev = vec![0.0_f32; TOTAL_PARAMS];
    // group 0: curr=prev=e1 → cos=+1, |c|=1
    grad_curr[0] = 1.0; grad_prev[0] = 1.0;
    // group 1: curr=e1, prev=e2 → cos=0, |c|=1
    grad_curr[4] = 1.0; grad_prev[5] = 1.0;
    // group 2: curr=e1, prev=-e1 → cos=-1, |c|=1
    grad_curr[8] = 1.0; grad_prev[8] = -1.0;
    // group 3: curr=prev=(3,4,0,0) → cos=+1, |c|=5
    grad_curr[12] = 3.0; grad_curr[13] = 4.0;
    grad_prev[12] = 3.0; grad_prev[13] = 4.0;
    // group 4: curr=0, prev=e1 → cos = 0/(0×1+EPS) = 0, |c|=0 (cold-start curr)
    grad_prev[17] = 1.0;
    // group 5: curr=e1, prev=0 → cos = 0/(1×0+EPS) = 0, |c|=1 (cold-start prev — Pearl A sentinel state)
    grad_curr[20] = 1.0;
    // group 6: same as group 0
    grad_curr[24] = 1.0; grad_prev[24] = 1.0;
    // group 7: same as group 0
    grad_curr[28] = 1.0; grad_prev[28] = 1.0;

    grad_curr_buf.write_from_slice(&grad_curr);
    grad_prev_buf.write_from_slice(&grad_prev);

    let curr_ptr        = grad_curr_buf.dev_ptr;
    let prev_ptr        = grad_prev_buf.dev_ptr;
    let offsets_dev     = offsets_buf.dev_ptr;
    let scratch_dev     = scratch_buf.dev_ptr;
    let cosine_base_i32 = SCRATCH_BASE_COSINE as i32;
    let l2_base_i32     = SCRATCH_BASE_L2 as i32;
    let writeback_ptr   = grad_prev_buf.dev_ptr;  // writeback target = grad_prev_buf itself

    unsafe {
        stream
            .launch_builder(&kernel)
            .arg(&curr_ptr)
            .arg(&prev_ptr)
            .arg(&offsets_dev)
            .arg(&scratch_dev)
            .arg(&cosine_base_i32)
            .arg(&l2_base_i32)
            .arg(&writeback_ptr)
            .launch(LaunchConfig {
                grid_dim:         (1, 1, 1),
                block_dim:        (8, 1, 1),
                shared_mem_bytes: 0,
            })
            .expect("launch grad_cosine_sim_update");
    }
    stream.synchronize().expect("sync after grad_cosine_sim_update");

    let scratch = scratch_buf.read_all();
    let prev_after = grad_prev_buf.read_all();

    let expected_cos = [1.0_f32, 0.0, -1.0, 1.0, 0.0, 0.0, 1.0, 1.0];
    let expected_l2  = [1.0_f32, 1.0,  1.0, 5.0, 0.0, 1.0, 1.0, 1.0];

    for g in 0..N_GROUPS {
        let cos = scratch[SCRATCH_BASE_COSINE + g];
        let l2  = scratch[SCRATCH_BASE_L2 + g];
        println!("group={g} cos={cos:.5} (expected {:.5}) l2={l2:.5} (expected {:.5})",
                 expected_cos[g], expected_l2[g]);

        assert!(
            (cos - expected_cos[g]).abs() < 1e-5,
            "group={g}: cos={cos:.5} expected={:.5}", expected_cos[g]
        );
        assert!(
            (l2 - expected_l2[g]).abs() < 1e-5,
            "group={g}: l2={l2:.5} expected={:.5}", expected_l2[g]
        );
    }

    // Writeback verification: grad_prev_buf must now equal grad_curr_buf.
    // This is load-bearing for cross-step cosine evolution — if writeback
    // is wrong, every subsequent step's cosine_sim is computed against a
    // stale or mis-aligned previous gradient.
    for i in 0..TOTAL_PARAMS {
        assert!(
            (prev_after[i] - grad_curr[i]).abs() < 1e-7,
            "writeback mismatch at i={i}: prev_after={} != grad_curr={}",
            prev_after[i], grad_curr[i]
        );
    }
}

// ── Test 10 + 11 helpers (Task A5) ────────────────────────────────────────────

fn load_q_skew_kurtosis_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
    let module = stream
        .context()
        .load_cubin(SP5_Q_SKEW_KURTOSIS_CUBIN.to_vec())
        .expect("load q_skew_kurtosis_kernel cubin");
    module
        .load_function("q_skew_kurtosis_update")
        .expect("load q_skew_kurtosis_update function")
}

fn load_pearl_5_iqn_tau_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
    let module = stream
        .context()
        .load_cubin(SP5_PEARL_5_IQN_TAU_CUBIN.to_vec())
        .expect("load pearl_5_iqn_tau_kernel cubin");
    module
        .load_function("pearl_5_iqn_tau_update")
        .expect("load pearl_5_iqn_tau_update function")
}

// ── Test 10: zero skew → symmetric default τ schedule ─────────────────────────

/// SP5 Task A5 unit test: constant Q-values (zero variance → zero skew) produce
/// the symmetric default τ schedule {0.05, 0.25, 0.5, 0.75, 0.95} on all branches.
///
/// Synthetic input: B=4 batch, q_out_buf [4 × 13] all set to constant 3.0.
///
/// Analytical ground truth for constant input:
///   mean = 3.0, all deviations = 0 → m2 = m3 = m4 = 0
///   m2_15 = 0 + EPS_DIV = 1e-12
///   m3 / m2_15 = 0 / 1e-12 = 0 → skew = 0 (before clamp, still 0)
///   τ[q] = tau_defaults[q] + 0 × SKEW_SHIFT = tau_defaults[q]
///   Expected τ per branch: {0.05, 0.25, 0.5, 0.75, 0.95}
///
/// This test chains q_skew_kurtosis_update → pearl_5_iqn_tau_update exactly as
/// the production launcher does (scratch[171..175) is the skew bridge slot).
/// No CPU reference computation per `feedback_no_cpu_test_fallbacks.md`.
#[test]
#[ignore = "requires GPU"]
fn pearl_5_iqn_tau_symmetric_q_yields_uniform_default_schedule() {
    const BATCH_SIZE: usize        = 4;
    const TOTAL_ACTIONS: usize     = 13;
    const SKEW_IDX_BASE: usize     = 171; // SCRATCH_PEARL_5_SKEW
    const EX_KURT_IDX_BASE: usize  = 175; // SCRATCH_PEARL_5_EX_KURT
    const TAU_IDX_BASE: usize      = 179; // SCRATCH_PEARL_5_TAU
    const SCRATCH_SIZE: usize      = TAU_IDX_BASE + 20; // 199

    let stream = make_test_stream();
    let skew_kernel = load_q_skew_kurtosis_kernel(&stream);
    let tau_kernel  = load_pearl_5_iqn_tau_kernel(&stream);

    // q_out_buf [B × 13]: all 3.0 — constant Q across all branches.
    let q_buf = unsafe { MappedF32Buffer::new(BATCH_SIZE * TOTAL_ACTIONS) }
        .expect("alloc q_buf");
    let q_data = vec![3.0_f32; BATCH_SIZE * TOTAL_ACTIONS];
    q_buf.write_from_slice(&q_data);

    // action_counts = {4, 3, 3, 3}
    let action_counts = [4_i32, 3, 3, 3];
    let action_counts_buf = unsafe { MappedI32Buffer::new(4) }
        .expect("alloc action_counts_buf");
    action_counts_buf.write_from_slice(&action_counts);

    // action_offsets = {0, 4, 7, 10}
    let action_offsets = [0_i32, 4, 7, 10];
    let action_offsets_buf = unsafe { MappedI32Buffer::new(4) }
        .expect("alloc action_offsets_buf");
    action_offsets_buf.write_from_slice(&action_offsets);

    let scratch_buf = unsafe { MappedF32Buffer::new(SCRATCH_SIZE) }
        .expect("alloc scratch_buf");

    let batch_size_i32       = BATCH_SIZE as i32;
    let skew_idx_i32         = SKEW_IDX_BASE as i32;
    let ex_kurt_idx_i32      = EX_KURT_IDX_BASE as i32;
    let tau_idx_i32          = TAU_IDX_BASE as i32;

    let q_dev              = q_buf.dev_ptr;
    let counts_dev         = action_counts_buf.dev_ptr;
    let offsets_dev        = action_offsets_buf.dev_ptr;
    let scratch_dev        = scratch_buf.dev_ptr;

    // Step 1: q_skew_kurtosis_update
    unsafe {
        stream
            .launch_builder(&skew_kernel)
            .arg(&q_dev)
            .arg(&batch_size_i32)
            .arg(&counts_dev)
            .arg(&offsets_dev)
            .arg(&scratch_dev)
            .arg(&skew_idx_i32)
            .arg(&ex_kurt_idx_i32)
            .launch(LaunchConfig {
                grid_dim:         (1, 1, 1),
                block_dim:        (4, 1, 1),
                shared_mem_bytes: 0,
            })
            .expect("launch q_skew_kurtosis_update (zero-skew test)");
    }
    stream.synchronize().expect("sync after q_skew_kurtosis_update (zero-skew test)");

    // Step 2: pearl_5_iqn_tau_update
    unsafe {
        stream
            .launch_builder(&tau_kernel)
            .arg(&scratch_dev)
            .arg(&skew_idx_i32)
            .arg(&scratch_dev)
            .arg(&tau_idx_i32)
            .launch(LaunchConfig {
                grid_dim:         (1, 1, 1),
                block_dim:        (4, 1, 1),
                shared_mem_bytes: 0,
            })
            .expect("launch pearl_5_iqn_tau_update (zero-skew test)");
    }
    stream.synchronize().expect("sync after pearl_5_iqn_tau_update (zero-skew test)");

    let scratch = scratch_buf.read_all();

    // Verify skew[4] ≈ 0 (constant input → zero skew)
    for b in 0..4_usize {
        let skew = scratch[SKEW_IDX_BASE + b];
        println!("branch={b} skew={skew:.8}");
        assert!(
            skew.abs() < 1e-5,
            "branch={b}: expected skew≈0 for constant Q, got {skew}"
        );
    }

    // Verify τ[4×5] ≈ symmetric defaults {0.05, 0.25, 0.5, 0.75, 0.95}
    let defaults = [0.05_f32, 0.25, 0.5, 0.75, 0.95];
    for b in 0..4_usize {
        for q in 0..5_usize {
            let tau = scratch[TAU_IDX_BASE + b * 5 + q];
            let expected = defaults[q];
            println!("branch={b} q={q} tau={tau:.6} expected={expected:.6}");
            assert!(
                (tau - expected).abs() < 1e-5,
                "branch={b} q={q}: tau={tau:.6} expected={expected:.6}"
            );
        }
    }
}

// ── Test 11: negative skew shifts τ schedule down ─────────────────────────────

/// SP5 Task A5 unit test: negative Q-distribution skew shifts the τ schedule
/// downward (toward lower quantiles), verifying the SKEW_SHIFT=0.05 mechanism.
///
/// Synthetic input: inject skew = -1.0 directly into the scratch buffer
/// (SCRATCH_PEARL_5_SKEW slot) and launch only pearl_5_iqn_tau_update.
/// This isolates the τ-schedule kernel from the skew-computation kernel.
///
/// Analytical ground truth for skew = -1.0:
///   τ[q] = tau_defaults[q] + (-1.0) × 0.05 = tau_defaults[q] - 0.05
///   Expected τ: {0.0, 0.20, 0.45, 0.70, 0.90}
///   Clamp check: 0.05 - 0.05 = 0.00 → clamped to 0.01 (structural envelope [0.01, 0.99])
///   So expected τ[0] = 0.01, τ[1..5] = {0.20, 0.45, 0.70, 0.90}
///
/// Checks that τ values are strictly less than the symmetric defaults (shifted down)
/// and that the 0.01 floor clamp fires on the first quantile.
/// No CPU reference computation per `feedback_no_cpu_test_fallbacks.md`.
#[test]
#[ignore = "requires GPU"]
fn pearl_5_iqn_tau_left_skew_shifts_quantiles_negatively() {
    const SKEW_IDX_BASE: usize = 171; // SCRATCH_PEARL_5_SKEW
    const TAU_IDX_BASE: usize  = 179; // SCRATCH_PEARL_5_TAU
    const SCRATCH_SIZE: usize  = TAU_IDX_BASE + 20; // 199
    const INJECTED_SKEW: f32   = -1.0;

    let stream = make_test_stream();
    let tau_kernel = load_pearl_5_iqn_tau_kernel(&stream);

    let scratch_buf = unsafe { MappedF32Buffer::new(SCRATCH_SIZE) }
        .expect("alloc scratch_buf");

    // Inject skew = -1.0 for all 4 branches directly into the scratch slot.
    let mut scratch_init = vec![0.0_f32; SCRATCH_SIZE];
    for b in 0..4_usize {
        scratch_init[SKEW_IDX_BASE + b] = INJECTED_SKEW;
    }
    scratch_buf.write_from_slice(&scratch_init);

    let scratch_dev = scratch_buf.dev_ptr;
    let skew_idx_i32 = SKEW_IDX_BASE as i32;
    let tau_idx_i32  = TAU_IDX_BASE as i32;

    unsafe {
        stream
            .launch_builder(&tau_kernel)
            .arg(&scratch_dev)
            .arg(&skew_idx_i32)
            .arg(&scratch_dev)
            .arg(&tau_idx_i32)
            .launch(LaunchConfig {
                grid_dim:         (1, 1, 1),
                block_dim:        (4, 1, 1),
                shared_mem_bytes: 0,
            })
            .expect("launch pearl_5_iqn_tau_update (left-skew test)");
    }
    stream.synchronize().expect("sync after pearl_5_iqn_tau_update (left-skew test)");

    let scratch = scratch_buf.read_all();

    // Analytical expected τ for skew = -1.0:
    // tau_defaults[q] + (-1.0) × 0.05, clamped to [0.01, 0.99]
    let expected = [0.01_f32, 0.20, 0.45, 0.70, 0.90];

    for b in 0..4_usize {
        for q in 0..5_usize {
            let tau = scratch[TAU_IDX_BASE + b * 5 + q];
            let exp = expected[q];
            println!("branch={b} q={q} tau={tau:.6} expected={exp:.6}");
            assert!(
                (tau - exp).abs() < 1e-5,
                "branch={b} q={q}: tau={tau:.6} expected={exp:.6}"
            );
        }
        // τ[1..5] must be strictly below their symmetric defaults (skew shifted them down).
        let defaults = [0.05_f32, 0.25, 0.5, 0.75, 0.95];
        for q in 1..5_usize {
            let tau = scratch[TAU_IDX_BASE + b * 5 + q];
            assert!(
                tau < defaults[q],
                "branch={b} q={q}: tau={tau:.6} should be below default={:.6}",
                defaults[q]
            );
        }
        // τ[0] must be at the floor (clamp fired: 0.05 - 0.05 = 0.00 → clamped to 0.01).
        let tau0 = scratch[TAU_IDX_BASE + b * 5];
        assert!(
            (tau0 - 0.01_f32).abs() < 1e-5,
            "branch={b} q=0: expected clamp floor 0.01, got {tau0:.6}"
        );
    }
}

// ── Pearl 6 helpers ───────────────────────────────────────────────────────────

fn load_pearl_6_kelly_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
    let ctx = stream.context();
    let module = ctx.load_cubin(SP5_PEARL_6_KELLY_CUBIN.to_vec())
        .expect("load pearl_6_kelly_kernel cubin");
    module.load_function("pearl_6_kelly_update")
        .expect("load pearl_6_kelly_update function")
}

/// Allocate a MappedF32Buffer for portfolio_state [n_envs * ps_stride] and ISV.
/// Returns (portfolio_buf, isv_buf). Both are mapped-pinned (no DtoH).
/// PS_STRIDE=43 matches state_layout.cuh.
const PS_STRIDE: usize = 43;
const PS_KELLY_WIN_COUNT:  usize = 14;
const PS_KELLY_LOSS_COUNT: usize = 15;
const PS_KELLY_SUM_WINS:   usize = 16;
const PS_KELLY_SUM_LOSSES: usize = 17;

// ISV slot indices for Pearl 6 outputs (match sp5_isv_slots.rs).
const KELLY_F_SMOOTH_IDX:      usize = 280;
const CONVICTION_SMOOTH_IDX:   usize = 281;
const TRADE_VAR_SMOOTH_IDX:    usize = 282;
const KELLY_SAMPLE_COUNT_IDX:  usize = 283;
const WIN_RATE_SMOOTH_IDX:     usize = 284;
const LOSS_RATE_SMOOTH_IDX:    usize = 285;

fn launch_pearl_6(
    stream:   &Arc<CudaStream>,
    kernel:   &CudaFunction,
    ps_buf:   &MappedF32Buffer,
    isv_buf:  &MappedF32Buffer,
    n_envs:   i32,
) {
    let ps_dev        = ps_buf.dev_ptr;
    let isv_dev       = isv_buf.dev_ptr;
    let ps_stride_i32 = PS_STRIDE as i32;
    let kelly_f_i32   = KELLY_F_SMOOTH_IDX     as i32;
    let conv_i32      = CONVICTION_SMOOTH_IDX  as i32;
    let tvar_i32      = TRADE_VAR_SMOOTH_IDX   as i32;
    let scount_i32    = KELLY_SAMPLE_COUNT_IDX as i32;
    let wr_i32        = WIN_RATE_SMOOTH_IDX    as i32;
    let lr_i32        = LOSS_RATE_SMOOTH_IDX   as i32;

    unsafe {
        stream
            .launch_builder(kernel)
            .arg(&ps_dev)
            .arg(&n_envs)
            .arg(&ps_stride_i32)
            .arg(&isv_dev)
            .arg(&kelly_f_i32)
            .arg(&conv_i32)
            .arg(&tvar_i32)
            .arg(&scount_i32)
            .arg(&wr_i32)
            .arg(&lr_i32)
            .launch(LaunchConfig {
                grid_dim:         (1, 1, 1),
                block_dim:        (6, 1, 1),
                shared_mem_bytes: 0,
            })
            .expect("launch pearl_6_kelly_update");
    }
    stream.synchronize().expect("sync after pearl_6_kelly_update");
}

/// SP5 Task A6 Test 12: within-fold EWMA blend.
///
/// Analytical derivation:
///   Single env with win_count=10, loss_count=5, sum_wins=0.5, sum_losses=0.25.
///   Bayesian priors: prior_wins=2, prior_losses=2, prior_sum_wins=0.01, prior_sum_losses=0.01.
///   eff_w=12, eff_l=7, wr=12/19≈0.6316, avg_w=(0.5+0.01)/12≈0.0425, avg_l=(0.25+0.01)/7≈0.0371
///   payoff=0.0425/0.0371≈1.1456, kelly_f=(1.1456×0.6316 - 0.3684)/1.1456≈(0.7234-0.3684)/1.1456≈0.3100
///   EWMA: new_isv = 0.99 × 0.5 + 0.01 × computed_kelly_f
///   ISV initial = 0.5 for all slots.
///
/// For sample_count (s==3): max(0.5, 15.0) = 15.0 (new_obs = total_trades = 15 > current=0.5).
/// For win_rate (s==4): new_obs = 10/15 ≈ 0.6667, new_isv = 0.99×0.5 + 0.01×0.6667 ≈ 0.5017.
/// For loss_rate (s==5): new_obs = 5/15 ≈ 0.3333, new_isv = 0.99×0.5 + 0.01×0.3333 ≈ 0.4983.
#[test]
#[ignore = "requires GPU"]
fn pearl_6_kelly_within_fold_ewma_blend() {
    let ctx    = CudaContext::new(0).expect("CUDA context");
    let stream = ctx.new_stream().expect("stream");
    let kernel = load_pearl_6_kelly_kernel(&stream);

    let n_envs = 1_i32;
    let isv_size = 290_usize; // covers slots [280..286)

    let mut ps_buf  = unsafe { MappedF32Buffer::new(n_envs as usize * PS_STRIDE) }.unwrap();
    let mut isv_buf = unsafe { MappedF32Buffer::new(isv_size) }.unwrap();

    // Set portfolio_state for env 0.
    {
        let ps_slice = ps_buf.host_slice_mut();
        ps_slice.fill(0.0_f32);
        ps_slice[PS_KELLY_WIN_COUNT]  = 10.0;
        ps_slice[PS_KELLY_LOSS_COUNT] =  5.0;
        ps_slice[PS_KELLY_SUM_WINS]   =  0.5;
        ps_slice[PS_KELLY_SUM_LOSSES] =  0.25;
    }

    // Initialise ISV to 0.5 for all Pearl 6 slots.
    {
        let isv_slice = isv_buf.host_slice_mut();
        isv_slice.fill(0.0_f32);
        isv_slice[KELLY_F_SMOOTH_IDX]     = 0.5;
        isv_slice[CONVICTION_SMOOTH_IDX]  = 0.5;
        isv_slice[TRADE_VAR_SMOOTH_IDX]   = 0.5;
        isv_slice[KELLY_SAMPLE_COUNT_IDX] = 0.5;
        isv_slice[WIN_RATE_SMOOTH_IDX]    = 0.5;
        isv_slice[LOSS_RATE_SMOOTH_IDX]   = 0.5;
    }

    launch_pearl_6(&stream, &kernel, &ps_buf, &isv_buf, n_envs);

    let isv = isv_buf.read_all();

    // slot 280 — kelly_f_smooth: EWMA blend of computed kelly_f.
    // Computed kelly_f: eff_w=12, eff_l=7, wr=12/19
    let eff_w = 12.0_f32; let eff_l = 7.0_f32;
    let wr = eff_w / (eff_w + eff_l);
    let avg_w = (0.5_f32 + 0.01) / eff_w;
    let avg_l = (0.25_f32 + 0.01) / eff_l;
    let pf = avg_w / avg_l.max(0.0001);
    let kf = ((pf * wr - (1.0 - wr)) / pf.max(0.0001)).clamp(0.0, 1.0);
    let expected_kelly_f = 0.99 * 0.5 + 0.01 * kf;
    println!("kf={kf:.6} expected_kelly_f={expected_kelly_f:.6} got={:.6}", isv[KELLY_F_SMOOTH_IDX]);
    assert!((isv[KELLY_F_SMOOTH_IDX] - expected_kelly_f).abs() < 2e-5,
        "kelly_f: expected {expected_kelly_f:.6}, got {:.6}", isv[KELLY_F_SMOOTH_IDX]);

    // slot 281 — conviction_smooth: identity (no-op update, current=0.5).
    // EWMA of (new_obs=current=0.5) with current=0.5 → still 0.5.
    assert!((isv[CONVICTION_SMOOTH_IDX] - 0.5).abs() < 1e-6,
        "conviction: expected 0.5 (identity), got {:.6}", isv[CONVICTION_SMOOTH_IDX]);

    // slot 283 — sample_count: cumulative via max(current=0.5, new_obs=15.0) = 15.0.
    assert!((isv[KELLY_SAMPLE_COUNT_IDX] - 15.0).abs() < 1e-4,
        "sample_count: expected 15.0, got {:.6}", isv[KELLY_SAMPLE_COUNT_IDX]);

    // slot 284 — win_rate_smooth: EWMA blend.
    let expected_wr = 0.99 * 0.5 + 0.01 * (10.0_f32 / 15.0);
    assert!((isv[WIN_RATE_SMOOTH_IDX] - expected_wr).abs() < 2e-5,
        "win_rate: expected {expected_wr:.6}, got {:.6}", isv[WIN_RATE_SMOOTH_IDX]);

    // slot 285 — loss_rate_smooth: EWMA blend.
    let expected_lr = 0.99 * 0.5 + 0.01 * (5.0_f32 / 15.0);
    assert!((isv[LOSS_RATE_SMOOTH_IDX] - expected_lr).abs() < 2e-5,
        "loss_rate: expected {expected_lr:.6}, got {:.6}", isv[LOSS_RATE_SMOOTH_IDX]);

    // slot 282 — trade_var_smooth: variance of per-env Kelly fractions across envs.
    // With n_envs=1, kelly_count=1, the kernel's `(kelly_count > 1) ? var : 0.0f`
    // branch returns 0.0 (single-env can't have cross-env variance).
    // EWMA blend: 0.99 × 0.5 + 0.01 × 0.0 = 0.495.
    let expected_tvar = 0.99 * 0.5 + 0.01 * 0.0_f32;
    assert!((isv[TRADE_VAR_SMOOTH_IDX] - expected_tvar).abs() < 2e-5,
        "trade_var: expected {expected_tvar:.6} (n_envs=1 → variance=0; EWMA blend with prior 0.5), got {:.6}",
        isv[TRADE_VAR_SMOOTH_IDX]);
}

/// SP5 Task A6 Test 13: cross-fold sample_count persists via max().
///
/// This test verifies the core cross-fold-persistence mechanism:
///   Step 1: portfolio_state total_trades=100. Run → ISV[283] = max(0, 100) = 100.
///   Step 2: Simulate fold boundary — reset portfolio_state to 0 (total_trades=0).
///           Run → new_obs=0, max(100, 0) = 100. ISV[283] preserved.
///   Step 3: portfolio_state total_trades=10 (recovery, below prior max).
///           Run → new_obs=10, max(100, 10) = 100. Still preserved.
///   Step 4: portfolio_state total_trades=150 (surpasses old max).
///           Run → new_obs=150, max(100, 150) = 150. Grows to new max.
#[test]
#[ignore = "requires GPU"]
fn pearl_6_kelly_cross_fold_sample_count_persists_via_max() {
    let ctx    = CudaContext::new(0).expect("CUDA context");
    let stream = ctx.new_stream().expect("stream");
    let kernel = load_pearl_6_kelly_kernel(&stream);

    let n_envs   = 1_i32;
    let isv_size = 290_usize;

    let mut ps_buf  = unsafe { MappedF32Buffer::new(n_envs as usize * PS_STRIDE) }.unwrap();
    let mut isv_buf = unsafe { MappedF32Buffer::new(isv_size) }.unwrap();

    // Helper: set portfolio_state wins/losses for env 0.
    let set_ps = |ps_buf: &mut MappedF32Buffer, wins: f32, losses: f32| {
        let sl = ps_buf.host_slice_mut();
        sl.fill(0.0_f32);
        sl[PS_KELLY_WIN_COUNT]  = wins;
        sl[PS_KELLY_LOSS_COUNT] = losses;
        sl[PS_KELLY_SUM_WINS]   = wins * 0.05;   // arbitrary; only count matters for this test
        sl[PS_KELLY_SUM_LOSSES] = losses * 0.05;
    };

    // Zero-init ISV (MappedF32Buffer::new already zeroes, but be explicit).
    isv_buf.host_slice_mut().fill(0.0_f32);

    // Step 1: total_trades=100, ISV[283] starts at 0.
    set_ps(&mut ps_buf, 60.0, 40.0);  // total=100
    launch_pearl_6(&stream, &kernel, &ps_buf, &isv_buf, n_envs);
    let v1 = isv_buf.read_all()[KELLY_SAMPLE_COUNT_IDX];
    println!("step1 ISV[283]={v1:.1} (expected 100)");
    assert!((v1 - 100.0).abs() < 0.5, "step1: expected 100.0, got {v1:.2}");

    // Step 2: fold boundary simulation — portfolio_state reset to 0.
    set_ps(&mut ps_buf, 0.0, 0.0);   // total=0
    launch_pearl_6(&stream, &kernel, &ps_buf, &isv_buf, n_envs);
    let v2 = isv_buf.read_all()[KELLY_SAMPLE_COUNT_IDX];
    println!("step2 ISV[283]={v2:.1} (expected 100 — preserved via max)");
    assert!((v2 - 100.0).abs() < 0.5,
        "step2: cross-fold persistence failed — expected 100.0, got {v2:.2}");

    // Step 3: portfolio_state recovering, total=10 (below old max).
    set_ps(&mut ps_buf, 6.0, 4.0);   // total=10
    launch_pearl_6(&stream, &kernel, &ps_buf, &isv_buf, n_envs);
    let v3 = isv_buf.read_all()[KELLY_SAMPLE_COUNT_IDX];
    println!("step3 ISV[283]={v3:.1} (expected 100 — still preserved)");
    assert!((v3 - 100.0).abs() < 0.5,
        "step3: max should still hold 100.0, got {v3:.2}");

    // Step 4: total_trades=150 — surpasses old max, should grow.
    set_ps(&mut ps_buf, 90.0, 60.0); // total=150
    launch_pearl_6(&stream, &kernel, &ps_buf, &isv_buf, n_envs);
    let v4 = isv_buf.read_all()[KELLY_SAMPLE_COUNT_IDX];
    println!("step4 ISV[283]={v4:.1} (expected 150 — grew to new max)");
    assert!((v4 - 150.0).abs() < 0.5,
        "step4: expected 150.0 (new max), got {v4:.2}");
}

// ── Pearl 8 trail helpers ─────────────────────────────────────────────────────

fn load_pearl_8_trail_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
    let module = stream
        .context()
        .load_cubin(SP5_PEARL_8_TRAIL_CUBIN.to_vec())
        .expect("load pearl_8_trail_kernel cubin");
    module
        .load_function("pearl_8_trail_update")
        .expect("load pearl_8_trail_update function")
}

/// Launch helper: writes 4 trail-distance floats to scratch_buf[base..base+4).
fn launch_pearl_8_trail(
    stream:      &Arc<CudaStream>,
    kernel:      &CudaFunction,
    features:    &MappedF32Buffer,  // features[bar_idx * market_dim + 9] = atr_norm
    bar_idx:     i32,
    market_dim:  i32,
    scratch_out: &MappedF32Buffer,
    base:        i32,
) {
    let feat_ptr  = features.dev_ptr;
    let scr_ptr   = scratch_out.dev_ptr;
    unsafe {
        stream
            .launch_builder(kernel)
            .arg(&feat_ptr)
            .arg(&bar_idx)
            .arg(&market_dim)
            .arg(&scr_ptr)
            .arg(&base)
            .launch(LaunchConfig {
                grid_dim:         (1, 1, 1),
                block_dim:        (4, 1, 1),
                shared_mem_bytes: 0,
            })
            .expect("launch pearl_8_trail_update");
    }
    stream.synchronize().expect("sync after pearl_8_trail_update");
}

/// SP5 Task A7 Test 14: per-direction ATR factor (2× vs floor=1.0).
///
/// Analytical derivation:
///   atr_norm = 0.5  →  log_atr = 0.5 × 16 - 7 = 1.0  →  atr_abs = e^1 ≈ 2.7183
///   Short[270] = Long[272] = atr_abs × 2.0 ≈ 5.4366
///   Hold[271]  = Flat[273] = EPS_CLAMP_FLOOR = 1.0
///
/// No CPU reference implementation (feedback_no_cpu_test_fallbacks.md).
/// Analytically-known inputs only — atr_norm=0.5 is chosen so that
/// log_atr=1.0 makes exp() trivial.
#[test]
#[ignore = "requires GPU"]
fn pearl_8_trail_per_direction_factors() {
    let ctx    = CudaContext::new(0).expect("CUDA context");
    let stream = ctx.new_stream().expect("stream");
    let kernel = load_pearl_8_trail_kernel(&stream);

    // atr_norm = 0.5 → atr_abs = exp(0.5 × 16 − 7) = exp(1.0) ≈ 2.7183
    let atr_norm: f32 = 0.5;
    let atr_abs_expected: f32 = (1.0_f32).exp(); // e^1

    // feature buffer: 1 bar × market_dim=10 (col 9 = atr_norm)
    let market_dim: i32  = 10;
    let bar_idx: i32     = 0;
    let mut feat_buf = unsafe { MappedF32Buffer::new((bar_idx as usize + 1) * market_dim as usize) }
        .expect("feat_buf alloc");
    feat_buf.host_slice_mut().fill(0.0_f32);
    feat_buf.host_slice_mut()[(bar_idx as usize) * (market_dim as usize) + 9] = atr_norm;

    let scratch_size: usize = 4;
    let base: i32           = 0;
    let mut scratch = unsafe { MappedF32Buffer::new(scratch_size) }.expect("scratch alloc");

    launch_pearl_8_trail(&stream, &kernel, &feat_buf, bar_idx, market_dim, &mut scratch, base);

    let out = scratch.read_all();

    let trail_long_short = atr_abs_expected * 2.0_f32;
    let trail_hold_flat  = 1.0_f32;

    println!("Short[0]={:.4} (expected {trail_long_short:.4})", out[0]);
    println!("Hold [1]={:.4} (expected {trail_hold_flat:.4})", out[1]);
    println!("Long [2]={:.4} (expected {trail_long_short:.4})", out[2]);
    println!("Flat [3]={:.4} (expected {trail_hold_flat:.4})", out[3]);

    let tol = 1e-4_f32;
    assert!((out[0] - trail_long_short).abs() < tol,
        "Short: expected {trail_long_short:.4}, got {:.4}", out[0]);
    assert!((out[1] - trail_hold_flat).abs() < tol,
        "Hold: expected {trail_hold_flat:.4}, got {:.4}", out[1]);
    assert!((out[2] - trail_long_short).abs() < tol,
        "Long: expected {trail_long_short:.4}, got {:.4}", out[2]);
    assert!((out[3] - trail_hold_flat).abs() < tol,
        "Flat: expected {trail_hold_flat:.4}, got {:.4}", out[3]);
}

/// SP5 Task A7 Test 15: ATR floor at quiet market (atr_norm very negative).
///
/// Analytical derivation:
///   atr_norm = -10.0  →  log_atr = -10.0 × 16 - 7 = -167.0  →  exp(-167) ≈ 0 << 0.01
///   atr_abs = max(exp(-167), 0.01) = 0.01  (Invariant 1 anchor: atr_abs floor)
///   Short[270] = Long[272] = 0.01 × 2.0 = 0.02
///   Hold[271]  = Flat[273] = 1.0
///
/// This tests the `fmaxf(expf(log_atr), 0.01f)` floor in the kernel — essential
/// for log-space underflow protection when the ATR feature is at very low values.
#[test]
#[ignore = "requires GPU"]
fn pearl_8_trail_atr_floor_at_quiet_market() {
    let ctx    = CudaContext::new(0).expect("CUDA context");
    let stream = ctx.new_stream().expect("stream");
    let kernel = load_pearl_8_trail_kernel(&stream);

    // atr_norm = -10.0 → exp(-167) ≈ 0 → clamped to 0.01 floor
    let atr_norm: f32 = -10.0;

    let market_dim: i32 = 10;
    let bar_idx: i32    = 0;
    let mut feat_buf = unsafe { MappedF32Buffer::new((bar_idx as usize + 1) * market_dim as usize) }
        .expect("feat_buf alloc");
    feat_buf.host_slice_mut().fill(0.0_f32);
    feat_buf.host_slice_mut()[(bar_idx as usize) * (market_dim as usize) + 9] = atr_norm;

    let scratch_size: usize = 4;
    let base: i32           = 0;
    let mut scratch = unsafe { MappedF32Buffer::new(scratch_size) }.expect("scratch alloc");

    launch_pearl_8_trail(&stream, &kernel, &feat_buf, bar_idx, market_dim, &mut scratch, base);

    let out = scratch.read_all();

    // atr_abs = 0.01 (floor), trail = 0.01 × 2.0 = 0.02 for Short and Long.
    let trail_long_short = 0.01_f32 * 2.0_f32; // = 0.02
    let trail_hold_flat  = 1.0_f32;

    println!("Short[0]={:.4} (expected {trail_long_short:.4} — ATR floor 0.01)", out[0]);
    println!("Hold [1]={:.4} (expected {trail_hold_flat:.4})", out[1]);
    println!("Long [2]={:.4} (expected {trail_long_short:.4} — ATR floor 0.01)", out[2]);
    println!("Flat [3]={:.4} (expected {trail_hold_flat:.4})", out[3]);

    let tol = 1e-5_f32;
    assert!((out[0] - trail_long_short).abs() < tol,
        "Short: ATR floor test failed — expected {trail_long_short:.5}, got {:.5}", out[0]);
    assert!((out[1] - trail_hold_flat).abs() < tol,
        "Hold: expected {trail_hold_flat:.5}, got {:.5}", out[1]);
    assert!((out[2] - trail_long_short).abs() < tol,
        "Long: ATR floor test failed — expected {trail_long_short:.5}, got {:.5}", out[2]);
    assert!((out[3] - trail_hold_flat).abs() < tol,
        "Flat: expected {trail_hold_flat:.5}, got {:.5}", out[3]);
}

// ── Task A8 (Pearl 1-ext num_atoms) ──────────────────────────────────────────

fn load_pearl_1_ext_num_atoms_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
    let module = stream
        .context()
        .load_cubin(SP5_PEARL_1_EXT_NUM_ATOMS_CUBIN.to_vec())
        .expect("load pearl_1_ext_num_atoms_kernel cubin");
    module
        .load_function("pearl_1_ext_num_atoms_update")
        .expect("load pearl_1_ext_num_atoms_update function")
}

/// SP5 Task A8 Test 16: pearl_1_ext_num_atoms threshold cascade.
///
/// Synthetic ISV[ATOM_V_HALF_BASE + b] for 4 branches:
///   b=0: v_half=0.05  → strictly less than 0.1  → 64 atoms (narrow)
///   b=1: v_half=0.5   → not < 0.1, but < 1.0    → 32 atoms (moderate)
///   b=2: v_half=2.0   → not < 1.0               → 16 atoms (wide)
///   b=3: v_half=0.099 → just below 0.1 boundary → 64 atoms (narrow)
///
/// Verifies the threshold cascade boundaries.
/// No CPU reference implementation (feedback_no_cpu_test_fallbacks.md).
/// Analytically-known: all expected values are structural constants
/// ATOMS_HIGH=64, ATOMS_MID=32, ATOMS_LOW=16 (Invariant 1 anchors).
#[test]
#[ignore = "requires GPU"]
fn pearl_1_ext_num_atoms_threshold_cascade() {
    let ctx    = CudaContext::new(0).expect("CUDA context");
    let stream = ctx.new_stream().expect("stream");
    let kernel = load_pearl_1_ext_num_atoms_kernel(&stream);

    // ISV buffer: place v_half values at ATOM_V_HALF_BASE=178..182.
    // Allocate enough slots to cover ISV[182) (4 slots from base 178 = indices 178..182).
    let isv_size: usize = 182;
    let mut isv_buf = unsafe { MappedF32Buffer::new(isv_size) }.expect("isv_buf alloc");
    isv_buf.host_slice_mut().fill(0.0_f32);
    // ATOM_V_HALF_BASE = 178
    isv_buf.host_slice_mut()[178] = 0.05_f32;   // b=0: < 0.1  → 64 atoms
    isv_buf.host_slice_mut()[179] = 0.5_f32;    // b=1: < 1.0  → 32 atoms
    isv_buf.host_slice_mut()[180] = 2.0_f32;    // b=2: >= 1.0 → 16 atoms
    isv_buf.host_slice_mut()[181] = 0.099_f32;  // b=3: < 0.1  → 64 atoms

    let scratch_size: usize = 4;
    let base: i32           = 0;
    let scratch = unsafe { MappedF32Buffer::new(scratch_size) }.expect("scratch alloc");

    let isv_dev   = isv_buf.dev_ptr;
    let scr_dev   = scratch.dev_ptr;
    let v_half_base_i32: i32 = 178;

    unsafe {
        stream
            .launch_builder(&kernel)
            .arg(&isv_dev)
            .arg(&v_half_base_i32)
            .arg(&scr_dev)
            .arg(&base)
            .launch(LaunchConfig {
                grid_dim:         (1, 1, 1),
                block_dim:        (4, 1, 1),
                shared_mem_bytes: 0,
            })
            .expect("launch pearl_1_ext_num_atoms_update");
    }
    stream.synchronize().expect("sync after pearl_1_ext_num_atoms_update");

    let out = scratch.read_all();

    println!("b=0 (v_half=0.05):  got {:.0} (expected 64 — narrow)", out[0]);
    println!("b=1 (v_half=0.5):   got {:.0} (expected 32 — moderate)", out[1]);
    println!("b=2 (v_half=2.0):   got {:.0} (expected 16 — wide)", out[2]);
    println!("b=3 (v_half=0.099): got {:.0} (expected 64 — just-narrow)", out[3]);

    let tol = 0.5_f32; // integer comparison via f32 — tolerate no rounding
    assert!((out[0] - 64.0_f32).abs() < tol,
        "b=0: expected 64 atoms (v_half=0.05 < 0.1), got {:.0}", out[0]);
    assert!((out[1] - 32.0_f32).abs() < tol,
        "b=1: expected 32 atoms (0.1 <= v_half=0.5 < 1.0), got {:.0}", out[1]);
    assert!((out[2] - 16.0_f32).abs() < tol,
        "b=2: expected 16 atoms (v_half=2.0 >= 1.0), got {:.0}", out[2]);
    assert!((out[3] - 64.0_f32).abs() < tol,
        "b=3: expected 64 atoms (v_half=0.099 < 0.1, just below narrow boundary), got {:.0}", out[3]);
}

/// SP5 Task A8 Test 17: pearl_1_ext_num_atoms exact-threshold boundary semantics.
///
/// At v_half = 0.1 exactly: `0.1 < 0.1` is false → falls to next branch → 32 atoms.
/// At v_half = 1.0 exactly: `1.0 < 1.0` is false → falls to else branch → 16 atoms.
///
/// Verifies the strictly-less-than threshold semantics (NOT <=).
/// No CPU reference implementation (feedback_no_cpu_test_fallbacks.md).
/// The two branches exercise: THRESHOLD_NARROW boundary (b=0) and
/// THRESHOLD_MODERATE boundary (b=1); the other two branches (b=2, b=3) are
/// set to unambiguous interior values to confirm they are not affected.
#[test]
#[ignore = "requires GPU"]
fn pearl_1_ext_num_atoms_exact_threshold_falls_to_next_branch() {
    let ctx    = CudaContext::new(0).expect("CUDA context");
    let stream = ctx.new_stream().expect("stream");
    let kernel = load_pearl_1_ext_num_atoms_kernel(&stream);

    let isv_size: usize = 182;
    let mut isv_buf = unsafe { MappedF32Buffer::new(isv_size) }.expect("isv_buf alloc");
    isv_buf.host_slice_mut().fill(0.0_f32);
    // ATOM_V_HALF_BASE = 178
    isv_buf.host_slice_mut()[178] = 0.1_f32;    // b=0: exactly 0.1 — NOT < 0.1 → 32 atoms
    isv_buf.host_slice_mut()[179] = 1.0_f32;    // b=1: exactly 1.0 — NOT < 1.0 → 16 atoms
    isv_buf.host_slice_mut()[180] = 0.05_f32;   // b=2: clearly narrow → 64 atoms (sanity)
    isv_buf.host_slice_mut()[181] = 5.0_f32;    // b=3: clearly wide   → 16 atoms (sanity)

    let scratch_size: usize = 4;
    let base: i32           = 0;
    let scratch = unsafe { MappedF32Buffer::new(scratch_size) }.expect("scratch alloc");

    let isv_dev   = isv_buf.dev_ptr;
    let scr_dev   = scratch.dev_ptr;
    let v_half_base_i32: i32 = 178;

    unsafe {
        stream
            .launch_builder(&kernel)
            .arg(&isv_dev)
            .arg(&v_half_base_i32)
            .arg(&scr_dev)
            .arg(&base)
            .launch(LaunchConfig {
                grid_dim:         (1, 1, 1),
                block_dim:        (4, 1, 1),
                shared_mem_bytes: 0,
            })
            .expect("launch pearl_1_ext_num_atoms_update (exact threshold)");
    }
    stream.synchronize().expect("sync after pearl_1_ext_num_atoms_update (exact threshold)");

    let out = scratch.read_all();

    println!("b=0 (v_half=0.1 exact):  got {:.0} (expected 32 — NOT < 0.1)", out[0]);
    println!("b=1 (v_half=1.0 exact):  got {:.0} (expected 16 — NOT < 1.0)", out[1]);
    println!("b=2 (v_half=0.05):       got {:.0} (expected 64 — sanity)", out[2]);
    println!("b=3 (v_half=5.0):        got {:.0} (expected 16 — sanity)", out[3]);

    let tol = 0.5_f32;
    assert!((out[0] - 32.0_f32).abs() < tol,
        "b=0 (v_half=0.1): expected 32 atoms (threshold is strictly <, not <=), got {:.0}", out[0]);
    assert!((out[1] - 16.0_f32).abs() < tol,
        "b=1 (v_half=1.0): expected 16 atoms (threshold is strictly <, not <=), got {:.0}", out[1]);
    assert!((out[2] - 64.0_f32).abs() < tol,
        "b=2 (v_half=0.05): expected 64 atoms (sanity check), got {:.0}", out[2]);
    assert!((out[3] - 16.0_f32).abs() < tol,
        "b=3 (v_half=5.0): expected 16 atoms (sanity check), got {:.0}", out[3]);
}

// ── Layer D Task D1 (PnL aggregation) ─────────────────────────────────────────

fn load_pnl_aggregation_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
    let module = stream
        .context()
        .load_cubin(SP5_PNL_AGGREGATION_CUBIN.to_vec())
        .expect("load pnl_aggregation_kernel cubin");
    module
        .load_function("pnl_aggregation_update")
        .expect("load pnl_aggregation_update function")
}

/// SP5 Layer D Task D1 (rewrite): `pnl_aggregation_update` correctness on a
/// synthetic per-bar return tape with one episode boundary, with
/// analytically-known total / mean / variance / max drawdown.
///
/// **Rewrite.** The original test (commit `5ee795f14`) targeted a per-trade
/// signature that doesn't exist in production; this test exercises the
/// rewritten kernel against the actual `compute_epoch_financials` (financials.rs)
/// formulas. Per `feedback_no_cpu_test_fallbacks.md` the expected outputs are
/// pre-computed analytical literals — no CPU reference implementation in the
/// test process.
///
/// Synthetic inputs:
///   step_returns = [+0.01, −0.005, +0.003, −0.002, +0.001]
///   done_flags   = [0,     0,      1,      0,      0]
///   num_bars     = 5
///   n_trades     = 5         (per-trade event count for mean / variance)
///   sum_returns  = 0.007     (per-trade Σ returns; arbitrary distinct from
///                             the per-bar sum so the kernel can't accidentally
///                             use the per-bar slice for the mean)
///   sum_sq_returns = 1.39e-4 (per-trade Σ returns²)
///   initial_capital = 100_000.0
///
/// Analytical ground truth (computed offline from the financials.rs formulas):
///
///   pnl_total = exp(Σ ln(max(1+r, 1e-10))) − 1
///             = exp( ln(1.01) + ln(0.995) + ln(1.003) + ln(0.998) + ln(1.001) ) − 1
///             = exp(0.0069307957) − 1 ≈ 0.0069548692
///
///   pnl_mean = sum_returns / n_trades = 0.007 / 5 = 0.0014
///
///   pnl_var  = max(0, sum_sq_returns/n_trades − mean²)
///            = max(0, 1.39e-4 / 5 − 1.96e-6)
///            = max(0, 2.78e-5 − 1.96e-6)
///            = 2.584e-5
///
///   pnl_max_dd: per-bar equity walk with episode-boundary reset.
///     Start: equity = peak = 100_000.
///     bar 0 (r=+0.0100): equity=101_000, peak=101_000, dd=0,        max_dd=0
///     bar 1 (r=−0.0050): equity=100_495, peak=101_000, dd=0.005,    max_dd=0.005
///     bar 2 (r=+0.0030): equity=100_796.485, peak=101_000, then done=1 →
///                        equity=peak=100_000, dd=0,                 max_dd=0.005
///     bar 3 (r=−0.0020): equity=99_800,  peak=100_000, dd=0.002,    max_dd=0.005
///     bar 4 (r=+0.0010): equity=99_899.8, peak=100_000, dd=0.001002, max_dd=0.005
///   max_dd = 0.005, capped at 1.0 ⇒ 0.005.
///
/// **Episode boundary reset semantics**: financials.rs:175-183 resets equity
/// and peak AFTER processing the done-bar's return — the loss on the done
/// bar belongs to the just-ending episode; the next bar starts fresh. This
/// test's bar 2 hits the boundary; without the reset the cumulative product
/// would continue to compound and bar 3/4's drawdowns would reference the
/// pre-reset peak of 101_000 instead of 100_000, changing max_dd.
#[test]
#[ignore = "requires GPU"]
fn pnl_aggregation_kernel_correctness() {
    let ctx    = CudaContext::new(0).expect("CUDA context");
    let stream = ctx.new_stream().expect("stream");
    let kernel = load_pnl_aggregation_kernel(&stream);

    // ── Synthetic input arrays ───────────────────────────────────────
    let returns_host: [f32; 5] = [0.01, -0.005, 0.003, -0.002, 0.001];
    let dones_host:   [f32; 5] = [0.0,  0.0,    1.0,   0.0,    0.0  ];
    let num_bars:        i32   = 5;
    let n_trades:        i32   = 5;
    let sum_returns:     f32   = 0.007;
    let sum_sq_returns:  f32   = 1.39e-4;
    let initial_capital: f32   = 100_000.0;

    let mut returns_buf =
        unsafe { MappedF32Buffer::new(num_bars as usize) }.expect("step_returns alloc");
    let mut dones_buf   =
        unsafe { MappedF32Buffer::new(num_bars as usize) }.expect("done_flags alloc");
    returns_buf.host_slice_mut().copy_from_slice(&returns_host);
    dones_buf.host_slice_mut().copy_from_slice(&dones_host);

    // ── Output scratch (4 slots) ─────────────────────────────────────
    let scratch = unsafe { MappedF32Buffer::new(4) }.expect("scratch alloc");

    // ── Launch ───────────────────────────────────────────────────────
    let returns_dev = returns_buf.dev_ptr;
    let dones_dev   = dones_buf.dev_ptr;
    let scratch_dev = scratch.dev_ptr;
    let total_idx:  i32 = 0;
    let mean_idx:   i32 = 1;
    let var_idx:    i32 = 2;
    let max_dd_idx: i32 = 3;

    let block_dim:        u32 = 256;
    let shared_mem_bytes: u32 = block_dim * std::mem::size_of::<f32>() as u32;

    unsafe {
        stream
            .launch_builder(&kernel)
            .arg(&returns_dev)
            .arg(&dones_dev)
            .arg(&num_bars)
            .arg(&n_trades)
            .arg(&sum_returns)
            .arg(&sum_sq_returns)
            .arg(&initial_capital)
            .arg(&scratch_dev)
            .arg(&total_idx)
            .arg(&mean_idx)
            .arg(&var_idx)
            .arg(&max_dd_idx)
            .launch(LaunchConfig {
                grid_dim:         (1, 1, 1),
                block_dim:        (block_dim, 1, 1),
                shared_mem_bytes,
            })
            .expect("launch pnl_aggregation_update");
    }
    stream.synchronize().expect("sync after pnl_aggregation_update");

    let out = scratch.read_all();

    // ── Analytical ground truth (literals; per feedback_no_cpu_test_fallbacks) ─
    let exp_total:  f32 = 0.0069548692;
    let exp_mean:   f32 = 0.0014;
    let exp_var:    f32 = 2.584e-5;
    let exp_max_dd: f32 = 0.005;

    println!("pnl_total  = {:.10} (expected {exp_total:.10})", out[0]);
    println!("pnl_mean   = {:.10} (expected {exp_mean:.10})",  out[1]);
    println!("pnl_var    = {:.10} (expected {exp_var:.10})",   out[2]);
    println!("pnl_max_dd = {:.10} (expected {exp_max_dd:.10})", out[3]);

    // FP tolerances:
    //   - pnl_total: log/exp roundtrip in f32 over 5 small returns; ~5 ulps.
    //   - pnl_mean: division by 5 of an exact f32; effectively exact.
    //   - pnl_var:  E[X²] − E[X]² catastrophic-cancellation prone; result is
    //              ~2.5e-5 so 1e-7 absolute = ~0.4% relative — comfortable.
    //   - pnl_max_dd: integer-arithmetic equity walk; effectively exact.
    let tol_total:  f32 = 1e-7;
    let tol_mean:   f32 = 1e-9;
    let tol_var:    f32 = 1e-7;
    let tol_max_dd: f32 = 1e-7;

    assert!((out[0] - exp_total).abs() < tol_total,
        "pnl_total: expected {exp_total:.10}, got {:.10} (delta {:.3e})",
        out[0], (out[0] - exp_total).abs());
    assert!((out[1] - exp_mean).abs()  < tol_mean,
        "pnl_mean: expected {exp_mean:.10}, got {:.10}", out[1]);
    assert!((out[2] - exp_var).abs()   < tol_var,
        "pnl_var: expected {exp_var:.10} (per-trade E[X²]−E[X]² form), got {:.10}",
        out[2]);
    assert!((out[3] - exp_max_dd).abs() < tol_max_dd,
        "pnl_max_dd: expected {exp_max_dd:.10} (episode reset at bar 2), got {:.10}",
        out[3]);
}

// ── Layer D Task D2 (Health composition) ──────────────────────────────────────

fn load_health_composition_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
    let module = stream
        .context()
        .load_cubin(SP5_HEALTH_COMPOSITION_CUBIN.to_vec())
        .expect("load health_composition_kernel cubin");
    module
        .load_function("health_composition_update")
        .expect("load health_composition_update function")
}

/// SP5 Layer D Task D2: `health_composition_update` formula-fidelity test.
///
/// Asserts the kernel reproduces `learning_health.rs::NormalizedComponents::
/// from_raw` + `compose()` bit-for-bit (within float precision). Per
/// `feedback_no_cpu_test_fallbacks.md` we do **not** programmatically
/// re-implement the formula in the test; instead the synthetic inputs are
/// chosen so each normalised component evaluates to a hand-computable
/// closed-form rational, and the expected outputs are checked against
/// literal pre-computed constants.
///
/// Synthetic raw inputs:
///
///   q_gap            = 0.05   → smoothstep(0.01, 0.5)
///                                 t = 0.04 / 0.49 = 0.08163265306...
///                                 ss = t²·(3 − 2t) = 0.018901683...
///   q_var            = 0.10   → smoothstep(0.001, 0.1)
///                                 t = (0.099)/(0.099) = 1   → ss = 1.0
///                                 (intentional saturation at upper edge)
///   atom_util        = 0.45   → smoothstep(0.2, 0.7) at the **midpoint**
///                                 t = 0.25/0.5 = 0.5 → ss = 0.5²·2 = 0.5
///   grad_norm        = 30.0   → 1 − smoothstep(10, 100)
///                                 t = 20/90 = 0.2̄
///                                 ss = t²·(3 − 2t) = (4/81)·(73/27)·…
///                                                   ≈ 0.126200274...
///                                 grad_stable = 1 − ss ≈ 0.873799725...
///   ens_disagreement = 0.20   → ens_agree = 1 − 0.20 = 0.80 (clamp pass-through)
///   grad_consistency = 0.15   → smoothstep(−0.2, 0.5) at the midpoint
///                                 t = 0.35/0.7 = 0.5 → ss = 0.5
///   spectral_gap     = 5.0    → 1 − smoothstep(2, 10)
///                                 t = 3/8 = 0.375
///                                 ss = t²·(3 − 2t) = 0.140625·2.25 = 0.316406250
///                                 spectral_gap_norm = 1 − 0.316406250 = 0.683593750
///
/// Two of the seven normalisations land at exact rational midpoints (output
/// 0.5) and one saturates exactly (output 1.0); the other four are
/// pre-computed analytically and pinned as literal constants below. The
/// composed score is the weighted sum.
///
///   health_score = 0.25 · 0.018901683
///                + 0.15 · 1.000000000
///                + 0.15 · 0.500000000
///                + 0.15 · 0.873799725
///                + 0.10 · 0.800000000
///                + 0.10 · 0.500000000
///                + 0.10 · 0.683593750
///                ≈ 0.559153878
///
/// **Formula fidelity is the load-bearing invariant** for D2: the kernel
/// performs a verbatim structural migration of the host-side composition,
/// no algorithmic change. The test pins both the cubic interpolation
/// (q_gap_norm interior point) and the composition arithmetic.
#[test]
#[ignore = "requires GPU"]
fn health_composition_kernel_correctness() {
    let ctx    = CudaContext::new(0).expect("CUDA context");
    let stream = ctx.new_stream().expect("stream");
    let kernel = load_health_composition_kernel(&stream);

    // ── Synthetic raw inputs ────────────────────────────────────────
    let q_gap:            f32 = 0.05;
    let q_var:            f32 = 0.10;
    let atom_util:        f32 = 0.45;
    let grad_norm:        f32 = 30.0;
    let ens_disagreement: f32 = 0.20;
    let grad_consistency: f32 = 0.15;
    let spectral_gap:     f32 = 5.0;

    // ── Output scratch (4 slots) ─────────────────────────────────────
    let scratch = unsafe { MappedF32Buffer::new(4) }.expect("scratch alloc");

    let scratch_dev             = scratch.dev_ptr;
    let health_score_idx:   i32 = 0;
    let q_gap_norm_idx:     i32 = 1;
    let q_var_norm_idx:     i32 = 2;
    let grad_norm_norm_idx: i32 = 3;

    unsafe {
        stream
            .launch_builder(&kernel)
            .arg(&q_gap)
            .arg(&q_var)
            .arg(&atom_util)
            .arg(&grad_norm)
            .arg(&ens_disagreement)
            .arg(&grad_consistency)
            .arg(&spectral_gap)
            .arg(&scratch_dev)
            .arg(&health_score_idx)
            .arg(&q_gap_norm_idx)
            .arg(&q_var_norm_idx)
            .arg(&grad_norm_norm_idx)
            .launch(LaunchConfig {
                grid_dim:         (1, 1, 1),
                block_dim:        (1, 1, 1),
                shared_mem_bytes: 0,
            })
            .expect("launch health_composition_update");
    }
    stream.synchronize().expect("sync after health_composition_update");

    let out = scratch.read_all();

    // ── Analytical ground truth (literals; per feedback_no_cpu_test_fallbacks) ─
    // q_gap_norm closed-form (exact rational):
    //   t = 4/49, t²(3 − 2t) = (16·139) / (49²·49) = 2224 / 117649 ≈ 0.0189036880
    let exp_q_gap_norm:    f32 = 0.018903688;
    // q_var_norm: smoothstep saturates at edge1 (0.10 == upper edge).
    let exp_q_var_norm:    f32 = 1.0;
    // grad_norm_norm = 1 − smoothstep(10, 100, 30) = 1 − (20/90)² · (3 − 40/90).
    //   t = 2/9, t²(3 − 2t) = (4/81) · (23/9) = 92 / 729 ≈ 0.1262002743
    //   grad_stable = 1 − 92/729 = 637 / 729 ≈ 0.8737997257
    let exp_grad_stable:   f32 = 0.873799725;
    // health_score: weighted sum of the 7 normalised values listed in the
    // docstring above. Using exact closed forms where possible:
    //   0.25 · (2224/117649)  ≈ 0.0047259220
    //   0.15 · 1.0             = 0.15
    //   0.15 · 0.5             = 0.075         (atom_util_norm midpoint)
    //   0.15 · (637/729)       ≈ 0.131069959
    //   0.10 · 0.80            = 0.08          (ens_agree)
    //   0.10 · 0.5             = 0.05          (grad_consistency_norm midpoint)
    //   0.10 · 0.68359375      = 0.068359375
    //   ────────────────────────────────────
    //   sum                    ≈ 0.5591552560
    let exp_health_score:  f32 = 0.559155256;

    println!("health_score   = {:.7} (expected {:.7})", out[0], exp_health_score);
    println!("q_gap_norm     = {:.7} (expected {:.7})", out[1], exp_q_gap_norm);
    println!("q_var_norm     = {:.7} (expected {:.7})", out[2], exp_q_var_norm);
    println!("grad_norm_norm = {:.7} (expected {:.7})", out[3], exp_grad_stable);

    // FP tolerances: composed score has 7 multiply-adds of small (≤1.0)
    // values; q_var_norm saturates exactly. q_gap_norm comes from cubic
    // interpolation of an irrational `t` in f32 — the literal we pinned is
    // truncated at the 9th digit, so the divergence between f32 evaluation
    // of the closed form and the literal is bounded by ~5e-9 * 3 (cubic).
    // 1e-6 absolute is safely above that.
    let tol_strict: f32 = 1e-6;
    let tol_score:  f32 = 1e-6;

    assert!((out[0] - exp_health_score).abs() < tol_score,
        "health_score: expected {exp_health_score:.7}, got {:.7}", out[0]);
    assert!((out[1] - exp_q_gap_norm).abs() < tol_strict,
        "q_gap_norm: expected {exp_q_gap_norm:.7}, got {:.7}", out[1]);
    assert!((out[2] - exp_q_var_norm).abs() < tol_strict,
        "q_var_norm: expected {exp_q_var_norm:.7}, got {:.7}", out[2]);
    assert!((out[3] - exp_grad_stable).abs() < tol_strict,
        "grad_norm_norm (= grad_stable): expected {exp_grad_stable:.7}, got {:.7}", out[3]);
}

// ── Layer D Task D3 (Training metrics EMA) ────────────────────────────────────

fn load_training_metrics_ema_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
    let module = stream
        .context()
        .load_cubin(SP5_TRAINING_METRICS_EMA_CUBIN.to_vec())
        .expect("load training_metrics_ema_kernel cubin");
    module
        .load_function("training_metrics_ema_update")
        .expect("load training_metrics_ema_update function")
}

/// SP5 Layer D Task D3: `training_metrics_ema_update` formula-fidelity test.
///
/// Asserts the kernel reproduces the host-side EMA arithmetic at
/// `training_loop.rs:5039-5113` bit-for-bit (within float precision). Per
/// `feedback_no_cpu_test_fallbacks.md` the test uses synthetic inputs whose
/// closed-form expected outputs are pre-computed and pinned as literals —
/// no programmatic re-implementation of the formula in the test body.
///
/// Synthetic inputs (deliberately exercise the steady-state branch of all
/// three threads — first-observation paths are tested in the second case
/// below by setting `sharpe_initialized=0` and `prev_max_dd_ema=0`):
///
///   epoch_sharpe         = 1.0
///   epoch_max_dd         = 0.05      (5% drawdown)
///   prev_sharpe_ema      = 0.5       (sharpe diverged from prev by |err|=0.5)
///   prev_max_dd_ema      = 0.10
///   prev_low_dd_ratio    = 0.4
///   sharpe_initialized   = true
///
/// Closed-form expected outputs:
///
///   tid=0 (training_sharpe_ema, adaptive α):
///     err   = |1.0 − 0.5| = 0.5
///     α     = clamp(0.05 + 0.3 · 0.5, 0, 0.5) = clamp(0.20, 0, 0.5) = 0.20
///     ema   = (1 − 0.20) · 0.5 + 0.20 · 1.0 = 0.40 + 0.20 = 0.60
///
///   tid=1 (max_dd_ema, fixed α=0.1, prev > 1e-12 so steady-state):
///     ema   = 0.9 · 0.10 + 0.1 · 0.05 = 0.090 + 0.005 = 0.095
///
///   tid=2 (low_dd_ratio, fixed α=0.15 over binary `is_low`):
///     thresh = new_max_dd_ema · 0.5 = 0.095 · 0.5 = 0.0475
///     is_low = (epoch_max_dd=0.05 < thresh=0.0475) ? 1.0 : 0.0 = 0.0
///              (epoch_max_dd is *just* above half the new EMA)
///     ema    = 0.85 · 0.4 + 0.15 · 0.0 = 0.340
///
/// **Formula fidelity is the load-bearing invariant** for D3: the kernel
/// performs a verbatim structural migration of the host-side EMA update,
/// no algorithmic change. The test pins the adaptive-α arithmetic, the
/// fixed-α recurrence, and the cross-thread `is_low` derivation
/// (tid=2 must read tid=1's *new* output via __syncthreads()).
#[test]
#[ignore = "requires GPU"]
fn training_metrics_ema_kernel_correctness() {
    let ctx    = CudaContext::new(0).expect("CUDA context");
    let stream = ctx.new_stream().expect("stream");
    let kernel = load_training_metrics_ema_kernel(&stream);

    // ── Case 1: steady-state branch on all three threads ────────────────
    {
        let epoch_sharpe:       f32  = 1.0;
        let epoch_max_dd:       f32  = 0.05;
        let prev_sharpe_ema:    f32  = 0.5;
        let prev_max_dd_ema:    f32  = 0.10;
        let prev_low_dd_ratio:  f32  = 0.4;
        let sharpe_initialized: i32  = 1;

        let scratch = unsafe { MappedF32Buffer::new(3) }.expect("scratch alloc");

        let scratch_dev             = scratch.dev_ptr;
        let sharpe_idx:       i32   = 0;
        let max_dd_idx:       i32   = 1;
        let low_dd_ratio_idx: i32   = 2;

        unsafe {
            stream
                .launch_builder(&kernel)
                .arg(&epoch_sharpe)
                .arg(&epoch_max_dd)
                .arg(&prev_sharpe_ema)
                .arg(&prev_max_dd_ema)
                .arg(&prev_low_dd_ratio)
                .arg(&sharpe_initialized)
                .arg(&scratch_dev)
                .arg(&sharpe_idx)
                .arg(&max_dd_idx)
                .arg(&low_dd_ratio_idx)
                .launch(LaunchConfig {
                    grid_dim:         (1, 1, 1),
                    block_dim:        (3, 1, 1),
                    shared_mem_bytes: 0,
                })
                .expect("launch training_metrics_ema_update steady-state");
        }
        stream.synchronize().expect("sync after training_metrics_ema_update steady-state");

        let out = scratch.read_all();

        // Analytical ground truth (closed-form literals):
        let exp_sharpe_ema:   f32 = 0.60;
        let exp_max_dd_ema:   f32 = 0.095;
        let exp_low_dd_ratio: f32 = 0.340;

        println!("[steady] training_sharpe_ema = {:.6} (expected {:.6})", out[0], exp_sharpe_ema);
        println!("[steady] max_dd_ema          = {:.6} (expected {:.6})", out[1], exp_max_dd_ema);
        println!("[steady] low_dd_ratio        = {:.6} (expected {:.6})", out[2], exp_low_dd_ratio);

        // FP tolerances: each output is 2-3 multiply-adds of small (≤1.0)
        // values; 1e-6 absolute is comfortable.
        let tol: f32 = 1e-6;
        assert!((out[0] - exp_sharpe_ema).abs() < tol,
            "training_sharpe_ema [steady]: expected {exp_sharpe_ema:.6}, got {:.6}", out[0]);
        assert!((out[1] - exp_max_dd_ema).abs() < tol,
            "max_dd_ema [steady]: expected {exp_max_dd_ema:.6}, got {:.6}", out[1]);
        assert!((out[2] - exp_low_dd_ratio).abs() < tol,
            "low_dd_ratio [steady]: expected {exp_low_dd_ratio:.6}, got {:.6}", out[2]);
    }

    // ── Case 2: cold-start branches on tid=0 (sharpe_initialized=0)
    //              and tid=1 (prev_max_dd_ema < 1e-12). tid=2 has no
    //              host-side sentinel; it always uses the EMA recurrence.
    {
        let epoch_sharpe:       f32  = 0.7;
        let epoch_max_dd:       f32  = 0.08;
        let prev_sharpe_ema:    f32  = 99.0;       // ignored when initialized=0
        let prev_max_dd_ema:    f32  = 0.0;        // < 1e-12 sentinel triggers replace
        let prev_low_dd_ratio:  f32  = 0.0;        // EMA from 0 is recurrence-safe
        let sharpe_initialized: i32  = 0;

        let scratch = unsafe { MappedF32Buffer::new(3) }.expect("scratch alloc");

        let scratch_dev             = scratch.dev_ptr;
        let sharpe_idx:       i32   = 0;
        let max_dd_idx:       i32   = 1;
        let low_dd_ratio_idx: i32   = 2;

        unsafe {
            stream
                .launch_builder(&kernel)
                .arg(&epoch_sharpe)
                .arg(&epoch_max_dd)
                .arg(&prev_sharpe_ema)
                .arg(&prev_max_dd_ema)
                .arg(&prev_low_dd_ratio)
                .arg(&sharpe_initialized)
                .arg(&scratch_dev)
                .arg(&sharpe_idx)
                .arg(&max_dd_idx)
                .arg(&low_dd_ratio_idx)
                .launch(LaunchConfig {
                    grid_dim:         (1, 1, 1),
                    block_dim:        (3, 1, 1),
                    shared_mem_bytes: 0,
                })
                .expect("launch training_metrics_ema_update cold-start");
        }
        stream.synchronize().expect("sync after training_metrics_ema_update cold-start");

        let out = scratch.read_all();

        // Closed-form for cold-start:
        //   tid=0: sharpe_initialized==0 ⇒ output replaces with epoch_sharpe = 0.7.
        //   tid=1: prev_max_dd_ema<1e-12 ⇒ output replaces with epoch_max_dd = 0.08.
        //   tid=2: thresh=0.04, is_low=(0.08<0.04)?1:0 = 0; ema = 0.85·0 + 0.15·0 = 0.
        let exp_sharpe_ema:   f32 = 0.7;
        let exp_max_dd_ema:   f32 = 0.08;
        let exp_low_dd_ratio: f32 = 0.0;

        println!("[cold]   training_sharpe_ema = {:.6} (expected {:.6})", out[0], exp_sharpe_ema);
        println!("[cold]   max_dd_ema          = {:.6} (expected {:.6})", out[1], exp_max_dd_ema);
        println!("[cold]   low_dd_ratio        = {:.6} (expected {:.6})", out[2], exp_low_dd_ratio);

        let tol: f32 = 1e-6;
        assert!((out[0] - exp_sharpe_ema).abs() < tol,
            "training_sharpe_ema [cold]: expected {exp_sharpe_ema:.6}, got {:.6}", out[0]);
        assert!((out[1] - exp_max_dd_ema).abs() < tol,
            "max_dd_ema [cold]: expected {exp_max_dd_ema:.6}, got {:.6}", out[1]);
        assert!((out[2] - exp_low_dd_ratio).abs() < tol,
            "low_dd_ratio [cold]: expected {exp_low_dd_ratio:.6}, got {:.6}", out[2]);
    }
}

// ── SP7 activation-flag fix unit test ─────────────────────────────────────────
//
// Verifies the activation-flag transitions in `loss_balance_controller_update`:
//   1. Cold start (zero ISV, zero grad_decomp_result): kernel takes the genuine
//      cold-start branch (h_n < EPS_DIV), writes 0.0 to scratch_active. The
//      apply_pearls_ad_kernel short-circuit on step_obs == 0.0 is observable
//      via the scratch buffer directly (we don't run apply_pearls in this test).
//   2. Active (non-zero IQN/CQL/C51 norms): kernel takes the active path,
//      writes 1.0 to scratch_active and a real budget to scratch_budget.
//   3. Resulting budget reflects controller compute (NOT bootstrap constant).
//
// Per `feedback_no_cpu_test_fallbacks.md`: assertions are on analytically-known
// kernel behaviour (active flag binary, budget structurally `cold_start_basis ·
// correction` from the seed-from-cold-start branch since old_budget=0).

fn load_loss_balance_controller_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
    let module = stream
        .context()
        .load_cubin(SP7_LOSS_BALANCE_CONTROLLER_CUBIN.to_vec())
        .expect("load loss_balance_controller_kernel cubin");
    module
        .load_function("loss_balance_controller_update")
        .expect("load loss_balance_controller_update function")
}

#[test]
#[ignore = "requires GPU"]
fn sp7_loss_balance_controller_activation_flag_transitions() {
    use ml::cuda_pipeline::sp5_isv_slots::{
        BUDGET_CQL_BASE, BUDGET_C51_BASE, FLATNESS_BASE,
        LB_DIFF_VAR_CQL_BASE, LB_SAMPLE_VAR_CQL_BASE,
        LB_DIFF_VAR_C51_BASE, LB_SAMPLE_VAR_C51_BASE,
        LB_CQL_ACTIVE_BASE, LB_C51_ACTIVE_BASE,
    };

    // ISV must cover up to SP5_SLOT_END (321) AND EPOCH_IDX_INDEX (39).
    const ISV_SIZE: usize = 321;
    const EPOCH_IDX_INDEX: usize = 39;

    // grad_decomp_result_pinned: 27 floats (9 components × 3 [mag, dir, trunk]).
    // Component element offsets per kernel header: iqn=0, cql_sx=6, c51=9.
    // We allocate the full 27-float layout to mirror production.
    const GRAD_DECOMP_TOTAL: usize = 27;
    const GRAD_OFFSET_IQN:    usize = 0;
    const GRAD_OFFSET_CQL_SX: usize = 6;
    const GRAD_OFFSET_C51:    usize = 9;

    // Scratch layout: 8 output slot-blocks × 4 branches = 32 floats.
    // We use small-but-distinct base offsets for the test.
    const SCRATCH_BUDGET_CQL: usize = 0;    // [0..4)
    const SCRATCH_BUDGET_C51: usize = 4;    // [4..8)
    const SCRATCH_DIFF_CQL:   usize = 8;    // [8..12)
    const SCRATCH_SAMPLE_CQL: usize = 12;   // [12..16)
    const SCRATCH_DIFF_C51:   usize = 16;   // [16..20)
    const SCRATCH_SAMPLE_C51: usize = 20;   // [20..24)
    const SCRATCH_ACTIVE_CQL: usize = 24;   // [24..28)
    const SCRATCH_ACTIVE_C51: usize = 28;   // [28..32)
    const SCRATCH_SIZE:       usize = 32;

    let stream = make_test_stream();
    let kernel = load_loss_balance_controller_kernel(&stream);

    let isv_buf = unsafe { MappedF32Buffer::new(ISV_SIZE) }.expect("alloc isv_buf");
    let grad_buf = unsafe { MappedF32Buffer::new(GRAD_DECOMP_TOTAL) }.expect("alloc grad_buf");
    let scratch_buf = unsafe { MappedF32Buffer::new(SCRATCH_SIZE) }.expect("alloc scratch_buf");

    let isv_dev = isv_buf.dev_ptr;
    let scratch_dev = scratch_buf.dev_ptr;
    let f32_size = std::mem::size_of::<f32>() as u64;
    let iqn_dev    = grad_buf.dev_ptr + (GRAD_OFFSET_IQN    as u64) * f32_size;
    let cql_sx_dev = grad_buf.dev_ptr + (GRAD_OFFSET_CQL_SX as u64) * f32_size;
    let c51_dev    = grad_buf.dev_ptr + (GRAD_OFFSET_C51    as u64) * f32_size;

    let flatness_isv_base_i32       = FLATNESS_BASE as i32;
    let budget_cql_isv_base_i32     = BUDGET_CQL_BASE as i32;
    let budget_c51_isv_base_i32     = BUDGET_C51_BASE as i32;
    let diff_var_cql_isv_base_i32   = LB_DIFF_VAR_CQL_BASE as i32;
    let sample_var_cql_isv_base_i32 = LB_SAMPLE_VAR_CQL_BASE as i32;
    let diff_var_c51_isv_base_i32   = LB_DIFF_VAR_C51_BASE as i32;
    let sample_var_c51_isv_base_i32 = LB_SAMPLE_VAR_C51_BASE as i32;
    let active_cql_isv_base_i32     = LB_CQL_ACTIVE_BASE as i32;
    let active_c51_isv_base_i32     = LB_C51_ACTIVE_BASE as i32;
    let epoch_idx_isv_index_i32     = EPOCH_IDX_INDEX as i32;
    let sb_budget_cql_i32           = SCRATCH_BUDGET_CQL as i32;
    let sb_budget_c51_i32           = SCRATCH_BUDGET_C51 as i32;
    let sb_diff_var_cql_i32         = SCRATCH_DIFF_CQL as i32;
    let sb_sample_var_cql_i32       = SCRATCH_SAMPLE_CQL as i32;
    let sb_diff_var_c51_i32         = SCRATCH_DIFF_C51 as i32;
    let sb_sample_var_c51_i32       = SCRATCH_SAMPLE_C51 as i32;
    let sb_active_cql_i32           = SCRATCH_ACTIVE_CQL as i32;
    let sb_active_c51_i32           = SCRATCH_ACTIVE_C51 as i32;

    let cfg = LaunchConfig {
        grid_dim:         (1, 1, 1),
        block_dim:        (8, 1, 1),
        shared_mem_bytes: 0,
    };

    // Helper: launch the controller kernel once with current ISV/grad state.
    let launch_once = |isv_data: &[f32], grad_data: &[f32], scratch_seed: f32| {
        isv_buf.write_from_slice(isv_data);
        grad_buf.write_from_slice(grad_data);
        // Seed scratch with a marker so we can detect "kernel didn't write".
        let scratch_init = vec![scratch_seed; SCRATCH_SIZE];
        scratch_buf.write_from_slice(&scratch_init);
        unsafe {
            stream
                .launch_builder(&kernel)
                .arg(&iqn_dev)
                .arg(&cql_sx_dev)
                .arg(&c51_dev)
                .arg(&isv_dev)
                .arg(&flatness_isv_base_i32)
                .arg(&budget_cql_isv_base_i32)
                .arg(&budget_c51_isv_base_i32)
                .arg(&diff_var_cql_isv_base_i32)
                .arg(&sample_var_cql_isv_base_i32)
                .arg(&diff_var_c51_isv_base_i32)
                .arg(&sample_var_c51_isv_base_i32)
                .arg(&active_cql_isv_base_i32)
                .arg(&active_c51_isv_base_i32)
                .arg(&epoch_idx_isv_index_i32)
                .arg(&scratch_dev)
                .arg(&sb_budget_cql_i32)
                .arg(&sb_budget_c51_i32)
                .arg(&sb_diff_var_cql_i32)
                .arg(&sb_sample_var_cql_i32)
                .arg(&sb_diff_var_c51_i32)
                .arg(&sb_sample_var_c51_i32)
                .arg(&sb_active_cql_i32)
                .arg(&sb_active_c51_i32)
                .launch(cfg)
                .expect("launch loss_balance_controller_update");
        }
        stream.synchronize().expect("sync after loss_balance_controller_update");
        scratch_buf.read_all()
    };

    // ── Step 1: cold start (all-zero ISV, all-zero grads). ───────────────────
    // Kernel hits cold-start branch: h_n=0 < EPS_DIV. was_active=0 (ISV is zero).
    // Expect: scratch_active = 0.0 for all 4 branches × 2 heads.
    // (apply_pearls_ad short-circuit on 0.0 means the consumer sees the
    //  bootstrap value, but we test the scratch directly here.)
    let isv_cold  = vec![0.0_f32; ISV_SIZE];
    let grad_cold = vec![0.0_f32; GRAD_DECOMP_TOTAL];

    let scratch_after_cold = launch_once(&isv_cold, &grad_cold, /* seed = */ -1.0);

    for b in 0..4_usize {
        let cql_active = scratch_after_cold[SCRATCH_ACTIVE_CQL + b];
        let c51_active = scratch_after_cold[SCRATCH_ACTIVE_C51 + b];
        assert!(
            cql_active.abs() < 1e-7,
            "cold start branch={b}: cql_active={cql_active:.6} should be 0.0 (genuine cold start)"
        );
        assert!(
            c51_active.abs() < 1e-7,
            "cold start branch={b}: c51_active={c51_active:.6} should be 0.0 (genuine cold start)"
        );
        // Budget slots also written 0.0 (no-op via apply_pearls_ad short-circuit).
        assert!(
            scratch_after_cold[SCRATCH_BUDGET_CQL + b].abs() < 1e-7,
            "cold start branch={b}: budget_cql={:.6} should be 0.0 (no-op marker)",
            scratch_after_cold[SCRATCH_BUDGET_CQL + b]
        );
        assert!(
            scratch_after_cold[SCRATCH_BUDGET_C51 + b].abs() < 1e-7,
            "cold start branch={b}: budget_c51={:.6} should be 0.0 (no-op marker)",
            scratch_after_cold[SCRATCH_BUDGET_C51 + b]
        );
    }

    // ── Step 2: active (synthetic non-zero IQN/CQL/C51 norms). ───────────────
    // grad_decomp pinned layout per component: [mag, dir, trunk]. We populate
    // all 3 slices so every branch has a non-zero h_n / iqn_n.
    //   IQN slice norms:    [1.0, 1.0, 1.0]
    //   CQL_SX slice norms: [0.5, 0.5, 0.5]
    //   C51 slice norms:    [0.3, 0.3, 0.3]
    // ISV: epoch_idx = 0 (first step) → welford_alpha = 1/max(1,0) = 1.0 → full
    // first-step update. Flatness defaults to 0 (clamped) → target_ratio for
    // CQL = 2.0, target_ratio for C51 = 0.0. With actual_ratio_CQL = 0.5/1.0 =
    // 0.5 → correction_CQL = 2.0/0.5 = 4.0. Old budget = 0 → seed-from-cold-start
    // branch: new_budget = COLD_START_FLOOR_CQL · correction = 0.02 · 4.0 = 0.08.
    // For C51: target=0 → correction=0 → new_budget = COLD_START_FLOOR_C51 · 0 = 0
    // (clamped at EPS_DIV = 1e-8).
    let mut grad_active = vec![0.0_f32; GRAD_DECOMP_TOTAL];
    for slice in 0..3_usize {
        grad_active[GRAD_OFFSET_IQN    + slice] = 1.0;
        grad_active[GRAD_OFFSET_CQL_SX + slice] = 0.5;
        grad_active[GRAD_OFFSET_C51    + slice] = 0.3;
    }

    let scratch_after_active = launch_once(&isv_cold, &grad_active, /* seed = */ -1.0);

    const EPS_DIV: f32 = 1e-8;
    const COLD_START_FLOOR_CQL: f32 = 0.02;

    for b in 0..4_usize {
        let cql_active = scratch_after_active[SCRATCH_ACTIVE_CQL + b];
        let c51_active = scratch_after_active[SCRATCH_ACTIVE_C51 + b];
        assert!(
            cql_active >= 0.5,
            "active branch={b}: cql_active={cql_active:.6} should be 1.0 after grads populated"
        );
        assert!(
            c51_active >= 0.5,
            "active branch={b}: c51_active={c51_active:.6} should be 1.0 after grads populated"
        );

        // Budget slots reflect controller's seed-from-cold-start compute, not
        // the bootstrap value. CQL: target_ratio=2.0, actual=0.5,
        // correction=4.0, seed = 0.02·4.0 = 0.08. C51: target_ratio=0.0
        // (flatness=0), so seed = 0.05·0 = 0 → clamped at EPS_DIV.
        let budget_cql = scratch_after_active[SCRATCH_BUDGET_CQL + b];
        let budget_c51 = scratch_after_active[SCRATCH_BUDGET_C51 + b];

        let expected_cql = COLD_START_FLOOR_CQL * 4.0;
        assert!(
            (budget_cql - expected_cql).abs() / expected_cql < 0.01,
            "active branch={b}: budget_cql={budget_cql:.6} expected {expected_cql:.6} (cold_start_basis · correction; not the {COLD_START_FLOOR_CQL:.6} bootstrap)"
        );
        // C51 target_ratio = ANCHOR_C51_RATIO · flatness = 1.0 · 0 = 0; the
        // seed becomes COLD_START_FLOOR_C51 · 0 = 0 → clamped at EPS_DIV.
        // The load-bearing assertion: budget != bootstrap (0.05).
        assert!(
            budget_c51 < 0.04,
            "active branch={b}: budget_c51={budget_c51:.6} should be ~EPS_DIV (target_ratio=0), not bootstrap 0.05"
        );
        assert!(
            budget_c51 >= EPS_DIV * 0.99,
            "active branch={b}: budget_c51={budget_c51:.6} should be at least EPS_DIV={EPS_DIV:.1e}"
        );
    }

    // ── Step 3: re-launch with prior was_active flags simulated by writing 1.0
    //   into LB_*_ACTIVE_BASE. Even with grads gated off (back to zero),
    //   transient cold-start should hold prior budget verbatim and re-assert
    //   active=1, NOT collapse back to consumer bootstrap.
    let mut isv_active = vec![0.0_f32; ISV_SIZE];
    for b in 0..4_usize {
        isv_active[LB_CQL_ACTIVE_BASE + b] = 1.0;     // was_active = 1
        isv_active[LB_C51_ACTIVE_BASE + b] = 1.0;
        isv_active[BUDGET_CQL_BASE + b]    = 0.07;    // prior budget != bootstrap
        isv_active[BUDGET_C51_BASE + b]    = 0.04;    // prior budget != bootstrap
    }

    let scratch_after_transient = launch_once(&isv_active, &grad_cold, /* seed = */ -1.0);
    for b in 0..4_usize {
        // Active flag held at 1.
        assert!(
            scratch_after_transient[SCRATCH_ACTIVE_CQL + b] >= 0.5,
            "transient grad-gate branch={b}: cql_active={:.6} should stay 1.0 (was_active=1)",
            scratch_after_transient[SCRATCH_ACTIVE_CQL + b]
        );
        assert!(
            scratch_after_transient[SCRATCH_ACTIVE_C51 + b] >= 0.5,
            "transient grad-gate branch={b}: c51_active={:.6} should stay 1.0 (was_active=1)",
            scratch_after_transient[SCRATCH_ACTIVE_C51 + b]
        );
        // Budget held verbatim at prior value (no drift on a measurement gap).
        assert!(
            (scratch_after_transient[SCRATCH_BUDGET_CQL + b] - 0.07).abs() < 1e-5,
            "transient grad-gate branch={b}: budget_cql={:.6} should hold prior 0.07",
            scratch_after_transient[SCRATCH_BUDGET_CQL + b]
        );
        assert!(
            (scratch_after_transient[SCRATCH_BUDGET_C51 + b] - 0.04).abs() < 1e-5,
            "transient grad-gate branch={b}: budget_c51={:.6} should hold prior 0.04",
            scratch_after_transient[SCRATCH_BUDGET_C51 + b]
        );
    }
}