4230 lines
189 KiB
Rust
4230 lines
189 KiB
Rust
// crates/ml/tests/sp15_phase1_oracle_tests.rs
|
||
//! SP15 Phase 1 oracle tests — honest numbers (sharpe, cost-net, DD,
|
||
//! baselines, dd_pct trunk grounding, test slice consumption).
|
||
|
||
#![allow(unsafe_code)] // CUDA kernel launch + mapped-pinned memory.
|
||
|
||
#[cfg(test)]
|
||
mod gpu {
|
||
// Per-task tests append here; see plan §Phase 1 task 1.X.
|
||
use std::sync::Arc;
|
||
|
||
use cudarc::driver::{CudaContext, CudaStream};
|
||
use ml::cuda_pipeline::gpu_dqn_trainer::{
|
||
launch_sp15_baseline_buyhold, launch_sp15_baseline_hold_only,
|
||
launch_sp15_baseline_naive_momentum, launch_sp15_baseline_naive_reversion,
|
||
launch_sp15_cost_net_sharpe, launch_sp15_dd_state, launch_sp15_dd_state_reduce,
|
||
launch_sp15_position_history_derivation, launch_sp15_sharpe_per_bar,
|
||
SP15_BASELINE_KERNELS_CUBIN, SP15_COST_NET_SHARPE_CUBIN,
|
||
// SP15 Wave 5 follow-up (2026-05-07): cubin re-exports for the
|
||
// pre-loaded `&CudaFunction` launcher migration. See
|
||
// `load_sp15_kernel` helper below.
|
||
SP15_DD_STATE_CUBIN, SP15_DD_STATE_REDUCE_CUBIN,
|
||
SP15_DD_TRAJECTORY_CUBIN, SP15_FINAL_REWARD_CUBIN,
|
||
SP15_PLASTICITY_INJECTION_CUBIN,
|
||
};
|
||
use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedI32Buffer};
|
||
use ml::cuda_pipeline::sp15_isv_slots::{
|
||
ALPHA_SPLIT_INDEX, COST_PER_BAR_AVG_INDEX, DD_CURRENT_INDEX, DD_MAX_INDEX,
|
||
DD_PCT_INDEX, DD_PERSISTENCE_INDEX, DD_PERSISTENCE_MAX_INDEX,
|
||
DD_RECOVERY_BARS_INDEX, OFI_IMPACT_LAMBDA_INDEX, SP15_SLOT_END,
|
||
};
|
||
|
||
/// Upper bound for any ISV index touched by SP15 oracle tests. Matches
|
||
/// `gpu_dqn_trainer::ISV_TOTAL_DIM` (which is `pub(crate)` and not
|
||
/// reachable from integration tests). The `sp15_isv_slots` regression
|
||
/// test `all_sp15_slots_fit_within_isv_total_dim` enforces that
|
||
/// `SP15_SLOT_END <= ISV_TOTAL_DIM`, so allocating an ISV buffer of
|
||
/// length `SP15_SLOT_END` is large enough for every SP15 slot the
|
||
/// kernel reads or writes.
|
||
const ISV_LEN: usize = SP15_SLOT_END;
|
||
|
||
/// Build a CUDA stream against device 0. Mirrors the helper in
|
||
/// `sp4_producer_unit_tests.rs` and `distributional_q_tests.rs`. If no
|
||
/// GPU is available the call panics — every test in this module is
|
||
/// `#[ignore]`-gated so CPU-only CI never reaches this path.
|
||
fn make_test_stream() -> Arc<CudaStream> {
|
||
let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
|
||
ctx.default_stream()
|
||
}
|
||
|
||
/// SP15 Wave 5 follow-up (2026-05-07): pre-load a SP15 kernel
|
||
/// handle on the test stream. The migrated launchers
|
||
/// (`launch_sp15_dd_state` / `launch_sp15_dd_state_reduce` /
|
||
/// `launch_sp15_dd_trajectory_decreasing` /
|
||
/// `launch_sp15_final_reward` /
|
||
/// `launch_sp15_plasticity_injection`) take a pre-loaded
|
||
/// `&CudaFunction` parameter — see field-doc on
|
||
/// `GpuExperienceCollector::exp_sp15_dd_state_kernel` for
|
||
/// rationale (graph-capture safety + CUDA_ERROR_ILLEGAL_ADDRESS
|
||
/// avoidance in subprocess child contexts). This helper inlines
|
||
/// the load at the call site so per-launch kernel resolution still
|
||
/// happens in tests (graph capture isn't a concern in oracle
|
||
/// scaffolds), while the production callers use struct-cached
|
||
/// handles via the collector.
|
||
fn load_sp15_kernel(
|
||
stream: &Arc<CudaStream>,
|
||
cubin: &'static [u8],
|
||
symbol: &str,
|
||
) -> cudarc::driver::CudaFunction {
|
||
let module = stream.context()
|
||
.load_cubin(cubin.to_vec())
|
||
.unwrap_or_else(|e| panic!("load sp15 cubin for {symbol}: {e}"));
|
||
module.load_function(symbol)
|
||
.unwrap_or_else(|e| panic!("load sp15 kernel symbol {symbol}: {e}"))
|
||
}
|
||
|
||
/// Test 1.1.a — synthetic zero-mean alternating PnL: mean=0, std=1, sharpe=0.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn unified_sharpe_kernel_zero_mean() {
|
||
let stream = make_test_stream();
|
||
|
||
let pnl: Vec<f32> = (0..1000)
|
||
.map(|i| if i % 2 == 0 { 1.0 } else { -1.0 })
|
||
.collect();
|
||
// Safety: CUDA context active via `make_test_stream` above.
|
||
let pnl_buf = unsafe { MappedF32Buffer::new(pnl.len()) }
|
||
.expect("alloc MappedF32Buffer for sharpe input");
|
||
pnl_buf.write_from_slice(&pnl);
|
||
let out_buf = unsafe { MappedF32Buffer::new(3) }
|
||
.expect("alloc MappedF32Buffer for sharpe output");
|
||
|
||
launch_sp15_sharpe_per_bar(
|
||
&stream,
|
||
pnl_buf.dev_ptr,
|
||
pnl.len() as i32,
|
||
out_buf.dev_ptr,
|
||
)
|
||
.expect("launch sharpe_per_bar_kernel");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after sharpe_per_bar_kernel launch");
|
||
|
||
let out = out_buf.read_all();
|
||
assert!(out[0].abs() < 1e-3, "mean={}, expected ~0", out[0]);
|
||
assert!((out[1] - 1.0).abs() < 1e-2, "std={}, expected ~1.0", out[1]);
|
||
assert!(out[2].abs() < 1e-3, "sharpe={}, expected ~0", out[2]);
|
||
}
|
||
|
||
/// Test 1.1.b — positive drift: mean=0.5, std≈1.0, sharpe≈0.5.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn unified_sharpe_kernel_positive_drift() {
|
||
let stream = make_test_stream();
|
||
|
||
let pnl: Vec<f32> = (0..1000)
|
||
.map(|i| {
|
||
let noise = if i % 2 == 0 { 1.0 } else { -1.0 };
|
||
0.5 + noise
|
||
})
|
||
.collect();
|
||
// Safety: CUDA context active via `make_test_stream` above.
|
||
let pnl_buf = unsafe { MappedF32Buffer::new(pnl.len()) }
|
||
.expect("alloc MappedF32Buffer for sharpe input");
|
||
pnl_buf.write_from_slice(&pnl);
|
||
let out_buf = unsafe { MappedF32Buffer::new(3) }
|
||
.expect("alloc MappedF32Buffer for sharpe output");
|
||
|
||
launch_sp15_sharpe_per_bar(
|
||
&stream,
|
||
pnl_buf.dev_ptr,
|
||
pnl.len() as i32,
|
||
out_buf.dev_ptr,
|
||
)
|
||
.expect("launch sharpe_per_bar_kernel");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after sharpe_per_bar_kernel launch");
|
||
|
||
let out = out_buf.read_all();
|
||
assert!((out[2] - 0.5).abs() < 0.05, "sharpe={}, expected ~0.5", out[2]);
|
||
}
|
||
|
||
/// Test 1.2 — single round-trip oracle for the cost-net sharpe kernel.
|
||
///
|
||
/// Setup: n=10 bars, position +1 contract held throughout, half_spread
|
||
/// 0.25, commission $1.50/RT, ofi=0 (zero out the OFI-impact term so
|
||
/// the spread + commission contributions are isolated).
|
||
/// Side events: bar 0 (entry, side_ind=1) and bar 5 (exit,
|
||
/// side_ind=1, rt_ind=1). Gross PnL: 10.0 booked at bar 5 (the
|
||
/// close-out bar).
|
||
///
|
||
/// Expected per spec §6.2:
|
||
/// total cost = 1.50 commission + 0.25 (entry half-spread)
|
||
/// + 0.25 (exit half-spread) = 2.00
|
||
/// cost_per_bar_avg = 2.00 / 10 = 0.20
|
||
/// mean(pnl) = 10.0 / 10 = 1.0
|
||
/// mean_pnl_net = 1.0 − 0.20 = 0.80
|
||
///
|
||
/// Validates:
|
||
/// • Per-side cost semantics (entry + exit = full spread per RT).
|
||
/// • Commission charged at close (rt_ind=1 only on bar 5).
|
||
/// • OFI-impact term scales with `|position|·|ofi|·side_ind`
|
||
/// (here zero because `ofi=0`).
|
||
/// • COST_PER_BAR_AVG_INDEX (slot 408) emitted to ISV.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn cost_net_sharpe_round_trip_charges_full_spread() {
|
||
let stream = make_test_stream();
|
||
let n = 10usize;
|
||
|
||
// Inputs.
|
||
let mut pnl = vec![0.0f32; n];
|
||
pnl[5] = 10.0;
|
||
let half_spread = vec![0.25f32; n];
|
||
let ofi = vec![0.0f32; n];
|
||
let position = vec![1.0f32; n];
|
||
// SP15 Wave 3b: side_ind / rt_ind are now f32 (0.0/1.0) so the
|
||
// cost_net_sharpe kernel signature matches the f32 streams emitted
|
||
// by the position_history_derivation_kernel; no u32→f32 adapter.
|
||
let mut side_ind = vec![0.0f32; n];
|
||
side_ind[0] = 1.0;
|
||
side_ind[5] = 1.0;
|
||
let mut rt_ind = vec![0.0f32; n];
|
||
rt_ind[5] = 1.0;
|
||
|
||
// Safety: CUDA context active via `make_test_stream` above.
|
||
let pnl_buf = unsafe { MappedF32Buffer::new(n) }
|
||
.expect("alloc MappedF32Buffer for pnl");
|
||
pnl_buf.write_from_slice(&pnl);
|
||
let half_spread_buf = unsafe { MappedF32Buffer::new(n) }
|
||
.expect("alloc MappedF32Buffer for half_spread");
|
||
half_spread_buf.write_from_slice(&half_spread);
|
||
let ofi_buf = unsafe { MappedF32Buffer::new(n) }
|
||
.expect("alloc MappedF32Buffer for ofi");
|
||
ofi_buf.write_from_slice(&ofi);
|
||
let position_buf = unsafe { MappedF32Buffer::new(n) }
|
||
.expect("alloc MappedF32Buffer for position");
|
||
position_buf.write_from_slice(&position);
|
||
let side_ind_buf = unsafe { MappedF32Buffer::new(n) }
|
||
.expect("alloc MappedF32Buffer for side_ind");
|
||
side_ind_buf.write_from_slice(&side_ind);
|
||
let rt_ind_buf = unsafe { MappedF32Buffer::new(n) }
|
||
.expect("alloc MappedF32Buffer for rt_ind");
|
||
rt_ind_buf.write_from_slice(&rt_ind);
|
||
|
||
// ISV bus seeded with OFI_IMPACT_LAMBDA=0 (mirrors the spec's
|
||
// zero-OFI test variable). The constructor seeds slot 407 to 2e-4
|
||
// in production; here we override to 0 so the spread term is the
|
||
// sole non-commission cost contributor.
|
||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc MappedF32Buffer for isv bus");
|
||
let mut isv_init = vec![0.0f32; ISV_LEN];
|
||
isv_init[OFI_IMPACT_LAMBDA_INDEX] = 0.0;
|
||
isv_buf.write_from_slice(&isv_init);
|
||
|
||
let out_buf = unsafe { MappedF32Buffer::new(3) }
|
||
.expect("alloc MappedF32Buffer for cost-net sharpe output");
|
||
|
||
// SP15 Wave 5 follow-up (2026-05-07): the launcher now consumes a
|
||
// pre-loaded `&CudaFunction` (graph-safe + cross-context-stable),
|
||
// so tests resolve the symbol once up-front rather than relying
|
||
// on the launcher itself.
|
||
let cost_net_module = stream
|
||
.context()
|
||
.load_cubin(SP15_COST_NET_SHARPE_CUBIN.to_vec())
|
||
.expect("load sp15_cost_net_sharpe cubin");
|
||
let cost_net_kernel = cost_net_module
|
||
.load_function("cost_net_sharpe_kernel")
|
||
.expect("load cost_net_sharpe_kernel symbol");
|
||
|
||
launch_sp15_cost_net_sharpe(
|
||
&stream,
|
||
pnl_buf.dev_ptr,
|
||
half_spread_buf.dev_ptr,
|
||
ofi_buf.dev_ptr,
|
||
side_ind_buf.dev_ptr,
|
||
rt_ind_buf.dev_ptr,
|
||
position_buf.dev_ptr,
|
||
n as i32,
|
||
/* commission_per_rt = */ 1.50,
|
||
isv_buf.dev_ptr,
|
||
out_buf.dev_ptr,
|
||
&cost_net_kernel,
|
||
)
|
||
.expect("launch cost_net_sharpe_kernel");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after cost_net_sharpe_kernel launch");
|
||
|
||
let isv_host = isv_buf.read_all();
|
||
let cost_per_bar_avg = isv_host[COST_PER_BAR_AVG_INDEX];
|
||
// Total cost = 1.50 commission + 0.25 + 0.25 = 2.00 over 10 bars
|
||
// → cost_per_bar_avg = 0.20.
|
||
assert!(
|
||
(cost_per_bar_avg - 0.20).abs() < 1e-3,
|
||
"cost_per_bar_avg = {}, expected 0.20",
|
||
cost_per_bar_avg
|
||
);
|
||
|
||
let out = out_buf.read_all();
|
||
// mean_pnl_net = mean(pnl) − mean(cost) = 1.0 − 0.20 = 0.80.
|
||
assert!(
|
||
(out[0] - 0.80).abs() < 1e-2,
|
||
"mean_pnl_net = {}, expected 0.80",
|
||
out[0]
|
||
);
|
||
}
|
||
|
||
/// Test 1.3 — synthetic equity curve drives DD slots correctly.
|
||
///
|
||
/// SP15 Phase 1.3.b contract (post Path A): the kernel is READ-ONLY
|
||
/// on `pos_state` and consumes the env-0 row's `PS_PREV_EQUITY` /
|
||
/// `PS_PEAK_EQUITY` slots. In production these are maintained by
|
||
/// `experience_env_step`; the oracle test simulates that role by
|
||
/// writing prev_equity and the running peak directly into the env-0
|
||
/// portfolio-state row before each `launch_sp15_dd_state` call. This
|
||
/// matches the live-system invariant that env_step writes equity
|
||
/// first, then dd_state reads.
|
||
///
|
||
/// Equity curve (cumulative): [+100, +10, -20, +15, -10, +5]
|
||
/// step 1: equity=100, peak=100, dd=0, rec=0
|
||
/// step 2: equity=110, peak=110, dd=0, rec=0
|
||
/// step 3: equity=90, peak=110, dd=(110-90)/110 ≈ 0.182, rec=1
|
||
/// step 4: equity=105, peak=110, dd=(110-105)/110 ≈ 0.045, rec=2
|
||
/// step 5: equity=95, peak=110, dd=(110-95)/110 ≈ 0.136, rec=3
|
||
/// step 6: equity=100, peak=110, dd=(110-100)/110 ≈ 0.091, rec=4
|
||
///
|
||
/// Expected: max_dd ≈ 0.182 (hit at step 3), current_dd ≈ 0.091
|
||
/// (never recovers), recovery_bars = 4 at end. dd_pct = 0.091 / 0.20
|
||
/// ≈ 0.455 ∈ (0, 1].
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn dd_state_kernel_tracks_drawdown_correctly() {
|
||
// PORTFOLIO_STRIDE / PS_PEAK_EQUITY / PS_PREV_EQUITY constants
|
||
// mirror state_layout.cuh — the kernel reads
|
||
// `pos_state[env * PORTFOLIO_STRIDE + PS_*]`. Update here if
|
||
// the CUDA-side layout shifts (the layout fingerprint check
|
||
// enforces synchrony in production).
|
||
const PORTFOLIO_STRIDE: usize = 43;
|
||
const PS_PEAK_EQUITY: usize = 7;
|
||
const PS_PREV_EQUITY: usize = 9;
|
||
|
||
let stream = make_test_stream();
|
||
|
||
// Equity at the END of each step (cumulative running sum starting
|
||
// from 0 prev_equity — same series the original test used, just
|
||
// pre-summed to match the new "env_step writes equity first" contract).
|
||
let equity_per_step: [f32; 6] = [100.0, 110.0, 90.0, 105.0, 95.0, 100.0];
|
||
let dd_budget: f32 = 0.20;
|
||
|
||
// SP15 Phase 1.3.b-followup contract: per-env tile + reduction.
|
||
// Single-env config (n_envs=1) reproduces the legacy Path A
|
||
// semantics — the mean-aggregate over a single env equals the
|
||
// single env's value, so the existing assertions still hold.
|
||
const N_ENVS: usize = 1;
|
||
|
||
// ISV bus sized to SP15_SLOT_END (slots 401-406 + 443 written
|
||
// by the reduction kernel).
|
||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc MappedF32Buffer for isv bus");
|
||
isv_buf.write_from_slice(&vec![0.0f32; ISV_LEN]);
|
||
|
||
// Per-env DD tile [n_envs * 6] (Phase 1.3.b-followup).
|
||
let dd_tile_buf = unsafe { MappedF32Buffer::new(N_ENVS * 6) }
|
||
.expect("alloc MappedF32Buffer for dd_state_per_env tile");
|
||
dd_tile_buf.write_from_slice(&vec![0.0f32; N_ENVS * 6]);
|
||
|
||
// Position state buffer sized for `N_ENVS` envs at PORTFOLIO_STRIDE
|
||
// plus headroom. The kernel reads `pos_state[env * STRIDE + ...]`.
|
||
let pos_len = (N_ENVS * PORTFOLIO_STRIDE).max(64);
|
||
let pos_state_buf = unsafe { MappedF32Buffer::new(pos_len) }
|
||
.expect("alloc MappedF32Buffer for pos_state");
|
||
|
||
// Step the env-0 equity curve manually (mirrors what env_step
|
||
// does in production); after each write the dd_state + reduce
|
||
// kernels run back-to-back.
|
||
let mut peak: f32 = 0.0;
|
||
for &equity in &equity_per_step {
|
||
peak = peak.max(equity);
|
||
let mut pos_init = vec![0.0f32; pos_len];
|
||
pos_init[PS_PREV_EQUITY] = equity;
|
||
pos_init[PS_PEAK_EQUITY] = peak;
|
||
pos_state_buf.write_from_slice(&pos_init);
|
||
|
||
launch_sp15_dd_state(
|
||
&stream,
|
||
&load_sp15_kernel(&stream, SP15_DD_STATE_CUBIN, "dd_state_kernel"),
|
||
dd_budget,
|
||
N_ENVS as i32,
|
||
dd_tile_buf.dev_ptr,
|
||
pos_state_buf.dev_ptr,
|
||
)
|
||
.expect("launch dd_state_kernel");
|
||
// SP15 Phase 1.3.b-followup-B: pass null for dd_trajectory_per_env
|
||
// (this test exercises only the dd_state path; the per-env
|
||
// trajectory tile is exercised by dedicated dd_trajectory_*
|
||
// tests below — null here keeps ISV[439] untouched).
|
||
launch_sp15_dd_state_reduce(
|
||
&stream,
|
||
&load_sp15_kernel(&stream, SP15_DD_STATE_REDUCE_CUBIN, "dd_state_reduce_kernel"),
|
||
N_ENVS as i32,
|
||
dd_tile_buf.dev_ptr,
|
||
0, // dd_trajectory_per_env (null)
|
||
isv_buf.dev_ptr,
|
||
)
|
||
.expect("launch dd_state_reduce_kernel");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after dd_state + reduce kernel launches");
|
||
}
|
||
|
||
// Mean-aggregated ISV scalars across envs (single-env so mean ==
|
||
// the env's value; the existing assertions on the original
|
||
// env-0-canonical contract still hold).
|
||
let isv = isv_buf.read_all();
|
||
let max_dd = isv[DD_MAX_INDEX];
|
||
let current_dd = isv[DD_CURRENT_INDEX];
|
||
let recovery_bars = isv[DD_RECOVERY_BARS_INDEX];
|
||
let dd_pct = isv[DD_PCT_INDEX];
|
||
|
||
// (110 − 90) / 110 ≈ 0.1818.
|
||
assert!(
|
||
(max_dd - 0.1818).abs() < 0.01,
|
||
"max_dd = {}, expected ~0.182",
|
||
max_dd
|
||
);
|
||
// (110 − 100) / 110 ≈ 0.0909 — never fully recovered.
|
||
assert!(
|
||
current_dd > 0.0,
|
||
"current_dd = {}, expected > 0 (never recovered)",
|
||
current_dd
|
||
);
|
||
// Last new high-water mark was at step 2 → 4 below-peak steps after.
|
||
assert!(
|
||
recovery_bars >= 4.0,
|
||
"recovery_bars = {}, expected >= 4",
|
||
recovery_bars
|
||
);
|
||
// current_dd / dd_budget ≈ 0.091 / 0.20 ≈ 0.455.
|
||
assert!(
|
||
dd_pct > 0.0 && dd_pct <= 1.0,
|
||
"dd_pct = {}, expected in (0, 1]",
|
||
dd_pct
|
||
);
|
||
|
||
// Per-env tile contents must match the ISV scalars (single-env
|
||
// mean degenerates to the env's value). Verifies the producer-
|
||
// consumer contract between `dd_state_kernel` and
|
||
// `dd_state_reduce_kernel`.
|
||
let tile = dd_tile_buf.read_all();
|
||
assert!(
|
||
(tile[0] - current_dd).abs() < 1e-5,
|
||
"tile[0] (DD_CURRENT) = {}, expected {}",
|
||
tile[0], current_dd
|
||
);
|
||
assert!(
|
||
(tile[1] - max_dd).abs() < 1e-5,
|
||
"tile[1] (DD_MAX) = {}, expected {}",
|
||
tile[1], max_dd
|
||
);
|
||
assert!(
|
||
(tile[5] - dd_pct).abs() < 1e-5,
|
||
"tile[5] (DD_PCT) = {}, expected {}",
|
||
tile[5], dd_pct
|
||
);
|
||
}
|
||
|
||
/// SP15 Phase 1.3.b-followup (2026-05-07) — per-env behavioral oracle.
|
||
///
|
||
/// Two-env config: env-0 is in recovery (DD shrinking from a peak),
|
||
/// env-1 is deepening into DD. After running a per-step series:
|
||
///
|
||
/// * The per-env tile must reflect each env's INDEPENDENT
|
||
/// trajectory (env-0's tile slot for DD_CURRENT shrinks bar-
|
||
/// over-bar; env-1's slot grows bar-over-bar).
|
||
/// * The mean-aggregated ISV scalars must equal the cross-env
|
||
/// mean (validates `dd_state_reduce_kernel`'s mean rule).
|
||
/// * The max-aggregated ISV[DD_PERSISTENCE_MAX_INDEX=443] must
|
||
/// equal the WORSE env's persistence (validates the max rule).
|
||
///
|
||
/// This is the canonical "envs in different DD contexts thread
|
||
/// their own DD context independently" guarantee that motivated
|
||
/// the Path A → Path B migration.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn dd_state_per_env_diverge_independently() {
|
||
const PORTFOLIO_STRIDE: usize = 43;
|
||
const PS_PEAK_EQUITY: usize = 7;
|
||
const PS_PREV_EQUITY: usize = 9;
|
||
const N_ENVS: usize = 2;
|
||
|
||
let stream = make_test_stream();
|
||
let dd_budget: f32 = 0.20;
|
||
|
||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc MappedF32Buffer for isv bus");
|
||
isv_buf.write_from_slice(&vec![0.0f32; ISV_LEN]);
|
||
|
||
let dd_tile_buf = unsafe { MappedF32Buffer::new(N_ENVS * 6) }
|
||
.expect("alloc MappedF32Buffer for dd_state_per_env tile");
|
||
dd_tile_buf.write_from_slice(&vec![0.0f32; N_ENVS * 6]);
|
||
|
||
let pos_len = (N_ENVS * PORTFOLIO_STRIDE).max(128);
|
||
let pos_state_buf = unsafe { MappedF32Buffer::new(pos_len) }
|
||
.expect("alloc MappedF32Buffer for pos_state");
|
||
|
||
// Six-step series:
|
||
// env-0 (recovery from a deep DD): 100 → 110 (peak) → 88 (DD~20%) → 92 → 96 → 99
|
||
// env-1 (deepening DD): 100 → 100 (peak) → 95 → 90 → 85 → 80
|
||
// env-0's DD trajectory shrinks (recovery); env-1's grows.
|
||
let env0_equity: [f32; 6] = [100.0, 110.0, 88.0, 92.0, 96.0, 99.0];
|
||
let env1_equity: [f32; 6] = [100.0, 100.0, 95.0, 90.0, 85.0, 80.0];
|
||
|
||
let mut env0_peak: f32 = 0.0;
|
||
let mut env1_peak: f32 = 0.0;
|
||
for step in 0..6 {
|
||
env0_peak = env0_peak.max(env0_equity[step]);
|
||
env1_peak = env1_peak.max(env1_equity[step]);
|
||
let mut pos_init = vec![0.0f32; pos_len];
|
||
pos_init[0 * PORTFOLIO_STRIDE + PS_PREV_EQUITY] = env0_equity[step];
|
||
pos_init[0 * PORTFOLIO_STRIDE + PS_PEAK_EQUITY] = env0_peak;
|
||
pos_init[1 * PORTFOLIO_STRIDE + PS_PREV_EQUITY] = env1_equity[step];
|
||
pos_init[1 * PORTFOLIO_STRIDE + PS_PEAK_EQUITY] = env1_peak;
|
||
pos_state_buf.write_from_slice(&pos_init);
|
||
|
||
launch_sp15_dd_state(
|
||
&stream,
|
||
&load_sp15_kernel(&stream, SP15_DD_STATE_CUBIN, "dd_state_kernel"),
|
||
dd_budget,
|
||
N_ENVS as i32,
|
||
dd_tile_buf.dev_ptr,
|
||
pos_state_buf.dev_ptr,
|
||
)
|
||
.expect("launch dd_state_kernel (per-env)");
|
||
launch_sp15_dd_state_reduce(
|
||
&stream,
|
||
&load_sp15_kernel(&stream, SP15_DD_STATE_REDUCE_CUBIN, "dd_state_reduce_kernel"),
|
||
N_ENVS as i32,
|
||
dd_tile_buf.dev_ptr,
|
||
0, // dd_trajectory_per_env (null) — per-env divergence test exercises only dd_state, not trajectory
|
||
isv_buf.dev_ptr,
|
||
)
|
||
.expect("launch dd_state_reduce_kernel (per-env)");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after per-env dd_state launches");
|
||
}
|
||
|
||
let tile = dd_tile_buf.read_all();
|
||
// env-0: prev=99, peak=110 → current_dd = (110-99)/110 ≈ 0.1.
|
||
let env0_dd_current = tile[0 * 6 + 0];
|
||
let env0_dd_max = tile[0 * 6 + 1];
|
||
let env0_persistence = tile[0 * 6 + 3];
|
||
let env0_dd_pct = tile[0 * 6 + 5];
|
||
// env-1: prev=80, peak=100 → current_dd = (100-80)/100 = 0.2.
|
||
let env1_dd_current = tile[1 * 6 + 0];
|
||
let env1_dd_max = tile[1 * 6 + 1];
|
||
let env1_persistence = tile[1 * 6 + 3];
|
||
let env1_dd_pct = tile[1 * 6 + 5];
|
||
|
||
// env-0 current_dd ≈ 0.1, max ≈ 0.2 (hit at step 2).
|
||
assert!(
|
||
(env0_dd_current - 0.1).abs() < 0.01,
|
||
"env-0 current_dd = {}, expected ~0.1 (recovering: 99 vs peak 110)",
|
||
env0_dd_current
|
||
);
|
||
assert!(
|
||
(env0_dd_max - 0.2).abs() < 0.01,
|
||
"env-0 max_dd = {}, expected ~0.2 (peak DD at step 2: 88 vs 110)",
|
||
env0_dd_max
|
||
);
|
||
// env-1 current_dd ≈ 0.2, max ≈ 0.2 (currently at the bottom).
|
||
assert!(
|
||
(env1_dd_current - 0.2).abs() < 1e-5,
|
||
"env-1 current_dd = {}, expected 0.2 (deepening: 80 vs peak 100)",
|
||
env1_dd_current
|
||
);
|
||
assert!(
|
||
(env1_dd_max - 0.2).abs() < 1e-5,
|
||
"env-1 max_dd = {}, expected 0.2",
|
||
env1_dd_max
|
||
);
|
||
// env-1 persistence ≥ 4 (4 consecutive below-peak bars: steps 2-5).
|
||
// env-0 persistence ≥ 4 (4 consecutive below-peak bars: steps 2-5
|
||
// since recovery to the new HWM never happened).
|
||
assert!(
|
||
env1_persistence >= 4.0,
|
||
"env-1 persistence = {}, expected ≥ 4",
|
||
env1_persistence
|
||
);
|
||
assert!(
|
||
env0_persistence >= 4.0,
|
||
"env-0 persistence = {}, expected ≥ 4",
|
||
env0_persistence
|
||
);
|
||
|
||
// Mean-aggregated ISV scalars must equal cross-env mean.
|
||
let isv = isv_buf.read_all();
|
||
let mean_dd_current = isv[DD_CURRENT_INDEX];
|
||
let mean_dd_pct = isv[DD_PCT_INDEX];
|
||
let mean_persistence = isv[DD_PERSISTENCE_INDEX];
|
||
let max_persistence = isv[DD_PERSISTENCE_MAX_INDEX];
|
||
|
||
let expected_mean_current = (env0_dd_current + env1_dd_current) / 2.0;
|
||
assert!(
|
||
(mean_dd_current - expected_mean_current).abs() < 1e-5,
|
||
"ISV[DD_CURRENT] = {}, expected mean of env-0 + env-1 = {}",
|
||
mean_dd_current, expected_mean_current
|
||
);
|
||
let expected_mean_dd_pct = (env0_dd_pct + env1_dd_pct) / 2.0;
|
||
assert!(
|
||
(mean_dd_pct - expected_mean_dd_pct).abs() < 1e-5,
|
||
"ISV[DD_PCT] = {}, expected mean of env-0 + env-1 = {}",
|
||
mean_dd_pct, expected_mean_dd_pct
|
||
);
|
||
let expected_mean_persistence = (env0_persistence + env1_persistence) / 2.0;
|
||
assert!(
|
||
(mean_persistence - expected_mean_persistence).abs() < 1e-5,
|
||
"ISV[DD_PERSISTENCE] = {}, expected mean of env-0 + env-1 = {}",
|
||
mean_persistence, expected_mean_persistence
|
||
);
|
||
// Max-aggregated persistence: must equal the worse env.
|
||
let expected_max_persistence = env0_persistence.max(env1_persistence);
|
||
assert!(
|
||
(max_persistence - expected_max_persistence).abs() < 1e-5,
|
||
"ISV[DD_PERSISTENCE_MAX] = {}, expected max(env-0, env-1) = {}",
|
||
max_persistence, expected_max_persistence
|
||
);
|
||
}
|
||
|
||
/// Test 1.4.a — buyhold on a positive-drift price series.
|
||
///
|
||
/// Build a synthetic price walk with mean-per-bar return +0.5 and
|
||
/// alternating ±1 noise (std≈1 per bar). Buyhold with position=+1
|
||
/// every bar collects the per-bar return + noise, so the sharpe
|
||
/// (mean/std of per-bar PnL) should be near 0.5 (drift / noise).
|
||
/// Bounds chosen to absorb the n=1000 sample-size noise: lower 0.2
|
||
/// rules out a sharpe-near-zero degenerate baseline; upper 1.0
|
||
/// catches accidental zero-cost or reduction bugs that would
|
||
/// inflate the sharpe.
|
||
///
|
||
/// Wave 3a (2026-05-06): assertion reads from caller-provided output
|
||
/// buffer `[mean, std, raw_sharpe]` instead of ISV slot 409 (kernel
|
||
/// signature change).
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn baseline_buyhold_positive_on_drift() {
|
||
let stream = make_test_stream();
|
||
|
||
// Synthetic +drift price series: per-bar diff = 0.5 + alternating
|
||
// ±1 noise → mean(diff)=0.5, std(diff)=1 → sharpe≈0.5 for
|
||
// buyhold (position=+1 every bar).
|
||
let n = 1000usize;
|
||
let mut prices = vec![4500.0f32; n];
|
||
for i in 1..n {
|
||
let dz: f32 = if i % 2 == 0 { 1.0 } else { -1.0 };
|
||
prices[i] = prices[i - 1] + 0.5 + dz;
|
||
}
|
||
let half_spread = vec![0.0f32; n]; // zero spread for clean test
|
||
let ofi = vec![0.0f32; n];
|
||
|
||
// Safety: CUDA context active via `make_test_stream` above.
|
||
let prices_buf = unsafe { MappedF32Buffer::new(n) }
|
||
.expect("alloc MappedF32Buffer for prices");
|
||
prices_buf.write_from_slice(&prices);
|
||
let half_spread_buf = unsafe { MappedF32Buffer::new(n) }
|
||
.expect("alloc MappedF32Buffer for half_spread");
|
||
half_spread_buf.write_from_slice(&half_spread);
|
||
let ofi_buf = unsafe { MappedF32Buffer::new(n) }
|
||
.expect("alloc MappedF32Buffer for ofi");
|
||
ofi_buf.write_from_slice(&ofi);
|
||
|
||
let out_buf = unsafe { MappedF32Buffer::new(3) }
|
||
.expect("alloc MappedF32Buffer for baseline output [mean, std, raw_sharpe]");
|
||
out_buf.write_from_slice(&[0.0f32; 3]);
|
||
|
||
// SP15 Wave 5 follow-up (2026-05-07): pre-load kernel handle.
|
||
let baseline_module = stream
|
||
.context()
|
||
.load_cubin(SP15_BASELINE_KERNELS_CUBIN.to_vec())
|
||
.expect("load sp15_baseline_kernels cubin");
|
||
let buyhold_kernel = baseline_module
|
||
.load_function("baseline_buyhold_kernel")
|
||
.expect("load baseline_buyhold_kernel symbol");
|
||
|
||
launch_sp15_baseline_buyhold(
|
||
&stream,
|
||
prices_buf.dev_ptr,
|
||
half_spread_buf.dev_ptr,
|
||
ofi_buf.dev_ptr,
|
||
n as i32,
|
||
/* commission_per_rt = */ 0.0,
|
||
out_buf.dev_ptr,
|
||
&buyhold_kernel,
|
||
)
|
||
.expect("launch baseline_buyhold_kernel");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after baseline_buyhold_kernel launch");
|
||
|
||
let out = out_buf.read_all();
|
||
let sharpe = out[2];
|
||
assert!(
|
||
sharpe > 0.2,
|
||
"buyhold sharpe = {}, expected > 0.2 on +drift series (mean={}, std={})",
|
||
sharpe, out[0], out[1]
|
||
);
|
||
assert!(
|
||
sharpe < 1.0,
|
||
"buyhold sharpe = {}, expected < 1.0 (sanity bound; mean={}, std={})",
|
||
sharpe, out[0], out[1]
|
||
);
|
||
}
|
||
|
||
/// Test 1.4.b — hold_only on any price series emits sharpe = 0.
|
||
///
|
||
/// HoldOnly never enters the market, so per-bar PnL is identically
|
||
/// zero on every bar. The kernel's reduction yields mean=0 std=0
|
||
/// and the `(std > 0) ? mean / std : 0` guard returns 0, the
|
||
/// spec-mandated sentinel for an undefined sharpe (no trades).
|
||
///
|
||
/// Wave 3a (2026-05-06): assertion reads from caller-provided output
|
||
/// buffer `[mean, std, raw_sharpe]` instead of ISV slot 410.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn baseline_hold_only_emits_zero() {
|
||
let stream = make_test_stream();
|
||
let n = 1000usize;
|
||
let prices = vec![4500.0f32; n];
|
||
let half_spread = vec![0.0f32; n];
|
||
let ofi = vec![0.0f32; n];
|
||
|
||
// Safety: CUDA context active via `make_test_stream` above.
|
||
let prices_buf = unsafe { MappedF32Buffer::new(n) }
|
||
.expect("alloc MappedF32Buffer for prices");
|
||
prices_buf.write_from_slice(&prices);
|
||
let half_spread_buf = unsafe { MappedF32Buffer::new(n) }
|
||
.expect("alloc MappedF32Buffer for half_spread");
|
||
half_spread_buf.write_from_slice(&half_spread);
|
||
let ofi_buf = unsafe { MappedF32Buffer::new(n) }
|
||
.expect("alloc MappedF32Buffer for ofi");
|
||
ofi_buf.write_from_slice(&ofi);
|
||
let out_buf = unsafe { MappedF32Buffer::new(3) }
|
||
.expect("alloc MappedF32Buffer for baseline output");
|
||
out_buf.write_from_slice(&[0.0f32; 3]);
|
||
|
||
// SP15 Wave 5 follow-up (2026-05-07): pre-load kernel handle.
|
||
let baseline_module = stream
|
||
.context()
|
||
.load_cubin(SP15_BASELINE_KERNELS_CUBIN.to_vec())
|
||
.expect("load sp15_baseline_kernels cubin");
|
||
let hold_only_kernel = baseline_module
|
||
.load_function("baseline_hold_only_kernel")
|
||
.expect("load baseline_hold_only_kernel symbol");
|
||
|
||
launch_sp15_baseline_hold_only(
|
||
&stream,
|
||
prices_buf.dev_ptr,
|
||
half_spread_buf.dev_ptr,
|
||
ofi_buf.dev_ptr,
|
||
n as i32,
|
||
/* commission_per_rt = */ 0.0,
|
||
out_buf.dev_ptr,
|
||
&hold_only_kernel,
|
||
)
|
||
.expect("launch baseline_hold_only_kernel");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after baseline_hold_only_kernel launch");
|
||
|
||
let out = out_buf.read_all();
|
||
let sharpe = out[2];
|
||
assert!(
|
||
sharpe.abs() < 1e-5,
|
||
"hold_only sharpe = {}, expected ~0 (no trades; mean={}, std={})",
|
||
sharpe, out[0], out[1]
|
||
);
|
||
// Also verify mean/std are themselves zero — hold_only's per-bar
|
||
// PnL stream is identically zero, not just zero-mean noise.
|
||
assert!(
|
||
out[0].abs() < 1e-5,
|
||
"hold_only mean = {}, expected ~0 (no trades)",
|
||
out[0]
|
||
);
|
||
}
|
||
|
||
/// Test 1.4.d + 1.4.h — symmetry: naive_momentum and naive_reversion
|
||
/// are sign-flipped policies, so on the same price series their
|
||
/// sharpe values should sum to ~0. The series is deliberately
|
||
/// mean-reverting (reversion outperforms) but the symmetry test is
|
||
/// invariant to which side wins.
|
||
///
|
||
/// Wave 3a (2026-05-06): each launcher writes to its own output
|
||
/// buffer `[mean, std, raw_sharpe]` (was ISV slots 412/416).
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn baseline_momentum_reversion_symmetry() {
|
||
let stream = make_test_stream();
|
||
|
||
// Mean-reverting series: AR(1) with negative coefficient anchors
|
||
// prices around 4500. On this series naive_momentum picks the
|
||
// wrong direction and naive_reversion picks the right one — the
|
||
// sum near zero is the sharpe-symmetry invariant we test.
|
||
let n = 1000usize;
|
||
let mut prices = vec![4500.0f32; n];
|
||
for i in 1..n {
|
||
let pull = (4500.0 - prices[i - 1]) * 0.3; // strong reversion
|
||
let dz: f32 = if i % 2 == 0 { 1.0 } else { -1.0 };
|
||
prices[i] = prices[i - 1] + pull + dz;
|
||
}
|
||
let half_spread = vec![0.0f32; n];
|
||
let ofi = vec![0.0f32; n];
|
||
|
||
// Safety: CUDA context active via `make_test_stream` above.
|
||
let prices_buf = unsafe { MappedF32Buffer::new(n) }
|
||
.expect("alloc MappedF32Buffer for prices");
|
||
prices_buf.write_from_slice(&prices);
|
||
let half_spread_buf = unsafe { MappedF32Buffer::new(n) }
|
||
.expect("alloc MappedF32Buffer for half_spread");
|
||
half_spread_buf.write_from_slice(&half_spread);
|
||
let ofi_buf = unsafe { MappedF32Buffer::new(n) }
|
||
.expect("alloc MappedF32Buffer for ofi");
|
||
ofi_buf.write_from_slice(&ofi);
|
||
|
||
let momentum_out = unsafe { MappedF32Buffer::new(3) }
|
||
.expect("alloc MappedF32Buffer for momentum output");
|
||
momentum_out.write_from_slice(&[0.0f32; 3]);
|
||
let reversion_out = unsafe { MappedF32Buffer::new(3) }
|
||
.expect("alloc MappedF32Buffer for reversion output");
|
||
reversion_out.write_from_slice(&[0.0f32; 3]);
|
||
|
||
// SP15 Wave 5 follow-up (2026-05-07): pre-load kernel handles.
|
||
let baseline_module = stream
|
||
.context()
|
||
.load_cubin(SP15_BASELINE_KERNELS_CUBIN.to_vec())
|
||
.expect("load sp15_baseline_kernels cubin");
|
||
let momentum_kernel = baseline_module
|
||
.load_function("baseline_naive_momentum_kernel")
|
||
.expect("load baseline_naive_momentum_kernel symbol");
|
||
let reversion_kernel = baseline_module
|
||
.load_function("baseline_naive_reversion_kernel")
|
||
.expect("load baseline_naive_reversion_kernel symbol");
|
||
|
||
launch_sp15_baseline_naive_momentum(
|
||
&stream,
|
||
prices_buf.dev_ptr,
|
||
half_spread_buf.dev_ptr,
|
||
ofi_buf.dev_ptr,
|
||
n as i32,
|
||
0.0,
|
||
momentum_out.dev_ptr,
|
||
&momentum_kernel,
|
||
)
|
||
.expect("launch baseline_naive_momentum_kernel");
|
||
launch_sp15_baseline_naive_reversion(
|
||
&stream,
|
||
prices_buf.dev_ptr,
|
||
half_spread_buf.dev_ptr,
|
||
ofi_buf.dev_ptr,
|
||
n as i32,
|
||
0.0,
|
||
reversion_out.dev_ptr,
|
||
&reversion_kernel,
|
||
)
|
||
.expect("launch baseline_naive_reversion_kernel");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after baseline_naive_{momentum,reversion}_kernel launches");
|
||
|
||
let momentum_sharpe = momentum_out.read_all()[2];
|
||
let reversion_sharpe = reversion_out.read_all()[2];
|
||
|
||
// Sign-flipped policies: their sum should be near zero (anti-
|
||
// correlated picks on the same per-bar return stream). Bound 0.5
|
||
// absorbs the asymmetry from cost charges on action-change bars
|
||
// (each policy hits trade events on different bars).
|
||
assert!(
|
||
(momentum_sharpe + reversion_sharpe).abs() < 0.5,
|
||
"momentum + reversion = {} + {} = {}, expected ~0 (anti-correlated)",
|
||
momentum_sharpe,
|
||
reversion_sharpe,
|
||
momentum_sharpe + reversion_sharpe
|
||
);
|
||
}
|
||
|
||
/// Test 1.4.j (Wave 3a, 2026-05-06) — position_history derivation kernel
|
||
/// reconstructs per-bar position/side_ind/rt_ind from a known factored
|
||
/// action sequence.
|
||
///
|
||
/// Fixture: 1 window, 8 bars, b1=b2=b3=3 (production sizes). The
|
||
/// factored-action encoding is `dir * 27 + mag * 9 + order * 3 + urg`.
|
||
/// Per `state_layout.cuh`: DIR_SHORT=0, DIR_HOLD=1, DIR_LONG=2,
|
||
/// DIR_FLAT=3. So:
|
||
/// action= 0 → DIR_SHORT (position −1)
|
||
/// action= 27 → DIR_HOLD (keep prev)
|
||
/// action= 54 → DIR_LONG (position +1)
|
||
/// action= 81 → DIR_FLAT (position 0)
|
||
///
|
||
/// The action sequence below exercises every transition: Short→Hold→
|
||
/// Long→Hold→Flat→Long→Flat→Short. Expected per-bar derived state:
|
||
/// position : [-1, -1, +1, +1, 0, +1, 0, -1]
|
||
/// side_ind : [ 1, 0, 1, 0, 1, 1, 1, 1] /* changed-from-prev */
|
||
/// rt_ind : [ 0, 0, 0, 0, 1, 0, 1, 0] /* transitioned to flat */
|
||
/// (initial prev_position=0 — every window starts flat by convention.)
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn position_history_derivation_walks_action_sequence() {
|
||
let stream = make_test_stream();
|
||
|
||
const N_WINDOWS: usize = 1;
|
||
const MAX_LEN: usize = 8;
|
||
const B1: i32 = 3;
|
||
const B2: i32 = 3;
|
||
const B3: i32 = 3;
|
||
|
||
// Action sequence: Short, Hold, Long, Hold, Flat, Long, Flat, Short.
|
||
// dir codes: 0, 1, 2, 1, 3, 2, 3, 0 → action = dir * 27.
|
||
let actions: [i32; MAX_LEN] = [0, 27, 54, 27, 81, 54, 81, 0];
|
||
|
||
// Safety: CUDA context active via `make_test_stream` above.
|
||
let actions_buf = unsafe { MappedI32Buffer::new(N_WINDOWS * MAX_LEN) }
|
||
.expect("alloc MappedI32Buffer for actions_history");
|
||
actions_buf.write_from_slice(&actions);
|
||
|
||
let position_buf = unsafe { MappedF32Buffer::new(N_WINDOWS * MAX_LEN) }
|
||
.expect("alloc MappedF32Buffer for position_history");
|
||
position_buf.write_from_slice(&vec![0.0f32; N_WINDOWS * MAX_LEN]);
|
||
let side_ind_buf = unsafe { MappedF32Buffer::new(N_WINDOWS * MAX_LEN) }
|
||
.expect("alloc MappedF32Buffer for side_ind");
|
||
side_ind_buf.write_from_slice(&vec![0.0f32; N_WINDOWS * MAX_LEN]);
|
||
let rt_ind_buf = unsafe { MappedF32Buffer::new(N_WINDOWS * MAX_LEN) }
|
||
.expect("alloc MappedF32Buffer for rt_ind");
|
||
rt_ind_buf.write_from_slice(&vec![0.0f32; N_WINDOWS * MAX_LEN]);
|
||
|
||
launch_sp15_position_history_derivation(
|
||
&stream,
|
||
actions_buf.dev_ptr,
|
||
N_WINDOWS as i32,
|
||
MAX_LEN as i32,
|
||
B1,
|
||
B2,
|
||
B3,
|
||
position_buf.dev_ptr,
|
||
side_ind_buf.dev_ptr,
|
||
rt_ind_buf.dev_ptr,
|
||
)
|
||
.expect("launch sp15_position_history_derivation_kernel");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after position_history_derivation launch");
|
||
|
||
let position = position_buf.read_all();
|
||
let side_ind = side_ind_buf.read_all();
|
||
let rt_ind = rt_ind_buf.read_all();
|
||
|
||
let expected_position = [-1.0_f32, -1.0, 1.0, 1.0, 0.0, 1.0, 0.0, -1.0];
|
||
let expected_side_ind = [1.0_f32, 0.0, 1.0, 0.0, 1.0, 1.0, 1.0, 1.0];
|
||
let expected_rt_ind = [0.0_f32, 0.0, 0.0, 0.0, 1.0, 0.0, 1.0, 0.0];
|
||
|
||
for t in 0..MAX_LEN {
|
||
assert!(
|
||
(position[t] - expected_position[t]).abs() < 1e-5,
|
||
"position[{}] = {}, expected {}",
|
||
t, position[t], expected_position[t]
|
||
);
|
||
assert!(
|
||
(side_ind[t] - expected_side_ind[t]).abs() < 1e-5,
|
||
"side_ind[{}] = {}, expected {}",
|
||
t, side_ind[t], expected_side_ind[t]
|
||
);
|
||
assert!(
|
||
(rt_ind[t] - expected_rt_ind[t]).abs() < 1e-5,
|
||
"rt_ind[{}] = {}, expected {}",
|
||
t, rt_ind[t], expected_rt_ind[t]
|
||
);
|
||
}
|
||
}
|
||
|
||
/// Test 1.5 — bottleneck-aware dd_pct trunk-input concat (Wave 4.1a)
|
||
/// (LAYOUT FINGERPRINT BREAK — fingerprint marker
|
||
/// `TRUNK_INPUT_DD_PCT=sp15_wave_4_1a`).
|
||
///
|
||
/// Verifies the new `bn_tanh_concat_dd_kernel` correctly produces a
|
||
/// `[B, concat_dd_dim]` buffer where
|
||
/// `concat_dd_dim = bn_dim + portfolio_dim + 1`. The kernel folds the
|
||
/// dd_pct append into the same launch as the bn_tanh + portfolio concat:
|
||
/// - columns `[0, bn_dim)` = `tanh(bn_hidden[b, 0..bn_dim))`
|
||
/// - columns `[bn_dim, bn_dim+pd)` = `states[b, market_dim..market_dim+pd)`
|
||
/// - column `bn_dim + pd` = `isv[DD_PCT_INDEX=406]` (broadcast)
|
||
///
|
||
/// Spec correction (Wave 4.1a, per `feedback_trust_code_not_docs`): the
|
||
/// spec's "state_dim 48→49" is stale terminology pre-STATE_DIM
|
||
/// 48→112→128 evolution. Production `s1_input_dim` is
|
||
/// `bn_dim + (STATE_DIM - market_dim)` (= 102 with current
|
||
/// bn_dim=16, market_dim=42, STATE_DIM=128). Wave 4.1b will bump this
|
||
/// to 103 by flipping the 4 launch sites + buffer allocations.
|
||
///
|
||
/// This is a kernel-level oracle test (matches the Phase 1.1-1.4
|
||
/// pattern) — the trunk-grounding behavioral test (KL divergence
|
||
/// between forward_trunk(102-input) and forward_trunk(103-input))
|
||
/// is added in Wave 4.1c using the helpers defined at the bottom of
|
||
/// this module.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn bn_tanh_concat_dd_kernel_writes_dd_pct_column() {
|
||
use ml::cuda_pipeline::gpu_dqn_trainer::{launch_sp15_bn_concat_dd, DQN_UTILITY_CUBIN};
|
||
|
||
let stream = make_test_stream();
|
||
// SP15 Wave 4.1b OOB fix (2026-05-07): the launcher now consumes a
|
||
// pre-loaded `&CudaFunction` (graph-safe), so tests resolve the
|
||
// symbol once up-front rather than relying on the launcher itself.
|
||
let module = stream.context().load_cubin(DQN_UTILITY_CUBIN.to_vec())
|
||
.expect("load dqn_utility_kernels cubin");
|
||
let kernel = module.load_function("bn_tanh_concat_dd_kernel")
|
||
.expect("load bn_tanh_concat_dd_kernel symbol");
|
||
|
||
// Synthetic batch matches production layout: bn_dim=16,
|
||
// market_dim=42, STATE_DIM=128 → state_dim_padded=128 (no
|
||
// sub-128 pad in the current production config), portfolio_dim
|
||
// = STATE_DIM - market_dim = 86, concat_dd_dim = bn_dim +
|
||
// portfolio_dim + 1 = 16 + 86 + 1 = 103. Wave 4.1b bumps the
|
||
// production `s1_input_dim` from 102 → 103 (matches the +1
|
||
// from the dd_pct append).
|
||
let batch_size: usize = 4;
|
||
const BN_DIM: usize = 16;
|
||
const MARKET_DIM: usize = 42;
|
||
const STATE_DIM_PADDED: usize = 128;
|
||
let portfolio_dim: usize = STATE_DIM_PADDED - MARKET_DIM;
|
||
let concat_dd_dim: usize = BN_DIM + portfolio_dim + 1;
|
||
|
||
// Build bn_hidden [B, BN_DIM]: pre-tanh values; row b column j =
|
||
// 0.1 * (b + 1) * (j + 1) so each cell is small enough to keep
|
||
// tanh in the linear regime (well under 1.0 saturation).
|
||
let mut bn_hidden_init: Vec<f32> = vec![0.0; batch_size * BN_DIM];
|
||
for b in 0..batch_size {
|
||
for j in 0..BN_DIM {
|
||
bn_hidden_init[b * BN_DIM + j] = 0.1_f32 * ((b + 1) as f32) * ((j + 1) as f32);
|
||
}
|
||
}
|
||
|
||
// Build states [B, STATE_DIM_PADDED]: row b column i =
|
||
// (b * 1000 + i) as f32 for the entire padded width — only
|
||
// the portfolio slice [MARKET_DIM, MARKET_DIM + portfolio_dim)
|
||
// is read by the kernel.
|
||
let mut states: Vec<f32> = vec![0.0; batch_size * STATE_DIM_PADDED];
|
||
for b in 0..batch_size {
|
||
for i in 0..STATE_DIM_PADDED {
|
||
states[b * STATE_DIM_PADDED + i] = (b * 1000 + i) as f32;
|
||
}
|
||
}
|
||
|
||
// Safety: CUDA context active via `make_test_stream` above.
|
||
let bn_hidden_buf = unsafe { MappedF32Buffer::new(bn_hidden_init.len()) }
|
||
.expect("alloc MappedF32Buffer for bn_hidden");
|
||
bn_hidden_buf.write_from_slice(&bn_hidden_init);
|
||
let states_buf = unsafe { MappedF32Buffer::new(states.len()) }
|
||
.expect("alloc MappedF32Buffer for states");
|
||
states_buf.write_from_slice(&states);
|
||
|
||
// ISV bus: write 0.05 to slot DD_PCT_INDEX=406.
|
||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc MappedF32Buffer for ISV");
|
||
let mut isv_init = vec![0.0f32; ISV_LEN];
|
||
isv_init[DD_PCT_INDEX] = 0.05;
|
||
isv_buf.write_from_slice(&isv_init);
|
||
|
||
let out_len: usize = batch_size * concat_dd_dim;
|
||
let concat_buf = unsafe { MappedF32Buffer::new(out_len) }
|
||
.expect("alloc MappedF32Buffer for concat output");
|
||
|
||
launch_sp15_bn_concat_dd(
|
||
&stream,
|
||
&kernel,
|
||
bn_hidden_buf.dev_ptr,
|
||
states_buf.dev_ptr,
|
||
isv_buf.dev_ptr,
|
||
concat_buf.dev_ptr,
|
||
batch_size as i32,
|
||
BN_DIM as i32,
|
||
MARKET_DIM as i32,
|
||
STATE_DIM_PADDED as i32,
|
||
concat_dd_dim as i32,
|
||
)
|
||
.expect("launch bn_tanh_concat_dd_kernel");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after bn_tanh_concat_dd_kernel launch");
|
||
|
||
let concat = concat_buf.read_all();
|
||
let bn_hidden_after = bn_hidden_buf.read_all();
|
||
|
||
// Check 1: bn_tanh slice — output[b, 0..BN_DIM) ==
|
||
// tanh(input bn_hidden), and the in-place write semantics put
|
||
// the same tanh'd value in bn_hidden after the launch.
|
||
for b in 0..batch_size {
|
||
for j in 0..BN_DIM {
|
||
let pre = bn_hidden_init[b * BN_DIM + j];
|
||
let expected = pre.tanh();
|
||
let out_val = concat[b * concat_dd_dim + j];
|
||
let bn_val = bn_hidden_after[b * BN_DIM + j];
|
||
assert!(
|
||
(out_val - expected).abs() < 1e-6,
|
||
"concat[b={}, j={}] = {} ≠ tanh({}) = {} (bn_tanh column mismatch)",
|
||
b, j, out_val, pre, expected
|
||
);
|
||
assert!(
|
||
(bn_val - expected).abs() < 1e-6,
|
||
"bn_hidden[b={}, j={}] = {} ≠ tanh({}) = {} (in-place write mismatch)",
|
||
b, j, bn_val, pre, expected
|
||
);
|
||
}
|
||
}
|
||
|
||
// Check 2: portfolio passthrough slice — output[b, BN_DIM..BN_DIM+pd)
|
||
// == states[b, MARKET_DIM..MARKET_DIM+pd).
|
||
for b in 0..batch_size {
|
||
for k in 0..portfolio_dim {
|
||
let expected = states[b * STATE_DIM_PADDED + MARKET_DIM + k];
|
||
let out_val = concat[b * concat_dd_dim + BN_DIM + k];
|
||
assert!(
|
||
(out_val - expected).abs() < 1e-6,
|
||
"concat[b={}, BN_DIM+{}] = {} ≠ states[b={}, MARKET_DIM+{}] = {} \
|
||
(portfolio passthrough mismatch)",
|
||
b, k, out_val, b, k, expected
|
||
);
|
||
}
|
||
}
|
||
|
||
// Check 3: dd_pct broadcast column — output[b, concat_dd_dim - 1]
|
||
// == isv[DD_PCT_INDEX] = 0.05 for every batch row.
|
||
for b in 0..batch_size {
|
||
let last = concat[b * concat_dd_dim + (concat_dd_dim - 1)];
|
||
assert!(
|
||
(last - 0.05).abs() < 1e-6,
|
||
"concat[b={}, last={}] = {} ≠ isv[DD_PCT_INDEX] = 0.05 (broadcast mismatch)",
|
||
b, concat_dd_dim - 1, last
|
||
);
|
||
}
|
||
}
|
||
|
||
// SP15 Wave 2 (2026-05-06) — old standalone-composer oracle tests
|
||
// (`r_split_uses_sentinel_alpha_at_cold_start` +
|
||
// `r_quality_subtracts_explicit_cost`) DELETED with the deletion of
|
||
// the scalar `r_quality_discipline_split_kernel`. The α-blend stage
|
||
// is now exercised end-to-end by the fused-kernel oracle tests at
|
||
// the end of this module (`final_reward_alpha_blend_*`,
|
||
// `final_reward_dd_penalty_*`, `final_reward_dd_asymmetric_*`,
|
||
// `final_reward_skips_early_return_slots`) which drive
|
||
// `compute_sp15_final_reward_kernel` directly with a hand-crafted
|
||
// `out_rewards` + `r_discipline_out` + ISV bundle.
|
||
|
||
// SP15 Wave 2 (2026-05-06) — `r_quality_subtracts_explicit_cost`
|
||
// DELETED with the scalar composer. The Phase 3.2 cost-on-trade-
|
||
// close subtraction is no longer a separate stage: SP11's existing
|
||
// shaping (`cost_anneal_ptr` + inventory / churn / spread-cost
|
||
// terms in experience_kernels.cu) already charges the model at
|
||
// evaluation-time costs, and re-inserting an explicit cost
|
||
// subtraction at the SP15 layer would double-count.
|
||
|
||
// SP15 Wave 2 (2026-05-06) — `dd_penalty_quadratic_above_threshold`
|
||
// + `dd_penalty_zero_below_threshold` DELETED with the deletion of
|
||
// the standalone `dd_penalty_kernel`. The quadratic DD penalty stage
|
||
// is now exercised by the fused-kernel oracle tests below
|
||
// (`final_reward_dd_penalty_above_threshold`,
|
||
// `final_reward_dd_penalty_zero_below_threshold`).
|
||
|
||
/// SP15 Phase 3.4 — regret signal: when policy holds AND the ideal
|
||
/// trade is past `cost_t + trail_dist`, `regret_bar = ideal_pnl_missed
|
||
/// − cost_t` per spec §8.2 (3.4). REGRET_EMA also bootstraps from the
|
||
/// sentinel 0 to the first non-zero observation per
|
||
/// `pearl_first_observation_bootstrap`. Validates BOTH the regret
|
||
/// computation AND the EMA first-observation bootstrap in one launch.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn regret_logged_on_missed_profitable_trade() {
|
||
use ml::cuda_pipeline::sp15_isv_slots::REGRET_EMA_INDEX;
|
||
|
||
let stream = make_test_stream();
|
||
|
||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc MappedF32Buffer for isv bus");
|
||
let mut isv_init = vec![0.0f32; ISV_LEN];
|
||
isv_init[REGRET_EMA_INDEX] = 0.0; // sentinel — first-observation bootstrap
|
||
isv_buf.write_from_slice(&isv_init);
|
||
|
||
let regret_buf = unsafe { MappedF32Buffer::new(1) }
|
||
.expect("alloc MappedF32Buffer for regret_bar out");
|
||
regret_buf.write_from_slice(&[0.0_f32]);
|
||
|
||
// Policy held (action=0=Hold); ideal next-bar PnL = 5.0; cost = 1.0;
|
||
// trail = 0.5. Ideal exceeds cost+trail (5.0 > 1.5) → gate fires →
|
||
// regret_bar = 5.0 − 1.0 = 4.0. EMA: prev=0 (sentinel), regret_bar
|
||
// > 0 → first-observation bootstrap replaces directly: 0 → 4.0.
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_regret_signal(
|
||
&stream,
|
||
/*action*/ 0, // 0 = Hold (regret-eligible)
|
||
/*ideal_pnl_missed*/ 5.0,
|
||
/*cost_t*/ 1.0,
|
||
/*trail_dist*/ 0.5,
|
||
isv_buf.dev_ptr,
|
||
regret_buf.dev_ptr,
|
||
)
|
||
.expect("launch regret_signal_kernel");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after regret_signal_kernel launch");
|
||
|
||
let regret = regret_buf.read_all()[0];
|
||
assert!(
|
||
(regret - 4.0).abs() < 1e-5,
|
||
"regret = {regret}, expected 4.0"
|
||
);
|
||
|
||
// EMA bootstrapped from sentinel 0 → first observation 4.0 directly
|
||
// (per pearl_first_observation_bootstrap), NOT decayed via α blend.
|
||
let isv = isv_buf.read_all();
|
||
let ema = isv[REGRET_EMA_INDEX];
|
||
assert!(
|
||
(ema - 4.0).abs() < 1e-5,
|
||
"regret_ema = {ema}, expected 4.0 (first-observation bootstrap)"
|
||
);
|
||
}
|
||
|
||
/// SP15 Phase 3.4 — asymmetric gate: regret = 0 when action != Hold OR
|
||
/// when the ideal trade did not clear the cost+trail floor. Both gate
|
||
/// paths validated in one test sharing an ISV buffer — sentinel 0 in
|
||
/// the regret_buf ensures the kernel actually overwrites with zero
|
||
/// rather than skipping the store. Per spec §8.2 (3.4).
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn regret_zero_when_action_active_or_ideal_below_threshold() {
|
||
let stream = make_test_stream();
|
||
|
||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc MappedF32Buffer for isv bus");
|
||
let isv_init = vec![0.0f32; ISV_LEN];
|
||
isv_buf.write_from_slice(&isv_init);
|
||
|
||
let regret_buf = unsafe { MappedF32Buffer::new(1) }
|
||
.expect("alloc MappedF32Buffer for regret_bar out");
|
||
|
||
// Case A: action != Hold (action=1 = active, e.g. Long) — gate
|
||
// fails on the action condition; regret = 0 regardless of ideal.
|
||
regret_buf.write_from_slice(&[1.0_f32]); // sentinel non-zero
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_regret_signal(
|
||
&stream,
|
||
/*action*/ 1,
|
||
/*ideal_pnl_missed*/ 5.0,
|
||
/*cost_t*/ 1.0,
|
||
/*trail_dist*/ 0.5,
|
||
isv_buf.dev_ptr,
|
||
regret_buf.dev_ptr,
|
||
)
|
||
.expect("launch regret_signal_kernel (case A)");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after regret_signal_kernel launch (case A)");
|
||
let regret_a = regret_buf.read_all()[0];
|
||
assert!(
|
||
regret_a.abs() < 1e-9,
|
||
"case A: regret = {regret_a}, expected 0 (action != Hold)"
|
||
);
|
||
|
||
// Case B: action = Hold (0) but ideal_pnl_missed (1.0) ≤
|
||
// cost+trail (1.0 + 0.5 = 1.5) — gate fails on the
|
||
// ideal-above-floor condition; regret = 0.
|
||
regret_buf.write_from_slice(&[1.0_f32]); // sentinel non-zero
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_regret_signal(
|
||
&stream,
|
||
/*action*/ 0,
|
||
/*ideal_pnl_missed*/ 1.0,
|
||
/*cost_t*/ 1.0,
|
||
/*trail_dist*/ 0.5,
|
||
isv_buf.dev_ptr,
|
||
regret_buf.dev_ptr,
|
||
)
|
||
.expect("launch regret_signal_kernel (case B)");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after regret_signal_kernel launch (case B)");
|
||
let regret_b = regret_buf.read_all()[0];
|
||
assert!(
|
||
regret_b.abs() < 1e-9,
|
||
"case B: regret = {regret_b}, expected 0 (ideal <= cost+trail)"
|
||
);
|
||
}
|
||
|
||
// SP15 Phase 3.5.b (2026-05-06): the standalone `hold_floor_kernel.cu`
|
||
// + `launch_sp15_hold_floor` were deleted when the consumer wiring
|
||
// landed inline inside `experience_action_select`. The previous
|
||
// `hold_floor_bounded_sigmoid` oracle test that drove the standalone
|
||
// kernel directly is removed in the same commit per
|
||
// `feedback_no_legacy_aliases.md`. The bounded-sigmoid math is now
|
||
// exercised end-to-end by the new
|
||
// `action_select_applies_hold_floor_inline` test below
|
||
// (full action_select kernel call with deterministic e_dir → known
|
||
// entropy → expected Q[Hold] bump).
|
||
|
||
// SP15 Wave 2 (2026-05-06) — `dd_asymmetric_reward_gain_amplified_by_dd_pct`
|
||
// + `dd_asymmetric_reward_loss_unchanged` + `dd_asymmetric_reward_no_op_at_ath`
|
||
// DELETED with the standalone `dd_asymmetric_reward_kernel`. The
|
||
// gain-only DD-recovery boost stage is now exercised by the fused-
|
||
// kernel oracle tests below (`final_reward_dd_asymmetric_gain`,
|
||
// `final_reward_dd_asymmetric_loss`, `final_reward_dd_asymmetric_ath`).
|
||
|
||
/// SP15 Phase 3.5.3 — cooldown trips after K=5 consecutive losses,
|
||
/// sets COOLDOWN_BARS_REMAINING to M=20 (decremented to 19 in the
|
||
/// same kernel call by the per-bar decrement that runs after the
|
||
/// trigger). Per spec §9.2 (3.5.3): after K consecutive losing
|
||
/// trades, force Hold for M bars. K and M are ISV-tracked
|
||
/// (slots 433/434); initial sentinels K=5, M=20 hardcoded in the
|
||
/// constructor (ISV-driven K via `MEDIAN_STREAK_LENGTH=442`
|
||
/// producer is a follow-up sub-task).
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn cooldown_trips_after_k_losses() {
|
||
use ml::cuda_pipeline::sp15_isv_slots::{
|
||
COOLDOWN_BARS_REMAINING_INDEX, COOLDOWN_K_THRESHOLD_INDEX,
|
||
COOLDOWN_M_BARS_INDEX,
|
||
};
|
||
|
||
let stream = make_test_stream();
|
||
|
||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc MappedF32Buffer for isv bus");
|
||
let mut isv_init = vec![0.0f32; ISV_LEN];
|
||
isv_init[COOLDOWN_K_THRESHOLD_INDEX] = 5.0;
|
||
isv_init[COOLDOWN_M_BARS_INDEX] = 20.0;
|
||
isv_init[COOLDOWN_BARS_REMAINING_INDEX] = 0.0;
|
||
isv_buf.write_from_slice(&isv_init);
|
||
|
||
// 32-bit signed counter buffer (consecutive_losses) — but
|
||
// `MappedF32Buffer` is f32-only. Use a float counter (a
|
||
// non-negative integer represented as f32) — the kernel only
|
||
// performs `+= 1.0f` increments and `= 0.0f` resets, so f32
|
||
// representation is exact for any reachable streak length.
|
||
let consec_losses_buf = unsafe { MappedF32Buffer::new(1) }
|
||
.expect("alloc MappedF32Buffer for consec_losses");
|
||
consec_losses_buf.write_from_slice(&[0.0_f32]);
|
||
|
||
// Drive 5 losing trade-close events.
|
||
for _ in 0..5 {
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_cooldown(
|
||
&stream,
|
||
/*last_trade_outcome_loss*/ 1,
|
||
/*trade_close_event*/ 1,
|
||
isv_buf.dev_ptr,
|
||
consec_losses_buf.dev_ptr,
|
||
)
|
||
.expect("launch cooldown_kernel (5 losses)");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after cooldown_kernel launch (5 losses)");
|
||
}
|
||
|
||
let isv_host = isv_buf.read_all();
|
||
let remaining = isv_host[COOLDOWN_BARS_REMAINING_INDEX];
|
||
// After 5th loss the kernel orders: trigger sets remaining=M=20,
|
||
// resets consec_losses=0, then the per-bar decrement subtracts
|
||
// 1 in the same call → 19. Acceptable range [18, 20] per
|
||
// implementation-order-flexibility comment in the task spec.
|
||
assert!(
|
||
(18.0..=20.0).contains(&remaining),
|
||
"cooldown_remaining = {remaining} after 5 losses, expected ~19-20"
|
||
);
|
||
}
|
||
|
||
/// SP15 Phase 3.5.3 — winning trade resets `consecutive_losses`,
|
||
/// no cooldown trip. 4 losses then 1 win then 4 losses → no trip
|
||
/// (the streak is broken by the win, so the second 4-loss run does
|
||
/// not reach the K=5 trigger). Per spec §9.2 (3.5.3) post-amendment-2
|
||
/// fix: streak counter only updates on `trade_close_event != 0`;
|
||
/// the win path writes `consec_losses = 0.0f` resetting the streak.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn cooldown_no_trip_when_streak_broken() {
|
||
use ml::cuda_pipeline::sp15_isv_slots::{
|
||
COOLDOWN_BARS_REMAINING_INDEX, COOLDOWN_K_THRESHOLD_INDEX,
|
||
COOLDOWN_M_BARS_INDEX,
|
||
};
|
||
|
||
let stream = make_test_stream();
|
||
|
||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc MappedF32Buffer for isv bus");
|
||
let mut isv_init = vec![0.0f32; ISV_LEN];
|
||
isv_init[COOLDOWN_K_THRESHOLD_INDEX] = 5.0;
|
||
isv_init[COOLDOWN_M_BARS_INDEX] = 20.0;
|
||
isv_init[COOLDOWN_BARS_REMAINING_INDEX] = 0.0;
|
||
isv_buf.write_from_slice(&isv_init);
|
||
|
||
let consec_losses_buf = unsafe { MappedF32Buffer::new(1) }
|
||
.expect("alloc MappedF32Buffer for consec_losses");
|
||
consec_losses_buf.write_from_slice(&[0.0_f32]);
|
||
|
||
// 4 losses, then 1 win (resets streak), then 4 more losses → no trip.
|
||
for outcome in [1, 1, 1, 1, 0, 1, 1, 1, 1] {
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_cooldown(
|
||
&stream,
|
||
outcome as i32,
|
||
/*trade_close_event*/ 1,
|
||
isv_buf.dev_ptr,
|
||
consec_losses_buf.dev_ptr,
|
||
)
|
||
.expect("launch cooldown_kernel (streak-broken)");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after cooldown_kernel launch (streak-broken)");
|
||
}
|
||
|
||
let isv_host = isv_buf.read_all();
|
||
let remaining = isv_host[COOLDOWN_BARS_REMAINING_INDEX];
|
||
assert!(
|
||
remaining.abs() < 1e-9,
|
||
"cooldown_remaining = {remaining}, expected 0 (streak broken by win)"
|
||
);
|
||
}
|
||
|
||
/// SP15 Phase 3.5.3 — non-trade-close bar decrements cooldown but
|
||
/// doesn't change `consecutive_losses`. Per spec §9.2 (3.5.3)
|
||
/// post-amendment-2 fix: streak counter is updated only on
|
||
/// trade-close events; per-bar non-close calls just decrement the
|
||
/// cooldown counter without touching the loss tracker. Starts
|
||
/// mid-cooldown (10 bars remaining) and consec_losses=2; 5
|
||
/// non-trade-close bars decrement remaining to 5 and leave
|
||
/// consec_losses untouched at 2.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn cooldown_decrements_on_non_trade_bars() {
|
||
use ml::cuda_pipeline::sp15_isv_slots::{
|
||
COOLDOWN_BARS_REMAINING_INDEX, COOLDOWN_K_THRESHOLD_INDEX,
|
||
COOLDOWN_M_BARS_INDEX,
|
||
};
|
||
|
||
let stream = make_test_stream();
|
||
|
||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc MappedF32Buffer for isv bus");
|
||
let mut isv_init = vec![0.0f32; ISV_LEN];
|
||
isv_init[COOLDOWN_K_THRESHOLD_INDEX] = 5.0;
|
||
isv_init[COOLDOWN_M_BARS_INDEX] = 20.0;
|
||
isv_init[COOLDOWN_BARS_REMAINING_INDEX] = 10.0; // mid-cooldown
|
||
isv_buf.write_from_slice(&isv_init);
|
||
|
||
let consec_losses_buf = unsafe { MappedF32Buffer::new(1) }
|
||
.expect("alloc MappedF32Buffer for consec_losses");
|
||
consec_losses_buf.write_from_slice(&[2.0_f32]);
|
||
|
||
// 5 non-trade-close bars: cooldown decrements 5×, consec_losses unchanged.
|
||
for _ in 0..5 {
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_cooldown(
|
||
&stream,
|
||
/*last_trade_outcome_loss*/ 0,
|
||
/*trade_close_event*/ 0,
|
||
isv_buf.dev_ptr,
|
||
consec_losses_buf.dev_ptr,
|
||
)
|
||
.expect("launch cooldown_kernel (non-trade-close)");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after cooldown_kernel launch (non-trade-close)");
|
||
}
|
||
|
||
let isv_host = isv_buf.read_all();
|
||
let remaining = isv_host[COOLDOWN_BARS_REMAINING_INDEX];
|
||
assert!(
|
||
(remaining - 5.0).abs() < 1e-9,
|
||
"cooldown_remaining = {remaining}, expected 5"
|
||
);
|
||
let consec = consec_losses_buf.read_all()[0];
|
||
assert!(
|
||
(consec - 2.0).abs() < 1e-9,
|
||
"consec_losses = {consec}, expected 2 (non-trade-close bars must not touch streak counter)"
|
||
);
|
||
}
|
||
|
||
/// SP15 Phase 3.5.4 — plasticity injection fires when
|
||
/// `DD_PERSISTENCE` crosses the threshold AND
|
||
/// `PLASTICITY_FIRED_THIS_FOLD == 0`. Sets the fired flag and the
|
||
/// warm-bars counter to M_warm. Per spec §9.2 (3.5.4) post-amendment-2
|
||
/// fix: TWO-STEP recovery — kernel sets fired + warm-bars counter
|
||
/// (this commit); action-selection layer (consumer wiring follow-up)
|
||
/// reads the fired flag transition AND emits Flat that bar, then
|
||
/// forces Hold while
|
||
/// `max(COOLDOWN_BARS_REMAINING, PLASTICITY_WARM_BARS_REMAINING) > 0`.
|
||
/// Wave 4.2 / Phase 3.5.4.b: kernel now also performs the
|
||
/// Kaiming-He weight-reset; this trigger-only test passes the
|
||
/// extra `fan_in` + `seed` launcher params and asserts the
|
||
/// trigger behaviour (fired flag + warm-bars). The dedicated
|
||
/// `plasticity_injection_kernel_resets_last_10pct_kaiming_he`
|
||
/// test covers the weight-reset numerics.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn plasticity_fires_when_persistence_exceeds_threshold() {
|
||
use ml::cuda_pipeline::sp15_isv_slots::{
|
||
DD_PERSISTENCE_MAX_INDEX, PLASTICITY_FIRED_THIS_FOLD_INDEX,
|
||
PLASTICITY_PERSISTENCE_THRESHOLD_INDEX, PLASTICITY_WARM_BARS_REMAINING_INDEX,
|
||
};
|
||
|
||
let stream = make_test_stream();
|
||
|
||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc MappedF32Buffer for isv bus");
|
||
let mut isv_init = vec![0.0f32; ISV_LEN];
|
||
isv_init[DD_PERSISTENCE_MAX_INDEX] = 150.0; // exceeds threshold
|
||
isv_init[PLASTICITY_PERSISTENCE_THRESHOLD_INDEX] = 100.0;
|
||
isv_init[PLASTICITY_FIRED_THIS_FOLD_INDEX] = 0.0; // not yet fired
|
||
isv_init[PLASTICITY_WARM_BARS_REMAINING_INDEX] = 0.0;
|
||
isv_buf.write_from_slice(&isv_init);
|
||
|
||
// Wave 4.2 weight-reset is exercised in detail by the
|
||
// dedicated Kaiming-He oracle below; here we just provide a
|
||
// 64-element buffer so the launcher can run (n_weights / 10 =
|
||
// 6 trailing slots get re-sampled).
|
||
let weights_buf = unsafe { MappedF32Buffer::new(64) }
|
||
.expect("alloc MappedF32Buffer for advantage_head_weights");
|
||
weights_buf.write_from_slice(&vec![1.0f32; 64]);
|
||
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_plasticity_injection(
|
||
&stream,
|
||
&load_sp15_kernel(&stream, SP15_PLASTICITY_INJECTION_CUBIN, "plasticity_injection_kernel"),
|
||
isv_buf.dev_ptr,
|
||
weights_buf.dev_ptr,
|
||
/*n_weights*/ 64,
|
||
/*m_warm*/ 200.0,
|
||
/*fan_in*/ 256,
|
||
/*seed*/ 0xCAFE_BABE_DEAD_BEEF_u64,
|
||
)
|
||
.expect("launch plasticity_injection_kernel (fires)");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after plasticity_injection_kernel launch (fires)");
|
||
|
||
let isv_host = isv_buf.read_all();
|
||
let fired = isv_host[PLASTICITY_FIRED_THIS_FOLD_INDEX];
|
||
assert!(
|
||
(fired - 1.0).abs() < 1e-9,
|
||
"PLASTICITY_FIRED_THIS_FOLD = {fired}, expected 1.0"
|
||
);
|
||
// After fire: warm = m_warm = 200 set, then per-bar decrement
|
||
// takes warm to 199 in the same call. Acceptable range
|
||
// [198, 200] mirrors the cooldown_kernel trigger-then-decrement
|
||
// tolerance.
|
||
let warm = isv_host[PLASTICITY_WARM_BARS_REMAINING_INDEX];
|
||
assert!(
|
||
(198.0..=200.0).contains(&warm),
|
||
"PLASTICITY_WARM_BARS_REMAINING = {warm}, expected ~199-200"
|
||
);
|
||
}
|
||
|
||
/// SP15 Phase 3.5.4 — plasticity injection does NOT re-fire within
|
||
/// the same fold (debounced by `PLASTICITY_FIRED_THIS_FOLD == 1`).
|
||
/// Per spec §9.2 (3.5.4) post-amendment-2 fix: the gate fires at
|
||
/// most once per fold; PLASTICITY_FIRED_THIS_FOLD resets to 0.0 at
|
||
/// fold boundary re-arms the gate for the next fold (handled by
|
||
/// `sp15_plasticity_fired_this_fold` registry-arm dispatch). Mid-
|
||
/// warm-up state: `fired = 1`, `warm_bars = 50`. After one call
|
||
/// with `DD_PERSISTENCE = 200` (way over threshold): fired stays
|
||
/// 1 (debounced), warm_bars decrements 50 → 49.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn plasticity_debounced_within_fold() {
|
||
use ml::cuda_pipeline::sp15_isv_slots::{
|
||
DD_PERSISTENCE_MAX_INDEX, PLASTICITY_FIRED_THIS_FOLD_INDEX,
|
||
PLASTICITY_PERSISTENCE_THRESHOLD_INDEX, PLASTICITY_WARM_BARS_REMAINING_INDEX,
|
||
};
|
||
|
||
let stream = make_test_stream();
|
||
|
||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc MappedF32Buffer for isv bus");
|
||
let mut isv_init = vec![0.0f32; ISV_LEN];
|
||
isv_init[DD_PERSISTENCE_MAX_INDEX] = 200.0; // way over threshold
|
||
isv_init[PLASTICITY_PERSISTENCE_THRESHOLD_INDEX] = 100.0;
|
||
isv_init[PLASTICITY_FIRED_THIS_FOLD_INDEX] = 1.0; // ALREADY fired this fold
|
||
isv_init[PLASTICITY_WARM_BARS_REMAINING_INDEX] = 50.0; // mid-warm-up
|
||
isv_buf.write_from_slice(&isv_init);
|
||
|
||
let weights_buf = unsafe { MappedF32Buffer::new(64) }
|
||
.expect("alloc MappedF32Buffer for advantage_head_weights");
|
||
weights_buf.write_from_slice(&vec![1.0f32; 64]);
|
||
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_plasticity_injection(
|
||
&stream,
|
||
&load_sp15_kernel(&stream, SP15_PLASTICITY_INJECTION_CUBIN, "plasticity_injection_kernel"),
|
||
isv_buf.dev_ptr,
|
||
weights_buf.dev_ptr,
|
||
64,
|
||
200.0,
|
||
/*fan_in*/ 256,
|
||
/*seed*/ 0xDEAD_BEEF_F00D_CAFE_u64,
|
||
)
|
||
.expect("launch plasticity_injection_kernel (debounced)");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after plasticity_injection_kernel launch (debounced)");
|
||
|
||
let isv_host = isv_buf.read_all();
|
||
// fired stays 1 (already fired, debounced); warm_bars decrements 50 → 49.
|
||
assert!(
|
||
(isv_host[PLASTICITY_FIRED_THIS_FOLD_INDEX] - 1.0).abs() < 1e-9,
|
||
"PLASTICITY_FIRED_THIS_FOLD = {}, expected 1.0 (debounced — must NOT re-fire)",
|
||
isv_host[PLASTICITY_FIRED_THIS_FOLD_INDEX]
|
||
);
|
||
let warm = isv_host[PLASTICITY_WARM_BARS_REMAINING_INDEX];
|
||
assert!(
|
||
(warm - 49.0).abs() < 1e-9,
|
||
"PLASTICITY_WARM_BARS_REMAINING = {warm}, expected 49 (decremented from 50, no re-fire)"
|
||
);
|
||
|
||
// Wave 4.2: when debounced, the kernel returns early before
|
||
// the weight-reset region, so all 64 weights must remain
|
||
// exactly 1.0 (no Kaiming-He sample written).
|
||
let weights_host = weights_buf.read_all();
|
||
for (i, w) in weights_host.iter().enumerate() {
|
||
assert!(
|
||
(w - 1.0).abs() < 1e-9,
|
||
"weights[{i}] = {w}, expected 1.0 (debounced — kernel must skip weight-reset)"
|
||
);
|
||
}
|
||
}
|
||
|
||
/// SP15 Phase 3.5.4 — plasticity injection does NOT fire when
|
||
/// `DD_PERSISTENCE` is below the threshold. `warm_bars` stays at 0
|
||
/// (no decrement when already 0 per the kernel's `if (warm > 0.0f)`
|
||
/// guard). Sentinel non-zero output verification: starting with
|
||
/// `fired = 0, warm = 0` and `persistence < threshold`, both flags
|
||
/// remain at 0 after the launch.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn plasticity_no_fire_below_threshold() {
|
||
use ml::cuda_pipeline::sp15_isv_slots::{
|
||
DD_PERSISTENCE_MAX_INDEX, PLASTICITY_FIRED_THIS_FOLD_INDEX,
|
||
PLASTICITY_PERSISTENCE_THRESHOLD_INDEX, PLASTICITY_WARM_BARS_REMAINING_INDEX,
|
||
};
|
||
|
||
let stream = make_test_stream();
|
||
|
||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc MappedF32Buffer for isv bus");
|
||
let mut isv_init = vec![0.0f32; ISV_LEN];
|
||
isv_init[DD_PERSISTENCE_MAX_INDEX] = 50.0; // BELOW threshold
|
||
isv_init[PLASTICITY_PERSISTENCE_THRESHOLD_INDEX] = 100.0;
|
||
isv_init[PLASTICITY_FIRED_THIS_FOLD_INDEX] = 0.0;
|
||
isv_init[PLASTICITY_WARM_BARS_REMAINING_INDEX] = 0.0;
|
||
isv_buf.write_from_slice(&isv_init);
|
||
|
||
let weights_buf = unsafe { MappedF32Buffer::new(64) }
|
||
.expect("alloc MappedF32Buffer for advantage_head_weights");
|
||
weights_buf.write_from_slice(&vec![1.0f32; 64]);
|
||
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_plasticity_injection(
|
||
&stream,
|
||
&load_sp15_kernel(&stream, SP15_PLASTICITY_INJECTION_CUBIN, "plasticity_injection_kernel"),
|
||
isv_buf.dev_ptr,
|
||
weights_buf.dev_ptr,
|
||
64,
|
||
200.0,
|
||
/*fan_in*/ 256,
|
||
/*seed*/ 0xFEED_FACE_BAAD_F00D_u64,
|
||
)
|
||
.expect("launch plasticity_injection_kernel (no-fire)");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after plasticity_injection_kernel launch (no-fire)");
|
||
|
||
let isv_host = isv_buf.read_all();
|
||
assert!(
|
||
isv_host[PLASTICITY_FIRED_THIS_FOLD_INDEX].abs() < 1e-9,
|
||
"PLASTICITY_FIRED_THIS_FOLD = {}, expected 0.0 (below-threshold must not fire)",
|
||
isv_host[PLASTICITY_FIRED_THIS_FOLD_INDEX]
|
||
);
|
||
assert!(
|
||
isv_host[PLASTICITY_WARM_BARS_REMAINING_INDEX].abs() < 1e-9,
|
||
"PLASTICITY_WARM_BARS_REMAINING = {}, expected 0.0 (no fire → no warm-up set; the `if (warm > 0)` guard prevents underflow when already 0)",
|
||
isv_host[PLASTICITY_WARM_BARS_REMAINING_INDEX]
|
||
);
|
||
|
||
// Wave 4.2: when no-fire, the kernel returns early before the
|
||
// weight-reset region, so all 64 weights must remain exactly
|
||
// 1.0 (no Kaiming-He sample written).
|
||
let weights_host = weights_buf.read_all();
|
||
for (i, w) in weights_host.iter().enumerate() {
|
||
assert!(
|
||
(w - 1.0).abs() < 1e-9,
|
||
"weights[{i}] = {w}, expected 1.0 (no-fire — kernel must skip weight-reset)"
|
||
);
|
||
}
|
||
}
|
||
|
||
/// SP15 Wave 4.2 / Phase 3.5.4.b — when the plasticity injection
|
||
/// gate fires, the trailing 10% of `advantage_head_weights` must
|
||
/// be re-initialised to Kaiming-He samples
|
||
/// (`Normal(0, sqrt(2/fan_in))`, gain=√2 for ReLU/GLU
|
||
/// activations) drawn via cuRAND `curand_normal()`. The first 90%
|
||
/// of the buffer must remain bit-identical to its pre-launch
|
||
/// value (1.0) — the kernel writes element `[reset_start + tid]`
|
||
/// only.
|
||
///
|
||
/// Distribution checks (with `n_weights = 1280` → `reset_count =
|
||
/// 128`, `fan_in = 256`, expected `stddev = sqrt(2/256) ≈
|
||
/// 0.08839`):
|
||
/// - Per-element diff from 1.0: every reset slot must have
|
||
/// changed (probability of an exact-1.0 sample is essentially
|
||
/// zero — Kaiming stddev ≈ 0.088 means 1.0 is ~11.3 σ out, and
|
||
/// anyway the kernel writes the *sample*, not the original).
|
||
/// - Mean of reset region: |mean| < 0.05 (sampling-noise
|
||
/// tolerance for n=128 from a Normal with stddev≈0.088 — the
|
||
/// true SE of the mean is `0.088 / sqrt(128) ≈ 0.0078`, so
|
||
/// `0.05` is a comfortable ~6σ envelope around 0).
|
||
/// - Std of reset region: within ±20% of `sqrt(2/fan_in)`
|
||
/// (n=128 sample std has ~6% relative SE; ±20% is comfortable
|
||
/// headroom).
|
||
/// - ISV[PLASTICITY_FIRED_THIS_FOLD_INDEX] flipped 0 → 1.
|
||
/// - ISV[PLASTICITY_WARM_BARS_REMAINING_INDEX] = m_warm − 1
|
||
/// (per-bar decrement runs in the same launch as the trigger
|
||
/// per the cooldown_kernel trigger-then-decrement convention).
|
||
///
|
||
/// Determinism: a fixed seed (`42`) drives `curand_init(seed,
|
||
/// tid, 0, &state)` per thread; re-running this test always
|
||
/// produces identical samples.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn plasticity_injection_kernel_resets_last_10pct_kaiming_he() {
|
||
use ml::cuda_pipeline::sp15_isv_slots::{
|
||
DD_PERSISTENCE_MAX_INDEX, PLASTICITY_FIRED_THIS_FOLD_INDEX,
|
||
PLASTICITY_PERSISTENCE_THRESHOLD_INDEX, PLASTICITY_WARM_BARS_REMAINING_INDEX,
|
||
};
|
||
|
||
let stream = make_test_stream();
|
||
|
||
// ISV: ensure the trigger fires — large persistence, small
|
||
// threshold, fired flag clear, warm-bars 0.
|
||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc MappedF32Buffer for isv bus");
|
||
let mut isv_init = vec![0.0f32; ISV_LEN];
|
||
isv_init[DD_PERSISTENCE_MAX_INDEX] = 1_000.0;
|
||
isv_init[PLASTICITY_PERSISTENCE_THRESHOLD_INDEX] = 1.0;
|
||
isv_init[PLASTICITY_FIRED_THIS_FOLD_INDEX] = 0.0;
|
||
isv_init[PLASTICITY_WARM_BARS_REMAINING_INDEX] = 0.0;
|
||
isv_buf.write_from_slice(&isv_init);
|
||
|
||
// Weight buffer: n_weights = 1280, reset_count = 128 (10%).
|
||
// Pre-fill every slot with 1.0 so we can detect both
|
||
// "unchanged" (first 90%) and "changed" (last 10%) by exact
|
||
// float comparison + sample-statistic checks.
|
||
const N_WEIGHTS: usize = 1280;
|
||
const FAN_IN: i32 = 256;
|
||
const SEED: u64 = 42;
|
||
const M_WARM: f32 = 200.0;
|
||
let weights_buf = unsafe { MappedF32Buffer::new(N_WEIGHTS) }
|
||
.expect("alloc MappedF32Buffer for advantage_head_weights");
|
||
weights_buf.write_from_slice(&vec![1.0f32; N_WEIGHTS]);
|
||
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_plasticity_injection(
|
||
&stream,
|
||
&load_sp15_kernel(&stream, SP15_PLASTICITY_INJECTION_CUBIN, "plasticity_injection_kernel"),
|
||
isv_buf.dev_ptr,
|
||
weights_buf.dev_ptr,
|
||
N_WEIGHTS as i32,
|
||
M_WARM,
|
||
FAN_IN,
|
||
SEED,
|
||
)
|
||
.expect("launch plasticity_injection_kernel (kaiming-he)");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after plasticity_injection_kernel launch (kaiming-he)");
|
||
|
||
// ISV side-effects: fired flipped to 1, warm = M_warm - 1.
|
||
let isv_host = isv_buf.read_all();
|
||
assert!(
|
||
(isv_host[PLASTICITY_FIRED_THIS_FOLD_INDEX] - 1.0).abs() < 1e-9,
|
||
"PLASTICITY_FIRED_THIS_FOLD = {}, expected 1.0 after fire",
|
||
isv_host[PLASTICITY_FIRED_THIS_FOLD_INDEX]
|
||
);
|
||
let warm = isv_host[PLASTICITY_WARM_BARS_REMAINING_INDEX];
|
||
assert!(
|
||
(warm - (M_WARM - 1.0)).abs() < 1e-6,
|
||
"PLASTICITY_WARM_BARS_REMAINING = {warm}, expected {} (M_warm - 1, per-bar decrement runs same call)",
|
||
M_WARM - 1.0
|
||
);
|
||
|
||
// Weight-reset side-effects.
|
||
let weights_host = weights_buf.read_all();
|
||
let reset_count: usize = N_WEIGHTS / 10; // 128
|
||
let reset_start: usize = N_WEIGHTS - reset_count; // 1152
|
||
|
||
// First 90% must be bit-identical to the pre-launch value.
|
||
for (i, w) in weights_host[..reset_start].iter().enumerate() {
|
||
assert!(
|
||
(w - 1.0).abs() < 1e-9,
|
||
"weights[{i}] = {w}, expected 1.0 (outside reset region — kernel must not touch)"
|
||
);
|
||
}
|
||
|
||
// Last 10%: every slot must have moved off 1.0 (a Kaiming
|
||
// stddev≈0.088 makes the sentinel 1.0 ~11σ out, so an exact
|
||
// 1.0 sample is effectively zero-probability).
|
||
let resets = &weights_host[reset_start..];
|
||
for (i, w) in resets.iter().enumerate() {
|
||
assert!(
|
||
(w - 1.0).abs() > 1e-6,
|
||
"weights[{}] = {w}, expected resampled value ≠ 1.0",
|
||
reset_start + i
|
||
);
|
||
}
|
||
|
||
// Mean of resets ≈ 0 within sampling-noise envelope.
|
||
let n = resets.len() as f64;
|
||
let mean: f64 = resets.iter().map(|&w| w as f64).sum::<f64>() / n;
|
||
assert!(
|
||
mean.abs() < 0.05,
|
||
"Kaiming-reset region mean = {mean:.6}, expected |mean| < 0.05 (n={n}, true SE ~ 0.0078)"
|
||
);
|
||
|
||
// Std of resets ≈ sqrt(2/fan_in) within ±20%.
|
||
let expected_std = (2.0_f64 / FAN_IN as f64).sqrt();
|
||
let var: f64 = resets
|
||
.iter()
|
||
.map(|&w| (w as f64 - mean).powi(2))
|
||
.sum::<f64>()
|
||
/ (n - 1.0);
|
||
let std = var.sqrt();
|
||
assert!(
|
||
(std - expected_std).abs() / expected_std < 0.20,
|
||
"Kaiming-reset region std = {std:.6}, expected {expected_std:.6} ± 20%"
|
||
);
|
||
|
||
// Determinism re-check: same (seed, tid) on a fresh buffer
|
||
// produces bit-identical samples (catches accidental
|
||
// entropy sources like clock-derived seeds inside the
|
||
// kernel). Re-launch with the same seed against a fresh
|
||
// 1.0-filled buffer — but note ISV state must also be reset
|
||
// (fired=0) so the trigger re-fires.
|
||
weights_buf.write_from_slice(&vec![1.0f32; N_WEIGHTS]);
|
||
isv_init[PLASTICITY_FIRED_THIS_FOLD_INDEX] = 0.0;
|
||
isv_init[PLASTICITY_WARM_BARS_REMAINING_INDEX] = 0.0;
|
||
isv_buf.write_from_slice(&isv_init);
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_plasticity_injection(
|
||
&stream,
|
||
&load_sp15_kernel(&stream, SP15_PLASTICITY_INJECTION_CUBIN, "plasticity_injection_kernel"),
|
||
isv_buf.dev_ptr,
|
||
weights_buf.dev_ptr,
|
||
N_WEIGHTS as i32,
|
||
M_WARM,
|
||
FAN_IN,
|
||
SEED, // same seed
|
||
)
|
||
.expect("launch plasticity_injection_kernel (determinism re-check)");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after plasticity_injection_kernel launch (determinism)");
|
||
let resets_2 = weights_buf.read_all();
|
||
for i in 0..reset_count {
|
||
let a = resets[i];
|
||
let b = resets_2[reset_start + i];
|
||
assert!(
|
||
(a - b).abs() < 1e-9,
|
||
"non-deterministic Kaiming sample at tid={i}: first={a}, second={b}"
|
||
);
|
||
}
|
||
}
|
||
|
||
/// SP15 Phase 3.5.5 — DD_TRAJECTORY_DECREASING fires when dd_pct
|
||
/// decreases from a non-trivial drawdown (above the floor). Per
|
||
/// spec §9.2 (3.5.5) post-amendment-2 fix: per-bar boolean
|
||
/// computed from `dd_pct(t) < dd_pct(t-1) AND dd_pct(t-1) > floor`.
|
||
/// SP15 Phase 1.3.b-followup-B (2026-05-07) — migrated to per-env
|
||
/// contract. The kernel runs a per-env grid; this test uses
|
||
/// n_envs=1 so the assertions match the original Wave 4.3
|
||
/// single-env behaviour bit-identically. Reads per-env DD_PCT
|
||
/// from `dd_state_per_env[env*6 + 5]` (the canonical producer
|
||
/// `dd_state_kernel` writes there); reads per-env prev_dd from
|
||
/// `prev_dd_pct_per_env[env]`; writes trajectory to
|
||
/// `dd_trajectory_per_env[env]`. Floor still read from ISV[441].
|
||
///
|
||
/// Sequence: (1) bootstrap call sets prev=0.10 (cold-start
|
||
/// trajectory=0 because initial prev=0); (2) update dd_pct=0.08
|
||
/// and re-launch — prev=0.10 > floor=0.02 AND now=0.08 < prev →
|
||
/// trajectory=1. Verifies BOTH the trajectory output AND the
|
||
/// persistent prev_dd_pct advance.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn dd_trajectory_decreasing_fires_during_recovery() {
|
||
use ml::cuda_pipeline::sp15_isv_slots::DD_TRAJECTORY_FLOOR_INDEX;
|
||
|
||
let stream = make_test_stream();
|
||
const N_ENVS: i32 = 1;
|
||
|
||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc MappedF32Buffer for isv bus");
|
||
let mut isv_init = vec![0.0f32; ISV_LEN];
|
||
isv_init[DD_TRAJECTORY_FLOOR_INDEX] = 0.02; // floor sentinel
|
||
isv_buf.write_from_slice(&isv_init);
|
||
|
||
// Per-env DD state tile (shared with dd_state_kernel layout):
|
||
// tile[env*6 + 5] = DD_PCT for env. The kernel reads slot 5.
|
||
let dd_tile_buf = unsafe { MappedF32Buffer::new((N_ENVS as usize) * 6) }
|
||
.expect("alloc MappedF32Buffer for dd_state_per_env tile");
|
||
let mut tile_init = vec![0.0f32; (N_ENVS as usize) * 6];
|
||
tile_init[0 * 6 + 5] = 0.10; // current dd_pct for env 0
|
||
dd_tile_buf.write_from_slice(&tile_init);
|
||
|
||
// Per-env persistent prev_dd_pct scratch — init to 0.0
|
||
// (mirrors collector's `sp15_dd_trajectory_prev_dd_per_env`).
|
||
let prev_dd_buf = unsafe { MappedF32Buffer::new(N_ENVS as usize) }
|
||
.expect("alloc MappedF32Buffer for prev_dd_pct_per_env");
|
||
prev_dd_buf.write_from_slice(&vec![0.0_f32; N_ENVS as usize]);
|
||
|
||
// Per-env trajectory output tile (mirrors collector's
|
||
// `sp15_dd_trajectory_per_env`).
|
||
let traj_buf = unsafe { MappedF32Buffer::new(N_ENVS as usize) }
|
||
.expect("alloc MappedF32Buffer for dd_trajectory_per_env");
|
||
traj_buf.write_from_slice(&vec![0.0_f32; N_ENVS as usize]);
|
||
|
||
// First call: prev=0, current=0.10 → cold-start bootstrap
|
||
// (prev=0 < floor=0.02 → floor gate false → trajectory=0); the
|
||
// kernel advances prev_dd_pct to 0.10.
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_dd_trajectory_decreasing(
|
||
&stream,
|
||
&load_sp15_kernel(&stream, SP15_DD_TRAJECTORY_CUBIN, "dd_trajectory_decreasing_kernel"),
|
||
N_ENVS,
|
||
dd_tile_buf.dev_ptr,
|
||
prev_dd_buf.dev_ptr,
|
||
traj_buf.dev_ptr,
|
||
isv_buf.dev_ptr,
|
||
)
|
||
.expect("launch dd_trajectory_decreasing_kernel (bootstrap)");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after dd_trajectory_decreasing_kernel launch (bootstrap)");
|
||
|
||
// Now reduce dd_pct: 0.10 → 0.08 (recovery from non-trivial DD
|
||
// above floor). Update tile slot 5 only.
|
||
let mut tile_now = dd_tile_buf.read_all();
|
||
tile_now[0 * 6 + 5] = 0.08;
|
||
dd_tile_buf.write_from_slice(&tile_now);
|
||
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_dd_trajectory_decreasing(
|
||
&stream,
|
||
&load_sp15_kernel(&stream, SP15_DD_TRAJECTORY_CUBIN, "dd_trajectory_decreasing_kernel"),
|
||
N_ENVS,
|
||
dd_tile_buf.dev_ptr,
|
||
prev_dd_buf.dev_ptr,
|
||
traj_buf.dev_ptr,
|
||
isv_buf.dev_ptr,
|
||
)
|
||
.expect("launch dd_trajectory_decreasing_kernel (recovery)");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after dd_trajectory_decreasing_kernel launch (recovery)");
|
||
|
||
let trajectory = traj_buf.read_all()[0];
|
||
assert!(
|
||
(trajectory - 1.0).abs() < 1e-9,
|
||
"dd_trajectory_per_env[0] = {trajectory}, expected 1.0 (recovery: prev=0.10 > floor=0.02 AND now=0.08 < prev)"
|
||
);
|
||
// Persistent state advance verification: prev_dd_pct now == current dd_pct = 0.08.
|
||
let prev_after = prev_dd_buf.read_all()[0];
|
||
assert!(
|
||
(prev_after - 0.08).abs() < 1e-9,
|
||
"prev_dd_pct_per_env[0] = {prev_after}, expected 0.08 (kernel advances persistent state to current dd_pct)"
|
||
);
|
||
}
|
||
|
||
/// SP15 Phase 3.5.5 — no fire when dd_pct increasing (DD getting
|
||
/// worse). Per spec §9.2 (3.5.5) post-amendment-2 fix: trajectory
|
||
/// requires `dd_pct(t) < dd_pct(t-1)` (strict). Setup: prev=0.05,
|
||
/// current=0.10 → DD getting worse → trajectory=0. Sentinel-write
|
||
/// semantics: ISV slot starts at 0.0 (default), kernel writes 0.0
|
||
/// (the sentinel itself) — verified separately by the recovery
|
||
/// test that the kernel WILL write a non-zero output when the
|
||
/// gate fires, so the 0.0 here is a real write rather than a
|
||
/// skipped store.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn dd_trajectory_no_fire_when_dd_increasing() {
|
||
use ml::cuda_pipeline::sp15_isv_slots::DD_TRAJECTORY_FLOOR_INDEX;
|
||
|
||
let stream = make_test_stream();
|
||
const N_ENVS: i32 = 1;
|
||
|
||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc MappedF32Buffer for isv bus");
|
||
let mut isv_init = vec![0.0f32; ISV_LEN];
|
||
isv_init[DD_TRAJECTORY_FLOOR_INDEX] = 0.02;
|
||
isv_buf.write_from_slice(&isv_init);
|
||
|
||
// Per-env DD state tile: env-0's DD_PCT = 0.10 (current).
|
||
let dd_tile_buf = unsafe { MappedF32Buffer::new((N_ENVS as usize) * 6) }
|
||
.expect("alloc MappedF32Buffer for dd_state_per_env tile");
|
||
let mut tile_init = vec![0.0f32; (N_ENVS as usize) * 6];
|
||
tile_init[0 * 6 + 5] = 0.10;
|
||
dd_tile_buf.write_from_slice(&tile_init);
|
||
|
||
let prev_dd_buf = unsafe { MappedF32Buffer::new(N_ENVS as usize) }
|
||
.expect("alloc MappedF32Buffer for prev_dd_pct_per_env");
|
||
prev_dd_buf.write_from_slice(&[0.05_f32]); // prev was 0.05
|
||
|
||
// Sentinel non-zero output ensures the kernel actually
|
||
// overwrites with the 0.0 result rather than skipping the
|
||
// store (mirrors the dd_penalty_kernel sentinel-write check).
|
||
let traj_buf = unsafe { MappedF32Buffer::new(N_ENVS as usize) }
|
||
.expect("alloc MappedF32Buffer for dd_trajectory_per_env");
|
||
traj_buf.write_from_slice(&[0.5_f32]);
|
||
|
||
// current dd_pct = 0.10 > prev = 0.05 → DD getting worse, no
|
||
// recovery fire (the strict `<` comparison rejects the >
|
||
// case).
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_dd_trajectory_decreasing(
|
||
&stream,
|
||
&load_sp15_kernel(&stream, SP15_DD_TRAJECTORY_CUBIN, "dd_trajectory_decreasing_kernel"),
|
||
N_ENVS,
|
||
dd_tile_buf.dev_ptr,
|
||
prev_dd_buf.dev_ptr,
|
||
traj_buf.dev_ptr,
|
||
isv_buf.dev_ptr,
|
||
)
|
||
.expect("launch dd_trajectory_decreasing_kernel (dd-increasing)");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after dd_trajectory_decreasing_kernel launch (dd-increasing)");
|
||
|
||
let trajectory = traj_buf.read_all()[0];
|
||
assert!(
|
||
trajectory.abs() < 1e-9,
|
||
"dd_trajectory_per_env[0] = {trajectory}, expected 0 (DD increasing: prev=0.05 < now=0.10; sentinel 0.5 must be overwritten with the real 0.0 result, not skipped)"
|
||
);
|
||
}
|
||
|
||
/// SP15 Phase 3.5.5 — no fire when prev_dd below the floor (trivial
|
||
/// DD, not "in recovery"). Per spec §9.2 (3.5.5) post-amendment-2
|
||
/// fix: trajectory requires `dd_pct(t-1) > DD_TRAJECTORY_FLOOR`
|
||
/// (strict). The non-trivial gate keeps "0.001 → 0.0005" PnL
|
||
/// flutter from inflating the recovery signal — only meaningful
|
||
/// drawdowns count. Setup: prev=0.01 (BELOW floor=0.02), current=
|
||
/// 0.005 → dd_pct decreasing but prev was below floor → trajectory=0.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn dd_trajectory_no_fire_when_prev_below_floor() {
|
||
use ml::cuda_pipeline::sp15_isv_slots::DD_TRAJECTORY_FLOOR_INDEX;
|
||
|
||
let stream = make_test_stream();
|
||
const N_ENVS: i32 = 1;
|
||
|
||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc MappedF32Buffer for isv bus");
|
||
let mut isv_init = vec![0.0f32; ISV_LEN];
|
||
isv_init[DD_TRAJECTORY_FLOOR_INDEX] = 0.02;
|
||
isv_buf.write_from_slice(&isv_init);
|
||
|
||
// Per-env DD state tile: env-0's DD_PCT = 0.005 (below floor).
|
||
let dd_tile_buf = unsafe { MappedF32Buffer::new((N_ENVS as usize) * 6) }
|
||
.expect("alloc MappedF32Buffer for dd_state_per_env tile");
|
||
let mut tile_init = vec![0.0f32; (N_ENVS as usize) * 6];
|
||
tile_init[0 * 6 + 5] = 0.005;
|
||
dd_tile_buf.write_from_slice(&tile_init);
|
||
|
||
let prev_dd_buf = unsafe { MappedF32Buffer::new(N_ENVS as usize) }
|
||
.expect("alloc MappedF32Buffer for prev_dd_pct_per_env");
|
||
prev_dd_buf.write_from_slice(&[0.01_f32]); // prev BELOW floor 0.02
|
||
|
||
// Sentinel non-zero output (see recovery test for rationale).
|
||
let traj_buf = unsafe { MappedF32Buffer::new(N_ENVS as usize) }
|
||
.expect("alloc MappedF32Buffer for dd_trajectory_per_env");
|
||
traj_buf.write_from_slice(&[0.7_f32]);
|
||
|
||
// dd decreasing (0.01 → 0.005) but prev was below floor → not
|
||
// "in recovery" per the non-trivial-DD gate.
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_dd_trajectory_decreasing(
|
||
&stream,
|
||
&load_sp15_kernel(&stream, SP15_DD_TRAJECTORY_CUBIN, "dd_trajectory_decreasing_kernel"),
|
||
N_ENVS,
|
||
dd_tile_buf.dev_ptr,
|
||
prev_dd_buf.dev_ptr,
|
||
traj_buf.dev_ptr,
|
||
isv_buf.dev_ptr,
|
||
)
|
||
.expect("launch dd_trajectory_decreasing_kernel (prev-below-floor)");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after dd_trajectory_decreasing_kernel launch (prev-below-floor)");
|
||
|
||
let trajectory = traj_buf.read_all()[0];
|
||
assert!(
|
||
trajectory.abs() < 1e-9,
|
||
"dd_trajectory_per_env[0] = {trajectory}, expected 0 (prev=0.01 below floor=0.02; sentinel 0.7 must be overwritten with the real 0.0 result, not skipped)"
|
||
);
|
||
}
|
||
|
||
/// SP15 Wave 4.3 / Phase 3.5.5.b — PER sampler recovery oversample
|
||
/// boost. Verifies that when ISV[DD_TRAJECTORY_DECREASING_INDEX=439]
|
||
/// = 1.0 (recovery active) and
|
||
/// ISV[RECOVERY_OVERSAMPLE_WEIGHT_INDEX=440] = 2.0 (sentinel), the
|
||
/// `per_insert_pa` kernel multiplies the inserted base priority by
|
||
/// `1 + ω × δ = 3.0`, so recovery transitions land in the priority
|
||
/// buffer at 3× the non-recovery baseline. This is the canonical
|
||
/// statistical bias the deferred Phase 3.5.5 PER consumer was
|
||
/// supposed to deliver — closes out
|
||
/// `feedback_wire_everything_up.md`.
|
||
///
|
||
/// We test the priority-buffer write directly (rather than running
|
||
/// the full prefix-sum + sample loop and counting empirical hits)
|
||
/// because the boost factor is deterministic at insert time —
|
||
/// `priorities[idx] = base × (1 + 2.0 × decreasing)` is bit-exact
|
||
/// against expected. The deterministic-stratified `per_sample`
|
||
/// kernel then weights samples by `priorities^alpha`, so the
|
||
/// downstream sampling-frequency bias is implied by the priority
|
||
/// values (which `gpu_per_integration_test::test_gpu_per_priorities_affect_sampling`
|
||
/// already covers for the general TD-error → priority → sampling
|
||
/// chain). Verifying the priority-buffer write here, paired with
|
||
/// the existing PER sampling test, gives the full statistical-bias
|
||
/// proof without re-running a 10000-sample empirical experiment in
|
||
/// this oracle suite.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn per_sampler_weights_recovery_transitions_higher() {
|
||
use ml::dqn::gpu_replay_buffer::{GpuReplayBuffer, GpuReplayBufferConfig};
|
||
use ml::cuda_pipeline::sp15_isv_slots::RECOVERY_OVERSAMPLE_WEIGHT_INDEX;
|
||
use cudarc::driver::CudaSlice;
|
||
|
||
let stream = make_test_stream();
|
||
|
||
// Build an ISV bus and seed the global oversample-weight scalar.
|
||
// `RECOVERY_OVERSAMPLE_WEIGHT_INDEX = 440` is the 2.0 sentinel
|
||
// until the deferred ISV-driven producer lands per
|
||
// `feedback_isv_for_adaptive_bounds`. Phase 1.3.b-followup-B
|
||
// (2026-05-07): the per-env recovery flag is now read from
|
||
// `dd_trajectory_per_env[env_id]` (NOT ISV[439] as in Wave
|
||
// 4.3). With n_envs=1 + lookback=1 the per-transition mapping
|
||
// `env_id = (j % (n_envs * lookback)) / lookback = (j % 1) / 1 = 0`
|
||
// collapses to "all transitions read env-0's flag", matching
|
||
// the original Wave 4.3 single-flag behaviour bit-identically
|
||
// — preserved for back-compat assertions while migrating the
|
||
// contract.
|
||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc MappedF32Buffer for isv bus");
|
||
let mut isv_init = vec![0.0_f32; ISV_LEN];
|
||
isv_init[RECOVERY_OVERSAMPLE_WEIGHT_INDEX] = 2.0; // Phase 3.5.5 sentinel
|
||
isv_buf.write_from_slice(&isv_init);
|
||
|
||
// Per-env trajectory tile: env-0's flag toggles between
|
||
// batches (batch 1 = recovery; batch 2 = baseline).
|
||
const N_ENVS: i32 = 1;
|
||
const LOOKBACK: i32 = 1;
|
||
let dd_traj_buf = unsafe { MappedF32Buffer::new(N_ENVS as usize) }
|
||
.expect("alloc MappedF32Buffer for dd_trajectory_per_env");
|
||
dd_traj_buf.write_from_slice(&[1.0_f32]); // recovery active for batch 1
|
||
|
||
// GPU PER replay buffer — small capacity, we'll insert tiny batches.
|
||
let cap = 64;
|
||
let cfg = GpuReplayBufferConfig {
|
||
capacity: cap,
|
||
alpha: 0.6,
|
||
beta_start: 0.4,
|
||
beta_max: 1.0,
|
||
beta_annealing_steps: 1000,
|
||
epsilon: 1e-6,
|
||
max_memory_bytes: 4 * 1024 * 1024,
|
||
max_batch_size: 64,
|
||
};
|
||
let mut buf = GpuReplayBuffer::new(cfg, &stream).expect("alloc GpuReplayBuffer");
|
||
|
||
// Wire the ISV bus + per-env trajectory tile + per-env dims —
|
||
// without these, per_insert_pa falls back to boost_factor=1.0
|
||
// (the explicit non-recovery baseline covered by the smoke
|
||
// tests / existing gpu_per_integration_test).
|
||
buf.set_isv_signals_ptr(isv_buf.dev_ptr);
|
||
buf.set_sp15_dd_trajectory_per_env_ptr(dd_traj_buf.dev_ptr);
|
||
buf.set_sp15_per_env_dims(N_ENVS, LOOKBACK);
|
||
assert_eq!(
|
||
buf.isv_signals_dev_ptr(),
|
||
isv_buf.dev_ptr,
|
||
"set_isv_signals_ptr round-trip — wired ptr must match"
|
||
);
|
||
assert_eq!(
|
||
buf.sp15_dd_trajectory_per_env_dev_ptr(),
|
||
dd_traj_buf.dev_ptr,
|
||
"set_sp15_dd_trajectory_per_env_ptr round-trip — wired ptr must match"
|
||
);
|
||
assert_eq!(
|
||
buf.sp15_per_env_dims(),
|
||
(N_ENVS, LOOKBACK),
|
||
"set_sp15_per_env_dims round-trip — n_envs + lookback must match"
|
||
);
|
||
|
||
// ── Batch 1: 8 transitions, recovery active (δ=1, ω=2) ────
|
||
// Expected priority = max_priority(1.0) × (1 + 2.0 × 1.0) = 3.0.
|
||
// With n_envs=1 + lookback=1, each transition's env_id = 0 →
|
||
// reads dd_trajectory_per_env[0] = 1.0.
|
||
let n1 = 8_usize;
|
||
let sd = ml::state_layout::STATE_DIM;
|
||
let states_host: Vec<f32> = vec![0.0; n1 * sd];
|
||
let next_states_host: Vec<f32> = vec![0.0; n1 * sd];
|
||
let actions_host: Vec<u32> = vec![0_u32; n1];
|
||
let rewards_host: Vec<f32> = vec![0.0; n1];
|
||
let dones_host: Vec<f32> = vec![0.0; n1];
|
||
let aux_sign_host: Vec<i32> = vec![0_i32; n1];
|
||
|
||
let sf: CudaSlice<f32> = stream.clone_htod(&states_host).expect("htod states");
|
||
let nsf: CudaSlice<f32> = stream.clone_htod(&next_states_host).expect("htod next_states");
|
||
let af: CudaSlice<u32> = stream.clone_htod(&actions_host).expect("htod actions");
|
||
let rf: CudaSlice<f32> = stream.clone_htod(&rewards_host).expect("htod rewards");
|
||
let df: CudaSlice<f32> = stream.clone_htod(&dones_host).expect("htod dones");
|
||
let ax: CudaSlice<i32> = stream.clone_htod(&aux_sign_host).expect("htod aux_sign");
|
||
let ac: CudaSlice<f32> = stream.alloc_zeros::<f32>(n1).expect("aux_conf alloc");
|
||
|
||
let ao_zero: CudaSlice<i32> = stream.alloc_zeros::<i32>(n1).expect("aux_outcome alloc");
|
||
buf.insert_batch(&sf, &nsf, &af, &rf, &df, &ax, &ac, &ao_zero, n1)
|
||
.expect("insert_batch (recovery=1)");
|
||
stream.synchronize().expect("sync after insert_batch (recovery=1)");
|
||
|
||
// ── Batch 2: 8 transitions, recovery inactive (δ=0) ───────
|
||
// Expected priority = 1.0 × (1 + 2.0 × 0.0) = 1.0 (baseline).
|
||
// Update env-0's per-env tile to 0.0 (NOT ISV[439] — the
|
||
// production read site changed in Phase 1.3.b-followup-B).
|
||
dd_traj_buf.write_from_slice(&[0.0_f32]);
|
||
|
||
let n2 = 8_usize;
|
||
let sf2: CudaSlice<f32> = stream.clone_htod(&vec![0.0_f32; n2 * sd]).expect("htod sf2");
|
||
let nsf2: CudaSlice<f32> = stream.clone_htod(&vec![0.0_f32; n2 * sd]).expect("htod nsf2");
|
||
let af2: CudaSlice<u32> = stream.clone_htod(&vec![0_u32; n2]).expect("htod af2");
|
||
let rf2: CudaSlice<f32> = stream.clone_htod(&vec![0.0_f32; n2]).expect("htod rf2");
|
||
let df2: CudaSlice<f32> = stream.clone_htod(&vec![0.0_f32; n2]).expect("htod df2");
|
||
let ax2: CudaSlice<i32> = stream.clone_htod(&vec![0_i32; n2]).expect("htod ax2");
|
||
let ac2: CudaSlice<f32> = stream.alloc_zeros::<f32>(n2).expect("aux_conf alloc 2");
|
||
|
||
let ao2_zero: CudaSlice<i32> = stream.alloc_zeros::<i32>(n2).expect("aux_outcome alloc");
|
||
buf.insert_batch(&sf2, &nsf2, &af2, &rf2, &df2, &ax2, &ac2, &ao2_zero, n2)
|
||
.expect("insert_batch (recovery=0)");
|
||
stream.synchronize().expect("sync after insert_batch (recovery=0)");
|
||
|
||
// ── Read back priorities buffer + priorities^alpha ─────────
|
||
// Slots [0..n1) ← batch 1 (boosted to 3.0); slots [n1..n1+n2)
|
||
// ← batch 2 (baseline 1.0). Both columns of the (priority,
|
||
// priority^alpha) row pair must agree:
|
||
// priorities[i] = 3.0 for i<n1 (recovery)
|
||
// = 1.0 for n1<=i<n1+n2 (baseline)
|
||
// priorities_pa[i] = 3.0^0.6 ≈ 1.93 for i<n1
|
||
// = 1.0^0.6 = 1.0 for n1<=i<n1+n2
|
||
let mut prio_host = vec![0.0_f32; cap];
|
||
let mut prio_pa_host = vec![0.0_f32; cap];
|
||
stream
|
||
.memcpy_dtoh(buf.priorities_slice(), &mut prio_host)
|
||
.expect("dtoh priorities");
|
||
stream
|
||
.memcpy_dtoh(buf.priorities_pa_slice(), &mut prio_pa_host)
|
||
.expect("dtoh priorities_pa");
|
||
stream.synchronize().expect("sync after dtoh");
|
||
|
||
// Boosted slots [0..8) — δ=1, ω=2 → factor 3.0, then ^0.6 ≈ 1.9332.
|
||
let expected_pa_boosted = 3.0_f32.powf(0.6);
|
||
for i in 0..n1 {
|
||
assert!(
|
||
(prio_host[i] - 3.0).abs() < 1e-5,
|
||
"priorities[{i}] = {} (recovery slot), expected 3.0 = max_priority(1.0) × (1 + ω(2.0) × δ(1.0))",
|
||
prio_host[i]
|
||
);
|
||
assert!(
|
||
(prio_pa_host[i] - expected_pa_boosted).abs() < 1e-4,
|
||
"priorities_pa[{i}] = {} (recovery slot), expected {expected_pa_boosted} = 3.0^0.6",
|
||
prio_pa_host[i]
|
||
);
|
||
}
|
||
// Baseline slots [8..16) — δ=0 → factor 1.0, then ^0.6 = 1.0.
|
||
for i in n1..(n1 + n2) {
|
||
assert!(
|
||
(prio_host[i] - 1.0).abs() < 1e-5,
|
||
"priorities[{i}] = {} (baseline slot), expected 1.0 = max_priority × (1 + ω × 0)",
|
||
prio_host[i]
|
||
);
|
||
assert!(
|
||
(prio_pa_host[i] - 1.0).abs() < 1e-4,
|
||
"priorities_pa[{i}] = {} (baseline slot), expected 1.0 = 1.0^0.6",
|
||
prio_pa_host[i]
|
||
);
|
||
}
|
||
|
||
// ── Statistical bias verification: ratio of boosted/baseline
|
||
// priorities equals exactly 3.0 (the recovery oversample
|
||
// factor). When the per_sample kernel's prefix-sum then
|
||
// weights by priorities_pa, recovery transitions are sampled
|
||
// 3.0^0.6 ≈ 1.93× more frequently than baseline (per the
|
||
// existing alpha=0.6 PER convention).
|
||
let boosted_mean: f32 = prio_host[..n1].iter().sum::<f32>() / n1 as f32;
|
||
let baseline_mean: f32 = prio_host[n1..(n1 + n2)].iter().sum::<f32>() / n2 as f32;
|
||
let ratio = boosted_mean / baseline_mean;
|
||
assert!(
|
||
(ratio - 3.0).abs() < 1e-5,
|
||
"boosted/baseline priority ratio = {ratio}, expected 3.0 (= 1 + ω(2) × δ(1)) — the recovery oversample factor"
|
||
);
|
||
}
|
||
|
||
/// SP15 Phase 1.3.b-followup-B (2026-05-07) — per-transition
|
||
/// recovery boost. Verifies that when two envs have DIFFERENT
|
||
/// trajectory flags (env-0 in recovery, env-1 not), and BOTH
|
||
/// envs' transitions land in the SAME insert batch, the kernel
|
||
/// applies env-0's 3.0× boost to env-0's transitions and env-1's
|
||
/// 1.0× baseline to env-1's transitions — NOT the uniform
|
||
/// "whichever flag was last" behaviour of Wave 4.3 (where ISV[439]
|
||
/// was read once per insert call and applied to all transitions).
|
||
///
|
||
/// This is the canonical statistical-bias proof: the test directly
|
||
/// fails the Wave 4.3 implementation (which would produce
|
||
/// uniform 3.0× OR uniform 1.0× across all transitions, depending
|
||
/// on which env's trajectory write landed last) and passes the
|
||
/// Phase 1.3.b-followup-B per-transition implementation.
|
||
///
|
||
/// Setup:
|
||
/// * n_envs=2, lookback=2, cf_mult=2 → batch_size=8
|
||
/// * Layout: [main_n1*l1, ..., main_n2*l2, cf_n1*l1, ..., cf_n2*l2]
|
||
/// j=0,1: env-0 main bars (env_id=(0%4)/2=0, (1%4)/2=0)
|
||
/// j=2,3: env-1 main bars (env_id=(2%4)/2=1, (3%4)/2=1)
|
||
/// j=4,5: env-0 cf bars (env_id=(4%4)/2=0, (5%4)/2=0)
|
||
/// j=6,7: env-1 cf bars (env_id=(6%4)/2=1, (7%4)/2=1)
|
||
/// * dd_trajectory_per_env[0]=1.0 (env-0 recovering)
|
||
/// * dd_trajectory_per_env[1]=0.0 (env-1 not in recovery)
|
||
/// * ω=ISV[440]=2.0
|
||
///
|
||
/// Expected per-transition boost:
|
||
/// * env-0's transitions: priority = 1.0 × (1 + 2.0 × 1.0) = 3.0
|
||
/// * env-1's transitions: priority = 1.0 × (1 + 2.0 × 0.0) = 1.0
|
||
/// In the same insert batch, the slot mapping above places:
|
||
/// priorities[0..2] = 3.0 (env-0 main)
|
||
/// priorities[2..4] = 1.0 (env-1 main)
|
||
/// priorities[4..6] = 3.0 (env-0 cf)
|
||
/// priorities[6..8] = 1.0 (env-1 cf)
|
||
///
|
||
/// This pattern is impossible under the Wave 4.3 implementation
|
||
/// (uniform-across-batch); the assertions below directly fail the
|
||
/// Wave 4.3 kernel and pass the per-transition kernel.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn per_sampler_weights_per_transition_not_uniform_across_batch() {
|
||
use ml::dqn::gpu_replay_buffer::{GpuReplayBuffer, GpuReplayBufferConfig};
|
||
use ml::cuda_pipeline::sp15_isv_slots::RECOVERY_OVERSAMPLE_WEIGHT_INDEX;
|
||
use cudarc::driver::CudaSlice;
|
||
|
||
let stream = make_test_stream();
|
||
|
||
// Per-env trajectory tile: env-0 in recovery, env-1 not.
|
||
const N_ENVS: i32 = 2;
|
||
const LOOKBACK: i32 = 2;
|
||
const CF_MULT: usize = 2;
|
||
let batch_size = (N_ENVS as usize) * (LOOKBACK as usize) * CF_MULT; // 8
|
||
|
||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc MappedF32Buffer for isv bus");
|
||
let mut isv_init = vec![0.0_f32; ISV_LEN];
|
||
isv_init[RECOVERY_OVERSAMPLE_WEIGHT_INDEX] = 2.0;
|
||
isv_buf.write_from_slice(&isv_init);
|
||
|
||
let dd_traj_buf = unsafe { MappedF32Buffer::new(N_ENVS as usize) }
|
||
.expect("alloc MappedF32Buffer for dd_trajectory_per_env");
|
||
dd_traj_buf.write_from_slice(&[1.0_f32, 0.0_f32]); // env-0 recovery, env-1 not
|
||
|
||
// GPU PER replay buffer.
|
||
let cap = 64;
|
||
let cfg = GpuReplayBufferConfig {
|
||
capacity: cap,
|
||
alpha: 0.6,
|
||
beta_start: 0.4,
|
||
beta_max: 1.0,
|
||
beta_annealing_steps: 1000,
|
||
epsilon: 1e-6,
|
||
max_memory_bytes: 4 * 1024 * 1024,
|
||
max_batch_size: 64,
|
||
};
|
||
let mut buf = GpuReplayBuffer::new(cfg, &stream).expect("alloc GpuReplayBuffer");
|
||
|
||
// Wire all three pointers / dims.
|
||
buf.set_isv_signals_ptr(isv_buf.dev_ptr);
|
||
buf.set_sp15_dd_trajectory_per_env_ptr(dd_traj_buf.dev_ptr);
|
||
buf.set_sp15_per_env_dims(N_ENVS, LOOKBACK);
|
||
|
||
// Insert one batch with 8 transitions; layout matches
|
||
// experience-collector contract `[main_N*L | cf_N*L]`. The
|
||
// host-side data values don't matter for the boost computation
|
||
// — only the per-transition env_id mapping does.
|
||
let sd = ml::state_layout::STATE_DIM;
|
||
let states_host: Vec<f32> = vec![0.0; batch_size * sd];
|
||
let next_states_host: Vec<f32> = vec![0.0; batch_size * sd];
|
||
let actions_host: Vec<u32> = vec![0_u32; batch_size];
|
||
let rewards_host: Vec<f32> = vec![0.0; batch_size];
|
||
let dones_host: Vec<f32> = vec![0.0; batch_size];
|
||
let aux_sign_host: Vec<i32> = vec![0_i32; batch_size];
|
||
|
||
let sf: CudaSlice<f32> = stream.clone_htod(&states_host).expect("htod states");
|
||
let nsf: CudaSlice<f32> = stream.clone_htod(&next_states_host).expect("htod next_states");
|
||
let af: CudaSlice<u32> = stream.clone_htod(&actions_host).expect("htod actions");
|
||
let rf: CudaSlice<f32> = stream.clone_htod(&rewards_host).expect("htod rewards");
|
||
let df: CudaSlice<f32> = stream.clone_htod(&dones_host).expect("htod dones");
|
||
let ax: CudaSlice<i32> = stream.clone_htod(&aux_sign_host).expect("htod aux_sign");
|
||
let ac: CudaSlice<f32> = stream.alloc_zeros::<f32>(batch_size).expect("aux_conf alloc");
|
||
|
||
let ao_zero: CudaSlice<i32> = stream.alloc_zeros::<i32>(batch_size).expect("aux_outcome alloc");
|
||
buf.insert_batch(&sf, &nsf, &af, &rf, &df, &ax, &ac, &ao_zero, batch_size)
|
||
.expect("insert_batch (mixed-env)");
|
||
stream.synchronize().expect("sync after insert_batch");
|
||
|
||
// Read back the priorities buffer.
|
||
let mut prio_host = vec![0.0_f32; cap];
|
||
stream
|
||
.memcpy_dtoh(buf.priorities_slice(), &mut prio_host)
|
||
.expect("dtoh priorities");
|
||
stream.synchronize().expect("sync after dtoh");
|
||
|
||
// Verify the per-transition boost pattern: env-0's 4
|
||
// transitions (j=0,1,4,5) get 3.0; env-1's 4 transitions
|
||
// (j=2,3,6,7) get 1.0.
|
||
let expected: [f32; 8] = [3.0, 3.0, 1.0, 1.0, 3.0, 3.0, 1.0, 1.0];
|
||
for (j, &exp) in expected.iter().enumerate() {
|
||
assert!(
|
||
(prio_host[j] - exp).abs() < 1e-5,
|
||
"priorities[{j}] = {} (env_id={}; trajectory={}), expected {} per per-transition lookup. Wave 4.3 uniform-batch bug would produce 3.0 OR 1.0 across ALL slots.",
|
||
prio_host[j],
|
||
(j % 4) / 2,
|
||
if (j % 4) / 2 == 0 { 1.0 } else { 0.0 },
|
||
exp
|
||
);
|
||
}
|
||
|
||
// Statistical-bias verification: the env-0 mean (3.0) must be
|
||
// strictly greater than the env-1 mean (1.0). Under Wave 4.3
|
||
// both means would be equal (all transitions sharing the
|
||
// same scalar boost factor).
|
||
let env0_mean: f32 = (prio_host[0] + prio_host[1] + prio_host[4] + prio_host[5]) / 4.0;
|
||
let env1_mean: f32 = (prio_host[2] + prio_host[3] + prio_host[6] + prio_host[7]) / 4.0;
|
||
assert!(
|
||
(env0_mean - 3.0).abs() < 1e-5 && (env1_mean - 1.0).abs() < 1e-5,
|
||
"env-0 mean = {env0_mean} (expected 3.0); env-1 mean = {env1_mean} (expected 1.0). \
|
||
Wave 4.3 uniform-batch bug would produce env0_mean == env1_mean."
|
||
);
|
||
assert!(
|
||
env0_mean > env1_mean + 1.0,
|
||
"env-0 mean ({env0_mean}) must exceed env-1 mean ({env1_mean}) by ≥1.0 per per-transition boost; \
|
||
Wave 4.3 would produce env0_mean - env1_mean = 0 (uniform-across-batch)."
|
||
);
|
||
}
|
||
|
||
/// Test 1.1.b — numerical equivalence of the SP15 Phase 1.1.b sharpe
|
||
/// split.
|
||
///
|
||
/// Before Phase 1.1.b, `compute_backtest_metrics` computed sharpe inline
|
||
/// via:
|
||
/// `out[sharpe] = (mean / sqrt(max(var, 1e-10))) * annualization`
|
||
/// where `var = sum_sq/n - mean^2` and `mean = sum/n`.
|
||
///
|
||
/// After the split:
|
||
/// `sharpe_per_bar_kernel` returns `out[2] = mean / max(std, 1e-12)`
|
||
/// where `std = sqrt(max(sum_dev_sq/n, 1e-12))` (two-pass: mean →
|
||
/// sum of (r-mean)²).
|
||
/// The host then multiplies by `annualization_factor`.
|
||
///
|
||
/// The two formulas are mathematically identical (one-pass vs two-pass
|
||
/// variance, same lower bound semantics within float precision). This
|
||
/// test launches the new kernel against a synthetic non-trivial PnL
|
||
/// series and asserts the output matches the closed-form expectation
|
||
/// `(mean / std) * annualization_factor` to within 1e-5 relative
|
||
/// tolerance — which is the same tolerance one would obtain comparing
|
||
/// the pre-split and post-split implementations directly (both
|
||
/// algorithms agree to ~f32 precision; reduction-order differences in
|
||
/// the block tree-reduce produce ~1e-6 relative drift).
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn unified_sharpe_kernel_equivalence_under_annualization() {
|
||
let stream = make_test_stream();
|
||
|
||
// Synthetic series: deterministic sin-wave + drift.
|
||
// n=512, mean ≈ 0.05, std ≈ ~0.7. Sharpe(raw) ≈ 0.07.
|
||
let n: usize = 512;
|
||
let pnl: Vec<f32> = (0..n)
|
||
.map(|i| {
|
||
let t = i as f32 / 32.0;
|
||
0.05 + (t * 0.5).sin()
|
||
})
|
||
.collect();
|
||
// Annualization factor matches `GpuBacktestConfig::default()`:
|
||
// sqrt(390 * 252) ≈ 313.5.
|
||
let annualization_factor: f32 = (390.0_f32 * 252.0).sqrt();
|
||
|
||
// Closed-form expectation (canonical f64 reference).
|
||
let n_f = n as f64;
|
||
let sum: f64 = pnl.iter().map(|&x| x as f64).sum();
|
||
let mean: f64 = sum / n_f;
|
||
let sum_dev_sq: f64 = pnl.iter().map(|&x| {
|
||
let d = x as f64 - mean;
|
||
d * d
|
||
}).sum();
|
||
let var: f64 = sum_dev_sq / n_f;
|
||
let std: f64 = var.sqrt();
|
||
let expected_raw: f64 = mean / std;
|
||
let expected_annualised: f64 = expected_raw * annualization_factor as f64;
|
||
|
||
// Safety: CUDA context active via `make_test_stream`.
|
||
let pnl_buf = unsafe { MappedF32Buffer::new(n) }
|
||
.expect("alloc MappedF32Buffer for sharpe input");
|
||
pnl_buf.write_from_slice(&pnl);
|
||
let out_buf = unsafe { MappedF32Buffer::new(3) }
|
||
.expect("alloc MappedF32Buffer for sharpe output");
|
||
|
||
launch_sp15_sharpe_per_bar(
|
||
&stream,
|
||
pnl_buf.dev_ptr,
|
||
n as i32,
|
||
out_buf.dev_ptr,
|
||
)
|
||
.expect("launch sharpe_per_bar_kernel");
|
||
stream
|
||
.synchronize()
|
||
.expect("synchronize after sharpe_per_bar_kernel launch");
|
||
|
||
let out = out_buf.read_all();
|
||
let raw_sharpe: f32 = out[2];
|
||
// Apply host-side annualisation — same code path as
|
||
// `gpu_backtest_evaluator::consume_metrics_after_event`.
|
||
let actual_annualised: f32 = raw_sharpe * annualization_factor;
|
||
|
||
// f32 reduction-order tolerance: relative 1e-5 (5 decimal digits).
|
||
let rel_err = ((actual_annualised as f64 - expected_annualised).abs()
|
||
/ expected_annualised.abs()) as f64;
|
||
assert!(
|
||
rel_err < 1e-5,
|
||
"annualised sharpe relative error {:.3e} exceeds 1e-5: \
|
||
expected {:.6}, got {:.6}",
|
||
rel_err,
|
||
expected_annualised,
|
||
actual_annualised
|
||
);
|
||
|
||
// Cross-check raw_sharpe alone (without annualisation) — kernel
|
||
// contract sanity. Same tolerance.
|
||
let rel_err_raw = ((raw_sharpe as f64 - expected_raw).abs()
|
||
/ expected_raw.abs()) as f64;
|
||
assert!(
|
||
rel_err_raw < 1e-5,
|
||
"raw sharpe relative error {:.3e} exceeds 1e-5: \
|
||
expected {:.6}, got {:.6}",
|
||
rel_err_raw,
|
||
expected_raw,
|
||
raw_sharpe
|
||
);
|
||
}
|
||
|
||
// ── SP15 Phase 3.5.b + 3.5.3.b oracle tests ──────────────────────────
|
||
//
|
||
// These tests exercise the inline hold_floor + cooldown-mask wiring
|
||
// inside the production `experience_action_select` kernel. The fixture
|
||
// mirrors `distributional_q_tests.rs::ProdActionSelectFixture` (which
|
||
// is `pub(crate)` in the trainer's test module and not reachable from
|
||
// this integration-test file) — we re-build a minimal launcher here
|
||
// to keep the oracle test self-contained. Branch sizes match the
|
||
// production 4-direction action space.
|
||
|
||
use cudarc::driver::{CudaFunction, CudaSlice, LaunchConfig, PushKernelArg};
|
||
use std::sync::OnceLock;
|
||
|
||
const ACTION_SELECT_PROD_EXP_CUBIN: &[u8] = include_bytes!(concat!(
|
||
env!("OUT_DIR"),
|
||
"/experience_kernels.cubin"
|
||
));
|
||
static ACTION_SELECT_KERNEL: OnceLock<CudaFunction> = OnceLock::new();
|
||
|
||
fn load_action_select(stream: &Arc<CudaStream>) -> &'static CudaFunction {
|
||
ACTION_SELECT_KERNEL.get_or_init(|| {
|
||
let module = stream
|
||
.context()
|
||
.load_cubin(ACTION_SELECT_PROD_EXP_CUBIN.to_vec())
|
||
.expect("load experience_kernels cubin");
|
||
module
|
||
.load_function("experience_action_select")
|
||
.expect("load experience_action_select")
|
||
})
|
||
}
|
||
|
||
/// Branch sizes for the 4-direction production action space.
|
||
const AS_B0: i32 = 4; // direction (Short=0, Hold=1, Long=2, Flat=3)
|
||
const AS_B1: i32 = 3; // magnitude
|
||
const AS_B2: i32 = 3; // order
|
||
const AS_B3: i32 = 3; // urgency
|
||
const AS_Q_STRIDE: usize = (AS_B0 + AS_B1 + AS_B2 + AS_B3) as usize;
|
||
const AS_DIR_HOLD: i32 = 1;
|
||
const AS_DIR_LONG: i32 = 2;
|
||
|
||
fn decode_dir(action_idx: i32) -> i32 {
|
||
action_idx / (AS_B1 * AS_B2 * AS_B3)
|
||
}
|
||
|
||
/// Build C51 logits that produce a deterministic per-direction E[Q]
|
||
/// equal to `target_v[d]`. Uses a high peak logit so softmax is
|
||
/// ~delta on the closest atom; with the inverse-CDF Thompson sample
|
||
/// the `q_sample` ≈ `e_dir` (jitter < linear-support delta_z),
|
||
/// keeping `q_eff = e_dir + temp·(q_sample-e_dir)` tightly anchored
|
||
/// at e_dir for any temp ∈ [0.5, 2.0].
|
||
fn fill_peaked_logits_action_select(
|
||
batch: usize,
|
||
n_atoms: i32,
|
||
v_min: f32,
|
||
delta_z: f32,
|
||
target_v: &[f32; 4],
|
||
peak_logit: f32,
|
||
) -> Vec<f32> {
|
||
let mut out = vec![0.0_f32; batch * (AS_B0 as usize) * (n_atoms as usize)];
|
||
for i in 0..batch {
|
||
for (d, tv) in target_v.iter().enumerate() {
|
||
let peak_a = (((*tv - v_min) / delta_z).round() as usize)
|
||
.min((n_atoms - 1) as usize);
|
||
let base =
|
||
i * (AS_B0 as usize) * (n_atoms as usize) + d * (n_atoms as usize);
|
||
out[base + peak_a] = peak_logit;
|
||
}
|
||
}
|
||
out
|
||
}
|
||
|
||
fn fill_linear_support_action_select(
|
||
batch: usize, v_min: f32, v_max: f32, delta_z: f32,
|
||
) -> Vec<f32> {
|
||
let mut out = vec![0.0_f32; batch * (AS_B0 as usize) * 3];
|
||
for i in 0..batch {
|
||
for d in 0..(AS_B0 as usize) {
|
||
let base = i * (AS_B0 as usize) * 3 + d * 3;
|
||
out[base] = v_min;
|
||
out[base + 1] = v_max;
|
||
out[base + 2] = delta_z;
|
||
}
|
||
}
|
||
out
|
||
}
|
||
|
||
/// Launch `experience_action_select` with deterministic eval-mode
|
||
/// inputs (eps_start = eps_end = 0). `isv_dev_ptr` MUST point to a
|
||
/// >= ISV_LEN f32 buffer (slot 339 = thompson_temp; slot 426 = α;
|
||
/// slot 427 = k; slot 428 = ε₀; slot 435 = cooldown_remaining).
|
||
/// Returns the per-sample factored action picks.
|
||
fn launch_action_select_for_test(
|
||
stream: &Arc<CudaStream>,
|
||
batch: usize,
|
||
n_atoms: i32,
|
||
q_values: &CudaSlice<f32>,
|
||
b_logits_dir: &CudaSlice<f32>,
|
||
per_sample_support: &CudaSlice<f32>,
|
||
actions_buf: &mut CudaSlice<i32>,
|
||
q_gaps_buf: &mut CudaSlice<f32>,
|
||
intent_buf: &mut CudaSlice<i32>,
|
||
conv_buf: &mut CudaSlice<f32>,
|
||
mag_conv_buf: &mut CudaSlice<f32>,
|
||
isv_dev_ptr: u64,
|
||
) -> Vec<i32> {
|
||
let kernel = load_action_select(stream);
|
||
let blocks = ((batch as u32).div_ceil(256)).max(1);
|
||
let cfg = LaunchConfig {
|
||
grid_dim: (blocks, 1, 1),
|
||
block_dim: (256, 1, 1),
|
||
shared_mem_bytes: 0,
|
||
};
|
||
let n_i32: i32 = batch as i32;
|
||
let null_ptr: u64 = 0;
|
||
// Eval mode: cosine schedule degenerate (eps_start = eps_end = 0).
|
||
let eps_start: f32 = 0.0;
|
||
let eps_end: f32 = 0.0;
|
||
unsafe {
|
||
stream
|
||
.launch_builder(kernel)
|
||
.arg(q_values)
|
||
.arg(&mut *actions_buf)
|
||
.arg(&mut *q_gaps_buf)
|
||
.arg(&eps_start)
|
||
.arg(&eps_end)
|
||
.arg(&0_i32) // current_epoch
|
||
.arg(&1_i32) // total_epochs
|
||
.arg(&n_i32)
|
||
.arg(&AS_B0)
|
||
.arg(&AS_B1)
|
||
.arg(&AS_B2)
|
||
.arg(&AS_B3)
|
||
.arg(&0.0_f32) // q_gap_threshold
|
||
.arg(&null_ptr) // portfolio_states (NULL)
|
||
.arg(&0_i32) // min_hold_bars
|
||
.arg(&0.0_f32) // max_position
|
||
.arg(&1.0_f32) // eps_exp_mult
|
||
.arg(&1.0_f32) // eps_ord_mult
|
||
.arg(&1.0_f32) // eps_urg_mult
|
||
.arg(&0_i32) // timestep
|
||
.arg(&null_ptr) // per_sample_epsilon (NULL → cosine)
|
||
.arg(&isv_dev_ptr) // isv_signals_ptr
|
||
.arg(&0_i32) // contrarian_active
|
||
.arg(&mut *intent_buf)
|
||
.arg(&mut *conv_buf)
|
||
.arg(&mut *mag_conv_buf)
|
||
.arg(b_logits_dir)
|
||
.arg(per_sample_support)
|
||
.arg(&null_ptr) // atom_positions = NULL → linear support
|
||
.arg(&n_atoms)
|
||
// SP15 Phase 3.5.4.c (2026-05-07): plasticity_m_warm.
|
||
// Default 0.0f → fire-bar predicate dead. The new
|
||
// `plasticity_fires_force_flat_then_cooldown_holds` test
|
||
// overrides this via `launch_action_select_for_test_with_m_warm`.
|
||
.arg(&0.0f32)
|
||
.launch(cfg)
|
||
.expect("launch experience_action_select");
|
||
}
|
||
stream.synchronize().expect("sync after action_select");
|
||
stream.clone_dtoh(actions_buf).expect("readback actions")
|
||
}
|
||
|
||
/// SP15 Phase 3.5.4.c (2026-05-07): variant of `launch_action_select_for_test`
|
||
/// that takes an explicit `m_warm` so the new behavioral test
|
||
/// `plasticity_fires_force_flat_then_cooldown_holds` can drive the
|
||
/// kernel's fire-bar / OR-gate consumer paths. All other args
|
||
/// match the m_warm-defaulted helper.
|
||
#[allow(clippy::too_many_arguments)]
|
||
fn launch_action_select_for_test_with_m_warm(
|
||
stream: &Arc<CudaStream>,
|
||
batch: usize,
|
||
n_atoms: i32,
|
||
q_values: &CudaSlice<f32>,
|
||
b_logits_dir: &CudaSlice<f32>,
|
||
per_sample_support: &CudaSlice<f32>,
|
||
actions_buf: &mut CudaSlice<i32>,
|
||
q_gaps_buf: &mut CudaSlice<f32>,
|
||
intent_buf: &mut CudaSlice<i32>,
|
||
conv_buf: &mut CudaSlice<f32>,
|
||
mag_conv_buf: &mut CudaSlice<f32>,
|
||
isv_dev_ptr: u64,
|
||
m_warm: f32,
|
||
) -> Vec<i32> {
|
||
let kernel = load_action_select(stream);
|
||
let blocks = ((batch as u32).div_ceil(256)).max(1);
|
||
let cfg = LaunchConfig {
|
||
grid_dim: (blocks, 1, 1),
|
||
block_dim: (256, 1, 1),
|
||
shared_mem_bytes: 0,
|
||
};
|
||
let n_i32: i32 = batch as i32;
|
||
let null_ptr: u64 = 0;
|
||
let eps_start: f32 = 0.0;
|
||
let eps_end: f32 = 0.0;
|
||
unsafe {
|
||
stream
|
||
.launch_builder(kernel)
|
||
.arg(q_values)
|
||
.arg(&mut *actions_buf)
|
||
.arg(&mut *q_gaps_buf)
|
||
.arg(&eps_start)
|
||
.arg(&eps_end)
|
||
.arg(&0_i32) // current_epoch
|
||
.arg(&1_i32) // total_epochs
|
||
.arg(&n_i32)
|
||
.arg(&AS_B0)
|
||
.arg(&AS_B1)
|
||
.arg(&AS_B2)
|
||
.arg(&AS_B3)
|
||
.arg(&0.0_f32) // q_gap_threshold
|
||
.arg(&null_ptr)
|
||
.arg(&0_i32)
|
||
.arg(&0.0_f32)
|
||
.arg(&1.0_f32)
|
||
.arg(&1.0_f32)
|
||
.arg(&1.0_f32)
|
||
.arg(&0_i32)
|
||
.arg(&null_ptr)
|
||
.arg(&isv_dev_ptr)
|
||
.arg(&0_i32)
|
||
.arg(&mut *intent_buf)
|
||
.arg(&mut *conv_buf)
|
||
.arg(&mut *mag_conv_buf)
|
||
.arg(b_logits_dir)
|
||
.arg(per_sample_support)
|
||
.arg(&null_ptr)
|
||
.arg(&n_atoms)
|
||
.arg(&m_warm)
|
||
.launch(cfg)
|
||
.expect("launch experience_action_select");
|
||
}
|
||
stream.synchronize().expect("sync after action_select");
|
||
stream.clone_dtoh(actions_buf).expect("readback actions")
|
||
}
|
||
|
||
/// SP15 Phase 3.5.b — confidence-aware Hold floor wired inline.
|
||
///
|
||
/// Construct a per-sample direction posterior with a near-uniform
|
||
/// shape so Shannon entropy is high (≈ ln(4) ≈ 1.386 nats > ε₀ = 1.0).
|
||
/// e_dir = [0.50, 0.45, 0.50, 0.45] yields softmax probs ≈
|
||
/// [0.255, 0.245, 0.255, 0.245] → entropy ≈ 1.386 nats.
|
||
/// With α = 0.5, k = 10, ε₀ = 1.0:
|
||
/// x = 10 × (1.386 − 1.0) ≈ 3.86, σ(x) ≈ 0.979, hold_floor ≈ 0.490.
|
||
///
|
||
/// Without hold_floor: argmax(e_dir) is Short(0) or Long(2) (ties to
|
||
/// the lower index → Short wins on the strictly-greater test). With
|
||
/// hold_floor: Hold's e_dir 0.45 + 0.49 = 0.94 dominates, so picked
|
||
/// dir = Hold.
|
||
///
|
||
/// Validates: (a) the inline hold_floor computation; (b) the
|
||
/// SP15 3.5.b architectural property that the floor shifts the
|
||
/// argmax to Hold under high-entropy uncertainty.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn action_select_applies_hold_floor_inline() {
|
||
use ml::cuda_pipeline::sp15_isv_slots::{
|
||
COOLDOWN_BARS_REMAINING_INDEX, HOLD_FLOOR_ALPHA_INDEX,
|
||
HOLD_FLOOR_EPS0_INDEX, HOLD_FLOOR_K_INDEX,
|
||
};
|
||
let stream = make_test_stream();
|
||
let batch: usize = 1;
|
||
let n_atoms: i32 = 51;
|
||
let v_min: f32 = -1.0;
|
||
let v_max: f32 = 1.0;
|
||
let delta_z = (v_max - v_min) / (n_atoms as f32 - 1.0);
|
||
|
||
// Near-uniform e_dir: [Short=0.50, Hold=0.45, Long=0.50, Flat=0.45].
|
||
// Softmax(near-uniform) → high entropy ≈ ln(4).
|
||
let target_v = [0.50_f32, 0.45_f32, 0.50_f32, 0.45_f32];
|
||
let b_logits_host = fill_peaked_logits_action_select(
|
||
batch, n_atoms, v_min, delta_z, &target_v, /*peak_logit*/ 50.0,
|
||
);
|
||
let support_host = fill_linear_support_action_select(batch, v_min, v_max, delta_z);
|
||
let q_values_host = vec![0.0_f32; batch * AS_Q_STRIDE];
|
||
|
||
// ISV bus with Thompson temp at floor (0.5) so q_eff stays close
|
||
// to e_dir; hold_floor anchors at α = 0.5, k = 10, ε₀ = 1.0;
|
||
// cooldown disabled (= 0).
|
||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc MappedF32Buffer for isv bus");
|
||
let mut isv_init = vec![0.0f32; ISV_LEN];
|
||
// ISV[ISV_EVAL_THOMPSON_TEMP_IDX = 339] — 0.5 (floor).
|
||
isv_init[339] = 0.5;
|
||
isv_init[HOLD_FLOOR_ALPHA_INDEX] = 0.5;
|
||
isv_init[HOLD_FLOOR_K_INDEX] = 10.0;
|
||
isv_init[HOLD_FLOOR_EPS0_INDEX] = 1.0;
|
||
isv_init[COOLDOWN_BARS_REMAINING_INDEX] = 0.0;
|
||
isv_buf.write_from_slice(&isv_init);
|
||
|
||
// GPU buffers.
|
||
let q_values = stream.clone_htod(&q_values_host).expect("upload q_values");
|
||
let b_logits_dir = stream.clone_htod(&b_logits_host).expect("upload logits");
|
||
let per_sample_support = stream
|
||
.clone_htod(&support_host)
|
||
.expect("upload support");
|
||
let mut actions_buf = stream.alloc_zeros::<i32>(batch).expect("alloc actions");
|
||
let mut q_gaps_buf = stream.alloc_zeros::<f32>(batch).expect("alloc q_gaps");
|
||
let mut intent_buf = stream.alloc_zeros::<i32>(batch).expect("alloc intent");
|
||
let mut conv_buf = stream.alloc_zeros::<f32>(batch).expect("alloc conviction");
|
||
let mut mag_conv_buf = stream.alloc_zeros::<f32>(batch).expect("alloc mag_conv");
|
||
|
||
let actions = launch_action_select_for_test(
|
||
&stream, batch, n_atoms,
|
||
&q_values, &b_logits_dir, &per_sample_support,
|
||
&mut actions_buf, &mut q_gaps_buf, &mut intent_buf,
|
||
&mut conv_buf, &mut mag_conv_buf,
|
||
isv_buf.dev_ptr,
|
||
);
|
||
|
||
let dir_picked = decode_dir(actions[0]);
|
||
assert_eq!(
|
||
dir_picked, AS_DIR_HOLD,
|
||
"hold_floor inline must shift argmax to Hold when entropy > ε₀; \
|
||
got dir = {dir_picked} (expected {AS_DIR_HOLD})"
|
||
);
|
||
}
|
||
|
||
/// SP15 Phase 3.5.3.b — cooldown mask: when COOLDOWN_BARS_REMAINING > 0,
|
||
/// the picked direction is forced to Hold regardless of e_dir
|
||
/// posterior. Sets up e_dir = [Short=0.0, Hold=0.0, Long=1.0, Flat=0.0]
|
||
/// (Long strongly preferred), then activates cooldown (= 5.0). The
|
||
/// mask short-circuits Pass 2 entirely → dir = Hold.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn action_select_forces_hold_during_cooldown() {
|
||
use ml::cuda_pipeline::sp15_isv_slots::{
|
||
COOLDOWN_BARS_REMAINING_INDEX, HOLD_FLOOR_ALPHA_INDEX,
|
||
HOLD_FLOOR_EPS0_INDEX, HOLD_FLOOR_K_INDEX,
|
||
};
|
||
let stream = make_test_stream();
|
||
let batch: usize = 1;
|
||
let n_atoms: i32 = 51;
|
||
let v_min: f32 = -1.0;
|
||
let v_max: f32 = 1.0;
|
||
let delta_z = (v_max - v_min) / (n_atoms as f32 - 1.0);
|
||
|
||
// Long strongly preferred over Hold (e_dir gap = 1.0); without
|
||
// cooldown the argmax would land on Long.
|
||
let target_v = [0.0_f32, 0.0_f32, 1.0_f32, 0.0_f32];
|
||
let b_logits_host = fill_peaked_logits_action_select(
|
||
batch, n_atoms, v_min, delta_z, &target_v, /*peak_logit*/ 50.0,
|
||
);
|
||
let support_host = fill_linear_support_action_select(batch, v_min, v_max, delta_z);
|
||
let q_values_host = vec![0.0_f32; batch * AS_Q_STRIDE];
|
||
|
||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc MappedF32Buffer for isv bus");
|
||
let mut isv_init = vec![0.0f32; ISV_LEN];
|
||
isv_init[339] = 0.5;
|
||
isv_init[HOLD_FLOOR_ALPHA_INDEX] = 0.5;
|
||
isv_init[HOLD_FLOOR_K_INDEX] = 10.0;
|
||
isv_init[HOLD_FLOOR_EPS0_INDEX] = 1.0;
|
||
// Cooldown active — kernel must force Hold.
|
||
isv_init[COOLDOWN_BARS_REMAINING_INDEX] = 5.0;
|
||
isv_buf.write_from_slice(&isv_init);
|
||
|
||
let q_values = stream.clone_htod(&q_values_host).expect("upload q_values");
|
||
let b_logits_dir = stream.clone_htod(&b_logits_host).expect("upload logits");
|
||
let per_sample_support = stream
|
||
.clone_htod(&support_host)
|
||
.expect("upload support");
|
||
let mut actions_buf = stream.alloc_zeros::<i32>(batch).expect("alloc actions");
|
||
let mut q_gaps_buf = stream.alloc_zeros::<f32>(batch).expect("alloc q_gaps");
|
||
let mut intent_buf = stream.alloc_zeros::<i32>(batch).expect("alloc intent");
|
||
let mut conv_buf = stream.alloc_zeros::<f32>(batch).expect("alloc conviction");
|
||
let mut mag_conv_buf = stream.alloc_zeros::<f32>(batch).expect("alloc mag_conv");
|
||
|
||
let actions = launch_action_select_for_test(
|
||
&stream, batch, n_atoms,
|
||
&q_values, &b_logits_dir, &per_sample_support,
|
||
&mut actions_buf, &mut q_gaps_buf, &mut intent_buf,
|
||
&mut conv_buf, &mut mag_conv_buf,
|
||
isv_buf.dev_ptr,
|
||
);
|
||
|
||
let dir_picked = decode_dir(actions[0]);
|
||
assert_eq!(
|
||
dir_picked, AS_DIR_HOLD,
|
||
"cooldown mask must force Hold regardless of e_dir; \
|
||
got dir = {dir_picked} (expected {AS_DIR_HOLD}, even though Long had +1.0 advantage)"
|
||
);
|
||
}
|
||
|
||
/// SP15 Phase 3.5.3.b — when cooldown counter is 0 the mask is
|
||
/// inactive; the picked direction follows the standard Thompson +
|
||
/// hold_floor argmax. With Long strongly preferred (e_dir Long=1.0,
|
||
/// rest=0.0) AND the entropy below ε₀ (one-hot ≈ 0 entropy) so
|
||
/// hold_floor ≈ 0, picked dir should be Long.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn action_select_no_force_hold_when_cooldown_zero() {
|
||
use ml::cuda_pipeline::sp15_isv_slots::{
|
||
COOLDOWN_BARS_REMAINING_INDEX, HOLD_FLOOR_ALPHA_INDEX,
|
||
HOLD_FLOOR_EPS0_INDEX, HOLD_FLOOR_K_INDEX,
|
||
};
|
||
let stream = make_test_stream();
|
||
let batch: usize = 1;
|
||
let n_atoms: i32 = 51;
|
||
let v_min: f32 = -1.0;
|
||
let v_max: f32 = 1.0;
|
||
let delta_z = (v_max - v_min) / (n_atoms as f32 - 1.0);
|
||
|
||
// Long strongly preferred; near-degenerate e_dir → entropy near
|
||
// 0 (well below ε₀ = 1.0) → hold_floor ≈ α × σ(-10) ≈ 0.5 × 4.5e-5
|
||
// ≈ 2e-5 → Hold gain negligible → argmax stays at Long.
|
||
let target_v = [0.0_f32, 0.0_f32, 1.0_f32, 0.0_f32];
|
||
let b_logits_host = fill_peaked_logits_action_select(
|
||
batch, n_atoms, v_min, delta_z, &target_v, /*peak_logit*/ 50.0,
|
||
);
|
||
let support_host = fill_linear_support_action_select(batch, v_min, v_max, delta_z);
|
||
let q_values_host = vec![0.0_f32; batch * AS_Q_STRIDE];
|
||
|
||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc MappedF32Buffer for isv bus");
|
||
let mut isv_init = vec![0.0f32; ISV_LEN];
|
||
isv_init[339] = 0.5;
|
||
isv_init[HOLD_FLOOR_ALPHA_INDEX] = 0.5;
|
||
isv_init[HOLD_FLOOR_K_INDEX] = 10.0;
|
||
isv_init[HOLD_FLOOR_EPS0_INDEX] = 1.0;
|
||
// Cooldown inactive.
|
||
isv_init[COOLDOWN_BARS_REMAINING_INDEX] = 0.0;
|
||
isv_buf.write_from_slice(&isv_init);
|
||
|
||
let q_values = stream.clone_htod(&q_values_host).expect("upload q_values");
|
||
let b_logits_dir = stream.clone_htod(&b_logits_host).expect("upload logits");
|
||
let per_sample_support = stream
|
||
.clone_htod(&support_host)
|
||
.expect("upload support");
|
||
let mut actions_buf = stream.alloc_zeros::<i32>(batch).expect("alloc actions");
|
||
let mut q_gaps_buf = stream.alloc_zeros::<f32>(batch).expect("alloc q_gaps");
|
||
let mut intent_buf = stream.alloc_zeros::<i32>(batch).expect("alloc intent");
|
||
let mut conv_buf = stream.alloc_zeros::<f32>(batch).expect("alloc conviction");
|
||
let mut mag_conv_buf = stream.alloc_zeros::<f32>(batch).expect("alloc mag_conv");
|
||
|
||
let actions = launch_action_select_for_test(
|
||
&stream, batch, n_atoms,
|
||
&q_values, &b_logits_dir, &per_sample_support,
|
||
&mut actions_buf, &mut q_gaps_buf, &mut intent_buf,
|
||
&mut conv_buf, &mut mag_conv_buf,
|
||
isv_buf.dev_ptr,
|
||
);
|
||
|
||
let dir_picked = decode_dir(actions[0]);
|
||
assert_eq!(
|
||
dir_picked, AS_DIR_LONG,
|
||
"with cooldown = 0 and Long strongly preferred, picked dir \
|
||
must be Long; got dir = {dir_picked} (expected {AS_DIR_LONG})"
|
||
);
|
||
}
|
||
|
||
/// SP15 Phase 3.5.4.c (2026-05-07) — full plasticity-injection two-step
|
||
/// recovery sequence end-to-end. Sets up the ISV state so the trigger
|
||
/// fires (DD_PERSISTENCE > threshold, fired flag clear, warm = 0),
|
||
/// then drives the per-step launch sequence
|
||
/// `launch_sp15_plasticity_injection` → `experience_action_select`
|
||
/// across `M_warm + 2` bars. Asserts:
|
||
///
|
||
/// - **Bar 0 (fire bar)**: dir_picked == DIR_FLAT — the trigger sets
|
||
/// warm = M_warm then decrements to (M_warm − 1), so action_select
|
||
/// observes warm ≥ M_warm − 1.5f and forces DIR_FLAT (step 1 of
|
||
/// the spec §9.2 (3.5.4) two-step). e_dir is set to strongly
|
||
/// prefer Long (gap = +1.0); without the fire-bar branch the
|
||
/// argmax would land on Long, so observing DIR_FLAT proves the
|
||
/// fire-bar consumer fired.
|
||
///
|
||
/// - **Bars 1..M_warm − 1 (warm-up)**: dir_picked == DIR_HOLD on
|
||
/// every bar. The trigger debounces (already fired this fold)
|
||
/// but continues to decrement warm; the action_select OR-gate
|
||
/// forces Hold while warm > 0 (step 2 of the two-step). Bar 1
|
||
/// observes warm = M_warm − 2; bar M_warm − 1 observes warm = 1.
|
||
///
|
||
/// - **Bar M_warm (warm expired)**: warm = 0 (the trigger's
|
||
/// decrement hit floor on the previous bar). action_select OR-gate
|
||
/// inactive (cooldown also 0); standard Thompson + hold_floor
|
||
/// argmax resumes; with Long strongly preferred and entropy near
|
||
/// 0 → hold_floor ≈ 0 → dir_picked == DIR_LONG.
|
||
///
|
||
/// Uses M_warm = 5 (down from production's 200) to keep the test fast
|
||
/// while preserving the invariant. The `fan_in` / `n_weights` choices
|
||
/// don't matter for the consumer-side assertions (the kernel writes
|
||
/// the trailing 10% of advantage_head_weights in addition to its ISV
|
||
/// updates, but this test doesn't read those weights). Driven via
|
||
/// `launch_action_select_for_test_with_m_warm` so the consumer sees
|
||
/// the same `m_warm` the trigger wrote.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn plasticity_fires_force_flat_then_cooldown_holds() {
|
||
use ml::cuda_pipeline::sp15_isv_slots::{
|
||
COOLDOWN_BARS_REMAINING_INDEX, DD_PERSISTENCE_MAX_INDEX,
|
||
HOLD_FLOOR_ALPHA_INDEX, HOLD_FLOOR_EPS0_INDEX, HOLD_FLOOR_K_INDEX,
|
||
PLASTICITY_FIRED_THIS_FOLD_INDEX,
|
||
PLASTICITY_PERSISTENCE_THRESHOLD_INDEX,
|
||
PLASTICITY_WARM_BARS_REMAINING_INDEX,
|
||
};
|
||
|
||
let stream = make_test_stream();
|
||
let batch: usize = 1;
|
||
let n_atoms: i32 = 51;
|
||
let v_min: f32 = -1.0;
|
||
let v_max: f32 = 1.0;
|
||
let delta_z = (v_max - v_min) / (n_atoms as f32 - 1.0);
|
||
|
||
// Long strongly preferred → without fire-bar / OR-gate, argmax = Long.
|
||
let target_v = [0.0_f32, 0.0_f32, 1.0_f32, 0.0_f32];
|
||
let b_logits_host = fill_peaked_logits_action_select(
|
||
batch, n_atoms, v_min, delta_z, &target_v, /*peak_logit*/ 50.0,
|
||
);
|
||
let support_host = fill_linear_support_action_select(batch, v_min, v_max, delta_z);
|
||
let q_values_host = vec![0.0_f32; batch * AS_Q_STRIDE];
|
||
|
||
// ISV bus pre-fire: persistence way above threshold, fired clear,
|
||
// warm 0, cooldown 0. hold_floor params identical to the cooldown
|
||
// tests above so any DIR_HOLD outcome is forced by the OR-gate, not
|
||
// by hold_floor (entropy near 0 with one-hot logits → hold_floor ≈ 0).
|
||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc MappedF32Buffer for isv bus");
|
||
let mut isv_init = vec![0.0f32; ISV_LEN];
|
||
isv_init[339] = 0.5; // ISV_EVAL_THOMPSON_TEMP_IDX (floor)
|
||
isv_init[HOLD_FLOOR_ALPHA_INDEX] = 0.5;
|
||
isv_init[HOLD_FLOOR_K_INDEX] = 10.0;
|
||
isv_init[HOLD_FLOOR_EPS0_INDEX] = 1.0;
|
||
isv_init[COOLDOWN_BARS_REMAINING_INDEX] = 0.0;
|
||
isv_init[DD_PERSISTENCE_MAX_INDEX] = 1_000.0;
|
||
isv_init[PLASTICITY_PERSISTENCE_THRESHOLD_INDEX] = 1.0;
|
||
isv_init[PLASTICITY_FIRED_THIS_FOLD_INDEX] = 0.0;
|
||
isv_init[PLASTICITY_WARM_BARS_REMAINING_INDEX] = 0.0;
|
||
isv_buf.write_from_slice(&isv_init);
|
||
|
||
// advantage_head_weights buffer for the trigger's Kaiming-He reset.
|
||
// Uses the same 64-elem skeleton the existing plasticity oracle
|
||
// tests use; n_weights / 10 = 6 → 6 trailing slots get rewritten.
|
||
const N_WEIGHTS: i32 = 64;
|
||
const FAN_IN: i32 = 128; // matches default config adv_h
|
||
const M_WARM: f32 = 5.0; // shrunk from production 200.0 for test runtime
|
||
let weights_buf = unsafe { MappedF32Buffer::new(N_WEIGHTS as usize) }
|
||
.expect("alloc MappedF32Buffer for advantage_head_weights");
|
||
weights_buf.write_from_slice(&vec![1.0f32; N_WEIGHTS as usize]);
|
||
|
||
// Per-step GPU buffers (allocated once, reused across bars).
|
||
let q_values = stream.clone_htod(&q_values_host).expect("upload q_values");
|
||
let b_logits_dir = stream.clone_htod(&b_logits_host).expect("upload logits");
|
||
let per_sample_support = stream
|
||
.clone_htod(&support_host)
|
||
.expect("upload support");
|
||
let mut actions_buf = stream.alloc_zeros::<i32>(batch).expect("alloc actions");
|
||
let mut q_gaps_buf = stream.alloc_zeros::<f32>(batch).expect("alloc q_gaps");
|
||
let mut intent_buf = stream.alloc_zeros::<i32>(batch).expect("alloc intent");
|
||
let mut conv_buf = stream.alloc_zeros::<f32>(batch).expect("alloc conv");
|
||
let mut mag_conv = stream.alloc_zeros::<f32>(batch).expect("alloc mag_conv");
|
||
|
||
let m_warm_i = M_WARM as usize;
|
||
|
||
// ── Bar 0: fire bar → expect DIR_FLAT ──────────────────────────────
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_plasticity_injection(
|
||
&stream,
|
||
&load_sp15_kernel(&stream, SP15_PLASTICITY_INJECTION_CUBIN, "plasticity_injection_kernel"),
|
||
isv_buf.dev_ptr,
|
||
weights_buf.dev_ptr,
|
||
N_WEIGHTS,
|
||
M_WARM,
|
||
FAN_IN,
|
||
/*seed*/ 0xFEED_BEEF_CAFE_BABEu64,
|
||
).expect("launch plasticity (fire bar)");
|
||
let actions = launch_action_select_for_test_with_m_warm(
|
||
&stream, batch, n_atoms,
|
||
&q_values, &b_logits_dir, &per_sample_support,
|
||
&mut actions_buf, &mut q_gaps_buf, &mut intent_buf,
|
||
&mut conv_buf, &mut mag_conv,
|
||
isv_buf.dev_ptr, M_WARM,
|
||
);
|
||
let dir_picked = decode_dir(actions[0]);
|
||
// DIR_FLAT = 3 per state_layout.cuh.
|
||
assert_eq!(
|
||
dir_picked, 3,
|
||
"fire bar must force DIR_FLAT (=3); got dir = {dir_picked} \
|
||
(expected 3, even though Long had +1.0 advantage). \
|
||
ISV[fired]={}, ISV[warm]={}",
|
||
isv_buf.read_all()[PLASTICITY_FIRED_THIS_FOLD_INDEX],
|
||
isv_buf.read_all()[PLASTICITY_WARM_BARS_REMAINING_INDEX],
|
||
);
|
||
|
||
// Sanity-check the ISV state after the fire bar:
|
||
// fired flipped 0 → 1, warm = M_warm − 1 (per the kernel's
|
||
// trigger-then-decrement convention).
|
||
let isv_after_fire = isv_buf.read_all();
|
||
assert!(
|
||
(isv_after_fire[PLASTICITY_FIRED_THIS_FOLD_INDEX] - 1.0).abs() < 1e-9,
|
||
"after fire-bar PLASTICITY_FIRED_THIS_FOLD must be 1.0; got {}",
|
||
isv_after_fire[PLASTICITY_FIRED_THIS_FOLD_INDEX]
|
||
);
|
||
let warm_after_fire = isv_after_fire[PLASTICITY_WARM_BARS_REMAINING_INDEX];
|
||
assert!(
|
||
(warm_after_fire - (M_WARM - 1.0)).abs() < 1e-4,
|
||
"after fire-bar PLASTICITY_WARM_BARS_REMAINING must be M_warm − 1 = {}; got {}",
|
||
M_WARM - 1.0, warm_after_fire
|
||
);
|
||
|
||
// ── Bars 1..M_warm − 1 (warm-up) → expect DIR_HOLD ──────────────────
|
||
// Trigger continues to run on every bar (debounced — fired = 1, no
|
||
// re-fire), per-bar warm decrement runs unconditionally; consumer
|
||
// sees warm decreasing from M_warm − 2 down to 1; OR-gate forces Hold.
|
||
//
|
||
// Trace (M_warm = 5):
|
||
// Bar 0 (fire): trigger sets warm=5 → decrements to 4. action_select
|
||
// reads warm=4 ≥ 3.5 → DIR_FLAT. (asserted above)
|
||
// Bar 1: trigger debounces, warm 4→3. action_select reads warm=3 → Hold.
|
||
// Bar 2: warm 3→2 → Hold.
|
||
// Bar 3: warm 2→1 → Hold.
|
||
// Bar 4: warm 1→0. action_select reads warm=0 — OR-gate inactive →
|
||
// normal Thompson (NOT Hold). This bar is the "post-warm"
|
||
// assertion below, not part of the warm-up loop.
|
||
//
|
||
// So the warm-up loop runs for bars [1, 2, 3] = bars 1..M_warm − 1
|
||
// = bars 1..4 in iter terms (3 iterations). The previous formulation
|
||
// `1..m_warm_i` overran by one bar.
|
||
for bar in 1..(m_warm_i - 1) {
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_plasticity_injection(
|
||
&stream,
|
||
&load_sp15_kernel(&stream, SP15_PLASTICITY_INJECTION_CUBIN, "plasticity_injection_kernel"),
|
||
isv_buf.dev_ptr,
|
||
weights_buf.dev_ptr,
|
||
N_WEIGHTS,
|
||
M_WARM,
|
||
FAN_IN,
|
||
/*seed*/ 0xFEED_BEEF_CAFE_BABEu64 ^ (bar as u64),
|
||
).expect("launch plasticity (warm-up)");
|
||
let actions = launch_action_select_for_test_with_m_warm(
|
||
&stream, batch, n_atoms,
|
||
&q_values, &b_logits_dir, &per_sample_support,
|
||
&mut actions_buf, &mut q_gaps_buf, &mut intent_buf,
|
||
&mut conv_buf, &mut mag_conv,
|
||
isv_buf.dev_ptr, M_WARM,
|
||
);
|
||
let dir_picked = decode_dir(actions[0]);
|
||
assert_eq!(
|
||
dir_picked, AS_DIR_HOLD,
|
||
"warm-up bar {bar} must force DIR_HOLD (={AS_DIR_HOLD}); \
|
||
got dir = {dir_picked}. ISV[warm]={}",
|
||
isv_buf.read_all()[PLASTICITY_WARM_BARS_REMAINING_INDEX]
|
||
);
|
||
}
|
||
|
||
// After (M_warm − 2) warm-up bars the warm counter is now 1 — the
|
||
// very next call's per-bar decrement takes it to 0 (the post-warm
|
||
// boundary). Sanity-check that we're at the boundary so the post-
|
||
// warm assertion below tests the ACTUAL transition, not some later
|
||
// already-decremented state.
|
||
let isv_at_boundary = isv_buf.read_all();
|
||
let warm_at_boundary = isv_at_boundary[PLASTICITY_WARM_BARS_REMAINING_INDEX];
|
||
assert!(
|
||
(warm_at_boundary - 1.0).abs() < 1e-4,
|
||
"after warm-up loop PLASTICITY_WARM_BARS_REMAINING must be 1 (= boundary); got {}",
|
||
warm_at_boundary
|
||
);
|
||
|
||
// ── Bar M_warm − 1 (warm transitions to 0) → normal action selection resumes ─
|
||
// The trigger still runs (debounced — `fired` stays 1 for the
|
||
// remainder of the fold) and decrements warm 1 → 0. action_select
|
||
// reads warm = 0: OR-gate inactive (cooldown also 0), fire-bar
|
||
// predicate false (warm=0 < 3.5). With Long strongly preferred and
|
||
// hold_floor ≈ 0 (one-hot logits → low entropy), Thompson argmax
|
||
// returns DIR_LONG.
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_plasticity_injection(
|
||
&stream,
|
||
&load_sp15_kernel(&stream, SP15_PLASTICITY_INJECTION_CUBIN, "plasticity_injection_kernel"),
|
||
isv_buf.dev_ptr,
|
||
weights_buf.dev_ptr,
|
||
N_WEIGHTS,
|
||
M_WARM,
|
||
FAN_IN,
|
||
/*seed*/ 0xFEED_BEEF_CAFE_BABEu64 ^ (m_warm_i as u64 + 1),
|
||
).expect("launch plasticity (post-warm)");
|
||
let actions = launch_action_select_for_test_with_m_warm(
|
||
&stream, batch, n_atoms,
|
||
&q_values, &b_logits_dir, &per_sample_support,
|
||
&mut actions_buf, &mut q_gaps_buf, &mut intent_buf,
|
||
&mut conv_buf, &mut mag_conv,
|
||
isv_buf.dev_ptr, M_WARM,
|
||
);
|
||
let dir_picked = decode_dir(actions[0]);
|
||
assert_eq!(
|
||
dir_picked, AS_DIR_LONG,
|
||
"post-warm bar must follow normal Thompson argmax → DIR_LONG; \
|
||
got dir = {dir_picked}. ISV[warm]={}, ISV[fired]={}",
|
||
isv_buf.read_all()[PLASTICITY_WARM_BARS_REMAINING_INDEX],
|
||
isv_buf.read_all()[PLASTICITY_FIRED_THIS_FOLD_INDEX]
|
||
);
|
||
|
||
// Post-test ISV state: fired stays 1 (debounced for remainder of
|
||
// fold — fold-reset registry arms re-arm to 0 on next fold
|
||
// boundary), warm at 0 (decrement guard prevents underflow).
|
||
let isv_final = isv_buf.read_all();
|
||
assert!(
|
||
(isv_final[PLASTICITY_FIRED_THIS_FOLD_INDEX] - 1.0).abs() < 1e-9,
|
||
"post-test PLASTICITY_FIRED_THIS_FOLD must remain 1 (debounced); got {}",
|
||
isv_final[PLASTICITY_FIRED_THIS_FOLD_INDEX]
|
||
);
|
||
assert!(
|
||
isv_final[PLASTICITY_WARM_BARS_REMAINING_INDEX].abs() < 1e-4,
|
||
"post-test PLASTICITY_WARM_BARS_REMAINING must be 0; got {}",
|
||
isv_final[PLASTICITY_WARM_BARS_REMAINING_INDEX]
|
||
);
|
||
}
|
||
|
||
// ====================================================================
|
||
// SP15 Wave 2 (2026-05-06) — fused post-SP11 reward-axis composer
|
||
// oracle tests for `compute_sp15_final_reward_kernel`.
|
||
//
|
||
// The fused kernel reads SP11's `out_rewards[N*2*L]` in place,
|
||
// applies α-blend → DD asymmetric reward → quadratic DD penalty →
|
||
// SP12 v3 cap (state_layout.cuh macros), and writes the result back.
|
||
// Tests drive the kernel directly with hand-crafted input buffers and
|
||
// a single-slot N=1, L=1 layout — the on-policy slot lives at
|
||
// `out_rewards[0]` and the CF slot at `out_rewards[N*L]=out_rewards[1]`.
|
||
// ====================================================================
|
||
|
||
/// Helper: build the bundle of input buffers for a single-slot
|
||
/// (N=1, L=1) fused-kernel oracle. Returns a tuple of mapped-pinned
|
||
/// buffers the test can mutate directly. Both on-policy + CF slots
|
||
/// in `out_rewards` start at the on_policy_r/cf_r values so tests
|
||
/// can assert which slot got modified independently.
|
||
fn build_final_reward_oracle_inputs(
|
||
on_policy_r: f32,
|
||
cf_r: f32,
|
||
r_discipline: f32,
|
||
slot_completed: i32,
|
||
) -> (
|
||
MappedF32Buffer, // out_rewards [N*2*L = 2]
|
||
MappedF32Buffer, // r_discipline_out [N*L = 1]
|
||
ml::cuda_pipeline::mapped_pinned::MappedI32Buffer, // slot_completed_normally [N*L = 1]
|
||
MappedF32Buffer, // isv [ISV_LEN]
|
||
MappedF32Buffer, // dd_state_per_env [N * 6 = 6] (Phase 1.3.b-followup)
|
||
) {
|
||
// Safety: tests calling this helper hold a CUDA context via
|
||
// `make_test_stream`.
|
||
let out_rewards = unsafe { MappedF32Buffer::new(2) }
|
||
.expect("alloc out_rewards");
|
||
out_rewards.write_from_slice(&[on_policy_r, cf_r]);
|
||
|
||
let r_disc = unsafe { MappedF32Buffer::new(1) }
|
||
.expect("alloc r_discipline_out");
|
||
r_disc.write_from_slice(&[r_discipline]);
|
||
|
||
let flag = unsafe {
|
||
ml::cuda_pipeline::mapped_pinned::MappedI32Buffer::new(1)
|
||
}.expect("alloc slot_completed_normally");
|
||
flag.write_from_slice(&[slot_completed]);
|
||
|
||
let isv = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc ISV bus");
|
||
let isv_init = vec![0.0f32; ISV_LEN];
|
||
isv.write_from_slice(&isv_init);
|
||
|
||
// SP15 Phase 1.3.b-followup contract: per-env DD tile [N * 6].
|
||
// Single-env (N=1) → 6 slots. Tests that exercise DD-aware
|
||
// shaping write the DD context into this tile (slots 0 = DD_CURRENT
|
||
// and 5 = DD_PCT) AFTER calling this helper, mirroring the
|
||
// pattern used for ISV[DD_CURRENT_INDEX] / ISV[DD_PCT_INDEX]
|
||
// in the pre-followup tests.
|
||
let dd_tile = unsafe { MappedF32Buffer::new(6) }
|
||
.expect("alloc dd_state_per_env tile");
|
||
dd_tile.write_from_slice(&[0.0f32; 6]);
|
||
|
||
(out_rewards, r_disc, flag, isv, dd_tile)
|
||
}
|
||
|
||
/// Test final_reward 1 — α-blend at cold start (α=0.5 sentinel).
|
||
/// On-policy r=1.0, r_discipline=-2.0, all DD slots zero (no penalty,
|
||
/// no boost), so SP12 cap is also no-op. Expected:
|
||
/// r_total = 0.5 × 1.0 + 0.5 × (-2.0) = -0.5
|
||
/// Validates Stage 1 (α-blend) of the fused composer in isolation.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn final_reward_alpha_blend_at_cold_start() {
|
||
let stream = make_test_stream();
|
||
let (out_rewards, r_disc, flag, isv, dd_tile) =
|
||
build_final_reward_oracle_inputs(1.0, 0.0, -2.0, 1);
|
||
// Mirror constructor's α=0.5 cold-start sentinel.
|
||
let mut isv_h = isv.read_all();
|
||
isv_h[ALPHA_SPLIT_INDEX] = 0.5;
|
||
isv.write_from_slice(&isv_h);
|
||
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_final_reward(
|
||
&stream,
|
||
&load_sp15_kernel(&stream, SP15_FINAL_REWARD_CUBIN, "compute_sp15_final_reward_kernel"),
|
||
out_rewards.dev_ptr,
|
||
r_disc.dev_ptr,
|
||
flag.dev_ptr,
|
||
isv.dev_ptr,
|
||
dd_tile.dev_ptr,
|
||
1,
|
||
1,
|
||
).expect("launch compute_sp15_final_reward_kernel");
|
||
stream.synchronize().expect("sync");
|
||
|
||
// r_split = 0.5 × 1.0 + 0.5 × (-2.0) = -0.5; DD stages all no-op
|
||
// (slots zero); SP12 cap fmaxf(-10, fminf(+5, -0.5)) = -0.5.
|
||
let r = out_rewards.read_all()[0];
|
||
assert!(
|
||
(r - (-0.5)).abs() < 1e-5,
|
||
"on-policy r = {r}, expected -0.5"
|
||
);
|
||
}
|
||
|
||
/// Test final_reward 2 — quadratic DD penalty above threshold.
|
||
/// On-policy r=1.0, α=1.0 (so r_discipline drops out), dd=0.10,
|
||
/// dd_threshold=0.05, λ_dd=10. r passes through stage 1 unchanged
|
||
/// (1.0). Stage 2 gain-only boost with λ=0/dd_pct=0 → no-op.
|
||
/// Stage 3: penalty = 10 × (0.10−0.05)² = 0.025; r −= 0.025 = 0.975.
|
||
/// SP12 cap no-op. Validates Stage 3.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn final_reward_dd_penalty_above_threshold() {
|
||
use ml::cuda_pipeline::sp15_isv_slots::{
|
||
DD_THRESHOLD_INDEX, LAMBDA_DD_INDEX,
|
||
};
|
||
|
||
let stream = make_test_stream();
|
||
let (out_rewards, r_disc, flag, isv, dd_tile) =
|
||
build_final_reward_oracle_inputs(1.0, 0.0, 99.0, 1);
|
||
let mut isv_h = isv.read_all();
|
||
isv_h[ALPHA_SPLIT_INDEX] = 1.0; // α=1 → r_discipline ignored
|
||
isv_h[DD_THRESHOLD_INDEX] = 0.05;
|
||
isv_h[LAMBDA_DD_INDEX] = 10.0;
|
||
isv.write_from_slice(&isv_h);
|
||
// Phase 1.3.b-followup: DD_CURRENT is per-env now — write to
|
||
// dd_tile[env*6 + 0] for env=0 (single-env layout).
|
||
let mut tile_h = dd_tile.read_all();
|
||
tile_h[0] = 0.10; // DD_CURRENT
|
||
dd_tile.write_from_slice(&tile_h);
|
||
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_final_reward(
|
||
&stream,
|
||
&load_sp15_kernel(&stream, SP15_FINAL_REWARD_CUBIN, "compute_sp15_final_reward_kernel"),
|
||
out_rewards.dev_ptr,
|
||
r_disc.dev_ptr,
|
||
flag.dev_ptr,
|
||
isv.dev_ptr,
|
||
dd_tile.dev_ptr,
|
||
1,
|
||
1,
|
||
).expect("launch");
|
||
stream.synchronize().expect("sync");
|
||
|
||
// r_split = 1.0 × 1.0 + 0 × 99.0 = 1.0
|
||
// Stage 2 (gain): λ_dd_asymm=0 → boost=1+0=1 → r=1.0
|
||
// Stage 3 (penalty): 1.0 − 10×(0.10-0.05)² = 1.0 − 0.025 = 0.975
|
||
// Stage 4: cap is no-op (0.975 ∈ [-10, 5])
|
||
let r = out_rewards.read_all()[0];
|
||
assert!(
|
||
(r - 0.975).abs() < 1e-5,
|
||
"on-policy r = {r}, expected 0.975 (penalty applied)"
|
||
);
|
||
}
|
||
|
||
/// Test final_reward 3 — DD asymmetric gain boost.
|
||
/// On-policy r=4.0, α=1, dd_pct=0.5, λ_asymm=0.5, λ_dd=0 (no penalty).
|
||
/// Stage 1: r=4. Stage 2: r > 0 → boost = 1 + 0.5×0.5 = 1.25 →
|
||
/// r = 4×1.25 = 5.0. Stage 3: λ_dd=0 → no penalty. Stage 4: cap caps
|
||
/// 5.0 → 5.0 (POS_CAP=5 inclusive). Validates Stage 2.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn final_reward_dd_asymmetric_gain() {
|
||
use ml::cuda_pipeline::sp15_isv_slots::{
|
||
DD_ASYMMETRY_LAMBDA_INDEX, R_GAIN_DD_BOOST_INDEX,
|
||
};
|
||
|
||
let stream = make_test_stream();
|
||
let (out_rewards, r_disc, flag, isv, dd_tile) =
|
||
build_final_reward_oracle_inputs(4.0, 0.0, 0.0, 1);
|
||
let mut isv_h = isv.read_all();
|
||
isv_h[ALPHA_SPLIT_INDEX] = 1.0;
|
||
isv_h[DD_ASYMMETRY_LAMBDA_INDEX] = 0.5;
|
||
isv.write_from_slice(&isv_h);
|
||
// Phase 1.3.b-followup: DD_PCT is per-env now — write to
|
||
// dd_tile[env*6 + 5] for env=0 (single-env layout).
|
||
let mut tile_h = dd_tile.read_all();
|
||
tile_h[5] = 0.5; // DD_PCT
|
||
dd_tile.write_from_slice(&tile_h);
|
||
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_final_reward(
|
||
&stream,
|
||
&load_sp15_kernel(&stream, SP15_FINAL_REWARD_CUBIN, "compute_sp15_final_reward_kernel"),
|
||
out_rewards.dev_ptr,
|
||
r_disc.dev_ptr,
|
||
flag.dev_ptr,
|
||
isv.dev_ptr,
|
||
dd_tile.dev_ptr,
|
||
1,
|
||
1,
|
||
).expect("launch");
|
||
stream.synchronize().expect("sync");
|
||
|
||
let r = out_rewards.read_all()[0];
|
||
assert!(
|
||
(r - 5.0).abs() < 1e-5,
|
||
"on-policy r = {r}, expected 5.0 (gain boost)"
|
||
);
|
||
|
||
// Diagnostic slot R_GAIN_DD_BOOST written by the on-policy
|
||
// representative thread (idx == 0).
|
||
let isv_after = isv.read_all();
|
||
let boost = isv_after[R_GAIN_DD_BOOST_INDEX];
|
||
assert!(
|
||
(boost - 1.25).abs() < 1e-5,
|
||
"R_GAIN_DD_BOOST = {boost}, expected 1.25"
|
||
);
|
||
}
|
||
|
||
/// Test final_reward 4 — loss-side asymmetry: r<0 unchanged through
|
||
/// Stage 2 (preserves loss aversion). On-policy r=-3.0, dd_pct=0.5,
|
||
/// λ_asymm=0.5. Stage 2 NO-OP on r<0; r stays at -3.0. SP12 cap
|
||
/// no-op. Validates Stage 2 asymmetric guard.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn final_reward_dd_asymmetric_loss() {
|
||
use ml::cuda_pipeline::sp15_isv_slots::DD_ASYMMETRY_LAMBDA_INDEX;
|
||
|
||
let stream = make_test_stream();
|
||
let (out_rewards, r_disc, flag, isv, dd_tile) =
|
||
build_final_reward_oracle_inputs(-3.0, 0.0, 0.0, 1);
|
||
let mut isv_h = isv.read_all();
|
||
isv_h[ALPHA_SPLIT_INDEX] = 1.0;
|
||
isv_h[DD_ASYMMETRY_LAMBDA_INDEX] = 0.5;
|
||
isv.write_from_slice(&isv_h);
|
||
// Phase 1.3.b-followup: DD_PCT is per-env now.
|
||
let mut tile_h = dd_tile.read_all();
|
||
tile_h[5] = 0.5; // DD_PCT
|
||
dd_tile.write_from_slice(&tile_h);
|
||
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_final_reward(
|
||
&stream,
|
||
&load_sp15_kernel(&stream, SP15_FINAL_REWARD_CUBIN, "compute_sp15_final_reward_kernel"),
|
||
out_rewards.dev_ptr,
|
||
r_disc.dev_ptr,
|
||
flag.dev_ptr,
|
||
isv.dev_ptr,
|
||
dd_tile.dev_ptr,
|
||
1,
|
||
1,
|
||
).expect("launch");
|
||
stream.synchronize().expect("sync");
|
||
|
||
let r = out_rewards.read_all()[0];
|
||
assert!(
|
||
(r - (-3.0)).abs() < 1e-5,
|
||
"on-policy r = {r}, expected -3.0 (loss unchanged)"
|
||
);
|
||
}
|
||
|
||
/// Test final_reward 5 — SP12 v3 asymmetric cap clamping.
|
||
/// On-policy r = +20 (above POS_CAP=5) → clamped to +5.
|
||
/// CF r = -20 (below NEG_CAP=-10) → clamped to -10.
|
||
/// Both with α=1, all DD slots zero. Validates Stage 4 macros.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn final_reward_sp12_cap_clamps_both_directions() {
|
||
let stream = make_test_stream();
|
||
let (out_rewards, r_disc, flag, isv, dd_tile) =
|
||
build_final_reward_oracle_inputs(20.0, -20.0, 0.0, 1);
|
||
let mut isv_h = isv.read_all();
|
||
isv_h[ALPHA_SPLIT_INDEX] = 1.0;
|
||
isv.write_from_slice(&isv_h);
|
||
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_final_reward(
|
||
&stream,
|
||
&load_sp15_kernel(&stream, SP15_FINAL_REWARD_CUBIN, "compute_sp15_final_reward_kernel"),
|
||
out_rewards.dev_ptr,
|
||
r_disc.dev_ptr,
|
||
flag.dev_ptr,
|
||
isv.dev_ptr,
|
||
dd_tile.dev_ptr,
|
||
1,
|
||
1,
|
||
).expect("launch");
|
||
stream.synchronize().expect("sync");
|
||
|
||
let r = out_rewards.read_all();
|
||
// On-policy capped at POS_CAP=5.0; CF capped at NEG_CAP=-10.0.
|
||
// (REWARD_NEG_CAP / REWARD_POS_CAP from state_layout.cuh.)
|
||
assert!(
|
||
(r[0] - 5.0).abs() < 1e-5,
|
||
"on-policy r = {} (expected +5 from POS_CAP)", r[0]
|
||
);
|
||
assert!(
|
||
(r[1] - (-10.0)).abs() < 1e-5,
|
||
"CF r = {} (expected -10 from NEG_CAP)", r[1]
|
||
);
|
||
}
|
||
|
||
/// Test final_reward 6 (Q3 NEW) — early-return slots preserved.
|
||
/// On-policy slot 0 has slot_completed=0 + r=-10.0 (blown-account
|
||
/// sentinel). On-policy slot 1 has slot_completed=1 + r=2.5 (real
|
||
/// reward). After fused kernel:
|
||
/// slot 0: UNCHANGED at -10.0 (early-return passthrough)
|
||
/// slot 1: SP15-shaped per the formula
|
||
/// Validates Q3 resolution: SP11's terminal-episode encoding passes
|
||
/// through unmodified. Uses N=2, L=1 layout for two on-policy slots
|
||
/// (CF slots 2/3 also processed but irrelevant here).
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn final_reward_skips_early_return_slots() {
|
||
let stream = make_test_stream();
|
||
|
||
// N=2, L=1 → out_rewards[0..4) = [op_0, op_1, cf_0, cf_1]
|
||
let out_rewards = unsafe { MappedF32Buffer::new(4) }
|
||
.expect("alloc out_rewards");
|
||
out_rewards.write_from_slice(&[-10.0, 2.5, 0.0, 0.0]);
|
||
|
||
// r_discipline_out[N*L=2]; both = 0 so α-blend gives r unchanged
|
||
// when α=1.
|
||
let r_disc = unsafe { MappedF32Buffer::new(2) }
|
||
.expect("alloc r_disc");
|
||
r_disc.write_from_slice(&[0.0, 0.0]);
|
||
|
||
// slot_completed_normally[2]: slot 0 = 0 (early-return), slot 1 = 1.
|
||
let flag = unsafe {
|
||
ml::cuda_pipeline::mapped_pinned::MappedI32Buffer::new(2)
|
||
}.expect("alloc flag");
|
||
flag.write_from_slice(&[0, 1]);
|
||
|
||
let isv = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc ISV");
|
||
let mut isv_h = vec![0.0f32; ISV_LEN];
|
||
// α=1 + zeroed DD slots → fused kernel composition is identity
|
||
// for slot 1's r. The interesting check is that slot 0 (r=-10
|
||
// sentinel) does NOT get touched.
|
||
isv_h[ALPHA_SPLIT_INDEX] = 1.0;
|
||
isv.write_from_slice(&isv_h);
|
||
|
||
// Phase 1.3.b-followup: per-env DD tile [N=2 * 6] = 12 slots,
|
||
// all zero — DD has no effect on the reward (α=1 + zero DD ⇒
|
||
// identity composition).
|
||
let dd_tile = unsafe { MappedF32Buffer::new(2 * 6) }
|
||
.expect("alloc dd_state_per_env tile");
|
||
dd_tile.write_from_slice(&vec![0.0f32; 2 * 6]);
|
||
|
||
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_final_reward(
|
||
&stream,
|
||
&load_sp15_kernel(&stream, SP15_FINAL_REWARD_CUBIN, "compute_sp15_final_reward_kernel"),
|
||
out_rewards.dev_ptr,
|
||
r_disc.dev_ptr,
|
||
flag.dev_ptr,
|
||
isv.dev_ptr,
|
||
dd_tile.dev_ptr,
|
||
2,
|
||
1,
|
||
).expect("launch");
|
||
stream.synchronize().expect("sync");
|
||
|
||
let r = out_rewards.read_all();
|
||
assert!(
|
||
(r[0] - (-10.0)).abs() < 1e-9,
|
||
"slot 0 (early-return) = {} != -10.0 — fused kernel modified \
|
||
a slot it should have skipped (Q3 violated)", r[0]
|
||
);
|
||
assert!(
|
||
(r[1] - 2.5).abs() < 1e-5,
|
||
"slot 1 (normal completion) = {} != 2.5 — fused kernel \
|
||
composition altered an unmodified path (alpha=1, all DD zero)", r[1]
|
||
);
|
||
// CF slots also flagged via `idx % (N*L)`: slot 2 uses flag[0]=0
|
||
// (skipped), slot 3 uses flag[1]=1 (processed but starts at 0
|
||
// and stays at 0 after the identity composition).
|
||
assert!(
|
||
r[2].abs() < 1e-9,
|
||
"CF slot 2 = {} != 0.0 — should have been skipped via flag[0]=0", r[2]
|
||
);
|
||
assert!(
|
||
r[3].abs() < 1e-9,
|
||
"CF slot 3 = {} != 0.0 — alpha-blend identity should keep at 0", r[3]
|
||
);
|
||
}
|
||
}
|
||
|
||
/// SP15 Wave 4.1a — shared test helpers used by the bn_tanh_concat_dd
|
||
/// oracle test (above) and Wave 4.1c's behavioral KL test (deferred).
|
||
///
|
||
/// Per the spec, three helpers are needed: `kl_divergence` (CPU math),
|
||
/// `minimal_trainer_for_tests` (thin wrapper around `DQNTrainer::new`),
|
||
/// and `forward_trunk_for_test` (runs the trunk forward path against
|
||
/// a `[B, s1_input_dim+1]` post-bn_concat_dd buffer and returns
|
||
/// `[B, shared_h1]` GRN output).
|
||
///
|
||
/// `forward_trunk_for_test` is NOT implemented in Wave 4.1a — its
|
||
/// exact contract is owned by Wave 4.1c (which lands the consumer
|
||
/// migration that bumps `s1_input_dim` 102→103 and reshapes GRN
|
||
/// w_s1). Implementing it here would either require a new
|
||
/// non-trivial public API on `DQNTrainer` (test-only forward
|
||
/// surface) or a stub that returns dummy data — both violate
|
||
/// `feedback_no_stubs`. The two helpers that ARE implementable
|
||
/// today (`kl_divergence`, `minimal_trainer_for_tests`) are
|
||
/// real-and-useful; Wave 4.1c will add `forward_trunk_for_test`
|
||
/// alongside the public `DQNTrainer` surface change it requires.
|
||
#[cfg(test)]
|
||
mod sp15_wave_4_1a_test_helpers {
|
||
use ml::trainers::dqn::{DQNHyperparameters, DQNTrainer};
|
||
|
||
/// KL divergence helper: KL(p || q) = sum_i p[i] * ln(p[i] / q[i]).
|
||
///
|
||
/// Used by Wave 4.1c's behavioral test to compare
|
||
/// `softmax(forward_trunk(102-input))` vs
|
||
/// `softmax(forward_trunk(103-input))` action distributions; the
|
||
/// dd_pct column should produce a non-trivial KL divergence
|
||
/// (i.e. the policy actually conditions on drawdown context).
|
||
///
|
||
/// Contract:
|
||
/// - `p` and `q` MUST be the same length (panics otherwise).
|
||
/// - Treats both inputs as proper probability distributions:
|
||
/// entries with `p[i] == 0.0` contribute 0 to the sum (limit
|
||
/// `x ln x → 0` as `x → 0`); entries with `q[i] <= 0.0` raise
|
||
/// an `f32::INFINITY` term — caller is responsible for either
|
||
/// adding a small epsilon or asserting q is strictly positive.
|
||
/// - Result is in nats (natural log). Multiply by `1.0 /
|
||
/// std::f32::consts::LN_2` to convert to bits if needed.
|
||
///
|
||
/// This is a generic helper — no SP15-specific behaviour. Standard
|
||
/// information-theoretic KL formula.
|
||
#[allow(dead_code)] // Consumer (Wave 4.1c behavioral KL test) lands later.
|
||
pub(super) fn kl_divergence(p: &[f32], q: &[f32]) -> f32 {
|
||
assert_eq!(
|
||
p.len(),
|
||
q.len(),
|
||
"kl_divergence: distribution length mismatch — p.len()={} q.len()={}",
|
||
p.len(),
|
||
q.len()
|
||
);
|
||
let mut acc = 0.0_f32;
|
||
for (pi, qi) in p.iter().zip(q.iter()) {
|
||
// Convention: 0 * ln(0) = 0; preserves continuity at p=0.
|
||
if *pi <= 0.0 {
|
||
continue;
|
||
}
|
||
// q <= 0 with p > 0 → KL is infinite. Surface as +inf rather
|
||
// than NaN so caller assertions like `kl > 0.05` see the
|
||
// diagnostic value.
|
||
if *qi <= 0.0 {
|
||
return f32::INFINITY;
|
||
}
|
||
acc += *pi * ((*pi / *qi).ln());
|
||
}
|
||
acc
|
||
}
|
||
|
||
/// Construct a minimal `DQNTrainer` for tests that don't need the
|
||
/// full data pipeline. Mirrors the existing
|
||
/// `set_test_data_from_slices_fires_observer_and_stashes` pattern
|
||
/// (see `test_slice_consumption` module below): synchronous
|
||
/// constructor, no GPU forward passes required. Returns
|
||
/// `Some(trainer)` on success or `None` if the underlying
|
||
/// `DQNTrainer::new` fails (e.g. CPU-only build configuration where
|
||
/// trainer construction can't proceed) — caller is expected to
|
||
/// `eprintln!` a skip notice and return early in that case (the
|
||
/// L40S smoke is the canonical verifier when trainer construction
|
||
/// is unavailable).
|
||
#[allow(dead_code)] // Consumer (Wave 4.1c behavioral KL test) lands later.
|
||
pub(super) fn minimal_trainer_for_tests() -> Option<DQNTrainer> {
|
||
let hyperparams = DQNHyperparameters::conservative();
|
||
match DQNTrainer::new(hyperparams) {
|
||
Ok(t) => Some(t),
|
||
Err(e) => {
|
||
eprintln!(
|
||
"minimal_trainer_for_tests: trainer construction failed ({}). \
|
||
Caller should skip and rely on L40S smoke verification.",
|
||
e
|
||
);
|
||
None
|
||
}
|
||
}
|
||
}
|
||
|
||
// Compile-time smoke for the `kl_divergence` helper. Lives in the
|
||
// helper module (not as a #[test]) so it doesn't add to the test
|
||
// suite; the runtime invariants (zero KL on identical
|
||
// distributions, positive KL on differing ones) are exercised by
|
||
// the dedicated unit tests below.
|
||
|
||
#[test]
|
||
fn kl_divergence_zero_on_identical_distributions() {
|
||
let p = [0.25_f32, 0.25, 0.25, 0.25];
|
||
let q = [0.25_f32, 0.25, 0.25, 0.25];
|
||
let kl = kl_divergence(&p, &q);
|
||
assert!(
|
||
kl.abs() < 1e-6,
|
||
"KL(p || p) should be 0 for identical distributions, got {}",
|
||
kl
|
||
);
|
||
}
|
||
|
||
#[test]
|
||
fn kl_divergence_positive_on_differing_distributions() {
|
||
let p = [0.7_f32, 0.1, 0.1, 0.1];
|
||
let q = [0.25_f32, 0.25, 0.25, 0.25];
|
||
let kl = kl_divergence(&p, &q);
|
||
assert!(
|
||
kl > 0.0,
|
||
"KL(p || q) should be strictly positive when p ≠ q, got {}",
|
||
kl
|
||
);
|
||
// Hand-checked: 0.7 ln(0.7/0.25) + 3*(0.1 ln(0.1/0.25))
|
||
// = 0.7 * 1.0296 + 3 * 0.1 * (-0.9163)
|
||
// ≈ 0.7207 - 0.2749 ≈ 0.4458 nats.
|
||
assert!(
|
||
(kl - 0.4458_f32).abs() < 1e-3,
|
||
"KL ≈ 0.4458 nats expected, got {}",
|
||
kl
|
||
);
|
||
}
|
||
|
||
#[test]
|
||
fn kl_divergence_handles_zero_probability_in_p() {
|
||
// Convention: p[i] = 0 contributes nothing (limit x ln x → 0).
|
||
let p = [0.5_f32, 0.5, 0.0, 0.0];
|
||
let q = [0.4_f32, 0.4, 0.1, 0.1];
|
||
let kl = kl_divergence(&p, &q);
|
||
// Manual check: 0.5 ln(0.5/0.4) + 0.5 ln(0.5/0.4) + 0 + 0
|
||
// = 1.0 * ln(1.25) ≈ 0.2231 nats.
|
||
assert!(
|
||
(kl - 0.2231_f32).abs() < 1e-3,
|
||
"KL with zero p entries ≈ 0.2231 nats expected, got {}",
|
||
kl
|
||
);
|
||
}
|
||
|
||
#[test]
|
||
fn minimal_trainer_for_tests_constructs_or_skips() {
|
||
// Either the trainer constructs (Some) or we get a clean None
|
||
// for the skip path — never panics. Wave 4.1c will use this
|
||
// helper and treat None the same way `test_slice_consumption`
|
||
// does (eprintln + early return).
|
||
let _trainer = minimal_trainer_for_tests();
|
||
}
|
||
}
|
||
|
||
/// SP15 Wave 4.1c — behavioral KL test that closes out the Wave 4.1
|
||
/// (Phase 1.5.b) state_dim 102→103 + dd_pct trunk integration.
|
||
///
|
||
/// Wave 4.1a (`a8da1cb9c`) landed `bn_tanh_concat_dd_kernel`. Wave 4.1b
|
||
/// (`eb9515e41`) wired `s1_input_dim` 102→103 through the trunk + 4
|
||
/// forward + 3 backward call sites. Wave 4.1c proves the wiring
|
||
/// actually shifts the policy distribution: a synthetic GPU-side
|
||
/// projection of the post-bn_concat_dd buffer with random Xavier-init
|
||
/// weights produces measurably different action distributions when
|
||
/// ISV[DD_PCT_INDEX=406] differs (0.0 at-ATH vs 0.10 in-DD).
|
||
///
|
||
/// Why a synthetic projection (vs the real GRN trunk):
|
||
/// `forward_trunk_for_test`'s ideal contract — a public API that runs
|
||
/// the real GRN encoder + advantage heads against caller-owned
|
||
/// post-bn_concat input — would require either (a) a public surface
|
||
/// change on `DQNTrainer` exposing internal cuBLAS handles + GRN scratch
|
||
/// buffers + weight pointers (the trainer's `fused_ctx` is `pub(crate)`
|
||
/// and only initialised inside the training loop, not by
|
||
/// `DQNTrainer::new`), or (b) duplicating the trunk's cuBLAS setup in a
|
||
/// test (≥200 lines of buffer + handle plumbing). Both options are
|
||
/// architecturally heavier than the test's purpose justifies. Per the
|
||
/// spec dispatch ("the test's purpose is 'non-zero KL proves the wire
|
||
/// is connected' not 'verifies trained behavior'"), we instead compose
|
||
/// `launch_sp15_bn_concat_dd` (the production kernel that injects
|
||
/// dd_pct into the trunk input) with a single-layer cuBLAS sgemm using
|
||
/// random Xavier-init weights `[proj_h, 103]`. The matmul exercises the
|
||
/// new column-103 weights on the dd_pct value — exactly the path the
|
||
/// real GRN's `Linear_a` first GEMM takes for `w_a_h_s1[:, 102]` (the
|
||
/// dd_pct column added by Wave 4.1b's reshape). Two ISV values produce
|
||
/// two `[B, proj_h]` linear outputs; softmax + KL on the host (post-
|
||
/// readback policy extraction, not trunk replication) confirms the
|
||
/// wiring shifts the proxy policy.
|
||
///
|
||
/// What this test buys: layout fingerprint + behavior coupling. If
|
||
/// Wave 4.1b's column-102 Xavier init gets reverted to zero, or if a
|
||
/// future migration silently truncates the post-concat buffer back to
|
||
/// 102 columns, KL collapses to 0 and this test fires.
|
||
///
|
||
/// What this test does NOT verify: the full GRN composition (ELU /
|
||
/// GLU / LayerNorm / residual) propagating dd_pct through h_s2 + the
|
||
/// branch advantage heads. That end-to-end behavior is exercised by
|
||
/// the L40S smoke + production training runs.
|
||
#[cfg(test)]
|
||
mod sp15_wave_4_1c_behavioral {
|
||
use std::sync::Arc;
|
||
|
||
use cudarc::cublas::sys as cublas_sys;
|
||
use cudarc::driver::{CudaContext, CudaStream};
|
||
use cudarc::driver::CudaFunction;
|
||
use ml::cuda_pipeline::gpu_dqn_trainer::{launch_sp15_bn_concat_dd, DQN_UTILITY_CUBIN};
|
||
use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer;
|
||
use ml::cuda_pipeline::shared_cublas_handle::PerStreamCublasHandles;
|
||
use ml::cuda_pipeline::sp15_isv_slots::{DD_PCT_INDEX, SP15_SLOT_END};
|
||
|
||
/// Match the kernel-level oracle test's ISV bus length.
|
||
const ISV_LEN: usize = SP15_SLOT_END;
|
||
|
||
fn make_test_stream() -> Arc<CudaStream> {
|
||
let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
|
||
ctx.default_stream()
|
||
}
|
||
|
||
/// CPU softmax — pure post-output policy extraction, NOT trunk
|
||
/// replication. Reads back a `[B, proj_h]` linear output from the
|
||
/// GPU and converts each row to a probability distribution.
|
||
/// Numerically stable: subtracts row max before exp.
|
||
fn row_softmax(logits: &[f32], batch: usize, dim: usize) -> Vec<f32> {
|
||
let mut out = vec![0.0_f32; logits.len()];
|
||
for b in 0..batch {
|
||
let row = &logits[b * dim..(b + 1) * dim];
|
||
let row_max = row.iter().copied().fold(f32::NEG_INFINITY, f32::max);
|
||
let mut sum = 0.0_f32;
|
||
for (i, &v) in row.iter().enumerate() {
|
||
let e = (v - row_max).exp();
|
||
out[b * dim + i] = e;
|
||
sum += e;
|
||
}
|
||
for i in 0..dim {
|
||
out[b * dim + i] /= sum;
|
||
}
|
||
}
|
||
out
|
||
}
|
||
|
||
/// One synthetic forward pass: inject `dd_pct` into ISV[DD_PCT_INDEX],
|
||
/// launch `bn_tanh_concat_dd_kernel` to produce `[B, concat_dd_dim]`
|
||
/// post-concat output, then run a single GPU SGEMM with random
|
||
/// Xavier-init weights `w[proj_h, concat_dd_dim]` against the
|
||
/// post-concat buffer to produce `[B, proj_h]` synthetic logits.
|
||
/// Returns the host-readable logits.
|
||
///
|
||
/// The SGEMM is computed entirely on GPU via cuBLAS classic
|
||
/// `cublasSgemm_v2` (the same primitive `gpu_tlob.rs` uses). No
|
||
/// CPU-side trunk replication: this is a single matmul, NOT the GRN
|
||
/// composition (Linear_a + ELU + Linear_b + GLU + Linear_residual +
|
||
/// LayerNorm) — the synthetic projection's purpose is to convert a
|
||
/// 103-dim post-concat buffer into a small action-distribution
|
||
/// proxy for KL comparison.
|
||
///
|
||
/// Layout: same column-major cuBLAS convention `gpu_tlob.rs:472-490`
|
||
/// uses: `op(A) = W^T [concat_dd_dim, proj_h]`, `op(B) = bn_concat`
|
||
/// stored row-major as `[B, concat_dd_dim]` (= column-major
|
||
/// `[concat_dd_dim, B]`), output `C = [proj_h, B]` column-major
|
||
/// (= row-major `[B, proj_h]`). m=proj_h, n=batch, k=concat_dd_dim.
|
||
#[allow(clippy::too_many_arguments)]
|
||
fn synthetic_forward(
|
||
stream: &Arc<CudaStream>,
|
||
cublas_handle: cublas_sys::cublasHandle_t,
|
||
bn_concat_kernel: &CudaFunction,
|
||
bn_hidden_dev: u64,
|
||
states_dev: u64,
|
||
isv_buf: &MappedF32Buffer,
|
||
weights_dev: u64,
|
||
dd_pct: f32,
|
||
batch_size: usize,
|
||
bn_dim: usize,
|
||
market_dim: usize,
|
||
state_dim_padded: usize,
|
||
concat_dd_dim: usize,
|
||
proj_h: usize,
|
||
) -> Vec<f32> {
|
||
// Inject dd_pct via mapped-pinned ISV bus. Mapped pages are
|
||
// device-visible after a stream-sync barrier — see
|
||
// `feedback_no_htod_htoh_only_mapped_pinned`.
|
||
isv_buf.write_from_slice(&{
|
||
let mut isv = vec![0.0_f32; ISV_LEN];
|
||
isv[DD_PCT_INDEX] = dd_pct;
|
||
isv
|
||
});
|
||
|
||
// Allocate post-concat output buffer for THIS forward pass.
|
||
// Each call gets its own buffer so the two passes' outputs are
|
||
// independently readable post-sync.
|
||
let concat_buf = unsafe { MappedF32Buffer::new(batch_size * concat_dd_dim) }
|
||
.expect("alloc post-concat buffer");
|
||
|
||
// Producer: bn_tanh_concat_dd writes [B, concat_dd_dim] with
|
||
// dd_pct broadcast at column (concat_dd_dim - 1). Reads
|
||
// ISV[DD_PCT_INDEX]; the mapped-pinned write above has
|
||
// device-side visibility before this launch by virtue of the
|
||
// single-stream ordering (the launch is enqueued AFTER the
|
||
// write returns).
|
||
launch_sp15_bn_concat_dd(
|
||
stream,
|
||
bn_concat_kernel,
|
||
bn_hidden_dev,
|
||
states_dev,
|
||
isv_buf.dev_ptr,
|
||
concat_buf.dev_ptr,
|
||
batch_size as i32,
|
||
bn_dim as i32,
|
||
market_dim as i32,
|
||
state_dim_padded as i32,
|
||
concat_dd_dim as i32,
|
||
)
|
||
.expect("launch bn_tanh_concat_dd_kernel");
|
||
|
||
// Consumer: synthetic SGEMM `logits = bn_concat @ W^T`.
|
||
//
|
||
// Row-major view: bn_concat is `[B, concat_dd_dim]`, weights are
|
||
// `[proj_h, concat_dd_dim]`, output is `[B, proj_h]`.
|
||
// Equivalent column-major call (cuBLAS native):
|
||
// A = W with logical shape `[proj_h, concat_dd_dim]`,
|
||
// column-major lda = concat_dd_dim → use op(A) = A^T
|
||
// (CUBLAS_OP_T) so the kernel reads W as
|
||
// `[concat_dd_dim, proj_h]`.
|
||
// B = bn_concat with logical shape `[B, concat_dd_dim]`,
|
||
// column-major view = `[concat_dd_dim, B]`, ldb =
|
||
// concat_dd_dim, op(B) = N (CUBLAS_OP_N).
|
||
// C = logits, column-major `[proj_h, B]` = row-major
|
||
// `[B, proj_h]`, ldc = proj_h.
|
||
//
|
||
// Same layout `gpu_tlob.rs:472-490` uses for the QKV
|
||
// projection — there `W^T @ X` produces `[OUT, B]` output from
|
||
// `[OUT, IN]` weights and `[IN, B]` input.
|
||
let m = proj_h as i32;
|
||
let n = batch_size as i32;
|
||
let k = concat_dd_dim as i32;
|
||
let alpha = 1.0_f32;
|
||
let beta = 0.0_f32;
|
||
let logits_buf = unsafe { MappedF32Buffer::new(batch_size * proj_h) }
|
||
.expect("alloc synthetic logits buffer");
|
||
// Safety: cublas_handle is a valid handle from the test's
|
||
// `PerStreamCublasHandles`; pointers are device pointers from
|
||
// mapped-pinned buffers (CUDA-side visible). m/n/k are
|
||
// non-negative i32 within range. The call has the same shape
|
||
// signature `gpu_tlob.rs::sgemm_v2` uses.
|
||
unsafe {
|
||
let status = cublas_sys::cublasSgemm_v2(
|
||
cublas_handle,
|
||
cublas_sys::cublasOperation_t::CUBLAS_OP_T,
|
||
cublas_sys::cublasOperation_t::CUBLAS_OP_N,
|
||
m,
|
||
n,
|
||
k,
|
||
&alpha as *const f32,
|
||
weights_dev as *const f32,
|
||
k,
|
||
concat_buf.dev_ptr as *const f32,
|
||
k,
|
||
&beta as *const f32,
|
||
logits_buf.dev_ptr as *mut f32,
|
||
m,
|
||
);
|
||
assert_eq!(
|
||
status,
|
||
cublas_sys::cublasStatus_t::CUBLAS_STATUS_SUCCESS,
|
||
"cublasSgemm_v2 (synthetic_forward dd_pct={}): {:?}",
|
||
dd_pct,
|
||
status
|
||
);
|
||
}
|
||
|
||
stream.synchronize().expect("synchronize after synthetic forward");
|
||
|
||
// logits_buf is column-major `[proj_h, B]`. Read back as flat
|
||
// and transpose to row-major `[B, proj_h]` for downstream
|
||
// softmax — this matches the ergonomics softmax expects.
|
||
let raw = logits_buf.read_all();
|
||
let mut row_major = vec![0.0_f32; batch_size * proj_h];
|
||
for b in 0..batch_size {
|
||
for h in 0..proj_h {
|
||
row_major[b * proj_h + h] = raw[h * batch_size + b];
|
||
}
|
||
}
|
||
row_major
|
||
}
|
||
|
||
/// Wave 4.1c — proves dd_pct trunk integration shifts policy
|
||
/// distribution under drawdown vs at-ATH conditions.
|
||
///
|
||
/// Construction:
|
||
/// - Random Xavier-uniform weights `W[proj_h=4, 103]` (production
|
||
/// uses `[shared_h1=256, 103]`; we shrink to 4 to match a
|
||
/// 4-action proxy for direct softmax + KL).
|
||
/// - Deterministic synthetic `bn_hidden [B=4, 16]` and `states
|
||
/// [B=4, 128]` inputs (same as the Wave 4.1a oracle test's
|
||
/// fixture).
|
||
/// - Two passes through `launch_sp15_bn_concat_dd` + cuBLAS sgemm
|
||
/// differing ONLY in `ISV[DD_PCT_INDEX]` (0.0 vs 0.10).
|
||
///
|
||
/// Assertion: `KL(P_DD || P_ATH) > epsilon` averaged over the
|
||
/// batch, where `epsilon = 1e-6` (large enough to reject pure
|
||
/// numerical noise; small enough to pass on random-init weights
|
||
/// where the dd_pct contribution is `0.10 × W[:, 102]` ≈ small but
|
||
/// non-zero).
|
||
///
|
||
/// If KL ≈ 0 the test fires — meaning either Wave 4.1b's column-102
|
||
/// Xavier init silently zeroed, or the post-concat buffer's last
|
||
/// column isn't the dd_pct value, or the SGEMM K-dimension hasn't
|
||
/// caught up to 103. Each of those would be a real wiring bug.
|
||
#[test]
|
||
#[ignore = "requires GPU"]
|
||
fn dd_pct_trunk_input_shifts_policy_distribution() {
|
||
let stream = make_test_stream();
|
||
|
||
// Production-shape constants (same as Wave 4.1a oracle test).
|
||
let batch_size: usize = 4;
|
||
const BN_DIM: usize = 16;
|
||
const MARKET_DIM: usize = 42;
|
||
const STATE_DIM_PADDED: usize = 128;
|
||
let portfolio_dim: usize = STATE_DIM_PADDED - MARKET_DIM;
|
||
let concat_dd_dim: usize = BN_DIM + portfolio_dim + 1; // = 103
|
||
const PROJ_H: usize = 4; // 4-action proxy
|
||
|
||
// ── Deterministic input fixture (same shape as Wave 4.1a) ──
|
||
let mut bn_hidden_init: Vec<f32> = vec![0.0; batch_size * BN_DIM];
|
||
for b in 0..batch_size {
|
||
for j in 0..BN_DIM {
|
||
bn_hidden_init[b * BN_DIM + j] = 0.1_f32 * ((b + 1) as f32) * ((j + 1) as f32);
|
||
}
|
||
}
|
||
let mut states: Vec<f32> = vec![0.0; batch_size * STATE_DIM_PADDED];
|
||
for b in 0..batch_size {
|
||
for i in 0..STATE_DIM_PADDED {
|
||
states[b * STATE_DIM_PADDED + i] = (b * 1000 + i) as f32 * 1e-3_f32;
|
||
}
|
||
}
|
||
|
||
// Safety: CUDA context is active via `make_test_stream` above.
|
||
let bn_hidden_buf = unsafe { MappedF32Buffer::new(bn_hidden_init.len()) }
|
||
.expect("alloc MappedF32Buffer for bn_hidden");
|
||
bn_hidden_buf.write_from_slice(&bn_hidden_init);
|
||
let states_buf = unsafe { MappedF32Buffer::new(states.len()) }
|
||
.expect("alloc MappedF32Buffer for states");
|
||
states_buf.write_from_slice(&states);
|
||
|
||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
|
||
.expect("alloc MappedF32Buffer for ISV");
|
||
|
||
// ── Random Xavier-uniform weights `W[proj_h, concat_dd_dim]` ──
|
||
// Xavier-uniform fan-in/out for a layer with input dim
|
||
// `concat_dd_dim` and output dim `proj_h`: bound =
|
||
// sqrt(6 / (fan_in + fan_out)).
|
||
// Same formula `gpu_dqn_trainer.rs::xavier_init_params_buf`
|
||
// applies to GRN's `w_a_h_s1`. Deterministic LCG (seed = 42)
|
||
// so the test is byte-reproducible across runs.
|
||
let fan_in = concat_dd_dim as f32;
|
||
let fan_out = PROJ_H as f32;
|
||
let bound = (6.0_f32 / (fan_in + fan_out)).sqrt();
|
||
let mut weights = vec![0.0_f32; PROJ_H * concat_dd_dim];
|
||
let mut seed: u64 = 42;
|
||
for w in weights.iter_mut() {
|
||
// Numerical Recipes LCG — same constants used in many GPU
|
||
// pseudo-RNG seed paths; deterministic and uniform enough
|
||
// for test fixtures (cryptographic quality not required).
|
||
seed = seed.wrapping_mul(6364136223846793005).wrapping_add(1442695040888963407);
|
||
let u = ((seed >> 11) as u32 as f32) / (u32::MAX as f32);
|
||
*w = (u * 2.0 - 1.0) * bound;
|
||
}
|
||
let weights_buf = unsafe { MappedF32Buffer::new(weights.len()) }
|
||
.expect("alloc MappedF32Buffer for weights");
|
||
weights_buf.write_from_slice(&weights);
|
||
|
||
// ── cuBLAS classic handle for the synthetic SGEMM ──
|
||
let cublas_handles =
|
||
PerStreamCublasHandles::new(&stream).expect("PerStreamCublasHandles::new");
|
||
let cublas_handle = cublas_handles
|
||
.classic_handle_for(&stream)
|
||
.expect("classic_handle_for default stream");
|
||
|
||
// ── Pre-load the bn_tanh_concat_dd_kernel handle (graph-safe pattern;
|
||
// SP15 Wave 4.1b OOB fix 2026-05-07: launcher consumes &CudaFunction). ──
|
||
let util_module = stream.context().load_cubin(DQN_UTILITY_CUBIN.to_vec())
|
||
.expect("load dqn_utility_kernels cubin");
|
||
let bn_concat_kernel = util_module.load_function("bn_tanh_concat_dd_kernel")
|
||
.expect("load bn_tanh_concat_dd_kernel symbol");
|
||
|
||
// ── At-ATH pass (dd_pct = 0.0) ──
|
||
let logits_ath = synthetic_forward(
|
||
&stream,
|
||
cublas_handle,
|
||
&bn_concat_kernel,
|
||
bn_hidden_buf.dev_ptr,
|
||
states_buf.dev_ptr,
|
||
&isv_buf,
|
||
weights_buf.dev_ptr,
|
||
0.0_f32,
|
||
batch_size,
|
||
BN_DIM,
|
||
MARKET_DIM,
|
||
STATE_DIM_PADDED,
|
||
concat_dd_dim,
|
||
PROJ_H,
|
||
);
|
||
let p_ath = row_softmax(&logits_ath, batch_size, PROJ_H);
|
||
|
||
// ── 10% DD pass (dd_pct = 0.10) ──
|
||
let logits_dd = synthetic_forward(
|
||
&stream,
|
||
cublas_handle,
|
||
&bn_concat_kernel,
|
||
bn_hidden_buf.dev_ptr,
|
||
states_buf.dev_ptr,
|
||
&isv_buf,
|
||
weights_buf.dev_ptr,
|
||
0.10_f32,
|
||
batch_size,
|
||
BN_DIM,
|
||
MARKET_DIM,
|
||
STATE_DIM_PADDED,
|
||
concat_dd_dim,
|
||
PROJ_H,
|
||
);
|
||
let p_dd = row_softmax(&logits_dd, batch_size, PROJ_H);
|
||
|
||
// ── Mean KL(P_DD || P_ATH) over the batch ──
|
||
// Per-row KL using the helper from Wave 4.1a's seeded module.
|
||
let mut total_kl = 0.0_f32;
|
||
let mut max_kl = 0.0_f32;
|
||
for b in 0..batch_size {
|
||
let p = &p_dd[b * PROJ_H..(b + 1) * PROJ_H];
|
||
let q = &p_ath[b * PROJ_H..(b + 1) * PROJ_H];
|
||
let kl = super::sp15_wave_4_1a_test_helpers::kl_divergence(p, q);
|
||
total_kl += kl;
|
||
if kl > max_kl {
|
||
max_kl = kl;
|
||
}
|
||
assert!(
|
||
kl.is_finite(),
|
||
"row {} KL is non-finite ({}); P_DD or P_ATH may have produced \
|
||
zero-probability rows that would imply softmax NaN/Inf in the \
|
||
synthetic projection",
|
||
b, kl
|
||
);
|
||
}
|
||
let mean_kl = total_kl / (batch_size as f32);
|
||
|
||
// The dd_pct contribution to the synthetic logits is
|
||
// `0.10 × W[h, 102]` for each output unit h. With Xavier-
|
||
// uniform W ~ U[-bound, bound] where bound ≈ 0.235 for
|
||
// (fan_in=103, fan_out=4), the pre-softmax logit shift is
|
||
// O(0.10 × 0.235) ≈ 0.024 per unit. Through softmax the
|
||
// resulting row-distribution shift produces KL > 1e-6 for
|
||
// generic random weights — the threshold is set 1000× below
|
||
// the expected magnitude so it rejects ONLY pure-noise
|
||
// (zero-init or wiring-broken) cases.
|
||
//
|
||
// If a future migration breaks the wiring (e.g. the column-102
|
||
// Xavier-init silently zeroes, the bn_concat_buf truncates to
|
||
// 102 columns, or W's K-dim falls back to 102), `mean_kl`
|
||
// drops to 0 and this assert fires — pinpointing the
|
||
// regression to the 4.1a/4.1b/4.1c chain.
|
||
const KL_EPSILON: f32 = 1e-6;
|
||
assert!(
|
||
mean_kl > KL_EPSILON,
|
||
"Wave 4.1c regression — dd_pct does NOT shift the synthetic policy: \
|
||
mean_kl={} max_kl={} threshold={}. Either: (a) bn_tanh_concat_dd \
|
||
stopped writing the dd_pct column, (b) the synthetic SGEMM K-dim \
|
||
fell back to 102, or (c) Xavier weights at column 102 are \
|
||
silently zero. Investigate before lowering the threshold — see \
|
||
feedback_no_quickfixes.",
|
||
mean_kl, max_kl, KL_EPSILON
|
||
);
|
||
|
||
// Diagnostic — print observed KL so the audit doc / commit
|
||
// message can record the actual sample magnitude (not just
|
||
// 'above threshold').
|
||
eprintln!(
|
||
"Wave 4.1c — dd_pct trunk wiring KL: mean={:.3e} max={:.3e} \
|
||
threshold={:.0e} batch={} proj_h={} concat_dd_dim={}",
|
||
mean_kl, max_kl, KL_EPSILON, batch_size, PROJ_H, concat_dd_dim
|
||
);
|
||
}
|
||
}
|
||
|
||
/// SP15 Phase 1.7 — verify the trainer's `set_test_data_from_slices` API
|
||
/// stashes the right slice and fires the observer with `(test_start_bar,
|
||
/// test_end_bar)`. This is the unit-level surface; the production verification
|
||
/// (every fold's test slice actually flows through the fold loop) is covered
|
||
/// by the L40S smoke and a future sealed-eval consumer that follows up the
|
||
/// deferred `evaluate_dqn_graphed` invocation per `feedback_no_partial_refactor`
|
||
/// (see `dqn-wire-up-audit.md` Phase 1.7 entry).
|
||
#[cfg(test)]
|
||
mod test_slice_consumption {
|
||
use ml::trainers::dqn::{DQNHyperparameters, DQNTrainer};
|
||
|
||
/// Test 1.7 — observer fires with the per-fold test slice range and the
|
||
/// stashed `test_features` / `test_targets` match the input slices.
|
||
/// Constructs a real trainer (sync init does not require CUDA forward
|
||
/// passes) and exercises only the new test-slice plumbing surface.
|
||
#[test]
|
||
fn set_test_data_from_slices_fires_observer_and_stashes() {
|
||
// Mirror `dqn_trainer_p1_tests.rs` — `DQNTrainer::new` is sync and
|
||
// does not require GPU forward passes for construction.
|
||
let hyperparams = DQNHyperparameters::conservative();
|
||
let trainer_result = DQNTrainer::new(hyperparams);
|
||
let mut trainer = match trainer_result {
|
||
Ok(t) => t,
|
||
// CPU trainer construction may not be supported in all build
|
||
// configurations (e.g. cuda-required paths). In that case the
|
||
// L40S smoke is the canonical verifier — skip with a printed
|
||
// notice rather than fail.
|
||
Err(e) => {
|
||
eprintln!(
|
||
"test_slice_consumption: skipping — trainer construction failed ({}). \
|
||
L40S smoke is the canonical verifier for this surface.",
|
||
e
|
||
);
|
||
return;
|
||
}
|
||
};
|
||
|
||
let calls = std::sync::Arc::new(std::sync::Mutex::new(Vec::<(usize, usize)>::new()));
|
||
let calls_clone = calls.clone();
|
||
trainer.set_test_data_observer(move |start, end| {
|
||
calls_clone.lock().unwrap().push((start, end));
|
||
});
|
||
|
||
// Synthetic 1000-bar test slice at [5000..6000) — typical fold-0 layout
|
||
// for an 8000-bar train range with default WF fractions.
|
||
let test_features: Vec<[f64; 42]> = (0..1000).map(|_| [0.0; 42]).collect();
|
||
let test_targets: Vec<[f64; 6]> = (0..1000).map(|_| [0.0; 6]).collect();
|
||
trainer.set_test_data_from_slices(&test_features, &test_targets, 5000, 6000);
|
||
|
||
let observed = calls.lock().unwrap();
|
||
assert_eq!(observed.len(), 1, "observer should fire exactly once");
|
||
assert_eq!(observed[0], (5000, 6000), "observer received wrong range");
|
||
|
||
// The stashed `test_features` / `test_targets` / `test_start_bar` /
|
||
// `test_end_bar` fields are `pub(crate)` and not reachable from this
|
||
// integration test; the observer firing with the correct range is the
|
||
// unit-level proof that the API wired the slice through. The deeper
|
||
// invariant (the fold loop calls this method with `[fold.test_start
|
||
// ..fold.test_end)`) is verified by the L40S smoke when a future
|
||
// commit lands the deferred `evaluate_dqn_graphed` consumer and the
|
||
// HEALTH_DIAG `test_slice fold=N` line appears once per fold.
|
||
}
|
||
}
|