//! SP14 Layer A + B oracle tests (post-Layer C Phase C.1 deletion).
//!
//! Layer A is the pre-Smoke-A stability sweep: small fixes addressing the
//! aux_w controller pathologies surfaced by Smoke A. Tests here run on CPU —
//! they exercise pure-math constants and the host-side `compute_aux_w_p0b`
//! controller without GPU dependencies.
//!
//! - **A.2**: `set_aux_weight` clamp lift from SP11-era `[0.05, 0.3]` to the
//!   SP13 P0b design range `[base × floor_ratio, base × ceil_ratio]` =
//!   `[0.15, 1.5]`. Verified here as a constants check; the trainer-side
//!   wiring is verified by Smoke A2-A end-to-end behaviour.
//!
//! Layer B retained tests (post-Phase C.1 atomic α-machinery deletion,
//! 2026-05-08): only the q_disagreement diagnostic chain (B.3) and the
//! Coupling A forward feature wire (B.6) survive. The α-machinery oracle
//! tests (alpha_grad_schmitt_hysteresis, alpha_grad_adaptive_beta,
//! gradient_hack_circuit_breaker_fires) were deleted atomically with the
//! kernel files — see
//! `docs/superpowers/plans/2026-05-07-sp14-layer-c-separate-aux-trunk.md` §C.1.
//!
//! - **B.3**: `q_disagreement_update_kernel` — K=4↔K=2 mapped argmax mismatch
//!   between Q-head direction (Short/Hold/Long/Flat) and aux head's K=2
//!   prediction (down/up); Hold/Flat masked. Pearl-A first-observation
//!   bootstrap; Welford variance EMA in same kernel. No atomicAdd per
//!   `feedback_no_atomicadd.md`. Post-C.1 the producer survives as a
//!   HEALTH_DIAG diagnostic only — no consumer reads slots 383/384/389.
//!
//! - **B.6**: `dir_concat_qaux_kernel` — Coupling A forward feature wire.
//!   Concatenates `[h_s2 ; aux_softmax_diff]` into `[B, SH2 + 1]` for the
//!   direction Q-head's first FC SGEMM input. Preserved through Phase C.1.
//!
//! Run GPU tests on a GPU host:
//!
//!     SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 \
//!       cargo test -p ml --test sp14_oracle_tests --features cuda \
//!         -- --ignored --nocapture

/// Test A.2: SP13 P0b aux_w bounds resolve to [0.15, 1.5] from the documented
/// constants. The pre-fix `set_aux_weight` clamp at gpu_dqn_trainer.rs:14722
/// was `[0.05, 0.3]` — the SP11-era cap that silently chopped the entire
/// upper half of the controller's intended range. This test verifies the
/// constants used by the post-fix clamp resolve to the design range.
#[test]
fn set_aux_weight_clamp_range() {
    use ml::cuda_pipeline::sp13_isv_slots::{
        AUX_W_BASE, AUX_W_HARD_FLOOR_RATIO, AUX_W_HARD_CEIL_RATIO,
    };
    // Direct math check; not a GPU test.
    let floor = AUX_W_BASE * AUX_W_HARD_FLOOR_RATIO;
    let ceil = AUX_W_BASE * AUX_W_HARD_CEIL_RATIO;
    assert_eq!(floor, 0.15);
    assert_eq!(ceil, 1.5);

    // Verify the trainer struct's clamp uses these bounds. We use a direct
    // f32 clamp here mirroring the production code; the test ensures the
    // production constants resolve to the expected values.
    let test_value: f32 = 1.0;
    let clamped = test_value.clamp(floor, ceil);
    assert_eq!(clamped, 1.0,
        "1.0 should NOT be clamped under [0.15, 1.5]; got {}", clamped);

    let too_low: f32 = 0.05;
    let clamped_low = too_low.clamp(floor, ceil);
    assert_eq!(clamped_low, 0.15);

    let too_high: f32 = 2.0;
    let clamped_high = too_high.clamp(floor, ceil);
    assert_eq!(clamped_high, 1.5);
}

// ═══════════════════════════════════════════════════════════════════════════
// SP14 Layer B GPU oracle tests — `#[cfg(feature = "cuda")]` gated.
//
// Each test is `#[ignore = "requires GPU"]` per the codebase convention so
// non-GPU runners skip them without `--ignored`. All CPU↔GPU buffers are
// `MappedF32Buffer` per `feedback_no_htod_htoh_only_mapped_pinned.md` (tests
// are not exempt from the rule). Cubins are `include_bytes!`'d from
// `OUT_DIR`, mirroring `sp13_layer_b_oracle_tests.rs`.
//
// SP14 Layer C Phase C.1 (2026-05-08): the α-machinery cubin loaders
// (`load_alpha_grad`, `load_grad_hack`) and tests
// (`alpha_grad_schmitt_hysteresis`, `alpha_grad_adaptive_beta`,
// `gradient_hack_circuit_breaker_fires`) were deleted atomically with
// the kernel files per `feedback_no_legacy_aliases`. Surviving:
// q_disagreement_* (3 tests) + dir_concat_qaux_correct (1 test).
// ═══════════════════════════════════════════════════════════════════════════

#[cfg(feature = "cuda")]
#[allow(unsafe_code)] // CUDA kernel launch + mapped-pinned memory.
mod gpu {
    use std::sync::Arc;

    use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg};
    use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer;
    use ml::cuda_pipeline::sp14_isv_slots::{
        Q_DISAGREEMENT_LONG_EMA_INDEX,
        Q_DISAGREEMENT_SHORT_EMA_INDEX,
        Q_DISAGREEMENT_VARIANCE_EMA_INDEX,
    };

    // ── B.3 cubin handle ────────────────────────────────────────────────────

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

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

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

    /// Shared-memory bytes for the kernel: `2 * blockDim.x * sizeof(float)`
    /// (the kernel uses `extern __shared__ float smem[]` split into
    /// `num_smem` and `cnt_smem` arrays of `blockDim.x` floats each).
    /// blockDim.x = 256 → 2 × 256 × 4 = 2048 bytes.
    const Q_DIS_BLOCK: u32 = 256;
    const Q_DIS_SMEM_BYTES: u32 = 2 * Q_DIS_BLOCK * std::mem::size_of::<f32>() as u32;

    // ── Test B.3a: K=4 → K=2 mapping correctness ───────────────────────────

    /// 8-row batch with deliberate mix of agreements, disagreements, and
    /// masked rows (Hold/Flat). With Q_DISAGREEMENT_SHORT_EMA at sentinel
    /// 0.5, Pearl-A first-observation replaces the EMA directly with the
    /// batch mean disagreement (2 disagreements / 4 contributing rows = 0.5).
    /// Both short and long EMAs land on 0.5 because the Pearl-A bootstrap
    /// applies to both EMAs simultaneously when prev = sentinel.
    #[test]
    #[ignore = "requires GPU"]
    fn q_disagreement_k4_k2_mapping() {
        let stream = make_test_stream();
        let kernel = load_q_disagreement(&stream);

        const B: usize = 8;
        const K_AUX: usize = 2;
        const K_DIR: usize = 4;
        const ISV_DIM: usize = 1024;
        let mut isv = vec![0.0_f32; ISV_DIM];
        // Pearl-A sentinels.
        isv[Q_DISAGREEMENT_SHORT_EMA_INDEX] = 0.5;
        isv[Q_DISAGREEMENT_LONG_EMA_INDEX]  = 0.5;
        isv[Q_DISAGREEMENT_VARIANCE_EMA_INDEX] = 0.0;

        // 8-row batch: rows 0,1,2,3 contributing (Short/Long), rows 4,5 are
        // Hold (masked), rows 6,7 are Flat (masked).
        // Agreements (aux matches Q): rows 0 (Short→down), 2 (Long→up).
        // Disagreements: rows 1 (Short→up), 3 (Long→down). 2 dis / 4 = 0.5.
        let mut aux_softmax = vec![0.0_f32; B * K_AUX];
        let mut q_logits    = vec![0.0_f32; B * K_DIR];
        // Row 0: Q=Short(0), aux=down(0) → agreement
        aux_softmax[0 * K_AUX + 0] = 0.9; aux_softmax[0 * K_AUX + 1] = 0.1;
        q_logits[0 * K_DIR + 0] = 5.0;
        // Row 1: Q=Short(0), aux=up(1) → disagreement
        aux_softmax[1 * K_AUX + 0] = 0.1; aux_softmax[1 * K_AUX + 1] = 0.9;
        q_logits[1 * K_DIR + 0] = 5.0;
        // Row 2: Q=Long(2), aux=up(1) → agreement
        aux_softmax[2 * K_AUX + 0] = 0.1; aux_softmax[2 * K_AUX + 1] = 0.9;
        q_logits[2 * K_DIR + 2] = 5.0;
        // Row 3: Q=Long(2), aux=down(0) → disagreement
        aux_softmax[3 * K_AUX + 0] = 0.9; aux_softmax[3 * K_AUX + 1] = 0.1;
        q_logits[3 * K_DIR + 2] = 5.0;
        // Row 4: Q=Hold(1), masked
        q_logits[4 * K_DIR + 1] = 5.0;
        // Row 5: Q=Hold(1), masked
        q_logits[5 * K_DIR + 1] = 5.0;
        // Row 6: Q=Flat(3), masked
        q_logits[6 * K_DIR + 3] = 5.0;
        // Row 7: Q=Flat(3), masked
        q_logits[7 * K_DIR + 3] = 5.0;

        let aux_buf = unsafe { MappedF32Buffer::new(B * K_AUX) }
            .expect("alloc aux softmax buf");
        aux_buf.write_from_slice(&aux_softmax);
        let q_buf = unsafe { MappedF32Buffer::new(B * K_DIR) }
            .expect("alloc q_logits buf");
        q_buf.write_from_slice(&q_logits);
        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }
            .expect("alloc isv buf");
        isv_buf.write_from_slice(&isv);

        let b_i32: i32 = B as i32;
        let q_stride: i32 = K_DIR as i32;
        let alpha_short: f32 = 0.3;
        let alpha_long: f32 = 0.05;
        let alpha_var: f32 = 0.05;

        unsafe {
            stream
                .launch_builder(&kernel)
                .arg(&aux_buf.dev_ptr)
                .arg(&q_buf.dev_ptr)
                .arg(&isv_buf.dev_ptr)
                .arg(&b_i32)
                .arg(&q_stride)
                .arg(&alpha_short)
                .arg(&alpha_long)
                .arg(&alpha_var)
                .launch(LaunchConfig {
                    grid_dim:         (1, 1, 1),
                    block_dim:        (Q_DIS_BLOCK, 1, 1),
                    shared_mem_bytes: Q_DIS_SMEM_BYTES,
                })
                .expect("launch q_disagreement_update_kernel (k4_k2)");
        }
        stream.synchronize().expect("sync after q_disagreement");

        let result = isv_buf.read_all();
        let short = result[Q_DISAGREEMENT_SHORT_EMA_INDEX];
        let long  = result[Q_DISAGREEMENT_LONG_EMA_INDEX];
        // Pearl-A: sentinel=0.5 + batch_mean=0.5 → exactly 0.5.
        assert!((short - 0.5).abs() < 1e-5, "short EMA: expected ~0.5, got {short}");
        assert!((long  - 0.5).abs() < 1e-5, "long EMA: expected ~0.5, got {long}");
    }

    /// All-Hold batch: every row is Hold (action=1), so every row is
    /// masked, contributing 0 to numerator AND 0 to count. Kernel must
    /// preserve the prior EMA bit-exactly (no division by zero).
    #[test]
    #[ignore = "requires GPU"]
    fn q_disagreement_all_hold_no_contribution() {
        let stream = make_test_stream();
        let kernel = load_q_disagreement(&stream);

        const B: usize = 8;
        const K_AUX: usize = 2;
        const K_DIR: usize = 4;
        const ISV_DIM: usize = 1024;
        const SEED_SHORT: f32 = 0.42;
        const SEED_LONG:  f32 = 0.31;
        const SEED_VAR:   f32 = 0.07;
        let mut isv = vec![0.0_f32; ISV_DIM];
        isv[Q_DISAGREEMENT_SHORT_EMA_INDEX] = SEED_SHORT;
        isv[Q_DISAGREEMENT_LONG_EMA_INDEX]  = SEED_LONG;
        isv[Q_DISAGREEMENT_VARIANCE_EMA_INDEX] = SEED_VAR;

        let aux_softmax = vec![0.5_f32; B * K_AUX];
        let mut q_logits = vec![0.0_f32; B * K_DIR];
        for b in 0..B {
            q_logits[b * K_DIR + 1] = 5.0; // all-Hold
        }

        let aux_buf = unsafe { MappedF32Buffer::new(B * K_AUX) }
            .expect("alloc aux buf");
        aux_buf.write_from_slice(&aux_softmax);
        let q_buf = unsafe { MappedF32Buffer::new(B * K_DIR) }
            .expect("alloc q buf");
        q_buf.write_from_slice(&q_logits);
        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }
            .expect("alloc isv buf");
        isv_buf.write_from_slice(&isv);

        let b_i32: i32 = B as i32;
        let q_stride: i32 = K_DIR as i32;
        unsafe {
            stream
                .launch_builder(&kernel)
                .arg(&aux_buf.dev_ptr)
                .arg(&q_buf.dev_ptr)
                .arg(&isv_buf.dev_ptr)
                .arg(&b_i32)
                .arg(&q_stride)
                .arg(&0.3_f32)
                .arg(&0.05_f32)
                .arg(&0.05_f32)
                .launch(LaunchConfig {
                    grid_dim: (1, 1, 1),
                    block_dim: (Q_DIS_BLOCK, 1, 1),
                    shared_mem_bytes: Q_DIS_SMEM_BYTES,
                })
                .expect("launch q_disagreement_update_kernel (all-hold)");
        }
        stream.synchronize().expect("sync after q_disagreement");

        let result = isv_buf.read_all();
        let short = result[Q_DISAGREEMENT_SHORT_EMA_INDEX];
        let long  = result[Q_DISAGREEMENT_LONG_EMA_INDEX];
        let var   = result[Q_DISAGREEMENT_VARIANCE_EMA_INDEX];
        assert_eq!(short, SEED_SHORT, "short EMA must be preserved on empty batch");
        assert_eq!(long,  SEED_LONG,  "long EMA must be preserved on empty batch");
        assert_eq!(var,   SEED_VAR,   "var EMA must be preserved on empty batch");
    }

    /// Empty batch (B=0): no contributions; EMAs preserved bit-exactly.
    #[test]
    #[ignore = "requires GPU"]
    fn q_disagreement_empty_batch_preserves_ema() {
        let stream = make_test_stream();
        let kernel = load_q_disagreement(&stream);

        const ISV_DIM: usize = 1024;
        const SEED_SHORT: f32 = 0.42;
        const SEED_LONG:  f32 = 0.31;
        const SEED_VAR:   f32 = 0.07;
        let mut isv = vec![0.0_f32; ISV_DIM];
        isv[Q_DISAGREEMENT_SHORT_EMA_INDEX] = SEED_SHORT;
        isv[Q_DISAGREEMENT_LONG_EMA_INDEX]  = SEED_LONG;
        isv[Q_DISAGREEMENT_VARIANCE_EMA_INDEX] = SEED_VAR;

        // Allocate non-zero-sized buffers (cudarc rejects size-0 allocs)
        // but pass B=0 so the kernel processes zero rows.
        let aux_buf = unsafe { MappedF32Buffer::new(2) }
            .expect("alloc aux buf");
        aux_buf.write_from_slice(&[0.0_f32; 2]);
        let q_buf = unsafe { MappedF32Buffer::new(4) }
            .expect("alloc q buf");
        q_buf.write_from_slice(&[0.0_f32; 4]);
        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }
            .expect("alloc isv buf");
        isv_buf.write_from_slice(&isv);

        let b_i32: i32 = 0;
        let q_stride: i32 = 4;
        unsafe {
            stream
                .launch_builder(&kernel)
                .arg(&aux_buf.dev_ptr)
                .arg(&q_buf.dev_ptr)
                .arg(&isv_buf.dev_ptr)
                .arg(&b_i32)
                .arg(&q_stride)
                .arg(&0.3_f32)
                .arg(&0.05_f32)
                .arg(&0.05_f32)
                .launch(LaunchConfig {
                    grid_dim: (1, 1, 1),
                    block_dim: (Q_DIS_BLOCK, 1, 1),
                    shared_mem_bytes: Q_DIS_SMEM_BYTES,
                })
                .expect("launch q_disagreement_update_kernel (empty batch)");
        }
        stream.synchronize().expect("sync after q_disagreement");

        let result = isv_buf.read_all();
        assert_eq!(result[Q_DISAGREEMENT_SHORT_EMA_INDEX], SEED_SHORT);
        assert_eq!(result[Q_DISAGREEMENT_LONG_EMA_INDEX],  SEED_LONG);
        let var_ema = result[Q_DISAGREEMENT_VARIANCE_EMA_INDEX];
        assert_eq!(
            var_ema, SEED_VAR,
            "Empty-batch must preserve variance EMA bit-exactly; expected {SEED_VAR}, got {var_ema}",
        );
    }

    // ── B.6 cubin handle ────────────────────────────────────────────────────

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

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

    // ── Test B.6: Pre-SGEMM concat correctness ─────────────────────────────

    /// 4-row batch, SH2=8. Rows 0-1 aux predicts "down" (softmax=[0.9, 0.1]),
    /// rows 2-3 aux predicts "up" (softmax=[0.1, 0.9]).
    ///
    /// Verifies:
    ///   - Positions 0..SH2 in each output row contain the original h_s2 row
    ///     values unchanged (contiguous copy, per-thread map).
    ///   - Position SH2 contains `softmax[b, 1] - softmax[b, 0]`: -0.8 for
    ///     "down" rows, +0.8 for "up" rows.
    ///
    /// aux_softmax_diff is structurally bounded to [-1, +1] by the softmax
    /// composition per `pearl_bounded_modifier_outputs_require_structural_
    /// activation`. No reduction; no atomicAdd.
    ///
    /// Coupling A (forward feature wire) — preserved through SP14 Layer C
    /// Phase C.1 deletion of the α-machinery (Coupling B).
    #[test]
    #[ignore = "requires GPU"]
    fn dir_concat_qaux_correct() {
        let stream = make_test_stream();
        let kernel = load_dir_concat(&stream);

        const B: usize = 4;
        const SH2: usize = 8;
        const K_AUX: usize = 2;

        // h_s2 input: sequential floats 0..B*SH2
        let h_s2: Vec<f32> = (0..B * SH2).map(|i| i as f32).collect();

        // aux_softmax: rows 0,1 → "down" ([0.9, 0.1]) → diff = 0.1 - 0.9 = -0.8
        //              rows 2,3 → "up"   ([0.1, 0.9]) → diff = 0.9 - 0.1 = +0.8
        let mut aux_softmax = vec![0.0_f32; B * K_AUX];
        for b in 0..2 {
            aux_softmax[b * K_AUX + 0] = 0.9;
            aux_softmax[b * K_AUX + 1] = 0.1;
        }
        for b in 2..4 {
            aux_softmax[b * K_AUX + 0] = 0.1;
            aux_softmax[b * K_AUX + 1] = 0.9;
        }

        let h_buf = unsafe { MappedF32Buffer::new(B * SH2) }.expect("alloc h_s2 buf");
        h_buf.write_from_slice(&h_s2);
        let aux_buf = unsafe { MappedF32Buffer::new(B * K_AUX) }.expect("alloc aux buf");
        aux_buf.write_from_slice(&aux_softmax);
        let out_buf = unsafe { MappedF32Buffer::new(B * (SH2 + 1)) }.expect("alloc out buf");

        let b_i32: i32 = B as i32;
        let sh2_i32: i32 = SH2 as i32;
        let total = (B * (SH2 + 1)) as u32;
        let blocks = ((total + 255) / 256).max(1);

        unsafe {
            stream
                .launch_builder(&kernel)
                .arg(&h_buf.dev_ptr)
                .arg(&aux_buf.dev_ptr)
                .arg(&out_buf.dev_ptr)
                .arg(&b_i32)
                .arg(&sh2_i32)
                .launch(LaunchConfig {
                    grid_dim:         (blocks, 1, 1),
                    block_dim:        (256, 1, 1),
                    shared_mem_bytes: 0,
                })
                .expect("launch dir_concat_qaux_kernel");
        }
        stream.synchronize().expect("sync after dir_concat_qaux");

        let out = out_buf.read_all();
        for b in 0..B {
            for s in 0..SH2 {
                let want = h_s2[b * SH2 + s];
                let got = out[b * (SH2 + 1) + s];
                assert_eq!(got, want, "row {b} col {s}: want {want}, got {got}");
            }
            let diff = out[b * (SH2 + 1) + SH2];
            let want_diff: f32 = if b < 2 { -0.8 } else { 0.8 };
            assert!((diff - want_diff).abs() < 1e-5,
                "row {b} wire col: want {want_diff}, got {diff}");
        }
    }
}

// ═══════════════════════════════════════════════════════════════════════════
// Class A P0-A (2026-05-08) — adaptive REWARD_POS/NEG_CAP producer tests.
//
// Verifies the `reward_cap_update_kernel` producer:
//   1. Cold-start Pearl-A bootstrap: ISV stays at sentinel until first valid
//      observation, then REPLACES (no blend).
//   2. 2:1 asymmetry preserved at producer-time: NEG_CAP = -2 × POS_CAP.
//   3. Bounds enforced: POS in [1, 50], NEG in [-100, -2].
//   4. Welford slow EMA after bootstrap (α=0.01).
//
// All tests are #[ignore = "requires GPU"]; gated under #[cfg(feature = "cuda")].
// ═══════════════════════════════════════════════════════════════════════════

#[cfg(feature = "cuda")]
#[allow(unsafe_code)] // CUDA kernel launch + mapped-pinned memory.
mod sp14_p0a_reward_cap_gpu {
    use std::sync::Arc;

    use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg};
    use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedI32Buffer};
    use ml::cuda_pipeline::sp14_isv_slots::{
        REWARD_CAP_EMA_ALPHA, REWARD_CAP_SAFETY_FACTOR, REWARD_NEG_CAP_ADAPTIVE_INDEX,
        REWARD_NEG_TO_POS_RATIO, REWARD_POS_CAP_ADAPTIVE_INDEX, REWARD_POS_CAP_MAX,
        REWARD_POS_CAP_MIN, SENTINEL_REWARD_NEG_CAP, SENTINEL_REWARD_POS_CAP,
    };

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

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

    fn load_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
        let module = stream
            .context()
            .load_cubin(REWARD_CAP_UPDATE_CUBIN.to_vec())
            .expect("load reward_cap_update_kernel cubin");
        module
            .load_function("reward_cap_update")
            .expect("load reward_cap_update function")
    }

    /// 4 arrays × 256 × 4 bytes = 4096 bytes shmem.
    const BLK_DIM: u32 = 256;
    const SMEM_BYTES: u32 = 4 * BLK_DIM * std::mem::size_of::<f32>() as u32;

    /// Helper: launch with a fixed config.
    #[allow(clippy::too_many_arguments)]
    fn launch_reward_cap(
        stream: &Arc<CudaStream>,
        kernel: &CudaFunction,
        step_ret_ptr: u64,
        trade_close_ptr: u64,
        total_samples: i32,
        isv_ptr: u64,
        pos_idx: i32,
        neg_idx: i32,
        sentinel_pos: f32,
        sentinel_neg: f32,
        safety_factor: f32,
        neg_to_pos_ratio: f32,
        pos_min: f32,
        pos_max: f32,
        alpha: f32,
    ) {
        unsafe {
            stream
                .launch_builder(kernel)
                .arg(&step_ret_ptr)
                .arg(&trade_close_ptr)
                .arg(&total_samples)
                .arg(&isv_ptr)
                .arg(&pos_idx)
                .arg(&neg_idx)
                .arg(&sentinel_pos)
                .arg(&sentinel_neg)
                .arg(&safety_factor)
                .arg(&neg_to_pos_ratio)
                .arg(&pos_min)
                .arg(&pos_max)
                .arg(&alpha)
                .launch(LaunchConfig {
                    grid_dim: (1, 1, 1),
                    block_dim: (BLK_DIM, 1, 1),
                    shared_mem_bytes: SMEM_BYTES,
                })
                .expect("launch reward_cap_update");
        }
        stream.synchronize().expect("sync after reward_cap_update");
    }

    /// Test 1 — Pearl-A first-observation bootstrap.
    ///
    /// 4 winning closes (all step_ret>0) with values [1.0, 2.0, 3.0, 4.0],
    /// all trade_close=1. Mean=2.5, sigma≈1.118, p99≈2.5+2.326×1.118≈5.10,
    /// max=4.0. Estimator picks `max(p99, max)=5.10`, then × 1.5 = 7.65.
    /// Clamped to [1, 50] → 7.65 stays.
    ///
    /// Cold-start (current = sentinel 5.0): Pearl-A REPLACES → POS=7.65,
    /// NEG=-2 × 7.65 = -15.3.
    #[test]
    #[ignore = "requires GPU"]
    fn reward_cap_pearl_a_bootstrap() {
        let stream = make_stream();
        let kernel = load_kernel(&stream);

        const ISV_DIM: usize = 1024;
        let mut isv = vec![0.0_f32; ISV_DIM];
        // Cold-start: sentinels (matches pre-P0-A constants).
        isv[REWARD_POS_CAP_ADAPTIVE_INDEX] = SENTINEL_REWARD_POS_CAP;
        isv[REWARD_NEG_CAP_ADAPTIVE_INDEX] = SENTINEL_REWARD_NEG_CAP;

        let step_ret: Vec<f32> = vec![1.0, 2.0, 3.0, 4.0];
        let trade_close: Vec<i32> = vec![1, 1, 1, 1];

        let n = step_ret.len() as i32;
        let ret_buf = unsafe { MappedF32Buffer::new(step_ret.len()) }
            .expect("alloc step_ret");
        ret_buf.write_from_slice(&step_ret);
        // i32 buffer via MappedF32Buffer — repurpose: the kernel reads
        // `int*`, so we need a u32-aligned i32 slice. We use a separate
        // raw allocation for the i32 close mask.
        let close_buf = unsafe { MappedI32Buffer::new(trade_close.len()) }
            .expect("alloc trade_close");
        close_buf.write_from_slice(&trade_close);
        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }
            .expect("alloc isv");
        isv_buf.write_from_slice(&isv);

        launch_reward_cap(
            &stream, &kernel,
            ret_buf.dev_ptr, close_buf.dev_ptr, n, isv_buf.dev_ptr,
            REWARD_POS_CAP_ADAPTIVE_INDEX as i32,
            REWARD_NEG_CAP_ADAPTIVE_INDEX as i32,
            SENTINEL_REWARD_POS_CAP, SENTINEL_REWARD_NEG_CAP,
            REWARD_CAP_SAFETY_FACTOR, REWARD_NEG_TO_POS_RATIO,
            REWARD_POS_CAP_MIN, REWARD_POS_CAP_MAX,
            REWARD_CAP_EMA_ALPHA,
        );

        let result = isv_buf.read_all();
        let pos = result[REWARD_POS_CAP_ADAPTIVE_INDEX];
        let neg = result[REWARD_NEG_CAP_ADAPTIVE_INDEX];

        // Pearl-A: sentinel → REPLACE with target. Target = max(p99, max) × 1.5.
        // mean=2.5, var=(1+4+9+16)/4 - 6.25 = 7.5 - 6.25 = 1.25, sigma≈1.118
        // p99 ≈ 2.5 + 2.326 × 1.118 ≈ 5.10. max=4.0. Picks 5.10 × 1.5 ≈ 7.65.
        // Clamped to [1, 50] → stays at ~7.65.
        assert!(
            (pos - 7.65).abs() < 0.05,
            "Pearl-A bootstrap: expected POS≈7.65, got {pos}"
        );
        // 2:1 asymmetry preserved at producer-time: NEG = -2 × POS.
        assert!(
            (neg - (-2.0 * pos)).abs() < 1e-4,
            "NEG must equal -2 × POS at producer-time (Kahneman 2:1 single source of truth); got POS={pos}, NEG={neg}"
        );
    }

    /// Test 2 — No winning trades → ISV slots preserved bit-exactly.
    ///
    /// All trade_close=0 (no closes). Kernel must skip the EMA update
    /// (count==0 guard) and leave both slots at their seeded values.
    #[test]
    #[ignore = "requires GPU"]
    fn reward_cap_no_winning_trades_preserves_isv() {
        let stream = make_stream();
        let kernel = load_kernel(&stream);

        const ISV_DIM: usize = 1024;
        const SEED_POS: f32 = 12.5;
        const SEED_NEG: f32 = -25.0;
        let mut isv = vec![0.0_f32; ISV_DIM];
        isv[REWARD_POS_CAP_ADAPTIVE_INDEX] = SEED_POS;
        isv[REWARD_NEG_CAP_ADAPTIVE_INDEX] = SEED_NEG;

        // Step returns present but no closes — none should contribute.
        let step_ret: Vec<f32> = vec![3.0, 5.0, -2.0];
        let trade_close: Vec<i32> = vec![0, 0, 0];

        let n = step_ret.len() as i32;
        let ret_buf = unsafe { MappedF32Buffer::new(step_ret.len()) }
            .expect("alloc step_ret");
        ret_buf.write_from_slice(&step_ret);
        let close_buf = unsafe { MappedI32Buffer::new(trade_close.len()) }
            .expect("alloc trade_close");
        close_buf.write_from_slice(&trade_close);
        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }
            .expect("alloc isv");
        isv_buf.write_from_slice(&isv);

        launch_reward_cap(
            &stream, &kernel,
            ret_buf.dev_ptr, close_buf.dev_ptr, n, isv_buf.dev_ptr,
            REWARD_POS_CAP_ADAPTIVE_INDEX as i32,
            REWARD_NEG_CAP_ADAPTIVE_INDEX as i32,
            SENTINEL_REWARD_POS_CAP, SENTINEL_REWARD_NEG_CAP,
            REWARD_CAP_SAFETY_FACTOR, REWARD_NEG_TO_POS_RATIO,
            REWARD_POS_CAP_MIN, REWARD_POS_CAP_MAX,
            REWARD_CAP_EMA_ALPHA,
        );

        let result = isv_buf.read_all();
        assert_eq!(
            result[REWARD_POS_CAP_ADAPTIVE_INDEX], SEED_POS,
            "POS slot must be preserved bit-exactly when no winning closes",
        );
        assert_eq!(
            result[REWARD_NEG_CAP_ADAPTIVE_INDEX], SEED_NEG,
            "NEG slot must be preserved bit-exactly when no winning closes",
        );
    }

    /// Test 3 — Bounds enforcement: huge winning return clamps POS to MAX=50.
    ///
    /// Single huge winning close at 1000.0. p99 estimate ≈ 1000 + 0 = 1000,
    /// max=1000, × safety_factor=1.5 = 1500. Must clamp to POS_MAX=50.
    /// NEG = -2 × 50 = -100 (also clamped at the lower bound).
    #[test]
    #[ignore = "requires GPU"]
    fn reward_cap_bounds_clamp_extreme_outlier() {
        let stream = make_stream();
        let kernel = load_kernel(&stream);

        const ISV_DIM: usize = 1024;
        let mut isv = vec![0.0_f32; ISV_DIM];
        isv[REWARD_POS_CAP_ADAPTIVE_INDEX] = SENTINEL_REWARD_POS_CAP;
        isv[REWARD_NEG_CAP_ADAPTIVE_INDEX] = SENTINEL_REWARD_NEG_CAP;

        // Single huge winning close — would explode the cap if unclamped.
        let step_ret: Vec<f32> = vec![1000.0];
        let trade_close: Vec<i32> = vec![1];

        let n = step_ret.len() as i32;
        let ret_buf = unsafe { MappedF32Buffer::new(step_ret.len()) }
            .expect("alloc step_ret");
        ret_buf.write_from_slice(&step_ret);
        let close_buf = unsafe { MappedI32Buffer::new(trade_close.len()) }
            .expect("alloc trade_close");
        close_buf.write_from_slice(&trade_close);
        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }
            .expect("alloc isv");
        isv_buf.write_from_slice(&isv);

        launch_reward_cap(
            &stream, &kernel,
            ret_buf.dev_ptr, close_buf.dev_ptr, n, isv_buf.dev_ptr,
            REWARD_POS_CAP_ADAPTIVE_INDEX as i32,
            REWARD_NEG_CAP_ADAPTIVE_INDEX as i32,
            SENTINEL_REWARD_POS_CAP, SENTINEL_REWARD_NEG_CAP,
            REWARD_CAP_SAFETY_FACTOR, REWARD_NEG_TO_POS_RATIO,
            REWARD_POS_CAP_MIN, REWARD_POS_CAP_MAX,
            REWARD_CAP_EMA_ALPHA,
        );

        let result = isv_buf.read_all();
        let pos = result[REWARD_POS_CAP_ADAPTIVE_INDEX];
        let neg = result[REWARD_NEG_CAP_ADAPTIVE_INDEX];

        // Bounds: POS clamped to MAX=50.
        assert!(
            (pos - REWARD_POS_CAP_MAX).abs() < 1e-4,
            "POS must clamp to MAX=50 on extreme outlier; got {pos}"
        );
        // NEG = -2 × 50 = -100 (lower bound by construction).
        let expected_neg = -REWARD_NEG_TO_POS_RATIO * REWARD_POS_CAP_MAX;
        assert!(
            (neg - expected_neg).abs() < 1e-4,
            "NEG must clamp to -2 × MAX = {expected_neg}; got {neg}"
        );
    }

    /// Test 4 — Welford EMA blend after bootstrap.
    ///
    /// Pre-seed POS=10.0 (NOT sentinel — bootstrap path NOT taken).
    /// Single winning close at 4.0. mean=4, sigma=0, p99=4, max=4.
    /// Target = max(p99, max) × 1.5 = 4 × 1.5 = 6. Clamped to [1, 50] → 6.
    /// EMA blend: (1 - α) × 10 + α × 6 with α=0.01:
    /// = 0.99 × 10 + 0.01 × 6 = 9.96. NEG = -2 × 9.96 = -19.92.
    #[test]
    #[ignore = "requires GPU"]
    fn reward_cap_welford_ema_after_bootstrap() {
        let stream = make_stream();
        let kernel = load_kernel(&stream);

        const ISV_DIM: usize = 1024;
        const SEEDED_POS: f32 = 10.0;
        const SEEDED_NEG: f32 = -20.0;
        let mut isv = vec![0.0_f32; ISV_DIM];
        // NOT sentinel — Pearl-A skipped, EMA path taken.
        isv[REWARD_POS_CAP_ADAPTIVE_INDEX] = SEEDED_POS;
        isv[REWARD_NEG_CAP_ADAPTIVE_INDEX] = SEEDED_NEG;

        let step_ret: Vec<f32> = vec![4.0];
        let trade_close: Vec<i32> = vec![1];

        let n = step_ret.len() as i32;
        let ret_buf = unsafe { MappedF32Buffer::new(step_ret.len()) }
            .expect("alloc step_ret");
        ret_buf.write_from_slice(&step_ret);
        let close_buf = unsafe { MappedI32Buffer::new(trade_close.len()) }
            .expect("alloc trade_close");
        close_buf.write_from_slice(&trade_close);
        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }
            .expect("alloc isv");
        isv_buf.write_from_slice(&isv);

        launch_reward_cap(
            &stream, &kernel,
            ret_buf.dev_ptr, close_buf.dev_ptr, n, isv_buf.dev_ptr,
            REWARD_POS_CAP_ADAPTIVE_INDEX as i32,
            REWARD_NEG_CAP_ADAPTIVE_INDEX as i32,
            SENTINEL_REWARD_POS_CAP, SENTINEL_REWARD_NEG_CAP,
            REWARD_CAP_SAFETY_FACTOR, REWARD_NEG_TO_POS_RATIO,
            REWARD_POS_CAP_MIN, REWARD_POS_CAP_MAX,
            REWARD_CAP_EMA_ALPHA,
        );

        let result = isv_buf.read_all();
        let pos = result[REWARD_POS_CAP_ADAPTIVE_INDEX];
        let neg = result[REWARD_NEG_CAP_ADAPTIVE_INDEX];

        // EMA blend: 0.99 × 10 + 0.01 × 6 = 9.96.
        let expected_pos = (1.0 - REWARD_CAP_EMA_ALPHA) * SEEDED_POS
            + REWARD_CAP_EMA_ALPHA * (4.0 * REWARD_CAP_SAFETY_FACTOR);
        assert!(
            (pos - expected_pos).abs() < 1e-3,
            "POS Welford EMA: expected {expected_pos}, got {pos}"
        );
        // 2:1 ratio preserved.
        assert!(
            (neg - (-REWARD_NEG_TO_POS_RATIO * pos)).abs() < 1e-4,
            "NEG must equal -2 × POS at producer; got POS={pos}, NEG={neg}"
        );
    }
}

// ═══════════════════════════════════════════════════════════════════════════
// Class A P1-Producer (2026-05-08) — adaptive Bayesian Kelly priors producer
// tests.
//
// Verifies the `kelly_bayesian_priors_update_kernel` producer:
//   1. Pearl-A first-observation bootstrap: ISV at sentinel → REPLACES with
//      the realized aggregated stats (no blend).
//   2. Slow EMA (α=0.005) blend after bootstrap.
//   3. Bounds enforced: counts in [0.5, 100], sums in [0.001, 1.0].
//   4. No realized trades (zero PS_KELLY_* fields) → ISV slots preserved
//      bit-exactly.
//
// All tests are #[ignore = "requires GPU"]; gated under #[cfg(feature = "cuda")].
// ═══════════════════════════════════════════════════════════════════════════

#[cfg(feature = "cuda")]
#[allow(unsafe_code)] // CUDA kernel launch + mapped-pinned memory.
mod sp14_p1_kelly_priors_gpu {
    use std::sync::Arc;

    use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg};
    use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer;
    use ml::cuda_pipeline::sp14_isv_slots::{
        KELLY_PRIOR_COUNT_MAX, KELLY_PRIOR_COUNT_MIN, KELLY_PRIOR_EMA_ALPHA,
        KELLY_PRIOR_LOSSES_INDEX, KELLY_PRIOR_SUM_LOSSES_INDEX, KELLY_PRIOR_SUM_MAX,
        KELLY_PRIOR_SUM_MIN, KELLY_PRIOR_SUM_WINS_INDEX, KELLY_PRIOR_WINS_INDEX,
        SENTINEL_KELLY_PRIOR_LOSSES, SENTINEL_KELLY_PRIOR_SUM_LOSSES,
        SENTINEL_KELLY_PRIOR_SUM_WINS, SENTINEL_KELLY_PRIOR_WINS,
    };

    // PS_STRIDE constant — matches state_layout.cuh::PS_STRIDE and the
    // value used by `launch_kelly_cap_update`. Grown 41→43 by Plan 3 D.4c.
    const PS_STRIDE: usize = 43;
    // PS_KELLY_* slot offsets within a single env's portfolio_state row.
    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;

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

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

    fn load_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
        let module = stream
            .context()
            .load_cubin(KELLY_BAYESIAN_PRIORS_UPDATE_CUBIN.to_vec())
            .expect("load kelly_bayesian_priors_update_kernel cubin");
        module
            .load_function("kelly_bayesian_priors_update")
            .expect("load kelly_bayesian_priors_update function")
    }

    /// 4 arrays × 256 × 4 bytes = 4096 bytes shmem.
    const BLK_DIM: u32 = 256;
    const SMEM_BYTES: u32 = 4 * BLK_DIM * std::mem::size_of::<f32>() as u32;

    /// Helper: launch with a fixed config.
    #[allow(clippy::too_many_arguments)]
    fn launch_kelly_priors(
        stream: &Arc<CudaStream>,
        kernel: &CudaFunction,
        portfolio_state_ptr: u64,
        n_envs: i32,
        ps_stride: i32,
        isv_ptr: u64,
        wins_idx: i32,
        losses_idx: i32,
        sum_wins_idx: i32,
        sum_losses_idx: i32,
        sentinel_wins: f32,
        sentinel_losses: f32,
        sentinel_sum_wins: f32,
        sentinel_sum_losses: f32,
        count_min: f32,
        count_max: f32,
        sum_min: f32,
        sum_max: f32,
        alpha: f32,
    ) {
        unsafe {
            stream
                .launch_builder(kernel)
                .arg(&portfolio_state_ptr)
                .arg(&n_envs)
                .arg(&ps_stride)
                .arg(&isv_ptr)
                .arg(&wins_idx)
                .arg(&losses_idx)
                .arg(&sum_wins_idx)
                .arg(&sum_losses_idx)
                .arg(&sentinel_wins)
                .arg(&sentinel_losses)
                .arg(&sentinel_sum_wins)
                .arg(&sentinel_sum_losses)
                .arg(&count_min)
                .arg(&count_max)
                .arg(&sum_min)
                .arg(&sum_max)
                .arg(&alpha)
                .launch(LaunchConfig {
                    grid_dim: (1, 1, 1),
                    block_dim: (BLK_DIM, 1, 1),
                    shared_mem_bytes: SMEM_BYTES,
                })
                .expect("launch kelly_bayesian_priors_update");
        }
        stream.synchronize().expect("sync after kelly_bayesian_priors_update");
    }

    /// Build a portfolio_state buffer of `n_envs * PS_STRIDE` entries with
    /// the four PS_KELLY_* fields populated per env from the parallel slices.
    /// All other PS_* fields are zero.
    fn build_portfolio_state(
        wins_per_env: &[f32],
        losses_per_env: &[f32],
        sum_wins_per_env: &[f32],
        sum_losses_per_env: &[f32],
    ) -> Vec<f32> {
        let n = wins_per_env.len();
        assert_eq!(losses_per_env.len(), n);
        assert_eq!(sum_wins_per_env.len(), n);
        assert_eq!(sum_losses_per_env.len(), n);
        let mut ps = vec![0.0_f32; n * PS_STRIDE];
        for e in 0..n {
            let row_off = e * PS_STRIDE;
            ps[row_off + PS_KELLY_WIN_COUNT]   = wins_per_env[e];
            ps[row_off + PS_KELLY_LOSS_COUNT]  = losses_per_env[e];
            ps[row_off + PS_KELLY_SUM_WINS]    = sum_wins_per_env[e];
            ps[row_off + PS_KELLY_SUM_LOSSES]  = sum_losses_per_env[e];
        }
        ps
    }

    /// Test 1 — Pearl-A first-observation bootstrap.
    ///
    /// 4 envs with realized stats. Total wins = 1+2+3+4 = 10, losses = 0+1+0+1 = 2,
    /// sum_wins = 0.05+0.10+0.15+0.20 = 0.50, sum_losses = 0+0.03+0+0.05 = 0.08.
    /// Pre-EMA clamp: counts to [0.5, 100], sums to [0.001, 1.0] —
    /// 10/2/0.5/0.08 all in range, pass through unchanged. Cold-start
    /// (sentinels 2.0/2.0/0.01/0.01): Pearl-A REPLACES → ISV =
    /// (10, 2, 0.5, 0.08).
    #[test]
    #[ignore = "requires GPU"]
    fn kelly_priors_pearl_a_bootstrap() {
        let stream = make_stream();
        let kernel = load_kernel(&stream);

        const ISV_DIM: usize = 1024;
        let mut isv = vec![0.0_f32; ISV_DIM];
        // Cold-start: sentinels (matches pre-P1-Producer constants).
        isv[KELLY_PRIOR_WINS_INDEX]       = SENTINEL_KELLY_PRIOR_WINS;
        isv[KELLY_PRIOR_LOSSES_INDEX]     = SENTINEL_KELLY_PRIOR_LOSSES;
        isv[KELLY_PRIOR_SUM_WINS_INDEX]   = SENTINEL_KELLY_PRIOR_SUM_WINS;
        isv[KELLY_PRIOR_SUM_LOSSES_INDEX] = SENTINEL_KELLY_PRIOR_SUM_LOSSES;

        let wins_per_env       = [1.0_f32, 2.0, 3.0, 4.0];
        let losses_per_env     = [0.0_f32, 1.0, 0.0, 1.0];
        let sum_wins_per_env   = [0.05_f32, 0.10, 0.15, 0.20];
        let sum_losses_per_env = [0.0_f32, 0.03, 0.0, 0.05];
        let n_envs = wins_per_env.len();

        let ps = build_portfolio_state(&wins_per_env, &losses_per_env,
                                        &sum_wins_per_env, &sum_losses_per_env);

        let ps_buf = unsafe { MappedF32Buffer::new(ps.len()) }.expect("alloc ps");
        ps_buf.write_from_slice(&ps);
        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
        isv_buf.write_from_slice(&isv);

        launch_kelly_priors(
            &stream, &kernel,
            ps_buf.dev_ptr, n_envs as i32, PS_STRIDE as i32, isv_buf.dev_ptr,
            KELLY_PRIOR_WINS_INDEX as i32,
            KELLY_PRIOR_LOSSES_INDEX as i32,
            KELLY_PRIOR_SUM_WINS_INDEX as i32,
            KELLY_PRIOR_SUM_LOSSES_INDEX as i32,
            SENTINEL_KELLY_PRIOR_WINS, SENTINEL_KELLY_PRIOR_LOSSES,
            SENTINEL_KELLY_PRIOR_SUM_WINS, SENTINEL_KELLY_PRIOR_SUM_LOSSES,
            KELLY_PRIOR_COUNT_MIN, KELLY_PRIOR_COUNT_MAX,
            KELLY_PRIOR_SUM_MIN,   KELLY_PRIOR_SUM_MAX,
            KELLY_PRIOR_EMA_ALPHA,
        );

        let result = isv_buf.read_all();
        let pw  = result[KELLY_PRIOR_WINS_INDEX];
        let pl  = result[KELLY_PRIOR_LOSSES_INDEX];
        let psw = result[KELLY_PRIOR_SUM_WINS_INDEX];
        let psl = result[KELLY_PRIOR_SUM_LOSSES_INDEX];

        // Pearl-A: sentinel → REPLACE with target. Targets are the
        // aggregated totals clamped to dimensional-safety bounds (all
        // within range here, no clamping applied).
        assert!((pw  - 10.0).abs() < 1e-4,
            "Pearl-A bootstrap: prior_wins expected 10.0, got {pw}");
        assert!((pl  -  2.0).abs() < 1e-4,
            "Pearl-A bootstrap: prior_losses expected 2.0, got {pl}");
        assert!((psw -  0.50).abs() < 1e-4,
            "Pearl-A bootstrap: prior_sum_wins expected 0.50, got {psw}");
        assert!((psl -  0.08).abs() < 1e-4,
            "Pearl-A bootstrap: prior_sum_losses expected 0.08, got {psl}");
    }

    /// Test 2 — No realized trades → ISV slots preserved bit-exactly.
    ///
    /// All envs have zero PS_KELLY_* fields → total_wins + total_losses == 0.
    /// Kernel guard skips the EMA update; both slots stay at the seeded
    /// values (NOT sentinels — to verify the guard preserves arbitrary
    /// pre-existing state).
    #[test]
    #[ignore = "requires GPU"]
    fn kelly_priors_no_realized_trades_preserves_isv() {
        let stream = make_stream();
        let kernel = load_kernel(&stream);

        const ISV_DIM: usize = 1024;
        // Seed with non-sentinel values to verify they survive the guard.
        const SEED_WINS:        f32 = 7.0;
        const SEED_LOSSES:      f32 = 5.0;
        const SEED_SUM_WINS:    f32 = 0.25;
        const SEED_SUM_LOSSES:  f32 = 0.15;

        let mut isv = vec![0.0_f32; ISV_DIM];
        isv[KELLY_PRIOR_WINS_INDEX]       = SEED_WINS;
        isv[KELLY_PRIOR_LOSSES_INDEX]     = SEED_LOSSES;
        isv[KELLY_PRIOR_SUM_WINS_INDEX]   = SEED_SUM_WINS;
        isv[KELLY_PRIOR_SUM_LOSSES_INDEX] = SEED_SUM_LOSSES;

        // 3 envs with no realized trades.
        let n_envs: usize = 3;
        let ps = vec![0.0_f32; n_envs * PS_STRIDE];

        let ps_buf = unsafe { MappedF32Buffer::new(ps.len()) }.expect("alloc ps");
        ps_buf.write_from_slice(&ps);
        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
        isv_buf.write_from_slice(&isv);

        launch_kelly_priors(
            &stream, &kernel,
            ps_buf.dev_ptr, n_envs as i32, PS_STRIDE as i32, isv_buf.dev_ptr,
            KELLY_PRIOR_WINS_INDEX as i32,
            KELLY_PRIOR_LOSSES_INDEX as i32,
            KELLY_PRIOR_SUM_WINS_INDEX as i32,
            KELLY_PRIOR_SUM_LOSSES_INDEX as i32,
            SENTINEL_KELLY_PRIOR_WINS, SENTINEL_KELLY_PRIOR_LOSSES,
            SENTINEL_KELLY_PRIOR_SUM_WINS, SENTINEL_KELLY_PRIOR_SUM_LOSSES,
            KELLY_PRIOR_COUNT_MIN, KELLY_PRIOR_COUNT_MAX,
            KELLY_PRIOR_SUM_MIN,   KELLY_PRIOR_SUM_MAX,
            KELLY_PRIOR_EMA_ALPHA,
        );

        let result = isv_buf.read_all();
        assert_eq!(result[KELLY_PRIOR_WINS_INDEX],       SEED_WINS,
            "prior_wins must be preserved bit-exactly when no realized trades");
        assert_eq!(result[KELLY_PRIOR_LOSSES_INDEX],     SEED_LOSSES,
            "prior_losses must be preserved bit-exactly when no realized trades");
        assert_eq!(result[KELLY_PRIOR_SUM_WINS_INDEX],   SEED_SUM_WINS,
            "prior_sum_wins must be preserved bit-exactly when no realized trades");
        assert_eq!(result[KELLY_PRIOR_SUM_LOSSES_INDEX], SEED_SUM_LOSSES,
            "prior_sum_losses must be preserved bit-exactly when no realized trades");
    }

    /// Test 3 — Bounds enforcement: extreme aggregate clamps to upper/lower bounds.
    ///
    /// One env with huge counts (1e6 wins, 1e6 losses) and tiny sums (1e-9).
    /// Pre-EMA clamp: wins/losses → 100 (count_max), sums → 0.001 (sum_min).
    /// Cold-start (sentinels): Pearl-A REPLACES with clamped targets.
    #[test]
    #[ignore = "requires GPU"]
    fn kelly_priors_bounds_clamp_extreme() {
        let stream = make_stream();
        let kernel = load_kernel(&stream);

        const ISV_DIM: usize = 1024;
        let mut isv = vec![0.0_f32; ISV_DIM];
        isv[KELLY_PRIOR_WINS_INDEX]       = SENTINEL_KELLY_PRIOR_WINS;
        isv[KELLY_PRIOR_LOSSES_INDEX]     = SENTINEL_KELLY_PRIOR_LOSSES;
        isv[KELLY_PRIOR_SUM_WINS_INDEX]   = SENTINEL_KELLY_PRIOR_SUM_WINS;
        isv[KELLY_PRIOR_SUM_LOSSES_INDEX] = SENTINEL_KELLY_PRIOR_SUM_LOSSES;

        // Single env with extreme aggregates.
        let wins_per_env       = [1.0e6_f32];
        let losses_per_env     = [1.0e6_f32];
        let sum_wins_per_env   = [1.0e-9_f32];
        let sum_losses_per_env = [1.0e-9_f32];

        let ps = build_portfolio_state(&wins_per_env, &losses_per_env,
                                        &sum_wins_per_env, &sum_losses_per_env);

        let ps_buf = unsafe { MappedF32Buffer::new(ps.len()) }.expect("alloc ps");
        ps_buf.write_from_slice(&ps);
        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
        isv_buf.write_from_slice(&isv);

        launch_kelly_priors(
            &stream, &kernel,
            ps_buf.dev_ptr, 1, PS_STRIDE as i32, isv_buf.dev_ptr,
            KELLY_PRIOR_WINS_INDEX as i32,
            KELLY_PRIOR_LOSSES_INDEX as i32,
            KELLY_PRIOR_SUM_WINS_INDEX as i32,
            KELLY_PRIOR_SUM_LOSSES_INDEX as i32,
            SENTINEL_KELLY_PRIOR_WINS, SENTINEL_KELLY_PRIOR_LOSSES,
            SENTINEL_KELLY_PRIOR_SUM_WINS, SENTINEL_KELLY_PRIOR_SUM_LOSSES,
            KELLY_PRIOR_COUNT_MIN, KELLY_PRIOR_COUNT_MAX,
            KELLY_PRIOR_SUM_MIN,   KELLY_PRIOR_SUM_MAX,
            KELLY_PRIOR_EMA_ALPHA,
        );

        let result = isv_buf.read_all();
        let pw  = result[KELLY_PRIOR_WINS_INDEX];
        let pl  = result[KELLY_PRIOR_LOSSES_INDEX];
        let psw = result[KELLY_PRIOR_SUM_WINS_INDEX];
        let psl = result[KELLY_PRIOR_SUM_LOSSES_INDEX];

        // 1e6 wins → clamped to count_max=100.
        assert!((pw  - KELLY_PRIOR_COUNT_MAX).abs() < 1e-4,
            "prior_wins must clamp to count_max=100 on extreme; got {pw}");
        assert!((pl  - KELLY_PRIOR_COUNT_MAX).abs() < 1e-4,
            "prior_losses must clamp to count_max=100 on extreme; got {pl}");
        // 1e-9 sums → clamped to sum_min=0.001.
        assert!((psw - KELLY_PRIOR_SUM_MIN).abs() < 1e-6,
            "prior_sum_wins must clamp to sum_min=0.001 on tiny; got {psw}");
        assert!((psl - KELLY_PRIOR_SUM_MIN).abs() < 1e-6,
            "prior_sum_losses must clamp to sum_min=0.001 on tiny; got {psl}");
    }

    /// Test 4 — Slow EMA blend after bootstrap (α=0.005).
    ///
    /// Pre-seed prior_wins=10.0 (NOT sentinel — bootstrap path NOT taken).
    /// Single env with 5 wins. Target = 5 (in range). EMA blend:
    /// (1 - 0.005) × 10 + 0.005 × 5 = 0.995 × 10 + 0.025 = 9.975.
    #[test]
    #[ignore = "requires GPU"]
    fn kelly_priors_slow_ema_after_bootstrap() {
        let stream = make_stream();
        let kernel = load_kernel(&stream);

        const ISV_DIM: usize = 1024;
        const SEEDED_WINS:       f32 = 10.0;
        // Other slots stay at sentinel — only test EMA on prior_wins.
        let mut isv = vec![0.0_f32; ISV_DIM];
        isv[KELLY_PRIOR_WINS_INDEX]       = SEEDED_WINS;
        isv[KELLY_PRIOR_LOSSES_INDEX]     = SENTINEL_KELLY_PRIOR_LOSSES;
        isv[KELLY_PRIOR_SUM_WINS_INDEX]   = SENTINEL_KELLY_PRIOR_SUM_WINS;
        isv[KELLY_PRIOR_SUM_LOSSES_INDEX] = SENTINEL_KELLY_PRIOR_SUM_LOSSES;

        // Single env with 5 wins.
        let wins_per_env       = [5.0_f32];
        let losses_per_env     = [1.0_f32];
        let sum_wins_per_env   = [0.10_f32];
        let sum_losses_per_env = [0.02_f32];

        let ps = build_portfolio_state(&wins_per_env, &losses_per_env,
                                        &sum_wins_per_env, &sum_losses_per_env);

        let ps_buf = unsafe { MappedF32Buffer::new(ps.len()) }.expect("alloc ps");
        ps_buf.write_from_slice(&ps);
        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
        isv_buf.write_from_slice(&isv);

        launch_kelly_priors(
            &stream, &kernel,
            ps_buf.dev_ptr, 1, PS_STRIDE as i32, isv_buf.dev_ptr,
            KELLY_PRIOR_WINS_INDEX as i32,
            KELLY_PRIOR_LOSSES_INDEX as i32,
            KELLY_PRIOR_SUM_WINS_INDEX as i32,
            KELLY_PRIOR_SUM_LOSSES_INDEX as i32,
            SENTINEL_KELLY_PRIOR_WINS, SENTINEL_KELLY_PRIOR_LOSSES,
            SENTINEL_KELLY_PRIOR_SUM_WINS, SENTINEL_KELLY_PRIOR_SUM_LOSSES,
            KELLY_PRIOR_COUNT_MIN, KELLY_PRIOR_COUNT_MAX,
            KELLY_PRIOR_SUM_MIN,   KELLY_PRIOR_SUM_MAX,
            KELLY_PRIOR_EMA_ALPHA,
        );

        let result = isv_buf.read_all();
        let pw = result[KELLY_PRIOR_WINS_INDEX];
        // EMA: 0.995 × 10 + 0.005 × 5 = 9.975.
        let expected = (1.0_f32 - KELLY_PRIOR_EMA_ALPHA) * SEEDED_WINS
                     + KELLY_PRIOR_EMA_ALPHA * 5.0_f32;
        assert!(
            (pw - expected).abs() < 1e-4,
            "Slow EMA blend (α={KELLY_PRIOR_EMA_ALPHA}): expected {expected}, got {pw}"
        );
    }
}

// ═══════════════════════════════════════════════════════════════════════════
// Test B.8: Layout-fingerprint regression for direction Q-head input bump.
//
// CPU-only test (no GPU required). Ensures the fingerprint hash changed when
// `w_b0fc`'s input dim grew from `SH2` to `SH2+1`, per Invariant 8 of the
// SP14 plan: any structural reshuffle MUST bump
// `LAYOUT_FINGERPRINT_CURRENT` so old checkpoints fail-fast on load instead
// of silently mapping into the new architecture. The bump is achieved by
// renaming the seed entry `PARAM_W_B0FC` → `PARAM_W_B0FC_AUX1`. Per
// `feedback_no_legacy_aliases`, no `_DEPRECATED` shim is added.
// ═══════════════════════════════════════════════════════════════════════════

/// Inline FNV-1a 64-bit hash for the regression test. Mirrors the
/// `const fn fnv1a_64` in `gpu_dqn_trainer.rs` so the test computes
/// the same hash the production fingerprint uses. (Mirrors the
/// equivalent helper in `sp13_layer_b_oracle_tests.rs`.)
const fn fnv1a_64(bytes: &[u8]) -> u64 {
    const OFFSET: u64 = 0xcbf29ce484222325;
    const PRIME:  u64 = 0x00000100000001b3;
    let mut h: u64 = OFFSET;
    let mut i = 0;
    while i < bytes.len() {
        h ^= bytes[i] as u64;
        h = h.wrapping_mul(PRIME);
        i += 1;
    }
    h
}

/// Pre-B.8 seed — hashing this constant gives the pre-B.8 fingerprint
/// that the bump must differ from. The test asserts the current
/// fingerprint is distinct from this baseline; SP14 Layer C Phase C.1
/// (2026-05-08) further shifted the fingerprint by deleting α-machinery
/// slot entries and adding aux-trunk slots, so the inequality holds even
/// more strongly post-C.1.
const PRE_B8_SEED: &[u8] = b"PRE_B8_SEED_PLACEHOLDER_LEGACY_PARAM_W_B0FC";

#[test]
fn layout_fingerprint_bumps_after_sp14_wire() {
    use ml::cuda_pipeline::gpu_dqn_trainer::LAYOUT_FINGERPRINT_CURRENT;
    let pre_b8_fp = fnv1a_64(PRE_B8_SEED);
    assert_ne!(
        LAYOUT_FINGERPRINT_CURRENT, pre_b8_fp,
        "B.8 fingerprint must differ from pre-B.8 placeholder value. \
         current = {:#018x}, pre-B.8 placeholder = {:#018x}",
        LAYOUT_FINGERPRINT_CURRENT, pre_b8_fp
    );
}

// ═══════════════════════════════════════════════════════════════════════════
// Class A audit-fix Batch 4-A (2026-05-08) — adaptive DD saturation floor
// producer tests.
//
// Verifies the `dd_saturation_floor_update_kernel` producer:
//   1. Cold-start Pearl-A bootstrap: ISV stays at sentinel until first valid
//      observation, then REPLACES (no blend).
//   2. No-DD guard: when all envs have DD_MAX=0, slot is preserved bit-exactly.
//   3. Bounds enforced: floor in [0.10, 0.50].
//   4. Welford slow EMA after bootstrap (α=0.01).
//
// All tests are #[ignore = "requires GPU"]; gated under #[cfg(feature = "cuda")].
// ═══════════════════════════════════════════════════════════════════════════

#[cfg(feature = "cuda")]
#[allow(unsafe_code)] // CUDA kernel launch + mapped-pinned memory.
mod sp14_audit_4a_dd_saturation_floor_gpu {
    use std::sync::Arc;

    use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg};
    use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer;
    use ml::cuda_pipeline::sp14_isv_slots::{
        DD_SATURATION_FLOOR_ADAPTIVE_INDEX, DD_SATURATION_FLOOR_EMA_ALPHA,
        DD_SATURATION_FLOOR_MAX, DD_SATURATION_FLOOR_MIN,
        DD_SATURATION_FLOOR_SAFETY_FACTOR, SENTINEL_DD_SATURATION_FLOOR,
    };

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

    /// Per-env DD state tile width (mirrors `dd_state_kernel` layout in
    /// `dd_state_kernel.cu`: DD_CURRENT, DD_MAX, DD_RECOVERY_BARS,
    /// DD_PERSISTENCE, CALMAR, DD_PCT).
    const DD_STATE_STRIDE: i32 = 6;
    /// DD_MAX offset within each env's tile.
    const DD_MAX_OFF: i32 = 1;

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

    fn load_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
        let module = stream
            .context()
            .load_cubin(DD_SATURATION_FLOOR_UPDATE_CUBIN.to_vec())
            .expect("load dd_saturation_floor_update_kernel cubin");
        module
            .load_function("dd_saturation_floor_update")
            .expect("load dd_saturation_floor_update function")
    }

    /// 4 arrays × 256 × 4 bytes = 4096 bytes shmem.
    const BLK_DIM: u32 = 256;
    const SMEM_BYTES: u32 = 4 * BLK_DIM * std::mem::size_of::<f32>() as u32;

    /// Helper: launch with a fixed config.
    #[allow(clippy::too_many_arguments)]
    fn launch_dd_floor(
        stream: &Arc<CudaStream>,
        kernel: &CudaFunction,
        dd_state_ptr: u64,
        n_envs: i32,
        dd_state_stride: i32,
        dd_max_off: i32,
        isv_ptr: u64,
        floor_idx: i32,
        sentinel_floor: f32,
        safety_factor: f32,
        floor_min: f32,
        floor_max: f32,
        alpha: f32,
    ) {
        unsafe {
            stream
                .launch_builder(kernel)
                .arg(&dd_state_ptr)
                .arg(&n_envs)
                .arg(&dd_state_stride)
                .arg(&dd_max_off)
                .arg(&isv_ptr)
                .arg(&floor_idx)
                .arg(&sentinel_floor)
                .arg(&safety_factor)
                .arg(&floor_min)
                .arg(&floor_max)
                .arg(&alpha)
                .launch(LaunchConfig {
                    grid_dim: (1, 1, 1),
                    block_dim: (BLK_DIM, 1, 1),
                    shared_mem_bytes: SMEM_BYTES,
                })
                .expect("launch dd_saturation_floor_update");
        }
        stream.synchronize().expect("sync after dd_saturation_floor_update");
    }

    /// Build a per-env DD state tile [n_envs * 6] with the given DD_MAX
    /// values at offset 1; other offsets zeroed.
    fn build_dd_state_tile(dd_max_per_env: &[f32]) -> Vec<f32> {
        let n_envs = dd_max_per_env.len();
        let mut tile = vec![0.0_f32; n_envs * DD_STATE_STRIDE as usize];
        for (env, dd_max) in dd_max_per_env.iter().enumerate() {
            tile[env * DD_STATE_STRIDE as usize + DD_MAX_OFF as usize] = *dd_max;
        }
        tile
    }

    /// Test 1 — Pearl-A first-observation bootstrap.
    ///
    /// 4 envs with DD_MAX=[0.10, 0.15, 0.20, 0.25]. mean=0.175, var=0.003125,
    /// sigma≈0.0559. p75 ≈ 0.175 + 0.6745 × 0.0559 ≈ 0.213. Robust target
    /// = max(0.213, 0.175) = 0.213. × safety_factor=1.5 = 0.3194. Clamped to
    /// [0.10, 0.50] — stays at 0.3194.
    ///
    /// Cold-start (current = sentinel 0.25): Pearl-A REPLACES → 0.3194.
    #[test]
    #[ignore = "requires GPU"]
    fn dd_floor_pearl_a_bootstrap() {
        let stream = make_stream();
        let kernel = load_kernel(&stream);

        const ISV_DIM: usize = 1024;
        let mut isv = vec![0.0_f32; ISV_DIM];
        // Cold-start: sentinel matches pre-fix hardcoded value.
        isv[DD_SATURATION_FLOOR_ADAPTIVE_INDEX] = SENTINEL_DD_SATURATION_FLOOR;

        let dd_max: [f32; 4] = [0.10, 0.15, 0.20, 0.25];
        let n_envs = dd_max.len() as i32;
        let tile = build_dd_state_tile(&dd_max);

        let tile_buf = unsafe { MappedF32Buffer::new(tile.len()) }.expect("alloc tile");
        tile_buf.write_from_slice(&tile);
        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
        isv_buf.write_from_slice(&isv);

        launch_dd_floor(
            &stream, &kernel,
            tile_buf.dev_ptr, n_envs, DD_STATE_STRIDE, DD_MAX_OFF,
            isv_buf.dev_ptr, DD_SATURATION_FLOOR_ADAPTIVE_INDEX as i32,
            SENTINEL_DD_SATURATION_FLOOR, DD_SATURATION_FLOOR_SAFETY_FACTOR,
            DD_SATURATION_FLOOR_MIN, DD_SATURATION_FLOOR_MAX,
            DD_SATURATION_FLOOR_EMA_ALPHA,
        );

        let result = isv_buf.read_all();
        let floor = result[DD_SATURATION_FLOOR_ADAPTIVE_INDEX];

        // Pearl-A: sentinel → REPLACE with target. Target = max(p75, mean) × 1.5.
        // mean=0.175, var=(0.01+0.0225+0.04+0.0625)/4 - 0.175² = 0.0337/4-...
        //   sumsq = 0.01+0.0225+0.04+0.0625 = 0.135, sumsq/n = 0.03375
        //   variance = 0.03375 - 0.030625 = 0.003125, sigma ≈ 0.0559
        // p75 ≈ 0.175 + 0.6745 × 0.0559 ≈ 0.2127. max(p75, mean)=0.2127.
        // × 1.5 = 0.319. Clamped to [0.10, 0.50] → stays at ~0.319.
        assert!(
            (floor - 0.319).abs() < 0.01,
            "Pearl-A bootstrap: expected floor≈0.319, got {floor}"
        );
    }

    /// Test 2 — No DD observed → ISV slot preserved bit-exactly.
    ///
    /// All envs have DD_MAX=0 (no DD observed yet). Kernel must skip the
    /// EMA update (count==0 guard) and leave the slot at its seeded value.
    #[test]
    #[ignore = "requires GPU"]
    fn dd_floor_no_dd_preserves_isv() {
        let stream = make_stream();
        let kernel = load_kernel(&stream);

        const ISV_DIM: usize = 1024;
        const SEED: f32 = 0.37;
        let mut isv = vec![0.0_f32; ISV_DIM];
        isv[DD_SATURATION_FLOOR_ADAPTIVE_INDEX] = SEED;

        // All envs at DD_MAX=0.0 — no DD observed.
        let dd_max: [f32; 8] = [0.0; 8];
        let n_envs = dd_max.len() as i32;
        let tile = build_dd_state_tile(&dd_max);

        let tile_buf = unsafe { MappedF32Buffer::new(tile.len()) }.expect("alloc tile");
        tile_buf.write_from_slice(&tile);
        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
        isv_buf.write_from_slice(&isv);

        launch_dd_floor(
            &stream, &kernel,
            tile_buf.dev_ptr, n_envs, DD_STATE_STRIDE, DD_MAX_OFF,
            isv_buf.dev_ptr, DD_SATURATION_FLOOR_ADAPTIVE_INDEX as i32,
            SENTINEL_DD_SATURATION_FLOOR, DD_SATURATION_FLOOR_SAFETY_FACTOR,
            DD_SATURATION_FLOOR_MIN, DD_SATURATION_FLOOR_MAX,
            DD_SATURATION_FLOOR_EMA_ALPHA,
        );

        let result = isv_buf.read_all();
        assert_eq!(
            result[DD_SATURATION_FLOOR_ADAPTIVE_INDEX], SEED,
            "DD floor slot must be preserved bit-exactly when no DD observed",
        );
    }

    /// Test 3 — Bounds enforcement: huge DD_MAX clamps floor to MAX=0.50.
    ///
    /// 2 envs with DD_MAX=[0.80, 0.90] — both well above the saturation
    /// floor MAX. Target = max(p75, mean) × 1.5 ≈ 0.85 × 1.5 = 1.275 → clamps
    /// to 0.50.
    #[test]
    #[ignore = "requires GPU"]
    fn dd_floor_bounds_clamp_extreme_dd() {
        let stream = make_stream();
        let kernel = load_kernel(&stream);

        const ISV_DIM: usize = 1024;
        let mut isv = vec![0.0_f32; ISV_DIM];
        isv[DD_SATURATION_FLOOR_ADAPTIVE_INDEX] = SENTINEL_DD_SATURATION_FLOOR;

        let dd_max: [f32; 2] = [0.80, 0.90];
        let n_envs = dd_max.len() as i32;
        let tile = build_dd_state_tile(&dd_max);

        let tile_buf = unsafe { MappedF32Buffer::new(tile.len()) }.expect("alloc tile");
        tile_buf.write_from_slice(&tile);
        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
        isv_buf.write_from_slice(&isv);

        launch_dd_floor(
            &stream, &kernel,
            tile_buf.dev_ptr, n_envs, DD_STATE_STRIDE, DD_MAX_OFF,
            isv_buf.dev_ptr, DD_SATURATION_FLOOR_ADAPTIVE_INDEX as i32,
            SENTINEL_DD_SATURATION_FLOOR, DD_SATURATION_FLOOR_SAFETY_FACTOR,
            DD_SATURATION_FLOOR_MIN, DD_SATURATION_FLOOR_MAX,
            DD_SATURATION_FLOOR_EMA_ALPHA,
        );

        let result = isv_buf.read_all();
        let floor = result[DD_SATURATION_FLOOR_ADAPTIVE_INDEX];

        // Pearl-A: sentinel → REPLACE with clamped target. Should hit MAX=0.50.
        assert!(
            (floor - DD_SATURATION_FLOOR_MAX).abs() < 1e-5,
            "DD floor must clamp to MAX=0.50 on extreme DD; got {floor}"
        );
    }

    /// Test 4 — Welford slow EMA after bootstrap.
    ///
    /// Seed slot at 0.30 (NOT sentinel — past the bootstrap). Run with
    /// 4 envs at DD_MAX=[0.10, 0.10, 0.10, 0.10] (mean=0.10, sigma=0).
    /// Robust target = max(0.10+0×Z_75, 0.10) × 1.5 = 0.15. Clamped to
    /// [0.10, 0.50] → stays at 0.15.
    /// EMA blend: blended = (1-0.01)×0.30 + 0.01×0.15 = 0.297 + 0.0015 = 0.2985.
    #[test]
    #[ignore = "requires GPU"]
    fn dd_floor_welford_ema_after_bootstrap() {
        let stream = make_stream();
        let kernel = load_kernel(&stream);

        const ISV_DIM: usize = 1024;
        const SEED: f32 = 0.30;  // NOT sentinel — Pearl-A guard fires α-blend.
        let mut isv = vec![0.0_f32; ISV_DIM];
        isv[DD_SATURATION_FLOOR_ADAPTIVE_INDEX] = SEED;

        let dd_max: [f32; 4] = [0.10, 0.10, 0.10, 0.10];
        let n_envs = dd_max.len() as i32;
        let tile = build_dd_state_tile(&dd_max);

        let tile_buf = unsafe { MappedF32Buffer::new(tile.len()) }.expect("alloc tile");
        tile_buf.write_from_slice(&tile);
        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
        isv_buf.write_from_slice(&isv);

        launch_dd_floor(
            &stream, &kernel,
            tile_buf.dev_ptr, n_envs, DD_STATE_STRIDE, DD_MAX_OFF,
            isv_buf.dev_ptr, DD_SATURATION_FLOOR_ADAPTIVE_INDEX as i32,
            SENTINEL_DD_SATURATION_FLOOR, DD_SATURATION_FLOOR_SAFETY_FACTOR,
            DD_SATURATION_FLOOR_MIN, DD_SATURATION_FLOOR_MAX,
            DD_SATURATION_FLOOR_EMA_ALPHA,
        );

        let result = isv_buf.read_all();
        let floor = result[DD_SATURATION_FLOOR_ADAPTIVE_INDEX];

        // EMA blend: (1-0.01) × 0.30 + 0.01 × 0.15 = 0.2985.
        let expected = (1.0_f32 - DD_SATURATION_FLOOR_EMA_ALPHA) * SEED
                     + DD_SATURATION_FLOOR_EMA_ALPHA * 0.15;
        assert!(
            (floor - expected).abs() < 1e-4,
            "Welford EMA blend: expected ≈ {expected:.5}, got {floor:.5}"
        );
    }
}

// ═══════════════════════════════════════════════════════════════════════════
// Class A audit-fix Batch 4-B (2026-05-08, Item 3) — adaptive plan_threshold
// floor inline producer (Design Y inside `plan_threshold_update_kernel`).
//
// Verifies:
//   1. Pearl-A bootstrap: sentinel 0.1 → REPLACE with `0.5 × ema`.
//   2. Slow EMA (α=0.005) blend after bootstrap.
//   3. Bilateral clamp: huge readiness clamps floor to MAX=0.50.
//
// All tests are #[ignore = "requires GPU"]; gated under #[cfg(feature = "cuda")].
// ═══════════════════════════════════════════════════════════════════════════

#[cfg(feature = "cuda")]
#[allow(unsafe_code)] // CUDA kernel launch + mapped-pinned memory.
mod sp14_audit_4b_plan_threshold_floor_gpu {
    use std::sync::Arc;

    use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg};
    use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer;
    use ml::cuda_pipeline::sp14_isv_slots::{
        PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX, PLAN_THRESHOLD_FLOOR_EMA_ALPHA,
        PLAN_THRESHOLD_FLOOR_MAX, PLAN_THRESHOLD_FLOOR_MIN, SENTINEL_PLAN_THRESHOLD_FLOOR,
    };

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

    /// ISV slot indices used by the kernel (mirror gpu_dqn_trainer.rs).
    const READINESS_EMA_INDEX: usize = 75;
    const PLAN_THRESHOLD_INDEX: usize = 49;
    const SHARPE_EMA_INDEX: usize = 22;
    const ALPHA_BASE: f32 = 0.05;

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

    fn load_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
        let module = stream
            .context()
            .load_cubin(PLAN_THRESHOLD_UPDATE_CUBIN.to_vec())
            .expect("load plan_threshold_update_kernel cubin");
        module
            .load_function("plan_threshold_update")
            .expect("load plan_threshold_update function")
    }

    fn launch_plan_threshold(
        stream: &Arc<CudaStream>,
        kernel: &CudaFunction,
        readiness_ptr: u64,
        n_samples: i32,
        isv_ptr: u64,
    ) {
        let readiness_slot = READINESS_EMA_INDEX as i32;
        let thr_slot       = PLAN_THRESHOLD_INDEX as i32;
        let sharpe_slot    = SHARPE_EMA_INDEX as i32;
        let alpha_base     = ALPHA_BASE;
        unsafe {
            stream
                .launch_builder(kernel)
                .arg(&readiness_ptr)
                .arg(&n_samples)
                .arg(&isv_ptr)
                .arg(&readiness_slot)
                .arg(&thr_slot)
                .arg(&sharpe_slot)
                .arg(&alpha_base)
                .launch(LaunchConfig {
                    grid_dim: (1, 1, 1),
                    block_dim: (1, 1, 1),
                    shared_mem_bytes: 0,
                })
                .expect("launch plan_threshold_update");
        }
        stream.synchronize().expect("sync after plan_threshold_update");
    }

    /// Test 1 — Pearl-A first-observation bootstrap.
    ///
    /// Seed `readiness_ema` at 0.4 and `floor` at sentinel 0.1; readiness
    /// samples uniform 0.5 → ema converges toward 0.5. The new
    /// readiness_ema after one launch with 4 samples of 0.5 and α=0.05
    /// (sharpe=0): ema = 0.95 × 0.4 + 0.05 × 0.5 = 0.405. Threshold
    /// target = 0.5 × 0.405 = 0.2025. Pearl-A: floor at sentinel 0.1
    /// → REPLACE with 0.2025 directly. Bilateral clamp [0.05, 0.50] →
    /// stays at 0.2025.
    #[test]
    #[ignore = "requires GPU"]
    fn plan_threshold_floor_pearl_a_bootstrap() {
        let stream = make_stream();
        let kernel = load_kernel(&stream);

        const ISV_DIM: usize = 1024;
        let mut isv = vec![0.0_f32; ISV_DIM];
        // Seed readiness_ema at 0.4 (NOT sentinel — past readiness bootstrap).
        isv[READINESS_EMA_INDEX] = 0.4;
        // Seed floor at sentinel 0.1.
        isv[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX] = SENTINEL_PLAN_THRESHOLD_FLOOR;
        // Sharpe = 0 → adaptive α stays at α_base.
        isv[SHARPE_EMA_INDEX] = 0.0;

        let readiness: [f32; 4] = [0.5, 0.5, 0.5, 0.5];
        let n_samples = readiness.len() as i32;

        let r_buf = unsafe { MappedF32Buffer::new(readiness.len()) }.expect("alloc readiness");
        r_buf.write_from_slice(&readiness);
        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
        isv_buf.write_from_slice(&isv);

        launch_plan_threshold(&stream, &kernel, r_buf.dev_ptr, n_samples, isv_buf.dev_ptr);

        let result = isv_buf.read_all();
        let floor = result[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX];

        // ema = 0.95 × 0.4 + 0.05 × 0.5 = 0.405; target = 0.5 × 0.405 = 0.2025.
        let expected_ema  = (1.0_f32 - ALPHA_BASE) * 0.4 + ALPHA_BASE * 0.5;
        let expected_target = 0.5 * expected_ema;
        assert!(
            (floor - expected_target).abs() < 1e-4,
            "Pearl-A bootstrap: expected floor ≈ {expected_target:.5}, got {floor:.5}"
        );
    }

    /// Test 2 — Slow EMA after bootstrap.
    ///
    /// Seed floor at 0.30 (NOT sentinel), readiness_ema at 0.20.
    /// Readiness samples uniform 0.20 → ema stays at 0.20. Threshold
    /// target = 0.5 × 0.20 = 0.10. Floor blend:
    ///   blended = (1 − 0.005) × 0.30 + 0.005 × 0.10 = 0.299.
    #[test]
    #[ignore = "requires GPU"]
    fn plan_threshold_floor_slow_ema_after_bootstrap() {
        let stream = make_stream();
        let kernel = load_kernel(&stream);

        const ISV_DIM: usize = 1024;
        const SEED_FLOOR: f32 = 0.30;
        let mut isv = vec![0.0_f32; ISV_DIM];
        // Seed readiness_ema at 0.20 (past bootstrap).
        isv[READINESS_EMA_INDEX] = 0.20;
        // Seed floor at 0.30 — Pearl-A guard fires α-blend (NOT sentinel).
        isv[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX] = SEED_FLOOR;
        isv[SHARPE_EMA_INDEX] = 0.0;

        let readiness: [f32; 4] = [0.20, 0.20, 0.20, 0.20];
        let n_samples = readiness.len() as i32;

        let r_buf = unsafe { MappedF32Buffer::new(readiness.len()) }.expect("alloc readiness");
        r_buf.write_from_slice(&readiness);
        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
        isv_buf.write_from_slice(&isv);

        launch_plan_threshold(&stream, &kernel, r_buf.dev_ptr, n_samples, isv_buf.dev_ptr);

        let result = isv_buf.read_all();
        let floor = result[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX];

        // ema = 0.95 × 0.20 + 0.05 × 0.20 = 0.20 (unchanged).
        // target = 0.5 × 0.20 = 0.10.
        // blended = (1 − 0.005) × 0.30 + 0.005 × 0.10 = 0.2990.
        let new_ema  = (1.0_f32 - ALPHA_BASE) * 0.20 + ALPHA_BASE * 0.20;
        let target   = 0.5 * new_ema;
        let expected = (1.0_f32 - PLAN_THRESHOLD_FLOOR_EMA_ALPHA) * SEED_FLOOR
                     +         PLAN_THRESHOLD_FLOOR_EMA_ALPHA  * target;
        assert!(
            (floor - expected).abs() < 1e-4,
            "Slow EMA blend: expected ≈ {expected:.5}, got {floor:.5}"
        );
    }

    /// Test 3 — Bilateral clamp on extreme readiness.
    ///
    /// Seed readiness_ema at 1.5 (out-of-spec but defends against
    /// malformed prior state). After one launch with readiness=1.5,
    /// new ema ≈ 1.5; target = 0.75 → clamps to MAX=0.50. Pearl-A:
    /// floor was at sentinel 0.1 → REPLACES with clamped 0.50.
    #[test]
    #[ignore = "requires GPU"]
    fn plan_threshold_floor_bounds_clamp_extreme_readiness() {
        let stream = make_stream();
        let kernel = load_kernel(&stream);

        const ISV_DIM: usize = 1024;
        let mut isv = vec![0.0_f32; ISV_DIM];
        isv[READINESS_EMA_INDEX] = 1.5;  // out-of-spec to force clamp
        isv[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX] = SENTINEL_PLAN_THRESHOLD_FLOOR;
        isv[SHARPE_EMA_INDEX] = 0.0;

        let readiness: [f32; 4] = [1.5, 1.5, 1.5, 1.5];
        let n_samples = readiness.len() as i32;

        let r_buf = unsafe { MappedF32Buffer::new(readiness.len()) }.expect("alloc readiness");
        r_buf.write_from_slice(&readiness);
        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
        isv_buf.write_from_slice(&isv);

        launch_plan_threshold(&stream, &kernel, r_buf.dev_ptr, n_samples, isv_buf.dev_ptr);

        let result = isv_buf.read_all();
        let floor = result[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX];

        assert!(
            (floor - PLAN_THRESHOLD_FLOOR_MAX).abs() < 1e-5,
            "Bounds clamp: floor must clamp to MAX=0.50 on extreme readiness; got {floor}"
        );
        // Defensive: also exercises the MIN bound — floor never below MIN
        // even if computed value is below it.
        assert!(
            floor >= PLAN_THRESHOLD_FLOOR_MIN,
            "Floor must respect MIN bound; got {floor}"
        );
    }
}

// ═══════════════════════════════════════════════════════════════════════════
// SP16 Phase 1 (revised, 2026-05-08) — adaptive MIN_HOLD_TEMPERATURE
// producer (`min_hold_temperature_update_kernel`).
//
// Driving signal swapped from `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]`
// (slot 373) to hold-rate overrun reading
// `ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382]` minus
// `ISV[HOLD_RATE_TARGET_INDEX=381]`. Per train-multi-seed-pfh9n
// post-mortem: slot 373 stayed at sentinel 0.5 in Fold 1 and the
// kernel's early-return-on-sentinel guard pinned slot 460 at 50 for
// the entire fold. Hold rates are measured per-epoch from realised
// actions and survive fold reset cleanly.
//
// Verifies:
//   1. Pearl-A bootstrap: sentinel 50.0 → REPLACE with target_temp.
//   2. observed=sentinel 0.0 → ISV preserved bit-exactly (no producer
//      update; consumer falls back to scalar).
//   3. Over-holding (overrun=1.0) → temp HIGH (permissive exit ramp).
//   4. At-target (no overrun) → temp LOW (strict).
//   5. Under-target → temp LOW (strict — no relaxation needed).
//   6. Mid-cadence EMA blend after bootstrap (α=0.05).
//   7. Slot 373 dependency removed — kernel responds even when slot 373
//      is at sentinel 0.5 (proves the swap is real).
//
// All tests are #[ignore = "requires GPU"]; gated under #[cfg(feature = "cuda")].
// ═══════════════════════════════════════════════════════════════════════════

#[cfg(feature = "cuda")]
#[allow(unsafe_code)] // CUDA kernel launch + mapped-pinned memory.
mod sp16_phase1_min_hold_temperature_gpu {
    use std::sync::Arc;

    use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg};
    use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer;
    use ml::cuda_pipeline::sp14_isv_slots::{
        MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX,
        MIN_HOLD_TEMPERATURE_MAX,
        MIN_HOLD_TEMPERATURE_MIN, SENTINEL_HOLD_RATE_OBSERVED,
        SENTINEL_MIN_HOLD_TEMPERATURE,
        MHT_TARGET_MEAN_INDEX, MHT_TARGET_M2_INDEX,
        MHT_DIFF_MEAN_INDEX, MHT_DIFF_M2_INDEX,
        MHT_PREV_TARGET_INDEX, MHT_SAMPLE_COUNT_INDEX,
    };
    use ml::cuda_pipeline::sp13_isv_slots::{
        AUX_DIR_ACC_SHORT_EMA_INDEX, HOLD_RATE_OBSERVED_EMA_INDEX,
        HOLD_RATE_TARGET_INDEX,
    };

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

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

    fn load_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
        let module = stream
            .context()
            .load_cubin(MIN_HOLD_TEMPERATURE_UPDATE_CUBIN.to_vec())
            .expect("load min_hold_temperature_update_kernel cubin");
        module
            .load_function("min_hold_temperature_update")
            .expect("load min_hold_temperature_update function")
    }

    /// SP16 T3 (2026-05-08): kernel takes Welford accumulator slot indices
    /// instead of a hardcoded α. The hardcoded α=0.05 is gone — α is
    /// computed inside the kernel from the Welford state per
    /// `pearl_wiener_optimal_adaptive_alpha`.
    #[allow(clippy::too_many_arguments)]
    fn launch_temp(
        stream: &Arc<CudaStream>,
        kernel: &CudaFunction,
        isv_ptr: u64,
        temp_idx: i32,
        hold_rate_observed_idx: i32,
        hold_rate_target_idx: i32,
        target_mean_idx: i32,
        target_m2_idx: i32,
        diff_mean_idx: i32,
        diff_m2_idx: i32,
        prev_target_idx: i32,
        sample_count_idx: i32,
        sentinel_temp: f32,
        sentinel_observed_hold_rate: f32,
        temp_min: f32,
        temp_max: f32,
    ) {
        unsafe {
            stream
                .launch_builder(kernel)
                .arg(&isv_ptr)
                .arg(&temp_idx)
                .arg(&hold_rate_observed_idx)
                .arg(&hold_rate_target_idx)
                .arg(&target_mean_idx)
                .arg(&target_m2_idx)
                .arg(&diff_mean_idx)
                .arg(&diff_m2_idx)
                .arg(&prev_target_idx)
                .arg(&sample_count_idx)
                .arg(&sentinel_temp)
                .arg(&sentinel_observed_hold_rate)
                .arg(&temp_min)
                .arg(&temp_max)
                .launch(LaunchConfig {
                    grid_dim: (1, 1, 1),
                    block_dim: (1, 1, 1),
                    shared_mem_bytes: 0,
                })
                .expect("launch min_hold_temperature_update");
        }
        stream.synchronize().expect("sync after min_hold_temperature_update");
    }

    /// Standard launch helper: takes the canonical SP14/SP13 constants,
    /// avoiding the verbose argument list at every test.
    fn launch_canonical(
        stream: &Arc<CudaStream>,
        kernel: &CudaFunction,
        isv_ptr: u64,
    ) {
        launch_temp(
            stream, kernel, isv_ptr,
            MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX as i32,
            HOLD_RATE_OBSERVED_EMA_INDEX as i32,
            HOLD_RATE_TARGET_INDEX as i32,
            MHT_TARGET_MEAN_INDEX as i32,
            MHT_TARGET_M2_INDEX as i32,
            MHT_DIFF_MEAN_INDEX as i32,
            MHT_DIFF_M2_INDEX as i32,
            MHT_PREV_TARGET_INDEX as i32,
            MHT_SAMPLE_COUNT_INDEX as i32,
            SENTINEL_MIN_HOLD_TEMPERATURE,
            SENTINEL_HOLD_RATE_OBSERVED,
            MIN_HOLD_TEMPERATURE_MIN,
            MIN_HOLD_TEMPERATURE_MAX,
        );
    }

    /// SP16 Phase 1 (revised) — Test 1: MIN_HOLD_TEMPERATURE producer
    /// responds to hold-rate overrun, NOT slot 373.
    ///
    /// Setup: observed_hold_rate=0.40, target_hold_rate=0.20 (100%
    /// overrun → overrun_norm=1.0). Pearl-A bootstrap from sentinel
    /// 50.0 → temp REPLACED directly with target_temp = TEMP_MAX = 50.
    /// Verifies the kernel climbs toward TEMP_MAX in the canonical
    /// over-holding regime.
    #[test]
    #[ignore = "requires GPU"]
    fn sp16_phase1_min_hold_temp_climbs_with_hold_overrun() {
        let stream = make_stream();
        let kernel = load_kernel(&stream);

        const ISV_DIM: usize = 1024;
        let mut isv = vec![0.0_f32; ISV_DIM];
        // 100% overrun: observed = 2× target.
        isv[HOLD_RATE_OBSERVED_EMA_INDEX] = 0.40;
        isv[HOLD_RATE_TARGET_INDEX]       = 0.20;
        // Temp at sentinel 50.0 → Pearl-A REPLACE with target_temp.
        isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SENTINEL_MIN_HOLD_TEMPERATURE;

        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
        isv_buf.write_from_slice(&isv);

        launch_canonical(&stream, &kernel, isv_buf.dev_ptr);

        let result = isv_buf.read_all();
        let temp = result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX];

        // overrun = max(0, 0.40 − 0.20) = 0.20
        // overrun_norm = clamp(0.20 / max(0.20, 0.01), 0, 1) = 1.0
        // target_temp = 5 + 45 × 1.0 = 50.0
        // Pearl-A bootstrap: sentinel → REPLACE → blended = 50.0.
        assert!(
            (temp - MIN_HOLD_TEMPERATURE_MAX).abs() < 1e-4,
            "Over-holding (observed=2×target) must drive temp HIGH (permissive exit ramp); \
             expected {}, got {temp}",
            MIN_HOLD_TEMPERATURE_MAX
        );
    }

    /// SP16 Phase 1 (revised) — Test 2: at-target → temp LOW (strict).
    ///
    /// Setup: observed_hold_rate=target_hold_rate=0.20 (no overrun).
    /// Pearl-A bootstrap from sentinel 50.0 → REPLACE with target_temp
    /// = TEMP_MIN = 5. Then run the kernel a few more times to
    /// confirm the EMA stays pinned at 5.0 (no overrun → no climb).
    #[test]
    #[ignore = "requires GPU"]
    fn sp16_phase1_min_hold_temp_strict_when_at_target() {
        let stream = make_stream();
        let kernel = load_kernel(&stream);

        const ISV_DIM: usize = 1024;
        let mut isv = vec![0.0_f32; ISV_DIM];
        // No overrun: observed = target.
        isv[HOLD_RATE_OBSERVED_EMA_INDEX] = 0.20;
        isv[HOLD_RATE_TARGET_INDEX]       = 0.20;
        isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SENTINEL_MIN_HOLD_TEMPERATURE;

        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
        isv_buf.write_from_slice(&isv);

        // First launch: Pearl-A bootstrap → REPLACE → temp = 5.
        launch_canonical(&stream, &kernel, isv_buf.dev_ptr);
        // Three additional launches to confirm EMA does not drift up.
        for _ in 0..3 {
            launch_canonical(&stream, &kernel, isv_buf.dev_ptr);
        }

        let result = isv_buf.read_all();
        let temp = result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX];

        // overrun = max(0, 0.20 − 0.20) = 0
        // target_temp = 5 + 45 × 0 = 5.0
        // After Pearl-A bootstrap, EMA blend with target=5 stays at 5.
        assert!(
            (temp - MIN_HOLD_TEMPERATURE_MIN).abs() < 1e-4,
            "At-target (observed = target) must drive temp LOW (strict); \
             expected {}, got {temp}",
            MIN_HOLD_TEMPERATURE_MIN
        );
    }

    /// SP16 Phase 1 (revised) — Test 3: under-target → temp LOW (strict).
    ///
    /// Setup: observed_hold_rate=0.10, target_hold_rate=0.20 (under).
    /// `max(0, observed − target)` clamps the overrun to 0, so target
    /// temp is TEMP_MIN. Pearl-A bootstrap → temp pinned to 5.
    /// Verifies the asymmetry: under-holding does not relax the
    /// penalty (no exit ramp needed when the model is already
    /// below target).
    #[test]
    #[ignore = "requires GPU"]
    fn sp16_phase1_min_hold_temp_strict_when_under_target() {
        let stream = make_stream();
        let kernel = load_kernel(&stream);

        const ISV_DIM: usize = 1024;
        let mut isv = vec![0.0_f32; ISV_DIM];
        // Under-holding: observed = 0.5× target.
        isv[HOLD_RATE_OBSERVED_EMA_INDEX] = 0.10;
        isv[HOLD_RATE_TARGET_INDEX]       = 0.20;
        isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SENTINEL_MIN_HOLD_TEMPERATURE;

        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
        isv_buf.write_from_slice(&isv);

        launch_canonical(&stream, &kernel, isv_buf.dev_ptr);

        let result = isv_buf.read_all();
        let temp = result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX];

        // overrun = max(0, 0.10 − 0.20) = 0 (clamped — no relaxation).
        // target_temp = 5 + 45 × 0 = 5.0
        // Pearl-A REPLACE → temp = 5.
        assert!(
            (temp - MIN_HOLD_TEMPERATURE_MIN).abs() < 1e-4,
            "Under-target (observed < target) must drive temp LOW (strict — no relaxation); \
             expected {}, got {temp}",
            MIN_HOLD_TEMPERATURE_MIN
        );
    }

    /// SP16 Phase 1 (revised) — Test 4: slot 373 dependency removed.
    ///
    /// Setup: slot 373 (AUX_DIR_ACC_SHORT_EMA) is at sentinel 0.5 —
    /// the OLD kernel's early-return guard would fire here and the
    /// temp slot would stay at sentinel 50.0. Setup observed_hold_rate
    /// =0.40, target=0.20 (significant overrun).
    /// The NEW kernel does NOT read slot 373; it must respond to the
    /// overrun and climb toward TEMP_MAX. This is the canonical
    /// proof that the signal swap is real.
    #[test]
    #[ignore = "requires GPU"]
    fn sp16_phase1_min_hold_temp_no_longer_reads_slot_373() {
        let stream = make_stream();
        let kernel = load_kernel(&stream);

        const ISV_DIM: usize = 1024;
        let mut isv = vec![0.0_f32; ISV_DIM];
        // Slot 373 at OLD-kernel sentinel 0.5. Pre-swap this would have
        // tripped the early-return guard and slot 460 would stay at 50.
        isv[AUX_DIR_ACC_SHORT_EMA_INDEX] = 0.5;
        // Significant overrun on the NEW signal slots.
        isv[HOLD_RATE_OBSERVED_EMA_INDEX] = 0.40;
        isv[HOLD_RATE_TARGET_INDEX]       = 0.20;
        isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SENTINEL_MIN_HOLD_TEMPERATURE;

        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
        isv_buf.write_from_slice(&isv);

        launch_canonical(&stream, &kernel, isv_buf.dev_ptr);

        let result = isv_buf.read_all();
        let temp = result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX];

        // The NEW kernel must respond to the hold-rate overrun even
        // though slot 373 is at the OLD sentinel. Pearl-A REPLACE
        // → target_temp = 50.0. The OLD kernel would have left slot
        // 460 at SENTINEL_MIN_HOLD_TEMPERATURE=50.0 too — but the
        // value-arrival path is different. To distinguish them
        // unambiguously, set up a non-sentinel under-target case
        // where OLD-kernel-sentinel-guard would preserve slot, but
        // NEW kernel writes target_temp=5.
        // (This sub-test is the strict version below.)
        assert!(
            (temp - MIN_HOLD_TEMPERATURE_MAX).abs() < 1e-4,
            "Over-holding with slot 373 at OLD sentinel must still drive temp HIGH; \
             expected {}, got {temp}",
            MIN_HOLD_TEMPERATURE_MAX
        );

        // Stronger version: under-target case where slot 373 sentinel
        // pre-swap would have preserved slot 460 (kernel returned
        // without writing). Post-swap the kernel writes target_temp=5
        // because under-target overrun_norm=0.
        //
        // SP16 T3 (2026-05-08): with Wiener-optimal α, cold-start
        // (sample_count=0 → 1 < 3) uses α=1.0 (REPLACE). The blend
        // path is exercised, but the formula `(1-α)×current + α×target`
        // collapses to `target` when α=1. So slot 460 lands at
        // target_temp=5 directly (NOT the OLD 28.75 result).
        let mut isv2 = vec![0.0_f32; ISV_DIM];
        isv2[AUX_DIR_ACC_SHORT_EMA_INDEX] = 0.5;       // OLD sentinel
        isv2[HOLD_RATE_OBSERVED_EMA_INDEX] = 0.10;     // under
        isv2[HOLD_RATE_TARGET_INDEX]       = 0.20;
        const SEED_TEMP: f32 = 30.0;
        isv2[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SEED_TEMP;
        let isv_buf2 = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv2");
        isv_buf2.write_from_slice(&isv2);
        launch_canonical(&stream, &kernel, isv_buf2.dev_ptr);
        let result2 = isv_buf2.read_all();
        let temp2 = result2[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX];

        // overrun = 0 → target_temp = 5.
        // Cold-start path (sample_count<3) → α=1.0 → blended ≈ target_temp = 5.
        // Slot moved off SEED_TEMP=30 (the OLD kernel's no-write outcome).
        assert!(
            (temp2 - MIN_HOLD_TEMPERATURE_MIN).abs() < 1e-4,
            "Slot 373 at OLD sentinel must NOT preserve slot 460; under-target → \
             cold-start REPLACE with target_temp = {}, got {temp2:.5} \
             (if {temp2:.5} ≈ {SEED_TEMP} the OLD kernel's slot-373-guard is still firing — signal swap not live)",
            MIN_HOLD_TEMPERATURE_MIN
        );
        assert!(
            (temp2 - SEED_TEMP).abs() > 0.5,
            "Slot 460 must have moved off seed {SEED_TEMP} (proves the swap removed the slot-373-sentinel-guard)"
        );
    }

    /// SP16 T3 (2026-05-08) — Wiener-optimal cold-start path.
    ///
    /// Pre-T3 this test verified `(1−0.05)×30 + 0.05×50 = 31.0`. Under
    /// Wiener-optimal α the kernel computes α from Welford state. With
    /// freshly-zeroed Welford accumulators (the canonical fold-reset
    /// state), the first launch sees `sample_count=0 → 1 < 3` and
    /// takes the cold-start path α=1.0 → blended = target_temp = 50.0.
    ///
    /// Setup: temp at 30.0 (NOT sentinel — past Pearl-A bootstrap).
    /// observed=0.40, target=0.20 → overrun_norm=1.0 → target_temp=50.
    /// Welford state at fold-reset zero. Cold-start path → α=1.0 →
    /// blend collapses to target_temp = MIN_HOLD_TEMPERATURE_MAX = 50.
    /// After 3 launches with this constant target, `sample_count = 3`
    /// reaches the Wiener threshold; with constant target_temp=50 every
    /// epoch, diff_var is exactly 0 and α floors at WELFORD_ALPHA_MIN.
    #[test]
    #[ignore = "requires GPU"]
    fn sp16_phase1_min_hold_temp_mid_cadence_ema_after_bootstrap() {
        let stream = make_stream();
        let kernel = load_kernel(&stream);

        const ISV_DIM: usize = 1024;
        const SEED_TEMP: f32 = 30.0;
        let mut isv = vec![0.0_f32; ISV_DIM];
        isv[HOLD_RATE_OBSERVED_EMA_INDEX] = 0.40;  // overrun_norm=1.0
        isv[HOLD_RATE_TARGET_INDEX]       = 0.20;
        isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SEED_TEMP;  // NOT sentinel

        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
        isv_buf.write_from_slice(&isv);

        launch_canonical(&stream, &kernel, isv_buf.dev_ptr);

        let result = isv_buf.read_all();
        let temp = result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX];

        // Cold-start path (sample_count=1 < 3) → α=1.0 → blend
        // collapses to target_temp = TEMP_MAX = 50.
        let expected = MIN_HOLD_TEMPERATURE_MAX;
        assert!(
            (temp - expected).abs() < 1e-4,
            "Wiener cold-start REPLACE: expected ≈ {expected:.5}, got {temp:.5}"
        );
        // Welford counter must have been incremented to 1.
        assert!(
            (result[MHT_SAMPLE_COUNT_INDEX] - 1.0).abs() < 1e-4,
            "Welford sample_count must be 1 after first launch, got {}",
            result[MHT_SAMPLE_COUNT_INDEX]
        );
    }

    /// Helper test (updated post-swap): observed-hold-rate at sentinel
    /// 0.0 (no per-step Hold observations yet) → kernel preserves
    /// slot 460 bit-exactly. Cold-start fold-1 protection mirroring
    /// the pre-swap dir_acc=0.5 sentinel test.
    #[test]
    #[ignore = "requires GPU"]
    fn sp16_phase1_min_hold_temp_sentinel_observed_preserves_isv() {
        let stream = make_stream();
        let kernel = load_kernel(&stream);

        const ISV_DIM: usize = 1024;
        const SEED: f32 = 32.0;
        let mut isv = vec![0.0_f32; ISV_DIM];
        // observed at SP13 fold-reset sentinel 0.0.
        isv[HOLD_RATE_OBSERVED_EMA_INDEX] = SENTINEL_HOLD_RATE_OBSERVED;
        isv[HOLD_RATE_TARGET_INDEX]       = 0.20;
        isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SEED;

        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
        isv_buf.write_from_slice(&isv);

        launch_canonical(&stream, &kernel, isv_buf.dev_ptr);

        let result = isv_buf.read_all();
        assert_eq!(
            result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX], SEED,
            "Temp slot must be preserved bit-exactly when observed_hold_rate is at sentinel"
        );
    }
}

// ═══════════════════════════════════════════════════════════════════════════
// SP16 Phase 0 (2026-05-08) — Q-by-action HEALTH_DIAG diagnostic.
//
// Per train-multi-seed-pfh9n post-mortem: structural-Q-bias hypothesis (Adam's
// `m/sqrt(v)` favours zero-variance Hold over noisy direction Q-targets) needs
// direct per-action Q observations to verify or refute. The production path is:
//
//   1. `update_q_dir_means_cached`: dtoh of `q_out_buf [B, total_actions]`,
//      then per-action averaging via `compute_q_dir_means_from_host_buf`.
//   2. `q_direction_action_means`: returns `[hold, long, short, flat]` cached
//      array (HEALTH_DIAG emit order).
//   3. training_loop.rs HEALTH_DIAG emit:
//      `q_by_action [hold=… long=… short=… flat=…]`.
//
// Action-index ordering verified at `state_layout.cuh:123-126`:
//   DIR_SHORT=0, DIR_HOLD=1, DIR_LONG=2, DIR_FLAT=3.
//
// This test exercises the actual averaging-path helper with known per-column
// values seeded into a [B, total_actions] row-major buffer. The slot order
// returned must match the HEALTH_DIAG `[hold, long, short, flat]` contract,
// independently of whichever dir_idx a given column corresponds to.
// ═══════════════════════════════════════════════════════════════════════════

/// SP16 Phase 0: verify the q_by_action diagnostic correctly extracts
/// per-direction-action Q-means from a `[B, total_actions]` row-major Q-out
/// buffer and emits in the contracted `[hold, long, short, flat]` slot order
/// matching the HEALTH_DIAG format. Validates the action-index ↔ slot
/// mapping against `state_layout.cuh:123-126` (DIR_SHORT=0, DIR_HOLD=1,
/// DIR_LONG=2, DIR_FLAT=3) so we can observe Q(Hold) trajectory vs
/// Q(direction) and falsify the structural-Q-bias hypothesis.
///
/// Construction:
///   * B=8, total_actions=12 (production layout: 4 dir + 3 mag + 3 ord + 3 urg).
///   * Each direction column i is filled with constant `dir_target[i]`
///     across all batch rows so the per-column mean equals the seeded
///     constant (no float averaging error). The remaining 8 mag/ord/urg
///     columns are filled with sentinel 99.0 to confirm the helper does
///     not bleed any non-direction column into the dir-action means.
///   * Targets chosen with 4 distinct values so a wrong index→slot mapping
///     is detected with 1e-3 tolerance.
///
/// CPU-only test — exercises pure-host averaging logic. The dtoh half of
/// the production updater is tested implicitly by smoke (HEALTH_DIAG
/// emits each epoch); the bug surface this test guards is the index→slot
/// mapping, which is data-independent.
#[test]
fn sp16_phase0_q_by_action_diagnostic_reads_four_action_means() {
    use ml::cuda_pipeline::gpu_dqn_trainer::compute_q_dir_means_from_host_buf;

    // Production direction-branch ordering, see state_layout.cuh:123-126.
    // dir_idx 0 → DIR_SHORT, 1 → DIR_HOLD, 2 → DIR_LONG, 3 → DIR_FLAT.
    const DIR_SHORT: usize = 0;
    const DIR_HOLD:  usize = 1;
    const DIR_LONG:  usize = 2;
    const DIR_FLAT:  usize = 3;

    // Returned-slot ordering, see compute_q_dir_means_from_host_buf docs.
    // HEALTH_DIAG `q_by_action [hold long short flat]`.
    const SLOT_HOLD:  usize = 0;
    const SLOT_LONG:  usize = 1;
    const SLOT_SHORT: usize = 2;
    const SLOT_FLAT:  usize = 3;

    const B: usize = 8;
    const B0: usize = 4;          // NUM_DIRECTIONS
    const TOTAL_ACTIONS: usize = 12; // 4 + 3 + 3 + 3 (production layout)

    // Distinct target Q-values per direction action.
    // Index by canonical dir_idx (matches the column the writer uses).
    let mut dir_target = [0.0_f32; 4];
    dir_target[DIR_SHORT] = 0.10;
    dir_target[DIR_HOLD ] = 0.50;
    dir_target[DIR_LONG ] = 0.20;
    dir_target[DIR_FLAT ] = 0.30;

    // Build a synthetic q_out_buf [B, total_actions] row-major. Direction
    // columns [0..4] are seeded with dir_target[col]; non-direction columns
    // [4..12] are sentinel 99.0 so any mis-indexing leaks visibly into the
    // returned means.
    let mut host = vec![0.0_f32; B * TOTAL_ACTIONS];
    for i in 0..B {
        let row_off = i * TOTAL_ACTIONS;
        for k in 0..B0 {
            host[row_off + k] = dir_target[k];
        }
        for k in B0..TOTAL_ACTIONS {
            host[row_off + k] = 99.0;
        }
    }

    let q_by_action = compute_q_dir_means_from_host_buf(&host, B, B0, TOTAL_ACTIONS);

    // Slot-order assertions: each slot must hold the canonical-dir-idx target.
    assert!(
        (q_by_action[SLOT_HOLD ] - dir_target[DIR_HOLD ]).abs() < 1e-3,
        "slot 0 (hold) must hold mean of DIR_HOLD column. got {:.6} expected {:.6}",
        q_by_action[SLOT_HOLD], dir_target[DIR_HOLD]
    );
    assert!(
        (q_by_action[SLOT_LONG ] - dir_target[DIR_LONG ]).abs() < 1e-3,
        "slot 1 (long) must hold mean of DIR_LONG column. got {:.6} expected {:.6}",
        q_by_action[SLOT_LONG], dir_target[DIR_LONG]
    );
    assert!(
        (q_by_action[SLOT_SHORT] - dir_target[DIR_SHORT]).abs() < 1e-3,
        "slot 2 (short) must hold mean of DIR_SHORT column. got {:.6} expected {:.6}",
        q_by_action[SLOT_SHORT], dir_target[DIR_SHORT]
    );
    assert!(
        (q_by_action[SLOT_FLAT ] - dir_target[DIR_FLAT ]).abs() < 1e-3,
        "slot 3 (flat) must hold mean of DIR_FLAT column. got {:.6} expected {:.6}",
        q_by_action[SLOT_FLAT], dir_target[DIR_FLAT]
    );

    // Defensive: any non-direction sentinel leakage would push the dir slot
    // past its dir_target. With sentinel=99.0 a one-column mis-index would
    // shift the slot by (99.0 - target) / B per leaked row — at least 12.4
    // for SLOT_HOLD's 0.50 target, far above the 1e-3 tolerance.
    for (slot, value) in q_by_action.iter().enumerate() {
        assert!(
            *value < 10.0,
            "slot {slot} contaminated by non-direction sentinel column \
             (got {value:.6}); index→slot mapping is wrong"
        );
    }

    // Empty-input guards: helper must return [0; 4] for degenerate cases.
    let empty = compute_q_dir_means_from_host_buf(&host, 0, B0, TOTAL_ACTIONS);
    assert_eq!(empty, [0.0_f32; 4], "B=0 must yield all zeros");
    let empty_b0 = compute_q_dir_means_from_host_buf(&host, B, 0, TOTAL_ACTIONS);
    assert_eq!(empty_b0, [0.0_f32; 4], "b0=0 must yield all zeros");
    // Truncated host (under-sized) must also yield zeros — defensive guard
    // that mirrors the production-path ModelError on under-sized q_out_buf.
    let small = vec![0.0_f32; 4];
    let truncated = compute_q_dir_means_from_host_buf(&small, B, B0, TOTAL_ACTIONS);
    assert_eq!(truncated, [0.0_f32; 4], "truncated host must yield all zeros");
}