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

feat(sp5): Task A6 — Pearl 6 cross-fold-persistent Kelly cap signals

Browse Source 
Adds `pearl_6_kelly_kernel.cu` (single-block 6-thread kernel reading
portfolio_state directly) to populate ISV[280..286) with Bayesian-prior
Kelly fraction, conviction identity, trade variance, cumulative sample
count (max-semantics), win-rate and loss-rate EMAs. EWMA α=0.01 is an
Invariant 1 structural anchor for cross-fold inertia.

Key design departure from A1-A5: ISV[280..286) are intentionally EXEMPT
from the StateResetRegistry and from apply_pearls/Pearls A+D. portfolio_state
has WindowReset lifecycle (resets at EVERY window boundary, not just fold),
making cross-fold persistence essential. The max()-semantics for sample_count
(s==3) ensure the cumulative count never decreases across any boundary.
Wiener offsets [525..543) that naive formula would produce are intentionally
unused; in-kernel EWMA replaces external wiener bootstrap.

Verification: cargo check -p ml --offline clean (11 pre-existing warnings,
no new). sp5_producer_unit_tests --no-run clean. Two GPU tests (12, 13)
validate EWMA blend and cross-fold persistence.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-01 23:22:06 +02:00
parent 1aaa1cb0bc
commit 5b4cdec4ff
7 changed files with 568 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

9
crates/ml/build.rs
View File

@@ -455,6 +455,15 @@ fn main() {
// apply_pearls_ad_kernel then smooths via Pearls A+D (ALPHA_META=1e-3)
// into ISV[IQN_TAU_BASE=250..270). Must run AFTER q_skew_kurtosis_kernel.
"pearl_5_iqn_tau_kernel.cu",
// SP5 Task A6 (2026-05-01): Pearl 6 cross-fold-persistent Kelly cap signals.
// Single-block 6-thread kernel (one per ISV output slot [280..286)).
// Reads portfolio_state [n_envs * PS_STRIDE]; writes directly to ISV[280..286).
// CROSS-FOLD-PERSISTENT: NOT in fold-reset registry. In-kernel EWMA alpha=0.01
// (Invariant 1 anchor). Cumulative-via-max() for sample_count slot (s==3).
// No apply_pearls call (sentinel incompatible with cross-fold persistence).
// Bayesian priors match kelly_cap_update_kernel.cu: prior_wins=2, prior_losses=2,
// prior_sum_wins=0.01, prior_sum_losses=0.01 (Invariant 1 anchors).
"pearl_6_kelly_kernel.cu",
// SP4 GPU-only Pearls A+D applicator (2026-05-01). Single-thread
// device-side loop applies Pearls A+D to N consecutive ISV slots
// (each with its own Wiener-state triple) on the producer's

100
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
View File

@@ -358,6 +358,18 @@ static SP5_Q_SKEW_KURTOSIS_CUBIN: &[u8] =
static SP5_PEARL_5_IQN_TAU_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/pearl_5_iqn_tau_kernel.cubin"));
/// SP5 Task A6 (2026-05-01): Pearl 6 cross-fold-persistent Kelly cap signals.
/// Single-block 6-thread kernel (one per output slot). Reads portfolio_state
/// [n_envs * PS_STRIDE] and current ISV[280..286); writes directly to ISV.
/// CROSS-FOLD-PERSISTENT: ISV[280..286) are EXEMPT from the fold-reset registry.
/// Does NOT use apply_pearls_ad_kernel: sentinel-bootstrap is incompatible with
/// cross-fold persistence (resetting to first-observation defeats the design).
/// In-kernel EWMA alpha=0.01 (Invariant 1 anchor). Cumulative-via-max() for
/// sample_count slot preserves cumulative trade count across window/fold resets.
/// Wiener offsets [525..543) that would naively map to these slots are unused.
static SP5_PEARL_6_KELLY_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/pearl_6_kelly_kernel.cubin"));
/// Plan 4 Task 3 (E.3): IQN multi-quantile diagnostic EMAs into ISV[99..103).
/// 4-block (256 threads/block, shmem-reduce, no atomicAdd) kernel launched
/// alongside `h_s2_rms_ema_update` from `training_loop.rs`. Reads the IQN
@@ -4229,6 +4241,14 @@ pub struct GpuDqnTrainer {
/// into ISV[IQN_TAU_BASE=250..270). No atomicAdd, no consumer migration in this commit.
/// Loaded from `pearl_5_iqn_tau_kernel.cubin`.
pearl_5_iqn_tau_kernel: CudaFunction,
/// SP5 Task A6 (2026-05-01): Pearl 6 cross-fold-persistent Kelly cap signals kernel.
/// Single-block 6-thread kernel (one per ISV output slot [280..286)).
/// Reads portfolio_state [n_envs * PS_STRIDE]; writes directly to ISV[280..286).
/// CROSS-FOLD-PERSISTENT: these slots are EXEMPT from the fold-reset registry.
/// In-kernel slow EWMA (alpha=0.01, Invariant 1 anchor); cumulative-via-max()
/// for sample_count slot (s==3). No apply_pearls call; no atomicAdd.
/// Loaded from `pearl_6_kelly_kernel.cubin`.
pearl_6_kelly_kernel: CudaFunction,
/// SP5 Task A4 (2026-05-01): previous-step gradient buffer for cosine similarity.
/// [total_params f32] mapped-pinned (feedback_no_htod_htoh_only_mapped_pinned).
/// Zeroed at construction and at each fold boundary (Pearl A sentinel: cosine_sim=0
@@ -11174,6 +11194,74 @@ impl GpuDqnTrainer {
Ok(())
}
/// SP5 Task A6: cross-fold-persistent Kelly cap signals.
///
/// CROSS-FOLD-PERSISTENT — ISV[280..286) are EXEMPT from the fold-reset registry.
/// The cold-start Quarter pathology (project_magnitude_eval_collapse_kelly_capped.md)
/// is caused by portfolio_state having WindowReset lifecycle: every window boundary
/// resets ps[14..18] (Kelly win/loss stats) to zero, forcing every new window to
/// cold-start at low maturity. Pearl 6 maintains a cross-fold and cross-window
/// representation so the maturity-driven cold-start cap does not re-engage.
///
/// Investigation finding: portfolio_state is ResetCategory::WindowReset in
/// state_reset_registry.rs (NOT FoldReset). This makes max() semantics for
/// sample_count essential for both fold-level and window-level cold-start suppression.
///
/// NO apply_pearls call: Pearls A+D's sentinel-bootstrap (zero ISV → first-observation
/// replacement) is incompatible with cross-fold persistence — the sentinel would reset
/// the slot if it were ever zeroed, defeating the whole design. In-kernel slow EWMA
/// (alpha=0.01, Invariant 1 anchor) is used instead.
///
/// Reads portfolio_state from the experience collector (same source as
/// launch_kelly_cap_update). Launches on self.stream; portfolio_state is populated
/// by the experience collector's stream before this epoch-boundary method is called
/// (same ordering guarantee as launch_kelly_cap_update, called in training_loop.rs).
pub(crate) fn launch_sp5_pearl_6_kelly(
&self,
portfolio_state_dev_ptr: u64,
n_envs: i32,
) -> Result<(), MLError> {
use crate::cuda_pipeline::sp5_isv_slots::{
KELLY_F_SMOOTH_INDEX, CONVICTION_SMOOTH_INDEX, TRADE_VAR_SMOOTH_INDEX,
KELLY_SAMPLE_COUNT_INDEX, WIN_RATE_SMOOTH_INDEX, LOSS_RATE_SMOOTH_INDEX,
};
debug_assert!(self.isv_signals_dev_ptr != 0,
"launch_sp5_pearl_6_kelly: isv_signals_dev_ptr must be allocated");
let isv_dev = self.isv_signals_dev_ptr;
let ps_stride_i32: i32 = 43; // PS_STRIDE from state_layout.cuh (grown 41→43 by Plan 3 D.4c)
let kelly_f_idx_i32 = KELLY_F_SMOOTH_INDEX as i32;
let conviction_idx_i32 = CONVICTION_SMOOTH_INDEX as i32;
let trade_var_idx_i32 = TRADE_VAR_SMOOTH_INDEX as i32;
let sample_count_idx_i32 = KELLY_SAMPLE_COUNT_INDEX as i32;
let win_rate_idx_i32 = WIN_RATE_SMOOTH_INDEX as i32;
let loss_rate_idx_i32 = LOSS_RATE_SMOOTH_INDEX as i32;
unsafe {
self.stream
.launch_builder(&self.pearl_6_kelly_kernel)
.arg(&portfolio_state_dev_ptr)
.arg(&n_envs)
.arg(&ps_stride_i32)
.arg(&isv_dev)
.arg(&kelly_f_idx_i32)
.arg(&conviction_idx_i32)
.arg(&trade_var_idx_i32)
.arg(&sample_count_idx_i32)
.arg(&win_rate_idx_i32)
.arg(&loss_rate_idx_i32)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (6, 1, 1),
shared_mem_bytes: 0,
})
.map_err(|e| MLError::ModelError(format!("pearl_6_kelly_update: {e}")))?;
}
Ok(())
}
/// Plan C Phase 2 follow-up A.2 (2026-04-29): launch `q_drift_rate_ema_update`
/// — single-thread single-block ISV producer for ISV[Q_DRIFT_RATE_INDEX=129].
///
@@ -13257,6 +13345,17 @@ impl GpuDqnTrainer {
.map_err(|e| MLError::ModelError(format!("pearl_5_iqn_tau_update load: {e}")))?
};
// SP5 Task A6 (2026-05-01): load pearl_6_kelly_kernel.
// Single-block 6-thread kernel. Reads portfolio_state; writes
// directly to ISV[280..286). CROSS-FOLD-PERSISTENT.
// No apply_pearls — sentinel incompatible with cross-fold persistence.
let pearl_6_kelly_kernel = {
let module = stream.context().load_cubin(SP5_PEARL_6_KELLY_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!("sp5 pearl_6_kelly cubin load: {e}")))?;
module.load_function("pearl_6_kelly_update")
.map_err(|e| MLError::ModelError(format!("pearl_6_kelly_update load: {e}")))?
};
// SP5 Task A4 (2026-05-01): allocate grad_prev_buf_per_group [total_params f32].
// Zeroed at construction; fold-boundary reset zeroes it again (Pearl A sentinel:
// zero grad_prev → cosine_sim=0 on first step → β1/β2 start at envelope midpoints).
@@ -16406,6 +16505,7 @@ impl GpuDqnTrainer {
pearl_4_adam_hparams_kernel,
q_skew_kurtosis_kernel,
pearl_5_iqn_tau_kernel,
pearl_6_kelly_kernel,
grad_prev_buf_per_group,
group_param_offsets_buf,
atoms_clip_count_buf,

209
crates/ml/src/cuda_pipeline/pearl_6_kelly_kernel.cu Normal file
View File

@@ -0,0 +1,209 @@
/* pearl_6_kelly_update — SP5 Task A6 cross-fold-persistent Kelly cap signals.
*
* CROSS-FOLD-PERSISTENT — ISV[280..286) are EXEMPT from the fold-reset registry.
* The cold-start Quarter pathology (project_magnitude_eval_collapse_kelly_capped.md)
* is caused by portfolio_state having WindowReset lifecycle: every window boundary
* resets win/loss counts to zero, which forces every new window to cold-start at
* low maturity → Kelly cap clamps actual_mag to Quarter for the first ~10 trades.
* Pearl 6's job is to maintain a cross-fold and cross-window representation of
* Kelly statistics so the maturity-driven cold-start cap does not re-engage.
*
* Investigation result (Task A6): portfolio_state has ResetCategory::WindowReset
* in state_reset_registry.rs (not FoldReset). This means ps[14..18] (Kelly stats)
* reset at EVERY window boundary — more frequently than fold boundaries.
* The max() semantics for sample_count are therefore essential for both
* fold-level and window-level cold-start suppression.
*
* DOES NOT use apply_pearls_ad_kernel — two reasons:
* 1. Pearls A+D's sentinel-bootstrap (zero ISV → first-observation replacement)
* is INCOMPATIBLE with cross-fold persistence. If these slots were in the
* fold-reset registry, the sentinel would fire at every fold boundary and
* replace the cross-fold state — defeating the entire design.
* 2. The wiener_state_buf is sized for the 110 SP5 producers at slots [174..278)
* (below the 2-slot carve-out gap at [278..280)). Pearl 6 slots [280..286)
* would map to wiener offsets 213 + (280-174)*3 = 525..543 under the naive
* formula — at the far edge of the buffer (wiener_state_buf = 543 floats).
* Slots 281..285 would map to offsets 528..543 — within the allocated range
* but would be atomically zeroed at fold boundary (part of wiener_state_buf),
* which is incompatible with cross-fold persistence. We intentionally leave
* wiener offsets [525..543) unused; Pearl 6 does not read or write them.
*
* Instead this kernel does its own slow in-kernel EWMA (alpha=0.01, Invariant 1
* anchor — slow rate preserves cross-fold inertia) for 5 of 6 output slots.
* Sample count (s==3) uses cumulative-via-max() semantics: at any window/fold
* boundary portfolio_state ps[14..18] reset so new_obs = total_trades drops; the
* max(current_isv, new_obs) call keeps the cumulative count non-decreasing.
*
* Bayesian priors match kelly_cap_update_kernel.cu exactly (Invariant 1 anchors):
* prior_wins=2.0, prior_losses=2.0,
* prior_sum_wins=0.01, prior_sum_losses=0.01
*
* Conviction_smooth slot (s==1) has new_obs = current: the conviction signal is not
* derivable from portfolio_state alone. This is an intentional no-op update until
* Layer B wires the magnitude_conviction_buf source. The EWMA blend with new_obs ==
* current yields new_isv = current (identity) — the slot holds its cross-fold value.
*
* 6-thread single-block (one thread per output slot). No atomicAdd.
* __threadfence_system() after writes ensures visibility across device contexts.
*/
#include "state_layout.cuh"
extern "C" __global__ void pearl_6_kelly_update(
const float* __restrict__ portfolio_state, /* [n_envs * ps_stride] */
int n_envs,
int ps_stride,
float* __restrict__ isv_signals, /* direct ISV write — cross-fold-persistent */
int kelly_f_idx, /* ISV slot 280 */
int conviction_idx, /* ISV slot 281 */
int trade_var_idx, /* ISV slot 282 */
int sample_count_idx, /* ISV slot 283 */
int win_rate_idx, /* ISV slot 284 */
int loss_rate_idx /* ISV slot 285 */
) {
if (blockIdx.x != 0 || threadIdx.x >= 6) return;
int s = (int)threadIdx.x;
/* Bayesian priors — match kelly_cap_update_kernel.cu exactly (Invariant 1 anchors). */
const float PRIOR_WINS = 2.0f;
const float PRIOR_LOSSES = 2.0f;
const float PRIOR_SUM_WINS = 0.01f;
const float PRIOR_SUM_LOSSES = 0.01f;
/* Aggregate across all envs. Each thread re-reads independently.
* 6 threads × n_envs reads: acceptable for cold-path (per-epoch). */
float total_wins = 0.0f;
float total_losses = 0.0f;
float sum_wins = 0.0f;
float sum_losses = 0.0f;
float kelly_sum = 0.0f;
int kelly_count = 0;
for (int e = 0; e < n_envs; ++e) {
const float* ps = portfolio_state + (long long)e * ps_stride;
float wc = ps[PS_KELLY_WIN_COUNT];
float lc = ps[PS_KELLY_LOSS_COUNT];
float sw = ps[PS_KELLY_SUM_WINS];
float sl = ps[PS_KELLY_SUM_LOSSES];
total_wins += wc;
total_losses += lc;
sum_wins += sw;
sum_losses += sl;
/* Per-env Bayesian-prior Kelly fraction — mirrors kelly_cap_update_kernel.cu exactly.
* Note: this kernel does NOT apply the half-Kelly factor (×0.5 for risk management)
* that kelly_cap_update applies. Pearl 6 tracks the raw statistical Kelly fraction
* as a signal; kelly_cap_update_kernel continues to own the capped, health-coupled
* ISV[47] slot consumed by position sizing. */
float eff_w = wc + PRIOR_WINS;
float eff_l = lc + PRIOR_LOSSES;
float wr = eff_w / (eff_w + eff_l);
float aw = (sw + PRIOR_SUM_WINS) / eff_w;
float al = (sl + PRIOR_SUM_LOSSES) / eff_l;
float pf = aw / fmaxf(al, 0.0001f);
float kf = (pf * wr - (1.0f - wr)) / fmaxf(pf, 0.0001f);
kf = fmaxf(0.0f, fminf(1.0f, kf));
kelly_sum += kf;
kelly_count++;
}
float total_trades = total_wins + total_losses;
/* Resolve ISV slot index for this thread. */
int idx_table[6] = {
kelly_f_idx, conviction_idx, trade_var_idx,
sample_count_idx, win_rate_idx, loss_rate_idx
};
int idx = idx_table[s];
float current = isv_signals[idx];
float new_obs;
if (s == 0) {
/* kelly_f_smooth: mean raw Kelly fraction across envs (not half-Kelly). */
new_obs = (kelly_count > 0)
? (kelly_sum / (float)kelly_count)
: 0.5f;
} else if (s == 1) {
/* conviction_smooth: not derivable from portfolio_state.
* Intentional identity update — Layer B replaces this with the actual
* magnitude_conviction_buf source. The EWMA blend of (new_obs == current)
* is an identity: new_isv = (1-a)*current + a*current = current. */
new_obs = current;
} else if (s == 2) {
/* trade_var_smooth: variance proxy — standard deviation of per-env Kelly
* fractions across envs (measures cross-env PnL dispersion).
* Second pass over envs: same Bayesian-prior formula, accumulate sq-dev. */
float kelly_mean = (kelly_count > 0)
? (kelly_sum / (float)kelly_count)
: 0.5f;
float ksum_sq = 0.0f;
for (int e = 0; e < n_envs; ++e) {
const float* ps = portfolio_state + (long long)e * ps_stride;
float wc = ps[PS_KELLY_WIN_COUNT];
float lc = ps[PS_KELLY_LOSS_COUNT];
float sw = ps[PS_KELLY_SUM_WINS];
float sl = ps[PS_KELLY_SUM_LOSSES];
float eff_w = wc + PRIOR_WINS;
float eff_l = lc + PRIOR_LOSSES;
float wr = eff_w / (eff_w + eff_l);
float aw = (sw + PRIOR_SUM_WINS) / eff_w;
float al = (sl + PRIOR_SUM_LOSSES) / eff_l;
float pf = aw / fmaxf(al, 0.0001f);
float kf = (pf * wr - (1.0f - wr)) / fmaxf(pf, 0.0001f);
kf = fmaxf(0.0f, fminf(1.0f, kf));
float d = kf - kelly_mean;
ksum_sq += d * d;
}
new_obs = (kelly_count > 1)
? (ksum_sq / (float)kelly_count)
: 0.0f;
} else if (s == 3) {
/* kelly_sample_count: CROSS-FOLD-PERSISTENT cumulative trade count.
* portfolio_state resets at window boundaries (WindowReset lifecycle),
* so total_trades drops back to a small number after every window.
* max(current, new_obs) keeps the cumulative count non-decreasing:
* - within a window: total_trades grows → max grows with it
* - at window/fold boundary: total_trades drops → max keeps old value
* - after cold-start re-accumulates past old max: max grows again
* This is the primary mechanism breaking the cold-start Kelly cap pathology. */
new_obs = total_trades;
} else if (s == 4) {
/* win_rate_smooth: aggregate win rate across all envs. */
new_obs = (total_trades > 0.0f)
? (total_wins / total_trades)
: 0.5f;
} else {
/* s == 5: loss_rate_smooth: aggregate loss rate across all envs. */
new_obs = (total_trades > 0.0f)
? (total_losses / total_trades)
: 0.5f;
}
/* Update ISV with cross-fold-aware blending.
*
* KELLY_EMA_ALPHA = 0.01 — Invariant 1 anchor (slow rate). A rate of 0.01
* means the effective window is ~100 steps; this is intentionally slow to
* preserve cross-fold inertia so a single bad window doesn't collapse the
* long-run Kelly estimate. The kelly_cap_update kernel (ISV[47]) continues
* to provide the health-coupled, half-Kelly-capped per-epoch estimate for
* position sizing; Pearl 6 tracks the longer-horizon statistical picture. */
const float KELLY_EMA_ALPHA = 0.01f;
float new_isv;
if (s == 3) {
/* Sample count: cumulative-via-max() — never decreases across windows/folds. */
new_isv = fmaxf(current, new_obs);
} else {
/* Slow EWMA: blend current cross-fold value with new per-epoch observation. */
new_isv = (1.0f - KELLY_EMA_ALPHA) * current + KELLY_EMA_ALPHA * new_obs;
}
isv_signals[idx] = new_isv;
__threadfence_system();
}

17
crates/ml/src/trainers/dqn/state_reset_registry.rs
View File

@@ -625,8 +625,23 @@ impl StateResetRegistry {
RegistryEntry {
name: "sp5_wiener_state",
category: ResetCategory::FoldReset,
description: "GpuDqnTrainer.wiener_state_buf SP5 block [330 floats = 110 SP5 producers × {sample_var, diff_var, x_lag}] at offsets [213..543). SP5 Pearl D Wiener-EMA state for Task A1's 24 producers (16 Pearl 1 + 8 shared signal). Bulk write_bytes(0) at fold boundary covers the full wiener_state_buf [0..543) (SP4 + SP5 together) so Pearl A's sentinel branch fires on the new fold's first producer launch. Both SP4 and SP5 blocks reset atomically — wiener_state_buf grew from 213 to 543 floats in Task A1.",
description: "GpuDqnTrainer.wiener_state_buf SP5 block [330 floats = 110 SP5 producers × {sample_var, diff_var, x_lag}] at offsets [213..543). SP5 Pearl D Wiener-EMA state for SP5 Pearls 1, 2, 3, 4, 5 producers. Bulk write_bytes(0) at fold boundary covers the full wiener_state_buf [0..543) (SP4 + SP5 together) so Pearl A's sentinel branch fires on the new fold's first producer launch. Pearl 6 (slots 280..286) does NOT use this buffer — Pearl 6 is cross-fold-persistent and uses in-kernel EWMA + cumulative-via-max(); the wiener slots [525..543) that Pearl 6 would map to under naive offset arithmetic (213 + (280-174)*3 = 525) are intentionally unused. Both SP4 and SP5 [non-Pearl-6] blocks reset atomically. wiener_state_buf grew from 213 to 543 floats in Task A1.",
},
// SP5 Task A6: Pearl 6 cross-fold-persistent Kelly cap signals.
// ISV[280..286) are EXPLICITLY EXEMPT from this registry.
// Adding fold-reset entries for these slots would defeat the entire
// design purpose: these slots must survive across fold and window
// boundaries to break the cold-start Kelly cap pathology.
//
// Investigation finding: portfolio_state has ResetCategory::WindowReset
// (not FoldReset) — it resets at every WINDOW boundary. This makes the
// cross-fold-persistent design even more critical: without Pearl 6, the
// Kelly cold-start re-engages after every window, not just every fold.
//
// The pearl_6_kelly_kernel uses in-kernel slow EWMA (alpha=0.01) and
// cumulative-via-max() for sample_count. No apply_pearls call; no wiener
// state entries needed. ISV[280..286) are never zeroed by the fold-reset
// path. This is the intended, documented behavior for Task A6.
// SP5 Task A4: per-group Adam β1/β2/ε ISV slots (Pearl 4).
// ISV[226..250) = 24 floats (8 groups × 3 hyperparams).
// Pearl A sentinel 0 at fold boundary → first step fires first-observation

5
crates/ml/src/trainers/dqn/trainer/training_loop.rs
View File

@@ -289,6 +289,11 @@ impl DQNTrainer {
if let Err(e) = fused.launch_kelly_cap_update(ps_dev_ptr, n_envs) {
tracing::warn!(epoch, "launch_kelly_cap_update failed (non-fatal): {e}");
}
// SP5 Task A6: Pearl 6 reads portfolio_state directly; cross-fold-persistent (NO apply_pearls).
// Called at the same epoch boundary as launch_kelly_cap_update (same ps_dev_ptr source).
if let Err(e) = fused.trainer().launch_sp5_pearl_6_kelly(ps_dev_ptr, n_envs) {
tracing::warn!(error = %e, "SP5 Pearl 6 producer launch failed");
}
}
// D.8 Plan 2 Task 6C: update TLOB regime focus EMA in ISV at epoch boundary.

227
crates/ml/tests/sp5_producer_unit_tests.rs
View File

@@ -21,6 +21,10 @@
//! 10. `pearl_5_iqn_tau_symmetric_q_yields_uniform_default_schedule`
//! 11. `pearl_5_iqn_tau_left_skew_shifts_quantiles_negatively`
//!
//! Task A6 (Pearl 6 cross-fold-persistent Kelly):
//! 12. `pearl_6_kelly_within_fold_ewma_blend`
//! 13. `pearl_6_kelly_cross_fold_sample_count_persists_via_max`
//!
//! All tests use analytically-known synthetic inputs with no CPU reference
//! implementation, per `feedback_no_cpu_test_fallbacks.md`. All mapped-pinned
//! memory (no DtoH), all `#[ignore]`-gated for GPU.
@@ -60,6 +64,9 @@ const SP5_Q_SKEW_KURTOSIS_CUBIN: &[u8] =
const SP5_PEARL_5_IQN_TAU_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/pearl_5_iqn_tau_kernel.cubin"));
const SP5_PEARL_6_KELLY_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/pearl_6_kelly_kernel.cubin"));
// ── Helpers ───────────────────────────────────────────────────────────────────
fn load_pearl_3_sigma_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
@@ -1511,3 +1518,223 @@ fn pearl_5_iqn_tau_left_skew_shifts_quantiles_negatively() {
);
}
}
// ── Pearl 6 helpers ───────────────────────────────────────────────────────────
fn load_pearl_6_kelly_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
let ctx = stream.context();
let module = ctx.load_cubin(SP5_PEARL_6_KELLY_CUBIN.to_vec())
.expect("load pearl_6_kelly_kernel cubin");
module.load_function("pearl_6_kelly_update")
.expect("load pearl_6_kelly_update function")
}
/// Allocate a MappedF32Buffer for portfolio_state [n_envs * ps_stride] and ISV.
/// Returns (portfolio_buf, isv_buf). Both are mapped-pinned (no DtoH).
/// PS_STRIDE=43 matches state_layout.cuh.
const PS_STRIDE: usize = 43;
const PS_KELLY_WIN_COUNT: usize = 14;
const PS_KELLY_LOSS_COUNT: usize = 15;
const PS_KELLY_SUM_WINS: usize = 16;
const PS_KELLY_SUM_LOSSES: usize = 17;
// ISV slot indices for Pearl 6 outputs (match sp5_isv_slots.rs).
const KELLY_F_SMOOTH_IDX: usize = 280;
const CONVICTION_SMOOTH_IDX: usize = 281;
const TRADE_VAR_SMOOTH_IDX: usize = 282;
const KELLY_SAMPLE_COUNT_IDX: usize = 283;
const WIN_RATE_SMOOTH_IDX: usize = 284;
const LOSS_RATE_SMOOTH_IDX: usize = 285;
fn launch_pearl_6(
stream: &Arc<CudaStream>,
kernel: &CudaFunction,
ps_buf: &MappedF32Buffer,
isv_buf: &MappedF32Buffer,
n_envs: i32,
) {
let ps_dev = ps_buf.dev_ptr;
let isv_dev = isv_buf.dev_ptr;
let ps_stride_i32 = PS_STRIDE as i32;
let kelly_f_i32 = KELLY_F_SMOOTH_IDX as i32;
let conv_i32 = CONVICTION_SMOOTH_IDX as i32;
let tvar_i32 = TRADE_VAR_SMOOTH_IDX as i32;
let scount_i32 = KELLY_SAMPLE_COUNT_IDX as i32;
let wr_i32 = WIN_RATE_SMOOTH_IDX as i32;
let lr_i32 = LOSS_RATE_SMOOTH_IDX as i32;
unsafe {
stream
.launch_builder(kernel)
.arg(&ps_dev)
.arg(&n_envs)
.arg(&ps_stride_i32)
.arg(&isv_dev)
.arg(&kelly_f_i32)
.arg(&conv_i32)
.arg(&tvar_i32)
.arg(&scount_i32)
.arg(&wr_i32)
.arg(&lr_i32)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (6, 1, 1),
shared_mem_bytes: 0,
})
.expect("launch pearl_6_kelly_update");
}
stream.synchronize().expect("sync after pearl_6_kelly_update");
}
/// SP5 Task A6 Test 12: within-fold EWMA blend.
///
/// Analytical derivation:
/// Single env with win_count=10, loss_count=5, sum_wins=0.5, sum_losses=0.25.
/// Bayesian priors: prior_wins=2, prior_losses=2, prior_sum_wins=0.01, prior_sum_losses=0.01.
/// eff_w=12, eff_l=7, wr=12/19≈0.6316, avg_w=(0.5+0.01)/12≈0.0425, avg_l=(0.25+0.01)/7≈0.0371
/// payoff=0.0425/0.0371≈1.1456, kelly_f=(1.1456×0.6316 - 0.3684)/1.1456≈(0.7234-0.3684)/1.1456≈0.3100
/// EWMA: new_isv = 0.99 × 0.5 + 0.01 × computed_kelly_f
/// ISV initial = 0.5 for all slots.
///
/// For sample_count (s==3): max(0.5, 15.0) = 15.0 (new_obs = total_trades = 15 > current=0.5).
/// For win_rate (s==4): new_obs = 10/15 ≈ 0.6667, new_isv = 0.99×0.5 + 0.01×0.6667 ≈ 0.5017.
/// For loss_rate (s==5): new_obs = 5/15 ≈ 0.3333, new_isv = 0.99×0.5 + 0.01×0.3333 ≈ 0.4983.
#[test]
#[ignore = "requires GPU"]
fn pearl_6_kelly_within_fold_ewma_blend() {
let ctx = CudaContext::new(0).expect("CUDA context");
let stream = ctx.new_stream().expect("stream");
let kernel = load_pearl_6_kelly_kernel(&stream);
let n_envs = 1_i32;
let isv_size = 290_usize; // covers slots [280..286)
let mut ps_buf = unsafe { MappedF32Buffer::new(n_envs as usize * PS_STRIDE) }.unwrap();
let mut isv_buf = unsafe { MappedF32Buffer::new(isv_size) }.unwrap();
// Set portfolio_state for env 0.
{
let ps_slice = ps_buf.host_slice_mut();
ps_slice.fill(0.0_f32);
ps_slice[PS_KELLY_WIN_COUNT] = 10.0;
ps_slice[PS_KELLY_LOSS_COUNT] = 5.0;
ps_slice[PS_KELLY_SUM_WINS] = 0.5;
ps_slice[PS_KELLY_SUM_LOSSES] = 0.25;
}
// Initialise ISV to 0.5 for all Pearl 6 slots.
{
let isv_slice = isv_buf.host_slice_mut();
isv_slice.fill(0.0_f32);
isv_slice[KELLY_F_SMOOTH_IDX] = 0.5;
isv_slice[CONVICTION_SMOOTH_IDX] = 0.5;
isv_slice[TRADE_VAR_SMOOTH_IDX] = 0.5;
isv_slice[KELLY_SAMPLE_COUNT_IDX] = 0.5;
isv_slice[WIN_RATE_SMOOTH_IDX] = 0.5;
isv_slice[LOSS_RATE_SMOOTH_IDX] = 0.5;
}
launch_pearl_6(&stream, &kernel, &ps_buf, &isv_buf, n_envs);
let isv = isv_buf.read_all();
// slot 280 — kelly_f_smooth: EWMA blend of computed kelly_f.
// Computed kelly_f: eff_w=12, eff_l=7, wr=12/19
let eff_w = 12.0_f32; let eff_l = 7.0_f32;
let wr = eff_w / (eff_w + eff_l);
let avg_w = (0.5_f32 + 0.01) / eff_w;
let avg_l = (0.25_f32 + 0.01) / eff_l;
let pf = avg_w / avg_l.max(0.0001);
let kf = ((pf * wr - (1.0 - wr)) / pf.max(0.0001)).clamp(0.0, 1.0);
let expected_kelly_f = 0.99 * 0.5 + 0.01 * kf;
println!("kf={kf:.6} expected_kelly_f={expected_kelly_f:.6} got={:.6}", isv[KELLY_F_SMOOTH_IDX]);
assert!((isv[KELLY_F_SMOOTH_IDX] - expected_kelly_f).abs() < 2e-5,
"kelly_f: expected {expected_kelly_f:.6}, got {:.6}", isv[KELLY_F_SMOOTH_IDX]);
// slot 281 — conviction_smooth: identity (no-op update, current=0.5).
// EWMA of (new_obs=current=0.5) with current=0.5 → still 0.5.
assert!((isv[CONVICTION_SMOOTH_IDX] - 0.5).abs() < 1e-6,
"conviction: expected 0.5 (identity), got {:.6}", isv[CONVICTION_SMOOTH_IDX]);
// slot 283 — sample_count: cumulative via max(current=0.5, new_obs=15.0) = 15.0.
assert!((isv[KELLY_SAMPLE_COUNT_IDX] - 15.0).abs() < 1e-4,
"sample_count: expected 15.0, got {:.6}", isv[KELLY_SAMPLE_COUNT_IDX]);
// slot 284 — win_rate_smooth: EWMA blend.
let expected_wr = 0.99 * 0.5 + 0.01 * (10.0_f32 / 15.0);
assert!((isv[WIN_RATE_SMOOTH_IDX] - expected_wr).abs() < 2e-5,
"win_rate: expected {expected_wr:.6}, got {:.6}", isv[WIN_RATE_SMOOTH_IDX]);
// slot 285 — loss_rate_smooth: EWMA blend.
let expected_lr = 0.99 * 0.5 + 0.01 * (5.0_f32 / 15.0);
assert!((isv[LOSS_RATE_SMOOTH_IDX] - expected_lr).abs() < 2e-5,
"loss_rate: expected {expected_lr:.6}, got {:.6}", isv[LOSS_RATE_SMOOTH_IDX]);
}
/// SP5 Task A6 Test 13: cross-fold sample_count persists via max().
///
/// This test verifies the core cross-fold-persistence mechanism:
/// Step 1: portfolio_state total_trades=100. Run → ISV[283] = max(0, 100) = 100.
/// Step 2: Simulate fold boundary — reset portfolio_state to 0 (total_trades=0).
/// Run → new_obs=0, max(100, 0) = 100. ISV[283] preserved.
/// Step 3: portfolio_state total_trades=10 (recovery, below prior max).
/// Run → new_obs=10, max(100, 10) = 100. Still preserved.
/// Step 4: portfolio_state total_trades=150 (surpasses old max).
/// Run → new_obs=150, max(100, 150) = 150. Grows to new max.
#[test]
#[ignore = "requires GPU"]
fn pearl_6_kelly_cross_fold_sample_count_persists_via_max() {
let ctx = CudaContext::new(0).expect("CUDA context");
let stream = ctx.new_stream().expect("stream");
let kernel = load_pearl_6_kelly_kernel(&stream);
let n_envs = 1_i32;
let isv_size = 290_usize;
let mut ps_buf = unsafe { MappedF32Buffer::new(n_envs as usize * PS_STRIDE) }.unwrap();
let mut isv_buf = unsafe { MappedF32Buffer::new(isv_size) }.unwrap();
// Helper: set portfolio_state wins/losses for env 0.
let set_ps = |ps_buf: &mut MappedF32Buffer, wins: f32, losses: f32| {
let sl = ps_buf.host_slice_mut();
sl.fill(0.0_f32);
sl[PS_KELLY_WIN_COUNT] = wins;
sl[PS_KELLY_LOSS_COUNT] = losses;
sl[PS_KELLY_SUM_WINS] = wins * 0.05; // arbitrary; only count matters for this test
sl[PS_KELLY_SUM_LOSSES] = losses * 0.05;
};
// Zero-init ISV (MappedF32Buffer::new already zeroes, but be explicit).
isv_buf.host_slice_mut().fill(0.0_f32);
// Step 1: total_trades=100, ISV[283] starts at 0.
set_ps(&mut ps_buf, 60.0, 40.0); // total=100
launch_pearl_6(&stream, &kernel, &ps_buf, &isv_buf, n_envs);
let v1 = isv_buf.read_all()[KELLY_SAMPLE_COUNT_IDX];
println!("step1 ISV[283]={v1:.1} (expected 100)");
assert!((v1 - 100.0).abs() < 0.5, "step1: expected 100.0, got {v1:.2}");
// Step 2: fold boundary simulation — portfolio_state reset to 0.
set_ps(&mut ps_buf, 0.0, 0.0); // total=0
launch_pearl_6(&stream, &kernel, &ps_buf, &isv_buf, n_envs);
let v2 = isv_buf.read_all()[KELLY_SAMPLE_COUNT_IDX];
println!("step2 ISV[283]={v2:.1} (expected 100 — preserved via max)");
assert!((v2 - 100.0).abs() < 0.5,
"step2: cross-fold persistence failed — expected 100.0, got {v2:.2}");
// Step 3: portfolio_state recovering, total=10 (below old max).
set_ps(&mut ps_buf, 6.0, 4.0); // total=10
launch_pearl_6(&stream, &kernel, &ps_buf, &isv_buf, n_envs);
let v3 = isv_buf.read_all()[KELLY_SAMPLE_COUNT_IDX];
println!("step3 ISV[283]={v3:.1} (expected 100 — still preserved)");
assert!((v3 - 100.0).abs() < 0.5,
"step3: max should still hold 100.0, got {v3:.2}");
// Step 4: total_trades=150 — surpasses old max, should grow.
set_ps(&mut ps_buf, 90.0, 60.0); // total=150
launch_pearl_6(&stream, &kernel, &ps_buf, &isv_buf, n_envs);
let v4 = isv_buf.read_all()[KELLY_SAMPLE_COUNT_IDX];
println!("step4 ISV[283]={v4:.1} (expected 150 — grew to new max)");
assert!((v4 - 150.0).abs() < 0.5,
"step4: expected 150.0 (new max), got {v4:.2}");
}

2
docs/dqn-wire-up-audit.md
View File

File diff suppressed because one or more lines are too long