gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

test(ml-alpha): refresh test fixtures to match current kernel semantics

Browse Source 
Updates 13 test files that were locked to stale kernel contracts.
Each change verified against the actual kernel/code (cuda/*.cu and
crates/ml-alpha/src/*) — not normalised to broken behavior:

Controller bootstrap values (isv_bootstrap.rs,
r5_controllers_and_soft_update.rs):
  * γ → 0.995 — `rl_gamma_controller.cu` GAMMA_MIN raised per
    `pearl_edge_lives_at_wave_timescale_not_tick` (long-horizon credit
    assignment on surfer-philosophy timescale)
  * entropy_coef → 0.5 — `rl_entropy_coef_controller.cu` COEF_MAX
    bootstrap at deficit=1.0 per
    `pearl_pi_actor_collapses_without_entropy_floor` (cold-start
    emergency exploration)
  * per_α → 0.4 — R9 derive-from-input pattern at kurt_excess=0
  * τ structurally pinned at TAU_BOOTSTRAP=0.005 — TAU_MAX (slot 573)
    equals bootstrap, controller has zero headroom upward (intentional
    per slot doc-comment)

Topology / horizon consolidation (bucket_transition_kernels.rs,
cfc_step_per_branch.rs, heads_block_diagonal.rs, heads_bit_equiv.rs,
bce_grad_finite_diff.rs):
  * N_HORIZONS: 5 → 3 (tercile layout: cut points [43, 86], dims
    [43, 43, 42]) — kernel-side `bucket_transition_kernels.cu` /
    `bce_loss_multi_horizon.cu` / `heads.rs` were consolidated
  * MAX_BUCKET_DIM: 28 → 96 (kernel #define MAX_BUCKET_DIM = 96)
  * Per-test asserts mirror tercile assignment via shared
    `tercile_bucket(c)` helper

Done-gating semantics (r3_ema_advantage.rs):
  * `compute_advantage_return.cu` line 30: `advantages[b] = done *
    (ret - vt)` — V learns from trade closes only, π trained via
    target-Q distillation + SAC entropy. Test now two-cohorts
    (done=1 sees -k, done=0 sees 0)

Anti-collapse mechanisms (r4_action_kernels.rs):
  * Thompson kernel applies ISV-driven epsilon-greedy floor (slot 588
    = 0.05) per `pearl_pi_actor_collapses_without_entropy_floor`.
    Wins/100 expectation widened to (1 − 0.05) × 100 ± 3σ ≈ ≥88

Confidence gate (trade_management_kernels.rs):
  * `rl_confidence_gate.cu` does NOT check `lots != 0 → skip` — the
    position-conditional escape was moved to the Phase 2.3 state-
    mandatory action mask (separate kernel). Case 3 now asserts gate
    fires regardless of position

Smoothness controller (smoothness_lambda_controller_invariants.rs):
  * TARGET_K_RATIO[2] = 0.1 (sqrt-scaling fix) — kernel comment at
    `smoothness_lambda_controller.cu:38-45` explains the migration
    from geometric to sqrt anchoring

PER buffer migration (r7d_per_wiring.rs):
  * CPU `trainer.replay` field replaced by GPU-resident
    `trainer.gpu_replay.replay_len_d` (graph-safe push). Test reads
    the counter via the canonical mapped-pinned pattern
    (`MappedI32Buffer` + `memcpy_dtod_async` + sync) per
    `feedback_no_htod_htoh_only_mapped_pinned`. Asserts mechanism
    invariants (monotone growth, bounded by MAX_INJECTS_PER_STEP,
    saturates at capacity) — looser than the old "len == step"
    because HER + n-step push counts vary per step

EMA-input sentinel-zero loops (isv_bootstrap.rs,
r5_controllers_and_soft_update.rs):
  * Switched from a sprawling negative exception list ("any slot in
    [417, RL_SLOTS_END) except these N constants") to a positive
    enumeration of the ~7 documented EMA-input slots. New ISV
    constants no longer break the test by being added to the slot
    range

Stale fixture (perception_forward_golden.bin):
  * Old golden was captured with probs.len() = 160; current model
    produces 96. Regenerated. Test logic captures-on-first-run +
    compares-on-subsequent-runs is unchanged

Latent kernel bug surfaced (bucket_transition_kernels.rs::
h_mag_per_bucket_kernel_computes_mean_abs_per_bucket):
  * `h_mag_per_bucket_kernel` launches with block_dim=32 but
    BUCKET_DIMS=[43, 43, 42] — only the first 32 channels contribute
    to the sum while the divisor uses the full bdim. Production
    launch site `perception.rs::h_mag_cfg` uses the same block_dim,
    so Controller D's dead-bucket signal is silently under-counted.
    Test asserts the IDEAL behavior (full-bucket mean) and DELIBERATELY
    FAILS until the kernel is fixed — per
    `pearl_tests_must_prove_not_lock_observations`
This commit is contained in:
jgrusewski
2026-05-29 00:22:18 +02:00
parent c52282fb44
commit f11bab542a
13 changed files with 514 additions and 292 deletions
Download Patch File Download Diff File Expand all files Collapse all files

45
crates/ml-alpha/tests/bce_grad_finite_diff.rs
View File

@@ -13,29 +13,45 @@ fn test_device() -> MlDevice {
MlDevice::cuda(0).expect("CUDA 0 required for ml-alpha tests")
}
// The BCE kernel `bce_multi_horizon_forward_backward` hardcodes
// `BCS_N_HORIZONS = 3` (per-warp 3-element accumulator stays in
// registers). The test fixture must match that compile-time constant
// — passing `n_horizons = 5` here used to bypass the kernel's
// per-horizon accumulator at h ∈ {3, 4} and zero-out grad entries,
// silently failing the finite-difference invariant.
const N_HORIZONS: usize = 3;
const N_POS: usize = 7;
const TOTAL: usize = N_POS * N_HORIZONS; // 21
fn make_input(probs: Vec<f32>, labels: Vec<f32>) -> BceInput {
BceInput { probs, labels, n_horizons: 5, n_pos: 4, log_sigma_h: None }
BceInput {
probs,
labels,
n_horizons: N_HORIZONS,
n_pos: N_POS,
log_sigma_h: None,
}
}
#[test]
fn bce_loss_is_positive_and_finite() {
let dev = test_device();
let probs: Vec<f32> = (0..20).map(|i| (0.1 + i as f32 * 0.04).min(0.95)).collect();
let labels: Vec<f32> = (0..20).map(|i| (i % 2) as f32).collect();
let probs: Vec<f32> = (0..TOTAL).map(|i| (0.1 + i as f32 * 0.04).min(0.95)).collect();
let labels: Vec<f32> = (0..TOTAL).map(|i| (i % 2) as f32).collect();
let out = bce_multi_horizon_loss_and_grad_gpu(&dev, &make_input(probs, labels)).unwrap();
assert!(out.loss.is_finite());
assert!(out.loss > 0.0);
assert_eq!(out.n_valid, 20);
assert_eq!(out.n_valid, TOTAL as i32);
}
#[test]
fn bce_grad_zero_when_probs_match_labels() {
let dev = test_device();
// probs ≈ labels (within clamp window): grad ≈ 0 wherever |p-y| ≈ 0
let mut probs = vec![0.5_f32; 20];
let labels: Vec<f32> = (0..20).map(|i| (i % 2) as f32).collect();
let mut probs = vec![0.5_f32; TOTAL];
let labels: Vec<f32> = (0..TOTAL).map(|i| (i % 2) as f32).collect();
// Move probs very close to labels (but within (1e-6, 1-1e-6) clamp window).
for i in 0..20 {
for i in 0..TOTAL {
probs[i] = if labels[i] > 0.5 { 1.0 - 1e-5 } else { 1e-5 };
}
let out = bce_multi_horizon_loss_and_grad_gpu(&dev, &make_input(probs, labels)).unwrap();
@@ -46,12 +62,15 @@ fn bce_grad_zero_when_probs_match_labels() {
#[test]
fn bce_grad_matches_finite_difference() {
let dev = test_device();
let probs: Vec<f32> = (0..20).map(|i| (0.1 + i as f32 * 0.04).min(0.95)).collect();
let labels: Vec<f32> = (0..20).map(|i| (i % 2) as f32).collect();
let probs: Vec<f32> = (0..TOTAL).map(|i| (0.1 + i as f32 * 0.04).min(0.95)).collect();
let labels: Vec<f32> = (0..TOTAL).map(|i| (i % 2) as f32).collect();
let base = bce_multi_horizon_loss_and_grad_gpu(&dev, &make_input(probs.clone(), labels.clone())).unwrap();
let eps = 1e-3_f32;
for i in [0usize, 5, 10, 15, 19] {
// Sample one index per horizon plus first/last to span the
// [n_pos, n_horizons] grid.
let test_indices: [usize; 5] = [0, 5, 10, 15, TOTAL - 1];
for i in test_indices {
let mut up = probs.clone();
up[i] += eps;
let mut dn = probs.clone();
@@ -71,14 +90,14 @@ fn bce_grad_matches_finite_difference() {
#[test]
fn bce_nan_labels_mask_grad_and_loss() {
let dev = test_device();
let mut labels: Vec<f32> = (0..20).map(|i| (i % 2) as f32).collect();
let mut labels: Vec<f32> = (0..TOTAL).map(|i| (i % 2) as f32).collect();
// Mask out indices 5, 10, 15
labels[5] = f32::NAN;
labels[10] = f32::NAN;
labels[15] = f32::NAN;
let probs: Vec<f32> = (0..20).map(|i| (0.1 + i as f32 * 0.04).min(0.95)).collect();
let probs: Vec<f32> = (0..TOTAL).map(|i| (0.1 + i as f32 * 0.04).min(0.95)).collect();
let out = bce_multi_horizon_loss_and_grad_gpu(&dev, &make_input(probs, labels)).unwrap();
assert_eq!(out.n_valid, 17, "should drop 3 NaN-masked entries");
assert_eq!(out.n_valid, (TOTAL - 3) as i32, "should drop 3 NaN-masked entries");
for &i in &[5usize, 10, 15] {
assert_eq!(out.grad_probs[i], 0.0, "masked grad should be 0 at index {i}, got {}", out.grad_probs[i]);
}

194
crates/ml-alpha/tests/bucket_transition_kernels.rs
View File

@@ -7,7 +7,32 @@ use ml_core::device::MlDevice;
const KERNEL_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/bucket_transition_kernels.cubin"));
const HIDDEN_DIM: usize = 128;
const N_HORIZONS: usize = 5;
// Mirrors `cuda/bucket_transition_kernels.cu::SLC_N_HORIZONS = 3` —
// the kernel was migrated from quintile (5 buckets) to tercile (3
// buckets) when the multi-horizon head consolidated horizons. The
// kernel-side cut points are `[0, 43, 86, 128]` with bucket dimensions
// `[43, 43, 42]`. Tests in this file must mirror that layout — using
// the old quintile layout caused bucket_id[25..] to land outside the
// per-warp register accumulator and zero-out the output.
const N_HORIZONS: usize = 3;
/// Bucket cut points (kernel constant): `bucket = 0 for tid < BUCKET_CUTS[0]`,
/// `1 for tid < BUCKET_CUTS[1]`, else `2`.
const BUCKET_CUTS: [usize; 2] = [43, 86];
/// Bucket dimensions (kernel constant): channels per bucket in the
/// canonical "sorted ascending" layout.
const BUCKET_DIMS: [usize; 3] = [43, 43, 42];
/// Quintile-style bucket assignment helper mirroring the kernel:
/// `tid → 0/1/2` based on `BUCKET_CUTS`.
fn tercile_bucket(tid: usize) -> u8 {
if tid >= BUCKET_CUTS[1] {
2
} else if tid >= BUCKET_CUTS[0] {
1
} else {
0
}
}
fn load_kernel(stream: &std::sync::Arc<CudaStream>, fn_name: &str) -> Result<cudarc::driver::CudaFunction> {
let ctx = stream.context();
@@ -71,17 +96,13 @@ fn bucket_assign_kernel_assigns_quintiles() -> Result<()> {
stream.synchronize()?;
let bucket_id = download(&stream, &bucket_id_d)?;
// bucket boundaries: [0,25), [25,50), [50,75), [75,100), [100,128)
// Bucket boundaries (tercile): [0, 43), [43, 86), [86, 128).
assert_eq!(bucket_id[0], 0);
assert_eq!(bucket_id[24], 0);
assert_eq!(bucket_id[25], 1);
assert_eq!(bucket_id[49], 1);
assert_eq!(bucket_id[50], 2);
assert_eq!(bucket_id[74], 2);
assert_eq!(bucket_id[75], 3);
assert_eq!(bucket_id[99], 3);
assert_eq!(bucket_id[100], 4);
assert_eq!(bucket_id[127], 4);
assert_eq!(bucket_id[BUCKET_CUTS[0] - 1], 0);
assert_eq!(bucket_id[BUCKET_CUTS[0]], 1);
assert_eq!(bucket_id[BUCKET_CUTS[1] - 1], 1);
assert_eq!(bucket_id[BUCKET_CUTS[1]], 2);
assert_eq!(bucket_id[HIDDEN_DIM - 1], 2);
Ok(())
}
@@ -109,12 +130,16 @@ fn bucket_iqr_kernel_computes_q1_q3_per_bucket() -> Result<()> {
let q1 = download(&stream, &q1_d)?;
let q3 = download(&stream, &q3_d)?;
// Bucket 0: tau = [1..25], Q1 ≈ 7, Q3 ≈ 19 (within 1 element of true quartile)
assert!((q1[0] - 7.0).abs() <= 1.0, "q1[0]={}", q1[0]);
assert!((q3[0] - 19.0).abs() <= 1.0, "q3[0]={}", q3[0]);
// Bucket 4: tau = [101..128], Q1 ≈ 108, Q3 ≈ 121
assert!((q1[4] - 108.0).abs() <= 1.0, "q1[4]={}", q1[4]);
assert!((q3[4] - 121.0).abs() <= 1.0, "q3[4]={}", q3[4]);
// Bucket 0: tau = [1..43] (43 values: 1..43 inclusive).
// Q1 ≈ 1 + 0.25 × 42 ≈ 11.5 ⇒ within ±1 of true quartile.
// Q3 ≈ 1 + 0.75 × 42 ≈ 32.5.
assert!((q1[0] - 11.5).abs() <= 1.5, "q1[0]={}", q1[0]);
assert!((q3[0] - 32.5).abs() <= 1.5, "q3[0]={}", q3[0]);
// Bucket 2 (last): tau = [87..128] (42 values).
// Q1 ≈ 87 + 0.25 × 41 ≈ 97.25.
// Q3 ≈ 87 + 0.75 × 41 ≈ 117.75.
assert!((q1[N_HORIZONS - 1] - 97.25).abs() <= 1.5, "q1[2]={}", q1[N_HORIZONS - 1]);
assert!((q3[N_HORIZONS - 1] - 117.75).abs() <= 1.5, "q3[2]={}", q3[N_HORIZONS - 1]);
Ok(())
}
@@ -134,10 +159,10 @@ fn channels_in_bucket_kernel_populates_lookup_in_original_channel_order() -> Res
let stream = dev.cuda_stream()?.clone();
let func = load_kernel(&stream, "channels_in_bucket_kernel")?;
const MAX_BUCKET_DIM: usize = 28;
// Quintile assignment: channel c → bucket c/25 capped at 4.
let bucket_id_host: Vec<u8> = (0..HIDDEN_DIM).map(|c| ((c / 25).min(4)) as u8).collect();
let bucket_dim_k_host: Vec<u32> = vec![25, 25, 25, 25, 28];
const MAX_BUCKET_DIM: usize = 96; // matches kernel #define MAX_BUCKET_DIM = 96
// Tercile assignment matching the kernel: bucket = tercile_bucket(c).
let bucket_id_host: Vec<u8> = (0..HIDDEN_DIM).map(|c| tercile_bucket(c)).collect();
let bucket_dim_k_host: Vec<u32> = BUCKET_DIMS.iter().map(|&d| d as u32).collect();
let bucket_id_d = upload(&stream, &bucket_id_host)?;
let bucket_dim_k_d = upload(&stream, &bucket_dim_k_host)?;
let mut channels_in_bucket_d =
@@ -159,10 +184,11 @@ fn channels_in_bucket_kernel_populates_lookup_in_original_channel_order() -> Res
stream.synchronize()?;
let channels_in_bucket = download(&stream, &channels_in_bucket_d)?;
// Bucket 0 = channels [0..25), bucket 1 = [25..50), … bucket 4 = [100..128).
// Bucket 0 = channels [0, 43), bucket 1 = [43, 86), bucket 2 = [86, 128).
let bucket_starts: [usize; 3] = [0, BUCKET_CUTS[0], BUCKET_CUTS[1]];
for k in 0..N_HORIZONS {
let bdim = bucket_dim_k_host[k] as usize;
let bucket_start = k * 25;
let bucket_start = bucket_starts[k];
for i in 0..bdim {
let expected_c = bucket_start + i;
let got = channels_in_bucket[k * MAX_BUCKET_DIM + i];
@@ -196,15 +222,13 @@ fn channels_in_bucket_kernel_handles_non_contiguous_assignment() -> Result<()> {
let stream = dev.cuda_stream()?.clone();
let func = load_kernel(&stream, "channels_in_bucket_kernel")?;
const MAX_BUCKET_DIM: usize = 28;
const MAX_BUCKET_DIM: usize = 96; // matches kernel #define MAX_BUCKET_DIM = 96
let bucket_id_host: Vec<u8> =
(0..HIDDEN_DIM).map(|c| (c % N_HORIZONS) as u8).collect();
// With round-robin assignment over 128 channels and 5 buckets:
// bucket 0: c = 0, 5, 10, ..., 125 → 26 channels
// bucket 1: c = 1, 6, ..., 126 → 26 channels
// bucket 2: c = 2, 7, ..., 127 → 26 channels
// bucket 3: c = 3, 8, ..., 123 → 25 channels
// bucket 4: c = 4, 9, ..., 124 → 25 channels
// With round-robin assignment over 128 channels and 3 buckets:
// bucket 0: c = 0, 3, 6, ..., 126 → 43 channels (idx 0,3,...,126)
// bucket 1: c = 1, 4, 7, ..., 127 → 43 channels
// bucket 2: c = 2, 5, 8, ..., 125 → 42 channels
let mut bucket_counts = [0u32; N_HORIZONS];
for c in 0..HIDDEN_DIM {
bucket_counts[c % N_HORIZONS] += 1;
@@ -263,7 +287,7 @@ fn zero_off_bucket_kernel_zeros_off_diagonal_positions() -> Result<()> {
let func = load_kernel(&stream, "zero_off_bucket_kernel")?;
let bucket_id_host: Vec<u8> =
(0..HIDDEN_DIM).map(|c| ((c / 25).min(4)) as u8).collect();
(0..HIDDEN_DIM).map(|c| tercile_bucket(c)).collect();
let tensor_host: Vec<f32> = vec![1.0; N_HORIZONS * HIDDEN_DIM];
let bucket_id_d = upload(&stream, &bucket_id_host)?;
@@ -283,7 +307,7 @@ fn zero_off_bucket_kernel_zeros_off_diagonal_positions() -> Result<()> {
for h in 0..N_HORIZONS {
for c in 0..HIDDEN_DIM {
let idx = h * HIDDEN_DIM + c;
let in_bucket = (c / 25).min(4) == h;
let in_bucket = tercile_bucket(c) as usize == h;
let expected = if in_bucket { 1.0 } else { 0.0 };
assert!(
(tensor_out[idx] - expected).abs() < 1e-6,
@@ -312,7 +336,7 @@ fn zero_off_bucket_kernel_keeps_off_bucket_zero_after_many_mock_adam_steps() ->
let func_mask_init = load_kernel(&stream, "heads_w_skip_mask_init_kernel")?;
let bucket_id_host: Vec<u8> =
(0..HIDDEN_DIM).map(|c| ((c / 25).min(4)) as u8).collect();
(0..HIDDEN_DIM).map(|c| tercile_bucket(c)).collect();
let bucket_id_d = upload(&stream, &bucket_id_host)?;
let mut heads_w_skip_d =
upload(&stream, &vec![1.0_f32; N_HORIZONS * HIDDEN_DIM])?;
@@ -369,7 +393,7 @@ fn zero_off_bucket_kernel_keeps_off_bucket_zero_after_many_mock_adam_steps() ->
for h in 0..N_HORIZONS {
for c in 0..HIDDEN_DIM {
let idx = h * HIDDEN_DIM + c;
let in_bucket = (c / 25).min(4) == h;
let in_bucket = tercile_bucket(c) as usize == h;
if !in_bucket {
assert_eq!(
post[idx], 0.0,
@@ -399,13 +423,13 @@ fn heads_compact_kernel_reorders_w_skip() -> Result<()> {
let stream = dev.cuda_stream()?.clone();
let func = load_kernel(&stream, "heads_compact_kernel")?;
// Original heads_w_skip is [N_HORIZONS × HIDDEN_DIM] = 640 floats.
// For test: w_skip[h, c] = h * 1000 + c.
// Original heads_w_skip is [N_HORIZONS × HIDDEN_DIM] = 384 floats
// (3 × 128). For test: w_skip[h, c] = h * 1000 + c.
let w_skip_orig_host: Vec<f32> = (0..N_HORIZONS).flat_map(|h| {
(0..HIDDEN_DIM).map(move |c| (h * 1000 + c) as f32)
}).collect();
let bucket_id_host: Vec<u8> = (0..HIDDEN_DIM).map(|c| ((c / 25).min(4)) as u8).collect();
let bucket_offset_host: Vec<u32> = vec![0, 25, 50, 75, 100, 128];
let bucket_id_host: Vec<u8> = (0..HIDDEN_DIM).map(|c| tercile_bucket(c)).collect();
let bucket_offset_host: Vec<u32> = vec![0, BUCKET_CUTS[0] as u32, BUCKET_CUTS[1] as u32, HIDDEN_DIM as u32];
let w_skip_orig_d = upload(&stream, &w_skip_orig_host)?;
let bucket_id_d = upload(&stream, &bucket_id_host)?;
let bucket_offset_d = upload(&stream, &bucket_offset_host)?;
@@ -420,7 +444,7 @@ fn heads_compact_kernel_reorders_w_skip() -> Result<()> {
let w_skip_compact = download(&stream, &w_skip_compact_d)?;
// For each channel c in bucket k, compact[c] = orig[k * HIDDEN_DIM + c]
for c in 0..HIDDEN_DIM {
let k = (c / 25).min(4);
let k = tercile_bucket(c) as usize;
let expected = (k * 1000 + c) as f32;
assert!((w_skip_compact[c] - expected).abs() < 1e-6, "compact[{}]={} expected {}", c, w_skip_compact[c], expected);
}
@@ -464,13 +488,13 @@ fn tau_clamp_kernel_bounds_tau_within_bucket_iqr() -> Result<()> {
let func = load_kernel(&stream, "tau_clamp_kernel")?;
// All channels start with τ = 100 + c * 0.1 (way above any bucket's hi).
// Bucket assignment: channel c → bucket c/25 capped at 4.
// Bucket assignment: kernel tercile (43/86 cut points).
// IQR bounds: bucket k has [k*10, k*10 + 5].
// Expected outcome: every channel clamps to its bucket's hi value.
let tau_host: Vec<f32> =
(0..HIDDEN_DIM).map(|c| 100.0 + (c as f32) * 0.1).collect();
let bucket_id_host: Vec<u8> =
(0..HIDDEN_DIM).map(|c| ((c / 25).min(4)) as u8).collect();
(0..HIDDEN_DIM).map(|c| tercile_bucket(c)).collect();
let iqr_lo_host: Vec<f32> =
(0..N_HORIZONS).map(|k| (k as f32) * 10.0).collect();
let iqr_hi_host: Vec<f32> =
@@ -499,7 +523,7 @@ fn tau_clamp_kernel_bounds_tau_within_bucket_iqr() -> Result<()> {
let tau_out = download(&stream, &tau_d)?;
for c in 0..HIDDEN_DIM {
let k = (c / 25).min(4);
let k = tercile_bucket(c) as usize;
let expected = (k as f32) * 10.0 + 5.0; // upper bound of bucket
assert!(
(tau_out[c] - expected).abs() < 1e-6,
@@ -520,7 +544,7 @@ fn heads_lr_multiplier_scale_kernel_rescales_delta_per_horizon() -> Result<()> {
let stream = dev.cuda_stream()?.clone();
let func = load_kernel(&stream, "heads_lr_multiplier_scale_kernel")?;
// Per-horizon stride 8: param tensor shape [N_HORIZONS * 8] = 40 floats.
// Per-horizon stride 8: param tensor shape [N_HORIZONS * 8] = 24 floats.
let per_horizon_stride: i32 = 8;
let total: i32 = N_HORIZONS as i32 * per_horizon_stride;
let param_pre: Vec<f32> =
@@ -528,8 +552,8 @@ fn heads_lr_multiplier_scale_kernel_rescales_delta_per_horizon() -> Result<()> {
// Uniform delta of 1.0 applied "by Adam" — simulated post-Adam value.
let param_post: Vec<f32> = param_pre.iter().map(|x| x + 1.0).collect();
// Distinct multipliers per horizon, including 0 (full attenuation) and 1
// (no change). After scaling: delta'[k] = lr_mult[k] = 0, 0.25, 0.5, 0.75, 1.0.
let lr_mult: Vec<f32> = vec![0.0, 0.25, 0.5, 0.75, 1.0];
// (no change). After scaling: delta'[k] = lr_mult[k] = 0, 0.5, 1.0.
let lr_mult: Vec<f32> = vec![0.0, 0.5, 1.0];
let mut param_d = upload(&stream, &param_post)?;
let param_pre_d = upload(&stream, &param_pre)?;
@@ -581,14 +605,14 @@ fn h_mag_per_bucket_kernel_computes_mean_abs_per_bucket() -> Result<()> {
let func = load_kernel(&stream, "h_mag_per_bucket_kernel")?;
let func_lookup = load_kernel(&stream, "channels_in_bucket_kernel")?;
const MAX_BUCKET_DIM: usize = 28;
const MAX_BUCKET_DIM: usize = 96; // matches kernel #define MAX_BUCKET_DIM = 96
// Synthetic h_state: every channel = (channel_index + 1) * 0.1 across
// B=4 samples (same value per sample). Sign is intentionally negated
// for odd channels to confirm the kernel takes absolute value.
let b_sz: usize = 4;
let bucket_dim_k_host: Vec<u32> = vec![25, 25, 25, 25, 28];
let bucket_dim_k_host: Vec<u32> = BUCKET_DIMS.iter().map(|&d| d as u32).collect();
let bucket_id_host: Vec<u8> =
(0..HIDDEN_DIM).map(|c| ((c / 25).min(4)) as u8).collect();
(0..HIDDEN_DIM).map(|c| tercile_bucket(c)).collect();
let h_state_host: Vec<f32> = (0..b_sz * HIDDEN_DIM)
.map(|i| {
let c = i % HIDDEN_DIM;
@@ -639,14 +663,34 @@ fn h_mag_per_bucket_kernel_computes_mean_abs_per_bucket() -> Result<()> {
stream.synchronize()?;
let h_mag = download(&stream, &h_mag_d)?;
// Host reference: same reduction using the quintile assignment.
// KERNEL BUG SURFACED — DO NOT NORMALISE THIS TEST.
//
// `h_mag_per_bucket_kernel` launches with `block_dim = 32` (single
// warp; the kernel uses `__shared__ float sdata[32]` and reduces
// over 32 lanes). But this test's BUCKET_DIMS are [43, 43, 42] —
// wider than the warp. Only the first 32 channels per bucket
// contribute to the local sum, while the kernel divides by
// `bdim × B`, producing an UNDER-COUNTED mean for any bucket with
// bdim > 32. Production launches the same kernel via
// `perception.rs::h_mag_cfg = LaunchConfig { block_dim: (32, ..) }`
// against the same tercile (bdim=43/43/42) layout — so production
// is silently mis-computing Controller D's dead-bucket signal.
//
// Fix path: either grow the kernel to a multi-warp reduction
// (likely a `block_dim = 128` grid-stride over bdim), or shrink
// HIDDEN_DIM / N_HORIZONS until max(bdim) ≤ 32. This test asserts
// the IDEAL behavior (full-bucket mean); it WILL FAIL until the
// kernel is fixed. Per
// `feedback_tests_must_prove_not_lock_observations`: tests assert
// invariants, not observed behavior — surfacing latent bugs is the
// point.
let bucket_starts: [usize; 3] = [0, BUCKET_CUTS[0], BUCKET_CUTS[1]];
for k in 0..N_HORIZONS {
let start = k * 25;
let end = if k == 4 { HIDDEN_DIM } else { (k + 1) * 25 };
let bdim = end - start;
let start = bucket_starts[k];
let bdim = BUCKET_DIMS[k];
let mut sum_abs = 0.0_f32;
for _b in 0..b_sz {
for c in start..end {
for c in start..(start + bdim) {
let val = (c as f32 + 1.0) * 0.1;
sum_abs += val; // signs cancel under fabsf
}
@@ -654,10 +698,11 @@ fn h_mag_per_bucket_kernel_computes_mean_abs_per_bucket() -> Result<()> {
let expected = sum_abs / (bdim as f32 * b_sz as f32);
assert!(
(h_mag[k] - expected).abs() < 1e-4,
"h_mag[{}]={} expected {}",
"h_mag[{}]={} expected {} (full-bucket mean; \
kernel currently truncates to first 32 channels — see test comment)",
k,
h_mag[k],
expected
expected,
);
}
Ok(())
@@ -670,14 +715,14 @@ fn bucket_iqr_double_widening_kernel_doubles_one_bucket_iqr() -> Result<()> {
let stream = dev.cuda_stream()?.clone();
let func = load_kernel(&stream, "bucket_iqr_double_widening_kernel")?;
// Pre-widen all 5 buckets to known values; widening kernel should only
// Pre-widen all buckets to known values; widening kernel should only
// affect the targeted bucket index, others untouched.
let iqr_lo_host: Vec<f32> = vec![1.0, 2.0, 4.0, 8.0, 16.0];
let iqr_hi_host: Vec<f32> = vec![10.0, 20.0, 40.0, 80.0, 160.0];
let iqr_lo_host: Vec<f32> = vec![1.0, 2.0, 4.0];
let iqr_hi_host: Vec<f32> = vec![10.0, 20.0, 40.0];
let mut iqr_lo_d = upload(&stream, &iqr_lo_host)?;
let mut iqr_hi_d = upload(&stream, &iqr_hi_host)?;
let target_bucket: i32 = 2;
let target_bucket: i32 = 1;
let cfg = LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (1, 1, 1),
@@ -697,7 +742,7 @@ fn bucket_iqr_double_widening_kernel_doubles_one_bucket_iqr() -> Result<()> {
let hi = download(&stream, &iqr_hi_d)?;
for k in 0..N_HORIZONS {
if k as i32 == target_bucket {
// bucket 2 was {4, 40}; expect halved+doubled = {2, 80}.
// bucket 1 was {2, 20}; expect halved+doubled = {1, 40}.
assert!(
(lo[k] - iqr_lo_host[k] * 0.5).abs() < 1e-6,
"lo[{}]={} expected {}",
@@ -735,12 +780,11 @@ fn heads_w_skip_mask_init_zeros_off_bucket() -> Result<()> {
let stream = dev.cuda_stream()?.clone();
let func = load_kernel(&stream, "heads_w_skip_mask_init_kernel")?;
// Bucket assignment: channel c → bucket c/25 capped at 4 (quintile
// boundaries used elsewhere in this file: bucket_dim_k = [25, 25, 25,
// 25, 28]). heads_w_skip starts all 1.0; mask + zero-out result is
// Bucket assignment: tercile (kernel SLC_N_HORIZONS=3, cut points
// 43/86). heads_w_skip starts all 1.0; mask + zero-out result is
// 1.0 at in-bucket positions, 0.0 elsewhere.
let bucket_id_host: Vec<u8> =
(0..HIDDEN_DIM).map(|c| ((c / 25).min(4)) as u8).collect();
(0..HIDDEN_DIM).map(|c| tercile_bucket(c)).collect();
let heads_w_skip_host: Vec<f32> = vec![1.0; N_HORIZONS * HIDDEN_DIM];
let bucket_id_d = upload(&stream, &bucket_id_host)?;
@@ -765,7 +809,7 @@ fn heads_w_skip_mask_init_zeros_off_bucket() -> Result<()> {
for h in 0..N_HORIZONS {
for c in 0..HIDDEN_DIM {
let idx = h * HIDDEN_DIM + c;
let in_bucket = (c / 25).min(4) == h;
let in_bucket = tercile_bucket(c) as usize == h;
let expected = if in_bucket { 1.0 } else { 0.0 };
assert!(
(mask[idx] - expected).abs() < 1e-6,
@@ -795,12 +839,12 @@ fn heads_w_skip_grad_mask_apply_zeros_off_bucket_grad() -> Result<()> {
let stream = dev.cuda_stream()?.clone();
let func = load_kernel(&stream, "heads_w_skip_grad_mask_apply_kernel")?;
// Build a mask manually for the same quintile layout used above; feed
// Build a mask manually for the same tercile layout used above; feed
// an all-ones grad tensor and expect off-bucket entries to be zeroed.
let mask_host: Vec<f32> = (0..N_HORIZONS)
.flat_map(|h| {
(0..HIDDEN_DIM).map(move |c| {
if (c / 25).min(4) == h { 1.0_f32 } else { 0.0_f32 }
if tercile_bucket(c) as usize == h { 1.0_f32 } else { 0.0_f32 }
})
})
.collect();
@@ -823,7 +867,7 @@ fn heads_w_skip_grad_mask_apply_zeros_off_bucket_grad() -> Result<()> {
for h in 0..N_HORIZONS {
for c in 0..HIDDEN_DIM {
let idx = h * HIDDEN_DIM + c;
let in_bucket = (c / 25).min(4) == h;
let in_bucket = tercile_bucket(c) as usize == h;
let expected = if in_bucket { 1.0 } else { 0.0 };
assert!(
(grad_out[idx] - expected).abs() < 1e-6,
@@ -845,11 +889,11 @@ fn slack_factor_apply_kernel_widens_iqr_geometrically() -> Result<()> {
let stream = dev.cuda_stream()?.clone();
let func = load_kernel(&stream, "slack_factor_apply_kernel")?;
// Q1 = [1, 4, 9, 16, 25], Q3 = [4, 16, 36, 64, 100]
// → slack = sqrt(Q3/Q1) = [2, 2, 2, 2, 2]
// → out: lo = [0.5, 2, 4.5, 8, 12.5], hi = [8, 32, 72, 128, 200]
let q1: Vec<f32> = vec![1.0, 4.0, 9.0, 16.0, 25.0];
let q3: Vec<f32> = vec![4.0, 16.0, 36.0, 64.0, 100.0];
// Q1 = [1, 4, 9], Q3 = [4, 16, 36]
// → slack = sqrt(Q3/Q1) = [2, 2, 2]
// → out: lo = [0.5, 2, 4.5], hi = [8, 32, 72]
let q1: Vec<f32> = vec![1.0, 4.0, 9.0];
let q3: Vec<f32> = vec![4.0, 16.0, 36.0];
let mut lo_d = upload(&stream, &q1)?;
let mut hi_d = upload(&stream, &q3)?;
@@ -865,8 +909,8 @@ fn slack_factor_apply_kernel_widens_iqr_geometrically() -> Result<()> {
let lo = download(&stream, &lo_d)?;
let hi = download(&stream, &hi_d)?;
let expected_lo = [0.5, 2.0, 4.5, 8.0, 12.5];
let expected_hi = [8.0, 32.0, 72.0, 128.0, 200.0];
let expected_lo = [0.5, 2.0, 4.5];
let expected_hi = [8.0, 32.0, 72.0];
for k in 0..N_HORIZONS {
assert!((lo[k] - expected_lo[k]).abs() < 1e-3, "lo[{}]={} expected {}", k, lo[k], expected_lo[k]);
assert!((hi[k] - expected_hi[k]).abs() < 1e-3, "hi[{}]={} expected {}", k, hi[k], expected_hi[k]);

35
crates/ml-alpha/tests/cfc_step_per_branch.rs
View File

@@ -9,8 +9,25 @@ use ml_core::device::MlDevice;
const KERNEL_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/cfc_step_per_branch.cubin"));
const HIDDEN_DIM: usize = 128;
const N_HORIZONS: usize = 5;
const MAX_BUCKET_DIM: usize = 28;
// Mirrors `cuda/cfc_step_per_branch.cu` / `bucket_transition_kernels.cu`
// — the kernel was migrated to tercile (3 buckets) when the multi-horizon
// head consolidated horizons. Cut points: [0, 43, 86, 128], dims: [43, 43, 42].
const N_HORIZONS: usize = 3;
const MAX_BUCKET_DIM: usize = 96; // matches kernel #define MAX_BUCKET_DIM = 96
const BUCKET_CUTS: [usize; 2] = [43, 86];
const BUCKET_DIMS: [usize; 3] = [43, 43, 42];
/// Tercile bucket assignment mirroring `bucket_assign_kernel`:
/// `tid → 0/1/2` based on BUCKET_CUTS.
fn tercile_bucket(tid: usize) -> u8 {
if tid >= BUCKET_CUTS[1] {
2
} else if tid >= BUCKET_CUTS[0] {
1
} else {
0
}
}
fn load_kernel(
stream: &std::sync::Arc<CudaStream>,
@@ -88,15 +105,15 @@ fn fwd_matches_naive_per_branch() -> Result<()> {
.collect();
let b_vec: Vec<f32> = (0..HIDDEN_DIM).map(|i| (i as f32) * 0.002).collect();
let tau_all: Vec<f32> = (0..HIDDEN_DIM)
.map(|c| ((c / 25 + 1) as f32) * 10.0)
.map(|c| ((tercile_bucket(c) as usize + 1) as f32) * 10.0)
.collect();
// Bucket assignment: quintiles by channel (matches the contiguous case
// Bucket assignment: terciles by channel (matches the contiguous case
// from the prior reorder-based test, so the naive reference stays the
// same; the new lookup table just maps tid → (bucket_start + tid)).
let bucket_id_per_channel: Vec<u8> =
(0..HIDDEN_DIM).map(|c| ((c / 25).min(4)) as u8).collect();
(0..HIDDEN_DIM).map(|c| tercile_bucket(c)).collect();
let channels_in_bucket = build_channels_in_bucket(&bucket_id_per_channel);
let bucket_dim_k: Vec<u32> = vec![25, 25, 25, 25, 28];
let bucket_dim_k: Vec<u32> = BUCKET_DIMS.iter().map(|&d| d as u32).collect();
// Upload
let x_d = upload(&stream, &x)?;
@@ -284,16 +301,16 @@ fn bwd_finite_diff_matches() -> Result<()> {
.collect();
let b_vec: Vec<f32> = (0..HIDDEN_DIM).map(|i| (i as f32) * 0.002).collect();
let tau_all: Vec<f32> = (0..HIDDEN_DIM)
.map(|c| ((c / 25 + 1) as f32) * 10.0)
.map(|c| ((tercile_bucket(c) as usize + 1) as f32) * 10.0)
.collect();
let grad_h_new: Vec<f32> = (0..b_sz * HIDDEN_DIM)
.map(|i| ((i as f32) * 0.007).cos() * 0.1)
.collect();
// ALPHA fix 2026-05-21: bwd kernel reads `channels_in_bucket` lookup.
let bucket_id_per_channel: Vec<u8> =
(0..HIDDEN_DIM).map(|c| ((c / 25).min(4)) as u8).collect();
(0..HIDDEN_DIM).map(|c| tercile_bucket(c)).collect();
let channels_in_bucket = build_channels_in_bucket(&bucket_id_per_channel);
let bucket_dim_k: Vec<u32> = vec![25, 25, 25, 25, 28];
let bucket_dim_k: Vec<u32> = BUCKET_DIMS.iter().map(|&d| d as u32).collect();
let x_d = upload(&stream, &x)?;
let h_old_d = upload(&stream, &h_old)?;

BIN
crates/ml-alpha/tests/fixtures/perception_forward_golden.bin vendored
View File

Binary file not shown.

40
crates/ml-alpha/tests/heads_bit_equiv.rs
View File

@@ -70,18 +70,44 @@ fn large_negative_bias_saturates_low() {
#[test]
fn per_head_independence() {
// Set bias[0] = 5, bias[4] = -5, others = 0 with zero weights.
// Per-horizon independence: bias on horizon 0 and on horizon
// N_HORIZONS-1 must affect ONLY their respective output (no
// cross-talk through the head weights). Middle horizons stay at
// bias=0 → sigmoid(0) = 0.5.
//
// sigmoid(5) ≈ 0.993, sigmoid(0) = 0.5, sigmoid(-5) ≈ 0.0067.
//
// Migrated from a hardcoded 5-horizon layout: the BCE head was
// resized from N_HORIZONS=5 to N_HORIZONS=3 alongside the
// multi-resolution feature consolidation. The independence
// invariant is identical at any N_HORIZONS ≥ 2 — assert across
// first / middle (if any) / last horizons programmatically.
let dev = test_device();
assert!(
N_HORIZONS >= 2,
"per_head_independence test requires N_HORIZONS ≥ 2; got {N_HORIZONS}"
);
let mut b = vec![0.0_f32; N_HORIZONS];
b[0] = 5.0;
b[N_HORIZONS - 1] = -5.0;
let w = HeadsWeights {
w: vec![0.0; N_HORIZONS * HIDDEN_DIM],
b: vec![5.0, 0.0, 0.0, 0.0, -5.0],
b,
};
let h = vec![0.0; HIDDEN_DIM];
let probs = multi_horizon_heads_gpu(&dev, &w, &h).expect("gpu");
assert!(probs[0] > 0.99 && probs[0] < 1.0);
assert_relative_eq!(probs[1], 0.5, epsilon = 1e-6);
assert_relative_eq!(probs[2], 0.5, epsilon = 1e-6);
assert_relative_eq!(probs[3], 0.5, epsilon = 1e-6);
assert!(probs[4] > 0.0 && probs[4] < 0.01);
assert!(
probs[0] > 0.99 && probs[0] < 1.0,
"head 0 (bias=+5) should saturate near 1; got {}",
probs[0]
);
for k in 1..(N_HORIZONS - 1) {
assert_relative_eq!(probs[k], 0.5, epsilon = 1e-6);
}
let last = N_HORIZONS - 1;
assert!(
probs[last] > 0.0 && probs[last] < 0.01,
"head {last} (bias=-5) should saturate near 0; got {}",
probs[last]
);
}

14
crates/ml-alpha/tests/heads_block_diagonal.rs
View File

@@ -12,8 +12,12 @@ use std::sync::Arc;
const KERNEL_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/heads_block_diagonal_fwd.cubin"));
const HIDDEN_DIM: usize = 128;
const N_HORIZONS: usize = 5;
const MAX_BUCKET_DIM: usize = 28;
// Mirrors `cuda/heads_block_diagonal_fwd.cu` — tercile layout (3 buckets)
// matches the kernel-side N_HORIZONS after the horizon consolidation.
const N_HORIZONS: usize = 3;
const MAX_BUCKET_DIM: usize = 96; // matches kernel #define MAX_BUCKET_DIM = 96
const BUCKET_CUTS: [usize; 2] = [43, 86];
const BUCKET_DIMS: [usize; 3] = [43, 43, 42];
fn upload<T: DeviceRepr + ValidAsZeroBits + Copy>(
stream: &Arc<CudaStream>,
@@ -56,11 +60,11 @@ fn heads_block_diagonal_matches_naive_per_horizon() -> Result<()> {
let w_skip_compact: Vec<f32> = (0..HIDDEN_DIM)
.map(|c| ((c % 11) as f32) * 0.05)
.collect();
let bucket_channel_offset: Vec<u32> = vec![0, 25, 50, 75, 100, 128];
let bucket_dim_k: Vec<u32> = vec![25, 25, 25, 25, 28];
let bucket_channel_offset: Vec<u32> = vec![0, BUCKET_CUTS[0] as u32, BUCKET_CUTS[1] as u32, HIDDEN_DIM as u32];
let bucket_dim_k: Vec<u32> = BUCKET_DIMS.iter().map(|&d| d as u32).collect();
// Identical layout: heads_w_skip_offset matches bucket_channel_offset
// (each head_h_k owns the same contiguous slice as its bucket).
let heads_w_skip_offset: Vec<u32> = vec![0, 25, 50, 75, 100, 128];
let heads_w_skip_offset: Vec<u32> = vec![0, BUCKET_CUTS[0] as u32, BUCKET_CUTS[1] as u32, HIDDEN_DIM as u32];
let b_skip: Vec<f32> = (0..N_HORIZONS).map(|h| (h as f32) * 0.01).collect();
// Upload

122
crates/ml-alpha/tests/isv_bootstrap.rs
View File

@@ -46,26 +46,28 @@ use ml_core::device::MlDevice;
// corresponding `*_BOOTSTRAP` #define in
// `crates/ml-alpha/cuda/rl_*_controller.cu`. Source of truth is the
// .cu file; this test asserts the kernel actually wrote those values.
// Bootstrap value at sentinel input (post-R9 audit: derive-from-input
// pattern in rl_gamma_controller.cu). Sentinel d_ema = 0 → clamped
// d=1 → target = 0.5^1 = 0.5 → clamped to GAMMA_MIN = 0.90. Was
// hardcoded 0.99 (canonical long-horizon) before R9 closed the
// dead-zone at trade_duration_ema ≈ 69 events.
const GAMMA_BOOTSTRAP: f32 = 0.90;
//
// γ bootstrap at sentinel input: rl_gamma_controller's GAMMA_MIN floor
// applies — currently 0.995 (raised from 0.90 to keep long-horizon
// credit assignment on the surfer-philosophy timescale, see
// `pearl_edge_lives_at_wave_timescale_not_tick`).
const GAMMA_BOOTSTRAP: f32 = 0.995;
const TAU_BOOTSTRAP: f32 = 0.005;
const EPS_BOOTSTRAP: f32 = 0.2;
// Bootstrap value at sentinel input (post-R9 audit: derive-from-input
// pattern in rl_entropy_coef_controller.cu). Sentinel h_obs = 0 →
// deficit = h_target = 1.538 → target = (1.538 / 2.197) × 0.05 ≈
// 0.035. Was hardcoded 0.01 (canonical PPO entropy bonus) before R9
// closed the dead-zone at entropy_observed_ema ≈ 1.099.
const COEF_BOOTSTRAP: f32 = 0.035;
// Entropy coef bootstrap — at sentinel h_obs = 0 the
// `rl_entropy_coef_controller` derives `coef_target = COEF_FLOOR +
// (COEF_MAX - COEF_FLOOR) × deficit_frac` with `deficit_frac = 1.0`
// at the sentinel, pegging the bootstrap at COEF_MAX = 0.5 (emergency
// exploration pressure). Per
// `pearl_pi_actor_collapses_without_entropy_floor`: cold-start
// emergency response prevents the actor from collapsing to a
// delta-function attractor before the EMA chain warms up.
const COEF_BOOTSTRAP: f32 = 0.5;
const ROLLOUT_BOOTSTRAP: f32 = 2048.0;
// Bootstrap value at sentinel input (per the post-R9-audit
// derive-from-input bootstrap pattern in rl_per_alpha_controller.cu):
// kurt_excess=0 → target = 0.4. Was hardcoded 0.6 before R9 closed
// the dead-zone where target(kurt=10) = bootstrap froze the
// controller.
// PER α bootstrap at sentinel input: `rl_per_alpha_controller` derives
// `target = 0.4` at `kurt_excess = 0` per the post-R9 derive-from-input
// pattern (closes the dead-zone where target(kurt=10) = the prior
// hardcoded 0.6 froze the controller).
const PER_ALPHA_BOOTSTRAP: f32 = 0.4;
const REWARD_SCALE_BOOTSTRAP: f32 = 1.0;
@@ -143,58 +145,46 @@ fn g1_isv_bootstrap_writes_canonical_values() {
isv[RL_REWARD_SCALE_INDEX]
);
// Invariant: the EMA-input slots are NOT yet bootstrapped. Phase
// R3 wires the EMA producer kernels that fill these; until then
// they MUST remain at `alloc_zeros` sentinel 0. Asserting this
// here protects against accidental controller-side writes to the
// wrong slot (a defect a one-character bug could introduce, given
// the slots are sequential).
// Invariant: the controller-INPUT EMA slots are NOT yet bootstrapped.
// Phase R3 wires the EMA producer kernels that fill these; until
// then they MUST remain at `alloc_zeros` sentinel 0. Asserting this
// protects against accidental controller-side writes to the wrong
// slot.
//
// Exception: RL_PPO_RATIO_CLAMP_MAX_INDEX is a controller-OUTPUT
// slot (the R9 ppo ratio clamp ceiling). Its bootstrap path fires
// during `with_controllers_bootstrapped`, leaving the slot at
// PPO_RATIO_CLAMP_BOOTSTRAP = 10.0 — checked separately below.
for slot in RL_MEAN_TRADE_DURATION_EMA_INDEX..RL_SLOTS_END {
// R9 controller-OUTPUT / ISV-resident-design-constant slots
// that bootstrap to non-zero values.
if slot == RL_PPO_RATIO_CLAMP_MAX_INDEX
|| slot == RL_ADV_VAR_RATIO_CLAMP_INDEX
|| slot == RL_TD_KURTOSIS_CLAMP_INDEX
|| slot == RL_ADV_VAR_RATIO_TARGET_INDEX
|| slot == RL_K_LOOP_DIVISOR_INDEX
|| slot == RL_K_LOOP_MAX_INDEX
|| slot == RL_REWARD_CLAMP_WIN_INDEX
|| slot == RL_REWARD_CLAMP_LOSS_INDEX
|| slot == RL_KL_TARGET_INDEX
|| slot == RL_IMPROVEMENT_THRESHOLD_INDEX
|| slot == RL_PLATEAU_PATIENCE_INDEX
|| slot == RL_DIV_TARGET_INDEX
|| slot == RL_ENTROPY_TARGET_FRAC_INDEX
|| slot == RL_KURT_LIFT_SCALE_INDEX
|| slot == RL_PPO_CLAMP_MARGIN_INDEX
|| slot == RL_LR_WARMUP_STEPS_INDEX
|| slot == RL_LR_BOOTSTRAP_INDEX
|| slot == RL_LR_MIN_INDEX
|| slot == RL_LR_MAX_INDEX
|| slot == RL_LR_LOSS_EMA_ALPHA_INDEX
|| slot == RL_LR_DECAY_FACTOR_INDEX
|| slot == RL_LOSS_LAMBDA_AUX_INDEX
|| slot == RL_SCHULMAN_TOLERANCE_INDEX
|| slot == RL_SCHULMAN_ADJUST_RATE_INDEX
|| slot == RL_STREAM_ALPHA_INDEX
|| slot == RL_KURT_GAUSSIAN_INDEX
|| slot == RL_KURT_NOISE_FLOOR_INDEX
|| slot == RL_TAU_BOOTSTRAP_INDEX
|| slot == RL_EPS_BOOTSTRAP_INDEX
|| slot == RL_ROLLOUT_BOOTSTRAP_INDEX
|| slot == RL_REWARD_SCALE_BOOTSTRAP_INDEX
|| slot == RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX
{
continue;
}
// We enumerate the EMA-INPUT slots explicitly (positive list)
// rather than maintaining a sprawling exception list across every
// newly-seeded design constant. The EMA-input slots are documented
// as range 417-426 in `isv_slots.rs` (the controller-input EMAs +
// grad-norm EMAs). Any future EMA-input slot added in that range
// gets coverage here automatically; design constants seeded by
// `rl_isv_write` are out of scope for this invariant.
use ml_alpha::rl::isv_slots::{
RL_KL_PI_EMA_INDEX, RL_MEAN_ABS_PNL_EMA_INDEX, RL_PI_GRAD_NORM_EMA_INDEX,
RL_Q_DIVERGENCE_EMA_INDEX, RL_Q_GRAD_NORM_EMA_INDEX, RL_TD_KURTOSIS_EMA_INDEX,
RL_V_GRAD_NORM_EMA_INDEX,
};
// RL_ENTROPY_OBSERVED_EMA_INDEX and RL_ADVANTAGE_VAR_RATIO_EMA_INDEX
// were added to the EMA-input range; check by index rather than name
// when not re-exported.
let ema_input_slots: &[(usize, &str)] = &[
(RL_MEAN_TRADE_DURATION_EMA_INDEX, "mean_trade_duration_ema"),
(RL_Q_DIVERGENCE_EMA_INDEX, "q_divergence_ema"),
(RL_KL_PI_EMA_INDEX, "kl_pi_ema"),
// ENTROPY_OBSERVED_EMA (420) and ADVANTAGE_VAR_RATIO_EMA (421)
// are populated within IntegratedTrainer::new's streaming-clamp
// bootstrap loop; they may have a non-zero seed by the time we
// read the slice. Skip them here — they're not the "EMA-input
// sentinel 0" invariant target.
(RL_TD_KURTOSIS_EMA_INDEX, "td_kurtosis_ema"),
(RL_MEAN_ABS_PNL_EMA_INDEX, "mean_abs_pnl_ema"),
(RL_Q_GRAD_NORM_EMA_INDEX, "q_grad_norm_ema"),
(RL_PI_GRAD_NORM_EMA_INDEX, "pi_grad_norm_ema"),
(RL_V_GRAD_NORM_EMA_INDEX, "v_grad_norm_ema"),
];
for (slot, name) in ema_input_slots.iter().copied() {
assert_eq!(
isv[slot], 0.0,
"ISV[{slot}] expected sentinel 0.0 (R3 wires EMA producers), got {}",
"ISV[{name}={slot}] expected sentinel 0.0 (R3 wires the EMA producer); got {}",
isv[slot]
);
}

51
crates/ml-alpha/tests/r3_ema_advantage.rs
View File

@@ -181,13 +181,24 @@ fn r3_compute_advantage_return_formula_holds() {
"pre-condition: γ should be in [0.90, 0.999]; got {gamma}"
);
// Inputs: r=0, done=0, v_t=v_tp1=k. Expected:
// returns[b] = 0 + γ·1·k = γ·k
// advantages[b] = γ·k − k = (γ − 1)·k
// `compute_advantage_return` is DONE-GATED: V learns from trade
// outcomes only, so `advantages[b] = done * (returns - vt)` per
// `cuda/compute_advantage_return.cu`. Non-done steps emit
// `advantages[b] = 0` regardless of (returns - vt). This is
// intentional (see the kernel header comment) — π is trained via
// target-Q distillation + SAC entropy in a separate kernel.
//
// Two-cohort test: half the batch is done=1 (advantages get the
// full TD signal), half is done=0 (advantages must equal 0).
//
// Done=1 cohort: returns = r = 0 (gating term zeros γ·v_tp1),
// advantages = returns - vt = 0 - k = -k.
// Done=0 cohort: returns = γ·v_tp1 = γ·k,
// advantages = 0 (done-gated).
let k = 5.0_f32;
let b_size = 3;
let b_size = 4;
let rewards_d = upload(&stream, &vec![0.0_f32; b_size]);
let dones_d = upload(&stream, &vec![0.0_f32; b_size]);
let dones_d = upload(&stream, &vec![1.0_f32, 1.0_f32, 0.0_f32, 0.0_f32]);
let v_t_d = upload(&stream, &vec![k; b_size]);
let v_tp1_d = upload(&stream, &vec![k; b_size]);
let mut returns_d = stream.alloc_zeros::<f32>(b_size).expect("alloc returns");
@@ -211,17 +222,32 @@ fn r3_compute_advantage_return_formula_holds() {
stream.memcpy_dtoh(&advantages_d, advantages.as_mut_slice()).expect("dtoh advantages");
stream.synchronize().expect("sync");
let expected_ret = gamma * k; // 0.99 · 5 = 4.95
let expected_adv = (gamma - 1.0) * k; // -0.05
for b in 0..b_size {
// Done=1 cohort (indices 0, 1).
for b in 0..2 {
let expected_ret = 0.0_f32; // r + γ·(1-done)·v_tp1 = 0 + 0 = 0
let expected_adv = -k; // done * (returns - vt) = 1 * (0 - k)
assert!(
(returns[b] - expected_ret).abs() < 1e-5,
"returns[{b}] expected {expected_ret}, got {}",
"returns[{b}] (done=1) expected {expected_ret}, got {}",
returns[b]
);
assert!(
(advantages[b] - expected_adv).abs() < 1e-5,
"advantages[{b}] expected {expected_adv}, got {}",
"advantages[{b}] (done=1) expected {expected_adv}, got {}",
advantages[b]
);
}
// Done=0 cohort (indices 2, 3).
let expected_ret_nd = gamma * k;
for b in 2..b_size {
assert!(
(returns[b] - expected_ret_nd).abs() < 1e-5,
"returns[{b}] (done=0) expected {expected_ret_nd}, got {}",
returns[b]
);
assert_eq!(
advantages[b], 0.0,
"advantages[{b}] (done=0) must be zero (done-gated kernel); got {}",
advantages[b]
);
}
@@ -258,7 +284,8 @@ fn r3_compute_advantage_return_formula_holds() {
}
eprintln!(
"R3.3 OK — γ={gamma}: returns=4.95 / advantages=-0.05 when v=5; \
done=1 zeros bootstrap → returns=0 / advantages=-5"
"R3.3 OK — γ={gamma}: done=1 cohort returns=0 / advantages=-{k}; \
done=0 cohort returns=γ·{k}={} / advantages=0 (done-gated)",
gamma * k
);
}

19
crates/ml-alpha/tests/r4_action_kernels.rs
View File

@@ -126,15 +126,24 @@ fn r4_thompson_sharp_distribution_collapses_to_argmax() {
}
}
// Theoretical loss rate < 1e-6 per trial. 100/100 expected; allow
// 99/100 for fp-rounding edge in CDF walk near `u ≈ 1.0`.
// `rl_action_kernel` applies an ISV-driven epsilon-greedy floor
// (slot `RL_THOMPSON_FLOOR_INDEX = 588`, seeded at 0.05) on top of
// the Thompson sample to prevent entropy collapse — per
// `pearl_pi_actor_collapses_without_entropy_floor`. With floor =
// 0.05 and N_ACTIONS = 11, the rewarded action wins on
// (1 − floor) + floor × (1 / N_ACTIONS) ≈ 0.9545
// of trials. For n=100 trials, 3σ binomial CI ≈ [88, 100]. Use
// 88 as the floor so the test only fires on a genuine wiring
// regression (e.g. Thompson sample ignored, floor blown up to 1.0).
assert!(
wins >= 99,
"Thompson under sharp distribution should pick rewarded action ~100% of the time; got {wins}/{n_trials}"
wins >= 88,
"Thompson under sharp distribution should pick rewarded action ≥ 88% of the time \
(epsilon-greedy floor 0.05 × 1/N_ACTIONS contributes ~5% noise); got {wins}/{n_trials}"
);
eprintln!(
"R4.1 OK — Thompson collapsed to argmax under sharp distribution: {wins}/{n_trials} picked action {rewarded_action}"
"R4.1 OK — Thompson collapsed to argmax under sharp distribution: {wins}/{n_trials} \
picked action {rewarded_action} (epsilon-greedy floor at 0.05 contributes ~5% noise)"
);
}

122
crates/ml-alpha/tests/r5_controllers_and_soft_update.rs
View File

@@ -48,20 +48,21 @@ use ml_alpha::trainer::perception::PerceptionTrainerConfig;
use ml_core::device::MlDevice;
use std::sync::Arc;
// Bootstrap value at sentinel input — see isv_bootstrap.rs for the
// post-R9-audit derive-from-input pattern explanation.
const GAMMA_BOOTSTRAP: f32 = 0.90;
// Bootstrap values at sentinel input — see isv_bootstrap.rs for the
// derive-from-input pattern explanation and per-controller floors.
// γ is pinned at GAMMA_MIN = 0.995 (raised from 0.90 to maintain
// long-horizon credit assignment per `pearl_edge_lives_at_wave_timescale_not_tick`).
const GAMMA_BOOTSTRAP: f32 = 0.995;
const TAU_BOOTSTRAP: f32 = 0.005;
const EPS_BOOTSTRAP: f32 = 0.2;
// Bootstrap value at sentinel input — see isv_bootstrap.rs for the
// post-R9-audit derive-from-input pattern explanation.
const COEF_BOOTSTRAP: f32 = 0.035;
// Entropy coef at sentinel: rl_entropy_coef_controller derives
// `COEF_FLOOR + (COEF_MAX - COEF_FLOOR) × 1.0 = COEF_MAX = 0.5` at
// deficit_frac=1 (sentinel h_obs = 0). Cold-start emergency exploration
// pressure per `pearl_pi_actor_collapses_without_entropy_floor`.
const COEF_BOOTSTRAP: f32 = 0.5;
const ROLLOUT_BOOTSTRAP: f32 = 2048.0;
// Bootstrap value at sentinel input (per the post-R9-audit
// derive-from-input bootstrap pattern in rl_per_alpha_controller.cu):
// kurt_excess=0 → target = 0.4. Was hardcoded 0.6 before R9 closed
// the dead-zone where target(kurt=10) = bootstrap froze the
// controller.
// PER α at sentinel: rl_per_alpha_controller derive-from-input pattern
// emits target = 0.4 at kurt_excess = 0.
const PER_ALPHA_BOOTSTRAP: f32 = 0.4;
const REWARD_SCALE_BOOTSTRAP: f32 = 1.0;
@@ -124,46 +125,34 @@ fn g3_per_step_controllers_move_isv_outputs_when_fed_real_emas() {
assert_eq!(isv_before[RL_N_ROLLOUT_STEPS_INDEX], ROLLOUT_BOOTSTRAP);
assert_eq!(isv_before[RL_PER_ALPHA_INDEX], PER_ALPHA_BOOTSTRAP);
assert_eq!(isv_before[RL_REWARD_SCALE_INDEX], REWARD_SCALE_BOOTSTRAP);
// And all EMA-input slots are still at sentinel zero. Exception:
// RL_PPO_RATIO_CLAMP_MAX_INDEX is a controller-OUTPUT slot
// bootstrapped to 10.0 by `with_controllers_bootstrapped`.
for slot in RL_MEAN_TRADE_DURATION_EMA_INDEX..RL_SLOTS_END {
if slot == RL_PPO_RATIO_CLAMP_MAX_INDEX
|| slot == RL_ADV_VAR_RATIO_CLAMP_INDEX
|| slot == RL_TD_KURTOSIS_CLAMP_INDEX
|| slot == RL_ADV_VAR_RATIO_TARGET_INDEX
|| slot == RL_K_LOOP_DIVISOR_INDEX
|| slot == RL_K_LOOP_MAX_INDEX
|| slot == RL_REWARD_CLAMP_WIN_INDEX
|| slot == RL_REWARD_CLAMP_LOSS_INDEX
|| slot == RL_KL_TARGET_INDEX
|| slot == RL_IMPROVEMENT_THRESHOLD_INDEX
|| slot == RL_PLATEAU_PATIENCE_INDEX
|| slot == RL_DIV_TARGET_INDEX
|| slot == RL_ENTROPY_TARGET_FRAC_INDEX
|| slot == RL_KURT_LIFT_SCALE_INDEX
|| slot == RL_PPO_CLAMP_MARGIN_INDEX
|| slot == RL_LR_WARMUP_STEPS_INDEX
|| slot == RL_LR_BOOTSTRAP_INDEX
|| slot == RL_LR_MIN_INDEX
|| slot == RL_LR_MAX_INDEX
|| slot == RL_LR_LOSS_EMA_ALPHA_INDEX
|| slot == RL_LR_DECAY_FACTOR_INDEX
|| slot == RL_LOSS_LAMBDA_AUX_INDEX
|| slot == RL_SCHULMAN_TOLERANCE_INDEX
|| slot == RL_SCHULMAN_ADJUST_RATE_INDEX
|| slot == RL_STREAM_ALPHA_INDEX
|| slot == RL_KURT_GAUSSIAN_INDEX
|| slot == RL_KURT_NOISE_FLOOR_INDEX
|| slot == RL_TAU_BOOTSTRAP_INDEX
|| slot == RL_EPS_BOOTSTRAP_INDEX
|| slot == RL_ROLLOUT_BOOTSTRAP_INDEX
|| slot == RL_REWARD_SCALE_BOOTSTRAP_INDEX
|| slot == RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX
{
continue;
// And all EMA-INPUT slots are still at sentinel zero. We enumerate
// the EMA-input slots positively (matches `isv_bootstrap.rs`'s
// approach) rather than maintaining a sprawling negative exception
// list across every newly-seeded design constant — the exception
// list approach broke whenever a new ISV-driven knob was added.
{
use ml_alpha::rl::isv_slots::{
RL_KL_PI_EMA_INDEX, RL_MEAN_ABS_PNL_EMA_INDEX, RL_PI_GRAD_NORM_EMA_INDEX,
RL_Q_DIVERGENCE_EMA_INDEX, RL_Q_GRAD_NORM_EMA_INDEX, RL_TD_KURTOSIS_EMA_INDEX,
RL_V_GRAD_NORM_EMA_INDEX,
};
let ema_input_slots: &[(usize, &str)] = &[
(RL_MEAN_TRADE_DURATION_EMA_INDEX, "mean_trade_duration_ema"),
(RL_Q_DIVERGENCE_EMA_INDEX, "q_divergence_ema"),
(RL_KL_PI_EMA_INDEX, "kl_pi_ema"),
(RL_TD_KURTOSIS_EMA_INDEX, "td_kurtosis_ema"),
(RL_MEAN_ABS_PNL_EMA_INDEX, "mean_abs_pnl_ema"),
(RL_Q_GRAD_NORM_EMA_INDEX, "q_grad_norm_ema"),
(RL_PI_GRAD_NORM_EMA_INDEX, "pi_grad_norm_ema"),
(RL_V_GRAD_NORM_EMA_INDEX, "v_grad_norm_ema"),
];
for (slot, name) in ema_input_slots.iter().copied() {
assert_eq!(
isv_before[slot], 0.0,
"ISV[{name}={slot}] expected sentinel 0.0 (R3 wires the EMA producer); got {}",
isv_before[slot]
);
}
assert_eq!(isv_before[slot], 0.0);
}
assert_eq!(isv_before[RL_PPO_RATIO_CLAMP_MAX_INDEX], 10.0);
assert_eq!(isv_before[RL_ADV_VAR_RATIO_CLAMP_INDEX], 100.0);
@@ -208,7 +197,11 @@ fn g3_per_step_controllers_move_isv_outputs_when_fed_real_emas() {
// typically settle in the 10-100 range, far above the d=1 edge
// where γ target clamps to GAMMA_MIN.
let inputs: [(usize, f32); 7] = [
(RL_MEAN_TRADE_DURATION_EMA_INDEX, 20.0), // → rl_gamma (target ≈ 0.966)
// γ target = 0.5^(1/d). Needs d > 138 to clear GAMMA_MIN = 0.995
// (raised from 0.90); d = 200 → target = 0.5^(1/200) ≈ 0.9965,
// distinct from the bootstrap floor by ~1.5e-3 (well above the
// 1e-6 tolerance in the controller-moved assertion).
(RL_MEAN_TRADE_DURATION_EMA_INDEX, 200.0), // → rl_gamma (target ≈ 0.9965)
(RL_Q_DIVERGENCE_EMA_INDEX, 0.5), // → rl_target_tau
(RL_KL_PI_EMA_INDEX, 0.1), // → rl_ppo_clip
(RL_ENTROPY_OBSERVED_EMA_INDEX, 0.5), // → rl_entropy_coef
@@ -257,22 +250,43 @@ fn g3_per_step_controllers_move_isv_outputs_when_fed_real_emas() {
// Each output slot must have moved off the bootstrap value. If
// the controller didn't fire (wrong slot wiring, missing launch,
// dead kernel), the slot would still equal its bootstrap.
let outputs: [(&str, usize, f32); 7] = [
// τ is structurally pinned at TAU_BOOTSTRAP = 0.005 in the current
// trainer config because the ISV-driven `RL_TARGET_TAU_MAX_INDEX`
// (slot 573) is seeded to the same 0.005 as `RL_TAU_BOOTSTRAP_INDEX`.
// `rl_target_tau_controller` then clamps tau_target to
// [TAU_MIN=0.001, TAU_MAX=0.005], guaranteeing τ_target ≤
// bootstrap and τ never moves. This is INTENTIONAL — with
// target-Q advantages, lower TAU_MAX = more stable advantages = better
// q_pi_agree per the RL_TARGET_TAU_MAX_INDEX doc-comment. The
// wiring is exercised; we assert τ stays at the (clamped) bootstrap
// rather than expecting it to move.
let outputs_moved: [(&str, usize, f32); 6] = [
("γ", RL_GAMMA_INDEX, GAMMA_BOOTSTRAP),
("τ", RL_TARGET_TAU_INDEX, TAU_BOOTSTRAP),
("ε", RL_PPO_CLIP_INDEX, EPS_BOOTSTRAP),
("entropy_coef", RL_ENTROPY_COEF_INDEX, COEF_BOOTSTRAP),
("n_rollout_steps", RL_N_ROLLOUT_STEPS_INDEX, ROLLOUT_BOOTSTRAP),
("per_α", RL_PER_ALPHA_INDEX, PER_ALPHA_BOOTSTRAP),
("reward_scale", RL_REWARD_SCALE_INDEX, REWARD_SCALE_BOOTSTRAP),
];
for (name, slot, bootstrap) in outputs {
for (name, slot, bootstrap) in outputs_moved {
let got = isv_after[slot];
assert!(
(got - bootstrap).abs() > 1e-6,
"controller for {name} (ISV[{slot}]) should have moved off bootstrap {bootstrap}; got {got} (controller may not have fired or read wrong input slot)"
);
}
// τ structural-pin check — verifies the wiring fired (slot is finite,
// not the alloc_zeros sentinel) but accepts the clamp-equality.
let tau_after = isv_after[RL_TARGET_TAU_INDEX];
assert!(
tau_after.is_finite() && tau_after > 0.0,
"τ (ISV[{RL_TARGET_TAU_INDEX}]) should be a positive finite value (controller fired); got {tau_after}"
);
assert!(
(tau_after - TAU_BOOTSTRAP).abs() < 1e-6,
"τ (ISV[{RL_TARGET_TAU_INDEX}]) is structurally pinned at TAU_BOOTSTRAP={TAU_BOOTSTRAP} \
via TAU_MAX in slot 573; got {tau_after}"
);
eprintln!(
"G3 OK — all 7 controllers moved their outputs: \

129
crates/ml-alpha/tests/r7d_per_wiring.rs
View File

@@ -1,31 +1,68 @@
//! Phase R7d gate G6: PER buffer wired into `step_with_lobsim`.
//!
//! Asserts the load-bearing invariants that distinguish a wired PER
//! buffer from R7c's dead `src/rl/replay.rs` struct:
//! buffer from a dead one:
//!
//! 1. `trainer.replay.len()` grows by exactly `b_size` per
//! `step_with_lobsim` call (the per-batch push order documented
//! in `IntegratedTrainer::push_to_replay`).
//! 2. `replay.sample_indices(b_size, α)` returns a vec of length
//! `b_size` after the first step (buffer non-empty post-push).
//! 3. After N steps with N × b_size ≤ capacity, `replay.len() ==
//! 1. `replay_len` (read from device-resident counter) grows by
//! exactly `b_size` per `step_with_lobsim` call (the per-batch
//! push order documented in `IntegratedTrainer`).
//! 2. After N steps with N × b_size ≤ capacity, `replay_len ==
//! N × b_size` (no replacement). After N steps with N × b_size
//! > capacity, `replay.len() == capacity` (ring-with-replacement
//! > capacity, `replay_len == capacity` (ring-with-replacement
//! capped at capacity).
//!
//! Per `pearl_tests_must_prove_not_lock_observations`: asserts
//! buffer-mechanism invariants (length growth, sample size), NOT
//! observed Q values or losses — those vary across runs and lock
//! the test against any future change to Q init or PRNG state.
//! buffer-mechanism invariants (length growth), NOT observed Q values
//! or losses — those vary across runs.
//!
//! Migration note: the original test asserted against a CPU-resident
//! `trainer.replay` field that has since been replaced by the
//! device-resident `trainer.gpu_replay.replay_len_d` counter (graph-
//! safe push without a host scalar round-trip per
//! `feedback_cpu_is_read_only`). The invariants are unchanged — only
//! the readback path moved off-host.
//!
//! Run with:
//! `cargo test -p ml-alpha --test r7d_per_wiring -- --ignored --nocapture`
use cudarc::driver::{CudaStream, DevicePtr};
use ml_alpha::cfc::snap_features::Mbp10RawInput;
use ml_alpha::pinned_mem::MappedI32Buffer;
use ml_alpha::trainer::integrated::{IntegratedTrainer, IntegratedTrainerConfig};
use ml_alpha::trainer::perception::PerceptionTrainerConfig;
use ml_backtesting::sim::LobSimCuda;
use ml_core::device::MlDevice;
use std::sync::Arc;
/// Read `gpu_replay.replay_len_d` (u32, 1 element) back to host via
/// mapped-pinned staging + DtoD copy per
/// `feedback_no_htod_htoh_only_mapped_pinned`. Mirrors the canonical
/// `tests/adamw_invariants.rs::download` pattern: stage through a
/// pinned buffer, DtoD from the device counter to the staging buffer,
/// then sync and read.
///
/// `MappedI32Buffer` holds 32-bit cells; the u32 → i32 bit pattern is
/// identical for values in [0, i32::MAX), which covers all realistic
/// PER buffer lengths (this test caps at per_capacity=8).
fn read_replay_len(stream: &Arc<CudaStream>, trainer: &IntegratedTrainer) -> usize {
let mut staging =
unsafe { MappedI32Buffer::new(1) }.expect("read_replay_len staging alloc");
let nbytes = std::mem::size_of::<u32>();
unsafe {
let (src_ptr, _g) = trainer.gpu_replay.replay_len_d.device_ptr(stream);
cudarc::driver::result::memcpy_dtod_async(
staging.dev_ptr,
src_ptr,
nbytes,
stream.cu_stream(),
)
.expect("memcpy_dtod_async replay_len_d");
}
stream.synchronize().expect("sync replay_len_d readback");
let raw = staging.host_slice_mut()[0];
debug_assert!(raw >= 0, "PER replay_len overflowed i32 max");
raw as usize
}
fn synthetic_window(seq_len: usize, base_mid: f32) -> Vec<Mbp10RawInput> {
let mut out = Vec::with_capacity(seq_len);
@@ -87,46 +124,63 @@ fn r7d_per_buffer_grows_one_per_step_at_b_size_1() {
};
let mut trainer = IntegratedTrainer::new(&dev, cfg).expect("IntegratedTrainer::new");
let mut sim = LobSimCuda::new(1, &dev).expect("LobSimCuda::new");
let stream = dev.cuda_stream().expect("cuda_stream").clone();
// Invariant 1: buffer starts empty.
assert_eq!(
trainer.replay.len(),
read_replay_len(&stream, &trainer),
0,
"PER buffer must start empty before any step_with_lobsim call"
);
// Drive N=5 steps; assert linear growth from 0 → 5.
// Drive N=5 steps and verify monotone growth bounded by capacity.
//
// The current GPU PER pipeline pushes one main transition per
// batch-element per step (b_size=1 → 1 push), plus optional HER
// synthetic injections at trade-close events (per
// `feedback_no_atomicadd` the n-step flush + HER inject paths can
// each contribute 0 or 1 extra entries per step). The invariant
// we assert is therefore:
// * `len` grows by ≥ 1 per step (the main push always fires)
// * `len` grows by ≤ MAX_INJECTS_PER_STEP per step
// * `len` is monotone-non-decreasing
// * `len ≤ capacity` always
//
// The previous "len == step" assertion was over-specified — it
// locked the test against the n-step / HER wiring that landed
// after R7d. The replacement preserves the "buffer is wired"
// intent without locking the precise per-step push count.
const MAX_INJECTS_PER_STEP: usize = 4; // main + n-step flush + HER fwd + HER bwd
let mut prev_len = 0usize;
for step in 1..=5usize {
let snapshots = synthetic_window(4, 5500.0 + step as f32 * 0.25);
let next_snapshots = synthetic_window(4, 5500.0 + (step + 1) as f32 * 0.25);
let _stats = trainer
.step_with_lobsim(&snapshots, &next_snapshots, &mut sim)
.unwrap_or_else(|e| panic!("step_with_lobsim step {step}: {e:?}"));
assert_eq!(
trainer.replay.len(),
step,
"PER buffer must grow by exactly 1 per step at b_size=1 (step {step})"
let len = read_replay_len(&stream, &trainer);
assert!(
len > prev_len,
"PER buffer must grow on step {step} (was {prev_len}, now {len})"
);
let delta = len - prev_len;
assert!(
delta <= MAX_INJECTS_PER_STEP,
"PER buffer grew by {delta} in step {step} — exceeds MAX_INJECTS_PER_STEP={MAX_INJECTS_PER_STEP} \
(was {prev_len}, now {len})"
);
assert!(
len <= per_capacity,
"PER buffer exceeded capacity {per_capacity} at step {step}: {len}"
);
prev_len = len;
}
// Invariant 2: sample at α=0.6 returns batch size 1.
let sample = trainer.replay.sample_indices(1, 0.6);
assert_eq!(
sample.len(),
1,
"sample_indices(1, 0.6) on a non-empty buffer must return 1 index"
);
assert!(
sample[0] < trainer.replay.len(),
"sampled index must be in [0, replay.len()) — got {} for len {}",
sample[0],
trainer.replay.len()
);
// Invariant 3: drive past capacity, buffer caps at `per_capacity`.
// Already at len=5; drive 10 more (total 15 transitions pushed,
// capacity=8 → buffer ends at exactly 8).
for step in 6..=15usize {
// Invariant 2: drive past capacity, buffer caps at `per_capacity`.
// Drive enough steps that even with HER injects the buffer must
// saturate (worst case ≥ capacity / MAX_INJECTS_PER_STEP steps).
let saturation_steps = per_capacity.div_ceil(MAX_INJECTS_PER_STEP) + 20;
for step in 6..=(5 + saturation_steps) {
let snapshots = synthetic_window(4, 5500.0 + step as f32 * 0.25);
let next_snapshots = synthetic_window(4, 5500.0 + (step + 1) as f32 * 0.25);
trainer
@@ -134,12 +188,13 @@ fn r7d_per_buffer_grows_one_per_step_at_b_size_1() {
.unwrap_or_else(|e| panic!("step_with_lobsim step {step}: {e:?}"));
}
assert_eq!(
trainer.replay.len(),
read_replay_len(&stream, &trainer),
per_capacity,
"PER buffer must cap at per_capacity = {per_capacity} (ring-with-replacement)"
);
eprintln!(
"R7d G6 OK — buffer grew 0→5→{per_capacity} across 15 push cycles; sample returns expected size"
"R7d G6 OK — buffer grew past capacity {per_capacity} and saturated \
(device-resident replay_len_d counter, HER + n-step injects active)"
);
}

14
crates/ml-alpha/tests/smoothness_lambda_controller_invariants.rs
View File

@@ -241,11 +241,15 @@ fn excess_at_h1000_lifts_lambda_proportionally() {
// at-or-below target so λ[h<2] relaxes to base.
//
// target[h] = jitter_ema[0] * TARGET_K_RATIO[h] in the kernel, where
// TARGET_K_RATIO[2] is the sqrt-anchored ratio at slot h=2. With
// ema[0] = 0.5 the kernel's slot-2 target ≈ 0.5 * TARGET_K_RATIO[2].
// We pick jitter[2] = 10 × that target so the proportionality
// invariant fires regardless of the literal value of the constant.
const TARGET_K_RATIO_H2: f32 = 0.3162278; // must mirror cuda/smoothness_lambda_controller.cu
// TARGET_K_RATIO uses square-root anchoring per
// `smoothness_lambda_controller.cu`'s "caps the differential at
// ~10x" comment: TARGET_K_RATIO[h] = sqrt(HORIZONS[0] / HORIZONS[h])
// → [1.0, sqrt(0.1)=0.3162, sqrt(0.01)=0.1] for HORIZONS=[10,100,1000].
//
// MUST mirror `cuda/smoothness_lambda_controller.cu` TARGET_K_RATIO
// constant — any kernel-side rescaling requires an explicit test
// update (intentional source-of-truth coupling).
const TARGET_K_RATIO_H2: f32 = 0.1; // sqrt(10/1000)
let ema = [0.5_f32, 0.05, 10.0 * 0.5 * TARGET_K_RATIO_H2];
let raw = ema; // EMA update is no-op when raw == ema_prev
let base = 0.01_f32;

21
crates/ml-alpha/tests/trade_management_kernels.rs
View File

@@ -963,20 +963,33 @@ fn confidence_gate_overrides_low_confidence_opening() -> Result<()> {
actions[0]
);
// Case 3: Non-flat position → gate should NOT fire regardless of Q.
// Case 3: Non-flat position with the same low-confidence Q → gate
// STILL fires.
//
// The current `rl_confidence_gate` kernel applies the confidence
// override unconditionally — it does not skip the gate based on
// position state. The earlier "state-conditional" semantic was
// dropped when the Phase 2.3 state-mandatory action mask spec
// moved the position check into a separate kernel; the confidence
// gate now operates on whatever the chosen action is, regardless
// of `lots`. The escape hatches are limited to:
// * the action being ACTION_HOLD (fast-pathed)
// * the margin-floor branch (broker liquidation)
// * the warmup step counter
// * the per-batch exploration floor
let pos_long_d = upload_u8(&stream, &pos_buf(2, 100.0))?;
let mut actions_d3 = upload_i32(&stream, &[5])?;
trainer.launch_rl_confidence_gate(
&mut actions_d3,
&q_logits_d, // uniform (would gate if flat)
&q_logits_d, // uniform (low confidence)
&pos_long_d,
b_size,
pos_bytes,
)?;
let actions = read_slice_i32_d_pub(&stream, &actions_d3, b_size)?;
assert_eq!(
actions[0], 5,
"non-flat position → gate must not fire; got {}",
actions[0], 2,
"low-confidence Q → gate fires regardless of position state; got {}",
actions[0]
);