From 5b4cdec4ffeae60a4fc647075b2f74a4fcc140f5 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 1 May 2026 23:22:06 +0200 Subject: [PATCH] =?UTF-8?q?feat(sp5):=20Task=20A6=20=E2=80=94=20Pearl=206?= =?UTF-8?q?=20cross-fold-persistent=20Kelly=20cap=20signals?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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 --- crates/ml/build.rs | 9 + .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 100 ++++++++ .../src/cuda_pipeline/pearl_6_kelly_kernel.cu | 209 ++++++++++++++++ .../src/trainers/dqn/state_reset_registry.rs | 17 +- .../src/trainers/dqn/trainer/training_loop.rs | 5 + crates/ml/tests/sp5_producer_unit_tests.rs | 227 ++++++++++++++++++ docs/dqn-wire-up-audit.md | 2 + 7 files changed, 568 insertions(+), 1 deletion(-) create mode 100644 crates/ml/src/cuda_pipeline/pearl_6_kelly_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index f631f52b7..85084e728 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -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 diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 57fcd993b..111b437a3 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -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, diff --git a/crates/ml/src/cuda_pipeline/pearl_6_kelly_kernel.cu b/crates/ml/src/cuda_pipeline/pearl_6_kelly_kernel.cu new file mode 100644 index 000000000..ed7cdbcfe --- /dev/null +++ b/crates/ml/src/cuda_pipeline/pearl_6_kelly_kernel.cu @@ -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(); +} diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 7192c8b1c..12b96edcb 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -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 diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 7d6a5de69..8a3bf4a9e 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -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. diff --git a/crates/ml/tests/sp5_producer_unit_tests.rs b/crates/ml/tests/sp5_producer_unit_tests.rs index 0db77f949..1be9c34bb 100644 --- a/crates/ml/tests/sp5_producer_unit_tests.rs +++ b/crates/ml/tests/sp5_producer_unit_tests.rs @@ -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) -> 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) -> 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, + 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}"); +} diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 4810e3683..1fa9165aa 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,8 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +SP5 Task A6 — Pearl 6 cross-fold-persistent Kelly cap signals GPU producer (Layer A, 2026-05-01): one new CUDA kernel (`pearl_6_kelly_kernel.cu`) lands as a Layer A additive producer feeding ISV slots [280..286) with cross-fold-persistent Kelly cap statistics. No consumer migration in this commit — Layer B wires ISV[280..286) into the Kelly cap controller once all Layer A producers are complete. `pearl_6_kelly_update`: single-block 6-thread kernel; thread `s ∈ [0,6)` handles ISV slot `KELLY_F_SMOOTH_IDX + s` (= 280 + s). Reads `portfolio_state[env × ps_stride + offset]` for all `n_envs` envs (cold-path, per-epoch, 6×n_envs reads acceptable). Bayesian priors `prior_wins=2, prior_losses=2, prior_sum_wins=0.01, prior_sum_losses=0.01` are Invariant 1 structural anchors (match `kelly_cap_update_kernel.cu`). In-kernel EWMA α=0.01 is an Invariant 1 anchor — slow rate to preserve cross-fold inertia. Per-slot semantics: s==0 (kelly_f_smooth): Bayesian Kelly fraction `kf = (payoff×wr - (1-wr)) / payoff`, clamped to [0,1], EWMA; s==1 (conviction_smooth): intentional identity update (`new_obs = current`) — Layer B will wire the actual conviction source; s==2 (trade_var_smooth): variance proxy (second pass over envs), EWMA; s==3 (sample_count): cumulative-via-max `new_isv = max(current, new_obs)` — count never decreases when portfolio_state resets; s==4 (win_rate_smooth): `new_obs = total_wins / total_trades`, EWMA; s==5 (loss_rate_smooth): `new_obs = total_losses / total_trades`, EWMA. `__threadfence_system()` after all writes. **Investigation finding**: `portfolio_state` has `WindowReset` lifecycle (resets at every WINDOW boundary, not just fold boundary — more aggressive than FoldReset). This makes cross-fold persistence even more critical: without the `max()` semantics for sample_count and the in-kernel EWMA carryover, every window boundary (not just fold boundary) would re-engage the Kelly cold-start Quarter-dominance pathology. **Registry carve-out**: ISV[280..286) are EXEMPT from the StateResetRegistry. This is an intentional, documented departure from the standard fold-reset discipline — the entire purpose of these slots is to survive fold and window resets. No `RegistryEntry` blocks added. State reset registry updated with an explicit audit comment block (near `sp5_wiener_state` entry) documenting: (a) ISV[280..286) exempt, (b) portfolio_state WindowReset lifecycle, (c) in-kernel design rationale. **Wiener state carve-out**: Pearl 6 does NOT use `apply_pearls` / `launch_apply_pearls` for write-back. Under the naive wiener offset formula `213 + (280-174)×3 = 525`, slots [280..282) would map to wiener_state_buf[525..534) (within the 543-float buffer), but slots [283..285) would map correctly. More importantly, wiener_state_buf is part of the fold-reset path — resetting Pearls A+D wiener state at fold boundary is incompatible with cross-fold persistence, so wiener/apply_pearls is correctly omitted entirely. In-kernel EWMA replaces the external Pearls A+D bootstrap discipline. **No apply_pearls**: sentinel-bootstrap (zero ISV → first-observation replacement) is also incompatible with cross-fold persistence (sentinel detection would fire after every fold reset, wiping accumulated history); replaced by in-kernel EWMA that naturally continues from the current ISV value. **Stream ordering**: Pearl 6 is launched at the same epoch-boundary site as `launch_kelly_cap_update`, reusing the same `ps_dev_ptr` and `n_envs` from the existing `if let (Some(ref fused), Some(ref collector))` block — natural stream ordering guarantees portfolio_state is populated before Pearl 6 reads it. No scratch buffer growth (Pearl 6 reads portfolio_state directly, no intermediate scratch). No wiener_state_buf growth. No StateResetRegistry entries added. Wire-up in training_loop.rs: `launch_sp5_pearl_6_kelly(ps_dev_ptr, n_envs)` called immediately after `launch_kelly_cap_update` in the existing epoch-boundary block with `tracing::warn!(error = %e, ...)` on error. Two GPU-only `#[ignore = "requires GPU"]` unit tests: (12) `pearl_6_kelly_within_fold_ewma_blend` — single env, known win/loss/sum_wins/sum_losses, verifies Bayesian Kelly EWMA for s==0, identity for s==1, cumulative-via-max for s==3, EWMA for s==4 and s==5; (13) `pearl_6_kelly_cross_fold_sample_count_persists_via_max` — 4-step test: 100 trades → count=100; reset ps to 0 → count still 100; 10 trades → still 100; 150 trades → count grows to 150. No CPU compute, no HtoD/HtoH, no atomicAdd, no stubs, no consumer migration. Touched: `cuda_pipeline/pearl_6_kelly_kernel.cu` (new), `build.rs` (+1 cubin entry), `cuda_pipeline/gpu_dqn_trainer.rs` (+1 static cubin + 1 struct field + cubin loader + struct initializer + launcher method `launch_sp5_pearl_6_kelly`), `trainers/dqn/state_reset_registry.rs` (audit comment block documenting registry carve-out — NO new RegistryEntry blocks), `trainers/dqn/trainer/training_loop.rs` (+5 LOC inside existing kelly epoch-boundary block), `tests/sp5_producer_unit_tests.rs` (+2 GPU-only tests + 1 cubin constant + 1 loader helper + module docstring). Hard rules: feedback_no_cpu_compute_strict, feedback_no_atomicadd, feedback_no_htod_htoh_only_mapped_pinned, feedback_no_stubs, feedback_no_partial_refactor (no consumer migration — Layer A additive only), feedback_adaptive_not_tuned (α=0.01 is Invariant 1 structural anchor not a tuned constant). + SP5 Task A5 fix-up — pearl_5_iqn_tau header comment corrected (2026-05-01): code-quality review caught that the kernel header described the τ shift as "toward the dense tail" but the actual math (`t = default + skew × 0.05`) shifts τ in the SAME direction as the skew sign — i.e. toward the LONG (sparse) tail, not the dense region. The math is correct and faithful to the plan's formula; only the prose was reversed. Replaced the misleading one-liner with a 7-line explanation: left-skewed Q (skew<0) shifts τ toward 0, right-skewed (skew>0) shifts τ toward 1, "leaning the quantile grid INTO the asymmetry direction" so the IQN sampler gets more resolution over the long tail (where the symmetric 5-tuple under-represents tail mass). Comment-only change; behavior unchanged. Touched: `crates/ml/src/cuda_pipeline/pearl_5_iqn_tau_kernel.cu` (header rewrite, +7/-1 LOC). Cubin rebuild clean. SP5 Task A5 — Pearl 5 per-branch IQN τ schedule GPU producer (Layer A, 2026-05-01): two new CUDA kernels (`q_skew_kurtosis_kernel.cu`, `pearl_5_iqn_tau_kernel.cu`) land as Layer A additive producers that feed ISV slots [250..270) with per-branch IQN quantile-τ schedules. No IQN τ-selection consumer migration in this commit — Layer B wires the τ consumer when distributional action-selection is migrated. `q_skew_kurtosis_update`: single-block 4-thread kernel (one thread per branch), reads `save_q_online[B × 13]` row-major; two-pass central moments (mean then m2/m3/m4); computes `skew = m3 / (m2^1.5 + EPS_DIV)` and `ex_kurt = m4 / (m2^2 + EPS_DIV) - 3`; EPS_DIV=1e-12 is an Invariant 1 anchor (numerical stability for degenerate constant-Q distribution); clamps skew to [-3,+3] and ex_kurt to [-3,+30] (structural envelopes); writes skew[4] to scratch[171..175) and ex_kurt[4] to scratch[175..179); `__threadfence_system()` after writes. No atomicAdd — one thread per branch with independent scratch writes. `pearl_5_iqn_tau_update`: single-block 4-thread kernel, reads scratch_buf[skew_idx_base + b]; shifts symmetric default schedule {0.05,0.25,0.5,0.75,0.95} by `skew × SKEW_SHIFT=0.05` (Invariant 1 anchor — modest schedule shift, max ±0.15 at clamped skew=±3); clamps every τ to structural envelope [0.01,0.99] (Invariant 1 anchor — IQN requires τ strictly in (0,1)); writes 20 τ floats to scratch_out[tau_idx_base + b×5 + q] at scratch[179..199). `launch_sp5_pearl_5_iqn_tau` fires both kernels sequentially then 20 `launch_apply_pearls` calls (4 branches × 5 quantiles → ISV[IQN_TAU_BASE=250..270)). ALPHA_META=1e-3 (same as SP4 default; Pearl 5's τ signal is direct per-step output, not an EMA-of-EMA, so slower smoothing relative to Pearl 4's 5e-4 not needed). Wiener offset formula: `base_wiener_offset + (isv_slot - SP5_SLOT_BASE) × 3` where `base_wiener_offset = SP4_PRODUCER_COUNT × 3 = 213`; for ISV slot 250 (first IQN_TAU slot): wiener_off = 213 + (250-174)×3 = 441. Buffer growth: `producer_step_scratch_buf` 171→199 slots (SP5_SCRATCH_TOTAL); 3 new scratch constants `SCRATCH_PEARL_5_SKEW=171`, `SCRATCH_PEARL_5_EX_KURT=175`, `SCRATCH_PEARL_5_TAU=179`. wiener_state_buf stays at 543 (SP5_WIENER_TOTAL_FLOATS — sized at Task A1 for all 110 SP5 slots, no growth). StateResetRegistry gains 1 new FoldReset entry: `sp5_iqn_tau` (ISV[250..270)) — Pearl A sentinel 0 at fold boundary triggers first-observation replacement to zero-skew symmetric default {0.05,0.25,0.5,0.75,0.95}. Wire-up in training_loop.rs: `launch_sp5_pearl_5_iqn_tau()` called immediately after `launch_sp5_pearl_4_adam_hparams()` with `tracing::warn!(error = %e, ...)` on error. Two GPU-only `#[ignore = "requires GPU"]` unit tests: (10) constant Q input → zero skew → symmetric default τ (chains both kernels); (11) injected skew=-1.0 → τ shifted down by 0.05, τ[0] clamped to 0.01 floor (exercises SKEW_SHIFT mechanism and structural envelope clamp). No CPU compute, no HtoD/HtoH, no atomicAdd, no stubs, no consumer migration. Touched: `cuda_pipeline/q_skew_kurtosis_kernel.cu` (new), `cuda_pipeline/pearl_5_iqn_tau_kernel.cu` (new), `build.rs` (+2 cubin entries), `cuda_pipeline/gpu_dqn_trainer.rs` (SP5_SCRATCH_TOTAL 171→199 + 3 new scratch constants + 2 static cubins + 2 struct fields + cubin loaders + struct initializer + launch method), `trainers/dqn/state_reset_registry.rs` (+1 FoldReset entry), `trainers/dqn/trainer/training_loop.rs` (+5 LOC after Pearl 4 launch), `tests/sp5_producer_unit_tests.rs` (+2 GPU-only tests + 2 cubin constants + 2 loader helpers + module docstring). Hard rules: feedback_no_cpu_compute_strict, feedback_no_atomicadd, feedback_no_htod_htoh_only_mapped_pinned, feedback_no_stubs, feedback_no_partial_refactor (no consumer migration — Layer A additive only).