From 316db416bb31799cd2f73be23c0c4f5e7c1e11f7 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 8 May 2026 08:15:53 +0200 Subject: [PATCH] =?UTF-8?q?fix(class-a-p0c):=20MIN=5FHOLD=5FTARGET=20?= =?UTF-8?q?=E2=86=92=20ISV[AVG=5FWIN=5FHOLD=5FTIME=5FBARS=5FINDEX=3D451]?= =?UTF-8?q?=20(adaptive)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Per Class A audit: MIN_HOLD_TARGET=30.0f hardcoded was creating a deterministic gradient pushing trades toward 30-bar holds regardless of edge expiry. User's trading frequency is between HFT-MFT and varies by regime; a 30-bar fixed target kills MFT-frequency alpha when the optimal hold for current data is shorter (or longer). The producer slot ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451] already exists from SP14 Layer C Phase C.4b (commit 3b71d2183) — Pearl-A-bootstrapped Welford EMA of observed winning trade hold times. Wiring fix only. Cold-start fallback: when slot still at sentinel (no winning trades observed yet), use MIN_HOLD_TARGET=30.0f as safety floor. Once a winning trade closes and the EMA bootstraps, the adaptive value takes over. Validity window: isv_hold_target > 0.0f && < 240.0f; outside window falls back to min_hold_target param (= MIN_HOLD_TARGET=30). Added #define AVG_WIN_HOLD_TIME_BARS_INDEX 451 to state_layout.cuh (C-side mirror of sp14_isv_slots.rs:97). Per feedback_isv_for_adaptive_bounds: every adaptive bound in ISV. Fixes the third Class A P0 hardcoded constant. Co-Authored-By: Claude Opus 4.7 (1M context) --- .../src/cuda_pipeline/experience_kernels.cu | 51 ++++++++++++++----- .../cuda_pipeline/gpu_experience_collector.rs | 27 ++++++---- crates/ml/src/cuda_pipeline/state_layout.cuh | 22 ++++++-- docs/dqn-wire-up-audit.md | 40 +++++++++++++++ 4 files changed, 113 insertions(+), 27 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 8000f1d87..1bbe48c43 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -2032,11 +2032,16 @@ extern "C" __global__ void experience_env_step( float* __restrict__ popart_component_per_sample, /* SP12 v3 (2026-05-04): per-trade event-driven reward composition. * - * `min_hold_target` (bars): patience requirement on voluntary trade - * exits. Phase 1 default 30 bars (Invariant-1 numerical anchor in - * state_layout.cuh::MIN_HOLD_TARGET). Voluntary exits with - * `hold_time < min_hold_target` get a soft penalty subtracted from - * `r_popart` via the deficit/(deficit + T) saturation formula. + * `min_hold_target` (bars): cold-start fallback for the patience + * requirement on voluntary trade exits. Class A P0-C (2026-05-08): + * the actual target used inside the kernel is ISV-driven from + * ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451] — EMA of winning-trade hold + * times. This parameter is the fallback when ISV[451] is still at + * sentinel (0.0 = no winning trades observed yet) or when + * isv_signals_ptr is NULL. Default 30 bars (MIN_HOLD_TARGET in + * state_layout.cuh). Voluntary exits with effective hold_time below + * the target get a soft penalty subtracted from `r_popart` via the + * deficit/(deficit + T) saturation formula. * * `min_hold_penalty_max`: max penalty magnitude when the deficit is * infinite (factor → 1). Phase 1 default 3.0 (60% of REWARD_POS_CAP). @@ -3111,7 +3116,7 @@ extern "C" __global__ void experience_env_step( r_trail = capped_pnl; } else { /* SP12 v3 Change 2: min-hold soft penalty for voluntary exits. - * Encourages MFT-leaning trade duration via patience requirement. + * Encourages patience-matched trade duration via a soft penalty. * Soft saturation formula: factor = deficit/(deficit + T) — bounded * [0,1], smooth transition (no cliff), monotone increasing in * deficit, monotone decreasing in T. Temperature anneals over @@ -3120,9 +3125,9 @@ extern "C" __global__ void experience_env_step( * enforces commitment to longer holds. * * Trail-fire exits exempted (the surrounding `else` already - * guarantees `!trail_triggered`). When `min_hold_target <= 0` - * (defensive guard against malformed configs) the entire block - * collapses to no-op since `deficit = max(0, -hold) <= 0`. + * guarantees `!trail_triggered`). When `effective_min_hold_target + * <= 0` (defensive guard against malformed configs) the entire + * block collapses to no-op since `deficit = max(0, -hold) <= 0`. * * Voluntary-exit only by construction: this lives in the * `else` arm of the `trail_triggered` branch, which itself is @@ -3138,12 +3143,32 @@ extern "C" __global__ void experience_env_step( * (oracle tests in `tests/sp12_reward_math_tests.rs`). The * device function returns the penalty MAGNITUDE (positive) and * subsumes the `hold_time >= target` early-exit; preserving the - * outer `if (segment_hold_time < min_hold_target)` guard would - * be redundant since the function already returns 0 in that - * case and `capped_pnl - shaping_scale * 0 == capped_pnl`. */ + * outer `if (segment_hold_time < target)` guard would be + * redundant since the function already returns 0 in that case + * and `capped_pnl - shaping_scale * 0 == capped_pnl`. + * + * Class A P0-C (2026-05-08): target is ISV-driven from + * ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451] — Pearl-A-bootstrapped + * Welford EMA of observed winning trade hold times. Adapts to + * whatever duration the model actually finds profitable in the + * current regime rather than a hardcoded MFT-leaning constant. + * Cold-start fallback: sentinel = 0.0 (no winning trades yet); + * guarded by the `> 0 && < 240` window below. Outside that + * window (cold-start or anomalous ISV) falls back to + * `min_hold_target` (= MIN_HOLD_TARGET=30 from Rust config). + * Per feedback_isv_for_adaptive_bounds. */ + const float isv_hold_target = (isv_signals_ptr != NULL) + ? isv_signals_ptr[AVG_WIN_HOLD_TIME_BARS_INDEX] + : 0.0f; + /* Validity window: > 0.0 excludes sentinel; < 240.0 caps at + * ~1 hour of bars (sanity guard against ISV corruption). */ + const float effective_min_hold_target = + (isv_hold_target > 0.0f && isv_hold_target < 240.0f) + ? isv_hold_target + : min_hold_target; /* cold-start / NULL fallback */ const float min_hold_penalty = compute_min_hold_penalty( segment_hold_time, - min_hold_target, + effective_min_hold_target, min_hold_temperature, min_hold_penalty_max ); diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index f75d526d8..b9ca7a8e1 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -299,10 +299,13 @@ pub struct ExperienceCollectorConfig { /// SP12 v3 (2026-05-04): per-trade event-driven reward composition. /// - /// Patience target on voluntary trade exits (in bars). Voluntary exits - /// with `hold_time < min_hold_target` get a soft penalty subtracted - /// from `r_popart` via the deficit/(deficit + T) saturation formula. - /// Phase 1 default = `MIN_HOLD_TARGET=30` (state_layout.cuh). + /// Cold-start fallback for the patience target on voluntary trade exits + /// (in bars). Class A P0-C (2026-05-08): the kernel uses + /// `ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451]` (EMA of winning-trade hold + /// times) as the effective target; this field is only used when ISV[451] + /// is still at sentinel (0.0 = no winning trades observed yet) or when + /// `isv_signals_ptr` is NULL. Default = `MIN_HOLD_TARGET=30` + /// (state_layout.cuh cold-start fallback). /// Trail-fire forced exits exempted — see kernel docstring. pub min_hold_target: f32, /// SP12 v3: max penalty magnitude when min-hold deficit is infinite @@ -5645,12 +5648,16 @@ impl GpuExperienceCollector { // SP12 v3 (2026-05-04): per-trade event-driven reward // composition parameters. The kernel applies a soft // min-hold penalty on voluntary exits (deficit/(deficit+T) - // saturation) using these three scalars; constants live - // in state_layout.cuh (MIN_HOLD_TARGET, MIN_HOLD_PENALTY_MAX, - // MIN_HOLD_TEMPERATURE_{START,END,DECAY}). The temperature - // is recomputed in Rust per epoch via the annealing - // schedule before this launch (see - // `training_loop.rs::min_hold_temperature_for_epoch`). + // saturation). Class A P0-C (2026-05-08): the effective + // min-hold target inside the kernel is ISV-driven from + // ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451]; `min_hold_target` + // here is the cold-start fallback only (used when ISV[451] + // is still at sentinel 0.0 or isv_signals_ptr is NULL). + // Constants in state_layout.cuh (MIN_HOLD_TARGET, + // MIN_HOLD_PENALTY_MAX, MIN_HOLD_TEMPERATURE_{START,END, + // DECAY}). Temperature recomputed per epoch via the + // annealing schedule in `training_loop.rs:: + // min_hold_temperature_for_epoch`. .arg(&config.min_hold_target) .arg(&config.min_hold_penalty_max) .arg(&config.min_hold_temperature) diff --git a/crates/ml/src/cuda_pipeline/state_layout.cuh b/crates/ml/src/cuda_pipeline/state_layout.cuh index 67986ba4e..f19c7ad2a 100644 --- a/crates/ml/src/cuda_pipeline/state_layout.cuh +++ b/crates/ml/src/cuda_pipeline/state_layout.cuh @@ -168,6 +168,17 @@ #define ISV_SEED_FRAC_EMA_IDX 84 // == SEED_FRAC_EMA_INDEX — adaptive EMA of (1 - done/target) ∈ [0, 1] (Plan 3 Task 8 B.3; consumed by Task 9 CQL ramp) #define ISV_EVAL_THOMPSON_TEMP_IDX 339 // == EVAL_THOMPSON_TEMP_INDEX — eval Thompson selector temperature (SP10 / Fix 38 2026-05-03; ISV-driven temperature blend on direction-branch Thompson sample) +// ──────────────────────────────────────────────────────────────────────────── +// === SP14-C.4b SLOT INDICES === (2026-05-05, aux trunk adaptive horizon) +// Mirror of crates/ml/src/cuda_pipeline/sp14_isv_slots.rs slot 451. +// Pearl-A-bootstrapped Welford EMA of observed winning trade hold times (bars). +// Sentinel = 0.0 (SENTINEL_AVG_WIN_HOLD_TIME_BARS). First winning trade close +// replaces sentinel directly (first-observation bootstrap). Consumer here +// (Class A P0-C): replaces MIN_HOLD_TARGET hardcoded constant in +// `compute_min_hold_penalty` call site — see experience_kernels.cu. +// ──────────────────────────────────────────────────────────────────────────── +#define AVG_WIN_HOLD_TIME_BARS_INDEX 451 // == AVG_WIN_HOLD_TIME_BARS_INDEX — EMA of winning-trade hold duration (bars); Pearl-A bootstrap sentinel 0.0 + // ──────────────────────────────────────────────────────────────────────────── // === SP13 SLOT INDICES === (2026-05-04, redefine success for predictive skill) // Mirror of crates/ml/src/cuda_pipeline/sp13_isv_slots.rs. Slots [372..383) @@ -229,10 +240,13 @@ // experience_kernels.cu line ~2758 (segment_complete vol-normalized P&L). // Per pearl_audit_unboundedness_for_implicit_asymmetry. // -// MIN_HOLD_TARGET — patience requirement on voluntary trade exits, in bars. -// 30 bars = MFT-leaning but not pure MFT (full MFT ≈ 100 bars, -// HFT-tolerant ≈ 10 bars). Matches the user-profile target band -// "between HFT and MFT". +// MIN_HOLD_TARGET — cold-start fallback for the patience requirement on +// voluntary trade exits, in bars. Used when ISV[AVG_WIN_HOLD_TIME_BARS_ +// INDEX=451] is still at sentinel (0.0 = no winning trades observed yet). +// Once the first winning trade closes and the EMA bootstraps, the +// adaptive ISV value takes over in experience_kernels.cu. Value 30 bars +// = MFT-leaning fallback ("between HFT and MFT" user-profile target band). +// Class A P0-C (2026-05-08): per feedback_isv_for_adaptive_bounds. // // MIN_HOLD_PENALTY_MAX — max penalty magnitude when voluntary exit at hold=0. // 3.0 = 60% of REWARD_POS_CAP. Meaningful gradient pressure but not diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 45d02125d..09a36d4b7 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -7770,3 +7770,43 @@ EMA blend logic directly in the kernel body and writes to ISV in-place. Fixed - C.7-C.8: ISV-driven aux trunk Adam β1/β2/ε/LR/grad-clip. - C.9: synthetic-data smoke + audit close-out + memory pearls. - C.10: L40S 30-epoch validation. + +--- + +## 2026-05-08 — Class A P0-C: MIN_HOLD_TARGET ISV-driven from AVG_WIN_HOLD_TIME_BARS_INDEX[451] + +### Problem + +`MIN_HOLD_TARGET=30.0f` (state_layout.cuh:254) was a hardcoded MFT-leaning constant applied to every voluntary exit's min-hold soft penalty. Creates a deterministic gradient pushing all trades toward 30-bar holds regardless of actual edge expiry or current regime. A profitable 12-bar scalp can be net-negative after penalty — kills MFT-frequency alpha. User's trading frequency is between HFT-MFT and varies by market regime. + +### Fix + +Wired `ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451]` — Pearl-A-bootstrapped Welford EMA of observed winning trade hold times — as the effective target in `compute_min_hold_penalty`. Added `#define AVG_WIN_HOLD_TIME_BARS_INDEX 451` to state_layout.cuh (C-side mirror of sp14_isv_slots.rs:97). The Rust `min_hold_target` parameter (default 30.0) becomes cold-start fallback only. + +### Cold-start behavior + +- Sentinel = `0.0f` (SENTINEL_AVG_WIN_HOLD_TIME_BARS in sp14_isv_slots.rs). +- Fallback condition: `isv_hold_target > 0.0f && isv_hold_target < 240.0f`. Outside window (or NULL ptr): uses `min_hold_target` = 30.0 from config. +- First winning trade close: Pearl-A bootstrap replaces sentinel with observation directly; adaptive path takes over immediately. +- Upper bound 240 bars = sanity cap (>= ~1hr of bars, protects against ISV corruption). + +### Sites modified + +| File | Lines | Change | +|------|-------|--------| +| `state_layout.cuh` | ~169 (add) | `#define AVG_WIN_HOLD_TIME_BARS_INDEX 451` block | +| `state_layout.cuh` | ~232-239 | MIN_HOLD_TARGET comment → cold-start fallback role | +| `experience_kernels.cu` | ~2034-2044 | `min_hold_target` param doc → cold-start fallback | +| `experience_kernels.cu` | ~3113-3165 | ISV read + fallback + pass `effective_min_hold_target` to `compute_min_hold_penalty` | +| `gpu_experience_collector.rs` | ~300-310 | struct field doc update | +| `gpu_experience_collector.rs` | ~5645-5658 | launch-site comment update | +| `docs/dqn-wire-up-audit.md` | here | this entry | + +### Verification + +- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --lib` — clean (19 pre-existing warnings, 0 new). +- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests --release -- --ignored --nocapture` — 4/4 pass. + +### Consumer audit + +Only one consumer of `MIN_HOLD_TARGET` was found in the CUDA call path (the `compute_min_hold_penalty` call in `experience_kernels.cu`). No additional consumers discovered. `trade_physics.cuh` takes `min_hold_target` as a parameter — no constant reads — no change needed there. Per `feedback_no_partial_refactor`: complete.