From f40ccc16a7ed88bcf044ceae7dd627ee0e0983e0 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 1 May 2026 23:31:37 +0200 Subject: [PATCH] =?UTF-8?q?fix(sp5):=20Task=20A6=20=E2=80=94=20close=20two?= =?UTF-8?q?=20minor=20review=20findings?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Combined spec/quality review caught two minor issues in the Pearl 6 commit. Both are mechanical fixes; no behavior change. 1. Test 12 (pearl_6_kelly_within_fold_ewma_blend) was missing an explicit assertion for slot 282 (TRADE_VAR_SMOOTH_IDX). The test setup initialized tvar_i32 and the launcher passed it through the kernel's parameter slot, but no assert! ever fired against the resulting ISV value. With n_envs=1, the kernel's `(kelly_count > 1) ? variance : 0.0f` branch returns 0 (no cross-env variance possible with 1 env), so EWMA blend yields 0.99 × 0.5 + 0.01 × 0.0 = 0.495. Added the missing assertion to close the within-fold coverage gap for s==2. 2. pearl_6_kelly_kernel.cu:136 doc comment said the slot computes "standard deviation of per-env Kelly fractions" but the code actually computes `ksum_sq / kelly_count` — i.e. the variance (second moment), not the standard deviation. The slot name TRADE_VAR_SMOOTH_INDEX correctly indicates variance; the comment was wrong. Updated comment to match: 'variance of per-env Kelly fractions' with explicit note that this is the second moment, NOT std-dev (no sqrtf applied). Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/src/cuda_pipeline/pearl_6_kelly_kernel.cu | 6 ++++-- crates/ml/tests/sp5_producer_unit_tests.rs | 9 +++++++++ docs/dqn-wire-up-audit.md | 2 ++ 3 files changed, 15 insertions(+), 2 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/pearl_6_kelly_kernel.cu b/crates/ml/src/cuda_pipeline/pearl_6_kelly_kernel.cu index ed7cdbcfe..e00add4a0 100644 --- a/crates/ml/src/cuda_pipeline/pearl_6_kelly_kernel.cu +++ b/crates/ml/src/cuda_pipeline/pearl_6_kelly_kernel.cu @@ -133,8 +133,10 @@ extern "C" __global__ void pearl_6_kelly_update( 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). + /* trade_var_smooth: variance of per-env Kelly fractions across envs + * (measures cross-env PnL dispersion). Computed as the second moment + * `ksum_sq / kelly_count` — i.e. variance, NOT standard deviation; + * matches the slot name TRADE_VAR_SMOOTH_INDEX. * Second pass over envs: same Bayesian-prior formula, accumulate sq-dev. */ float kelly_mean = (kelly_count > 0) ? (kelly_sum / (float)kelly_count) diff --git a/crates/ml/tests/sp5_producer_unit_tests.rs b/crates/ml/tests/sp5_producer_unit_tests.rs index 1be9c34bb..a56b90462 100644 --- a/crates/ml/tests/sp5_producer_unit_tests.rs +++ b/crates/ml/tests/sp5_producer_unit_tests.rs @@ -1669,6 +1669,15 @@ fn pearl_6_kelly_within_fold_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]); + + // slot 282 — trade_var_smooth: variance of per-env Kelly fractions across envs. + // With n_envs=1, kelly_count=1, the kernel's `(kelly_count > 1) ? var : 0.0f` + // branch returns 0.0 (single-env can't have cross-env variance). + // EWMA blend: 0.99 × 0.5 + 0.01 × 0.0 = 0.495. + let expected_tvar = 0.99 * 0.5 + 0.01 * 0.0_f32; + assert!((isv[TRADE_VAR_SMOOTH_IDX] - expected_tvar).abs() < 2e-5, + "trade_var: expected {expected_tvar:.6} (n_envs=1 → variance=0; EWMA blend with prior 0.5), got {:.6}", + isv[TRADE_VAR_SMOOTH_IDX]); } /// SP5 Task A6 Test 13: cross-fold sample_count persists via max(). diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 1fa9165aa..5e46aa7ce 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 fix-up — close two minor review findings (2026-05-01): combined spec-quality review caught two minor issues in the Pearl 6 commit. (1) Test 12 (`pearl_6_kelly_within_fold_ewma_blend`) was missing an explicit assertion for slot 282 (TRADE_VAR_SMOOTH_IDX), even though the test setup initialized `tvar_i32` and the launcher passed it. With n_envs=1, the kernel's `(kelly_count > 1) ? variance : 0.0f` branch returns 0, so EWMA blend yields `0.99 × 0.5 + 0.01 × 0.0 = 0.495`. Added `assert!((isv[TRADE_VAR_SMOOTH_IDX] - 0.495).abs() < 2e-5)` to close the within-fold test coverage gap for s==2. (2) `pearl_6_kelly_kernel.cu:136` doc comment said the slot is "standard deviation of per-env Kelly fractions" but the code computes `ksum_sq / kelly_count` (variance), and the slot is named `TRADE_VAR_SMOOTH_INDEX` — the comment was wrong, the name and code are right. Updated comment to "variance of per-env Kelly fractions" and added clarifying note that this is the second moment, NOT std-dev. Comment-only and test-only changes; no behavior change. cargo build --release + cargo test --no-run both clean. + 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.