From 7e9a8f6ef170fff4d66217dc26f33a8c2cb649de Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 8 May 2026 10:41:07 +0200 Subject: [PATCH] fix(class-a-audit-batch-4a): DD saturation floor adaptive + legacy DD path Case A MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Per Class A audit-fix Batch 4-A (deferred from P1-wiring/P1-producer due to audit-doc errors). Fixes 2 of 4 deferred items; Batch 4-B handles plan_threshold floor + MIN_HOLD_TEMPERATURE in a separate commit. Item 1: DD saturation floor (the upper end of the DD ramp at trade_physics.cuh:154 in apply_margin_cap, NOT line 548 as the audit doc claimed — that line is a magnitude action constant; the actual saturation floor lives in apply_margin_cap) - NEW slot DD_SATURATION_FLOOR_ADAPTIVE_INDEX=458 - Producer dd_saturation_floor_update_kernel.cu — p75(per-env DD_MAX) × 1.5 via Welford `mean + Z_75 × sigma` estimator with `max(p75, mean)` robustness guard, mirrors P0-A REWARD_POS_CAP producer pattern (Pearl-A bootstrap + Welford α=0.01) - Cold-start fallback: 0.25f (DD_SATURATION_FLOOR_DEFAULT in state_layout.cuh) - Bounds: [0.10, 0.50] (Category-1 dimensional safety) - Distinct from SP15_DD_THRESHOLD_INDEX=421 (the SP15 quadratic DD-penalty *trigger* threshold, a *lower* bound; this slot is the *upper* end of the linear position-size scaling ramp dd_scale = max(0.05, 1.0 − dd_frac/floor)) - Threaded `isv_signals_ptr` into `apply_margin_cap` with NULL-tolerant cold-start fallback to DD_SATURATION_FLOOR_DEFAULT - 4 oracle tests (Pearl-A bootstrap, no-DD guard, bounds clamp, Welford EMA) Item 2: Legacy compute_drawdown_penalty path → Case A (DELETED) - Decision rationale: SP15's quadratic asymmetric DD penalty (compute_sp15_final_reward_kernel.cu:154 via sp15_dd_penalty helper) runs unconditionally as a post-modifier on the SP11-composed reward with ISV-driven λ_dd (slot 420) and DD threshold (slot 421). Layering the legacy linear-ramp penalty inside the SP11 composer on top of the SP15 quadratic creates double-counting of DD shaping — exactly the code-smell the Class A audit was designed to eliminate. Per `feedback_no_legacy_aliases.md` and `feedback_no_partial_refactor.md`. - Atomic deletion across: - `compute_drawdown_penalty` device function (trade_physics.cuh) - Single call site at experience_kernels.cu:3822 - `dd_threshold` and `w_dd` kernel arguments - `w_dd` Rust config field (gpu_experience_collector.rs + trainers/dqn/config.rs DQNHyperparameters) - `w_dd` profile section + dispatch (training_profile.rs RewardSection, OptimizableParameterRanges, FixedRewardParameters, ParamLookup dispatch, profile→hyperparam mapping, test assertion) - `w_dd *= rki` risk-intensity multiplier (config.rs) - `w_dd` TOML keys (dqn-hyperopt.toml × 2, dqn-localdev.toml, dqn-production.toml, dqn-smoketest.toml) - Stale doc comments on hyperopt/adapters/dqn.rs + config.rs risk_intensity field - `config.dd_threshold` SURVIVES (still consumed by `launch_sp15_dd_state` as the dd_budget for DD_PCT scaling). Documented in field comment. ISV_TOTAL_DIM: 458 → 459 (Item 1 adds 1 slot; Item 2 is pure deletion) Cumulative WR-plateau fix series (this is commit 7): - Class C bug 1 + P0-B (8f218cab2) - P0-C (316db416b) - P0-A (394de7d43) - P1 wiring (c4b6d6ef2) — 1 of 4 wireable - P0-A downstream (657972a4b) - P1 producer (87d597d5d) - audit-fix 4-A (this commit) Verification: 16 sp14_oracle_tests pass (incl. 4 new), 36 sp15_phase1_oracle_tests pass, 12 sp14_isv_slots layout tests pass, 4 state_reset_registry tests pass (every-FoldReset-arm-has-dispatch contract holds), workspace cargo check clean. Per feedback_isv_for_adaptive_bounds + feedback_no_partial_refactor + feedback_no_legacy_aliases. Co-Authored-By: Claude Opus 4.7 (1M context) --- config/training/dqn-hyperopt.toml | 8 +- config/training/dqn-localdev.toml | 2 +- config/training/dqn-production.toml | 2 +- config/training/dqn-smoketest.toml | 2 +- crates/ml/build.rs | 29 ++ .../dd_saturation_floor_update_kernel.cu | 243 +++++++++++++++ .../src/cuda_pipeline/experience_kernels.cu | 54 ++-- crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs | 112 ++++++- .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 99 +++++- .../cuda_pipeline/gpu_experience_collector.rs | 37 ++- crates/ml/src/cuda_pipeline/sp14_isv_slots.rs | 92 ++++++ crates/ml/src/cuda_pipeline/state_layout.cuh | 28 ++ crates/ml/src/cuda_pipeline/trade_physics.cuh | 84 +++-- crates/ml/src/hyperopt/adapters/dqn.rs | 4 +- crates/ml/src/trainers/dqn/config.rs | 19 +- .../src/trainers/dqn/state_reset_registry.rs | 27 ++ .../src/trainers/dqn/trainer/training_loop.rs | 44 ++- crates/ml/src/training_profile.rs | 30 +- crates/ml/tests/sp14_oracle_tests.rs | 293 ++++++++++++++++++ docs/dqn-wire-up-audit.md | 65 ++++ docs/isv-slots.md | 10 + 21 files changed, 1191 insertions(+), 93 deletions(-) create mode 100644 crates/ml/src/cuda_pipeline/dd_saturation_floor_update_kernel.cu diff --git a/config/training/dqn-hyperopt.toml b/config/training/dqn-hyperopt.toml index bd75cdd6e..51735bb11 100644 --- a/config/training/dqn-hyperopt.toml +++ b/config/training/dqn-hyperopt.toml @@ -75,8 +75,10 @@ c51_alpha_max = [0.3, 0.9] her_ratio = [0.0, 0.8] # Composite reward weights +# `w_dd` removed (Class A audit-fix Batch 4-A, 2026-05-08) atomically with the +# legacy compute_drawdown_penalty path. SP15's λ_dd (ISV slot 420) is the +# producer-driven replacement and is not exposed as a hyperopt search dim. w_pnl = [0.0, 1.0] -w_dd = [0.0, 5.0] w_idle = [0.0, 0.1] dd_threshold = [0.005, 0.03] # HFT: tight drawdown tolerance (0.5%-3%) loss_aversion = [1.0, 1.0] @@ -121,7 +123,7 @@ max_trace_length = 7 hindsight_fraction = 0.1 hindsight_lookahead = 10 w_pnl = 0.3 -w_dd = 1.0 +# `w_dd` removed (Class A audit-fix Batch 4-A, 2026-05-08) atomically with legacy compute_drawdown_penalty. w_idle = 0.01 dd_threshold = 0.01 loss_aversion = 1.0 @@ -155,7 +157,7 @@ td_lambda = 0.9 hindsight_fraction = 0.1 hindsight_lookahead = 10 w_pnl = 0.3 -w_dd = 1.0 +# `w_dd` removed (Class A audit-fix Batch 4-A, 2026-05-08) atomically with legacy compute_drawdown_penalty. w_idle = 0.01 dd_threshold = 0.01 loss_aversion = 1.0 diff --git a/config/training/dqn-localdev.toml b/config/training/dqn-localdev.toml index ea1c322e5..151db8ebf 100644 --- a/config/training/dqn-localdev.toml +++ b/config/training/dqn-localdev.toml @@ -122,7 +122,7 @@ hindsight_fraction = 0.1 hindsight_lookahead = 10 loss_aversion = 1.0 q_gap_threshold = 0.1 -w_dd = 1.0 +# `w_dd` removed (Class A audit-fix Batch 4-A, 2026-05-08) atomically with legacy compute_drawdown_penalty. dd_threshold = 0.02 cea_weight = 0.3 b3_size = 3 diff --git a/config/training/dqn-production.toml b/config/training/dqn-production.toml index f7c821b6e..1a81fd2d2 100644 --- a/config/training/dqn-production.toml +++ b/config/training/dqn-production.toml @@ -149,7 +149,7 @@ hindsight_fraction = 0.1 hindsight_lookahead = 10 loss_aversion = 1.0 q_gap_threshold = 0.1 -w_dd = 1.0 +# `w_dd` removed (Class A audit-fix Batch 4-A, 2026-05-08) atomically with legacy compute_drawdown_penalty. dd_threshold = 0.02 cea_weight = 0.3 # Dense micro-reward component weights (tunable via hyperopt) diff --git a/config/training/dqn-smoketest.toml b/config/training/dqn-smoketest.toml index 49b183e11..daeebd395 100644 --- a/config/training/dqn-smoketest.toml +++ b/config/training/dqn-smoketest.toml @@ -140,7 +140,7 @@ hindsight_fraction = 0.1 hindsight_lookahead = 10 loss_aversion = 1.0 q_gap_threshold = 0.1 -w_dd = 1.0 +# `w_dd` removed (Class A audit-fix Batch 4-A, 2026-05-08) atomically with legacy compute_drawdown_penalty. dd_threshold = 0.02 cea_weight = 0.3 b3_size = 3 diff --git a/crates/ml/build.rs b/crates/ml/build.rs index c40e94105..8dcabbe8d 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -827,6 +827,35 @@ fn main() { // Block-tree-reduce (no atomicAdd) per // `feedback_no_atomicadd.md`. Per-epoch boundary launch. "kelly_bayesian_priors_update_kernel.cu", + // Class A audit-fix Batch 4-A (2026-05-08): adaptive DD saturation + // floor producer. Single-block 256-thread kernel that sweeps the + // per-env DD trajectory tile `dd_state_per_env[n_envs * 6]` + // (populated by `dd_state_kernel`), aggregates DD_MAX (offset 1) + // across envs via Welford `mean + Z_75 × sigma` p75 estimator + // with `max(p75, mean)` robustness guard × 1.5 safety factor, + // and slow-EMA blends into ISV[DD_SATURATION_FLOOR_ADAPTIVE_INDEX + // =458]. Pearl-A first-observation bootstrap (sentinel 0.25 + // matches pre-fix hardcoded `0.25f` in + // `trade_physics.cuh::apply_margin_cap` for bit-identical cold- + // start); α=0.01 slow EMA thereafter (DD distribution is a slow- + // moving fold-volatility property — same cadence as P0-A + // REWARD_POS_CAP). Bounds [0.10, 0.50] are Category-1 dimensional + // safety floors per `feedback_isv_for_adaptive_bounds`. Replaces + // the hardcoded saturation floor (the upper end of the linear + // position-size scaling ramp `dd_scale = max(0.05, 1.0 − + // dd_frac/saturation_floor)`). Distinct from + // `SP15_DD_THRESHOLD_INDEX=421` (the SP15 quadratic DD-penalty + // *trigger* threshold, a *lower* bound). Atomically deleted + // alongside this slot addition: legacy `compute_drawdown_penalty` + // device function in trade_physics.cuh + its single caller at + // experience_kernels.cu:3822 + the `w_dd` Rust config field + + // the `dd_threshold` and `w_dd` kernel args (Item 2 Case A — + // SP15's quadratic asymmetric DD penalty in + // `compute_sp15_final_reward_kernel.cu:154` is the production- + // grade replacement; layering both creates double-counting). + // Block-tree-reduce (no atomicAdd) per `feedback_no_atomicadd.md`. + // Per-epoch boundary launch. + "dd_saturation_floor_update_kernel.cu", // SP14 Layer C Phase C.6 (2026-05-08): h_s2_aux RMS EMA producer. // Single-block 256-thread kernel computing RMS = sqrt(mean(x²)) // over `h_s2_aux [B, SH2]` (aux trunk final output, no activation) diff --git a/crates/ml/src/cuda_pipeline/dd_saturation_floor_update_kernel.cu b/crates/ml/src/cuda_pipeline/dd_saturation_floor_update_kernel.cu new file mode 100644 index 000000000..79cdc63ce --- /dev/null +++ b/crates/ml/src/cuda_pipeline/dd_saturation_floor_update_kernel.cu @@ -0,0 +1,243 @@ +/* ══════════════════════════════════════════════════════════════════════════ + * Class A audit-fix Batch 4-A — adaptive DD saturation floor producer (2026-05-08). + * + * Replaces the hardcoded `0.25f` saturation floor in + * `trade_physics.cuh::apply_margin_cap` (line ~154) with an ISV-driven + * adaptive bound derived from the realized per-env DD_MAX distribution. + * + * Why this matters (per Class A audit deferred-item batch 4-A): + * + * apply_margin_cap implements a linear position-size scaling ramp: + * + * dd_scale = max(0.05, 1.0 − dd_frac / saturation_floor) + * + * The 0.25f literal sets the upper end of that ramp — the DD level at + * which dd_scale floors out at 5%. With a hardcoded 25% saturation + * floor, the ramp is calibrated for one specific DD distribution. + * Production training across 11 SPs has revealed that DDs vary widely + * (median fold DD_MAX=8% on calm folds, 35% on volatile folds), and + * the static 25% breakpoint either over-throttles (calm folds where + * small DDs unnecessarily clamp positions) or under-throttles + * (volatile folds where 35% DDs hit the 5% floor on most envs). + * + * Adaptive saturation floor: drive the breakpoint from the realized + * per-env DD_MAX distribution at fold/epoch boundary. The floor lands + * at p75(DD_MAX) × safety_factor=1.5, ensuring 25% of envs are still + * well within the linear ramp regardless of fold volatility. + * + * Algorithm (mirrors reward_cap_update_kernel and kelly_bayesian_priors + * patterns): + * + * Phase 1 (block-tree-reduce in shmem): each thread strides over envs, + * reading `dd_state_per_env[env*DD_STATE_STRIDE + DD_STATE_DD_MAX_OFF]` + * (the per-env running max DD ratio in the current fold). Block reduce + * accumulates sum, sum-of-squares, count, and max across all envs. + * + * Phase 2 (Pearl-A + EMA, single thread): if no envs experienced a + * non-zero DD this epoch (all DD_MAX zero), keep ISV slot unchanged — + * the sentinel persists and consumers fall back to the static 0.25f + * default via the NULL-tolerant guard pattern. + * + * p75 estimator (Welford rule of thumb): + * + * p75(DD_MAX) ≈ mean + Z_75 × sigma (Z_75 ≈ 0.6745) + * + * for a roughly-normal positive tail. We do NOT histogram (no + * atomicAdd; would also need full env-tile materialisation in shmem, + * plus a sort/select step). Per `feedback_no_atomicadd.md`. The mean+ + * sigma estimator is conservative: small-N regimes where sigma is + * inflated by a single deep DD env will pull the p75 up, which is the + * safe direction for a saturation floor (better to scale positions + * slowly than aggressively). + * + * target = max(p75_estimate, mean) × safety_factor (1.5×) + * + * The `max(p75, mean)` guard handles the degenerate case of all envs + * having identical DD_MAX (sigma=0 → p75=mean) — picks the larger + * value robustly. + * + * Pearl-A first-observation bootstrap: if `current == sentinel_floor` + * (0.25 within EPS), REPLACE with target directly. Otherwise EMA + * blend with α (slow, 0.01 — DD distribution is a slow-moving + * property of the fold's volatility regime). Per + * `pearl_first_observation_bootstrap.md`. + * + * Bounds: clamped to [DD_SATURATION_FLOOR_MIN=0.10, DD_SATURATION_FLOOR_MAX=0.50] + * (Category-1 dimensional safety per `feedback_isv_for_adaptive_bounds`). + * Below 10% the ramp would clip positions on minor noise; above 50% + * the ramp barely activates before the capital-floor circuit-breaker + * triggers at 25% PDT. + * + * Pearls + invariants: + * + * - `feedback_no_atomicadd.md` — single-block kernel; block-tree-reduce + * in shmem; no atomics across blocks. + * + * - `pearl_no_host_branches_in_captured_graph.md` — `CudaFunction` + * pre-loaded at construction; on-device guards for "no DD" and + * Pearl-A sentinel detection. + * + * - `pearl_first_observation_bootstrap.md` — first valid sample + * replaces sentinel directly; no blend. + * + * - `feedback_no_stubs.md` — full body, no return-zero placeholders. + * + * - `feedback_isv_for_adaptive_bounds.md` — DD floor bounds [0.10, 0.50] + * are dimensional safety, NOT tuning. + * + * - `pearl_symmetric_clamp_audit.md` — bilateral `fmaxf(lo, fminf(x, hi))` + * clamp on target before writing the ISV slot. + * + * Args: + * dd_state_per_env — `[n_envs * dd_state_stride]` device f32; + * per-env DD trajectory tile written by + * `dd_state_kernel`. We read the DD_MAX + * field at offset `dd_max_off`. + * n_envs — number of envs (alloc_episodes from + * the collector). + * dd_state_stride — DD_STATE_STRIDE_PER_ENV (== 6, mirrors + * the per-env tile in `dd_state_kernel`). + * dd_max_off — offset within each env's tile for + * DD_MAX (== 1). + * isv — `[ISV_TOTAL_DIM]` device f32, modified + * in place at index `floor_idx`. + * floor_idx — DD_SATURATION_FLOOR_ADAPTIVE_INDEX (458). + * sentinel_floor — SENTINEL_DD_SATURATION_FLOOR (0.25). + * safety_factor — DD_SATURATION_FLOOR_SAFETY_FACTOR (1.5). + * floor_min — DD_SATURATION_FLOOR_MIN (0.10). + * floor_max — DD_SATURATION_FLOOR_MAX (0.50). + * alpha — DD_SATURATION_FLOOR_EMA_ALPHA (0.01, slow). + * + * Launch: grid=(1, 1, 1), block=(BLK_DIM=256, 1, 1). + * Shared memory: 4 * BLK_DIM floats: + * sh_sum — sum of DD_MAX across envs + * sh_sumsq — sum-of-squares of DD_MAX + * sh_count — count of envs with DD_MAX > 0 + * sh_max — running max of DD_MAX (defensive bound on the estimator) + * The count is summed as float for unified shmem layout — values are + * bounded by n_envs ≤ 2^16 so f32 representation is exact. + * ══════════════════════════════════════════════════════════════════════════ */ + +#include + +#define BLK_DIM 256 +#define EPS_F 1e-6f +/* Z-score for ~p75 of a normal distribution (one-sided). Used as a + * lightweight estimator: p75 ≈ mean + Z_75 × sigma. */ +#define Z_75 0.6745f + +extern "C" __global__ +void dd_saturation_floor_update( + const float* __restrict__ dd_state_per_env, + int n_envs, + int dd_state_stride, + int dd_max_off, + float* __restrict__ isv, + int floor_idx, + float sentinel_floor, + float safety_factor, + float floor_min, + float floor_max, + float alpha) +{ + /* Single-block kernel — only block 0 does anything. */ + if (blockIdx.x != 0) return; + + extern __shared__ float shmem[]; + float* sh_sum = shmem; /* [BLK_DIM] */ + float* sh_sumsq = shmem + BLK_DIM; /* [BLK_DIM] */ + float* sh_count = shmem + 2 * BLK_DIM; /* [BLK_DIM] */ + float* sh_max = shmem + 3 * BLK_DIM; /* [BLK_DIM] */ + + int tid = threadIdx.x; + + /* Phase 1a: stride-based sweep — accumulate stats on per-env DD_MAX + * values. Skip envs with DD_MAX <= 0 (no DD observed yet). */ + float local_sum = 0.0f; + float local_sumsq = 0.0f; + float local_count = 0.0f; + float local_max = 0.0f; + for (int e = tid; e < n_envs; e += BLK_DIM) { + const long long base = (long long)e * dd_state_stride; + float dd_max = dd_state_per_env[base + dd_max_off]; + if (dd_max > 0.0f) { + local_sum += dd_max; + local_sumsq += dd_max * dd_max; + local_count += 1.0f; + if (dd_max > local_max) local_max = dd_max; + } + } + sh_sum[tid] = local_sum; + sh_sumsq[tid] = local_sumsq; + sh_count[tid] = local_count; + sh_max[tid] = local_max; + __syncthreads(); + + /* Phase 1b: block-tree-reduce (no atomicAdd) — halve every iter. + * Sum/sumsq/count fold via addition; max folds via fmaxf. */ + for (int s = BLK_DIM >> 1; s > 0; s >>= 1) { + if (tid < s) { + sh_sum[tid] += sh_sum[tid + s]; + sh_sumsq[tid] += sh_sumsq[tid + s]; + sh_count[tid] += sh_count[tid + s]; + sh_max[tid] = fmaxf(sh_max[tid], sh_max[tid + s]); + } + __syncthreads(); + } + + /* Phase 2: thread 0 finalises EMA + Pearl-A + clamp + writes slot. */ + if (tid != 0) return; + + float total_sum = sh_sum[0]; + float total_sumsq = sh_sumsq[0]; + float total_count = sh_count[0]; + float total_max = sh_max[0]; + + /* Guard: no envs with non-zero DD_MAX — keep ISV slot unchanged + * (cold-start sentinel persists; consumer falls back to the static + * 0.25f default via the sentinel-detection guard pattern). */ + if (total_count <= 0.0f) { + return; + } + + /* Welford-rule p75 estimator: mean + Z_75 × sigma. Sigma is + * computed as `sqrt(max(0, sumsq/n − mean²))` — the inner max + * absorbs floating-point under-flow when all observations are + * (near-)identical. */ + float mean = total_sum / total_count; + float variance = fmaxf(0.0f, total_sumsq / total_count - mean * mean); + float sigma = sqrtf(variance); + float p75_estimate = mean + Z_75 * sigma; + + /* Robust target: max(p75_estimate, mean) — handles the degenerate + * case of zero variance (all envs same DD) where p75=mean exactly, + * and the small-N case where sigma can drift slightly negative + * from f32 round-off. The `total_max` is NOT included here because + * a single deep-DD outlier env should NOT pull the saturation floor + * up — that would over-relax position scaling for the other 99% of + * envs. (Contrast with reward_cap_update where max() takeover IS + * used: there we explicitly want the cap to never clip an observed + * winning return; here we want the typical fold's behavior.) */ + (void)total_max; /* tracked for diagnostics only; not used by estimator */ + float target = fmaxf(p75_estimate, mean) * safety_factor; + + /* Clamp to dimensional-safety bounds [floor_min, floor_max]. + * Bilateral per `pearl_symmetric_clamp_audit.md`. */ + target = fmaxf(floor_min, fminf(target, floor_max)); + + /* Pearl-A first-observation bootstrap. The sentinel is exactly + * SENTINEL_DD_SATURATION_FLOOR=0.25; any value within EPS_F of + * the sentinel is treated as sentinel. */ + float current = isv[floor_idx]; + float blended; + if (fabsf(current - sentinel_floor) < EPS_F) { + blended = target; + } else { + blended = (1.0f - alpha) * current + alpha * target; + } + /* Re-clamp post-blend (defensive — handles malformed prior state + * outside [floor_min, floor_max]). */ + blended = fmaxf(floor_min, fminf(blended, floor_max)); + + isv[floor_idx] = blended; +} diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 791e4e169..89f3294af 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -1903,8 +1903,14 @@ extern "C" __global__ void experience_env_step( float spread_cost, /* bid-ask spread cost per unit (matches backtest) */ float contract_multiplier, /* e.g. 50.0 for ES, 20.0 for NQ */ float margin_pct, /* e.g. 0.06 (6% initial margin) */ - float dd_threshold, /* drawdown fraction before penalty (0.02 = 2%) */ - float w_dd, /* drawdown penalty weight (1.0 = full) */ + /* `dd_threshold` and `w_dd` parameters DELETED (Class A audit-fix + * Batch 4-A, 2026-05-08) atomically with the legacy + * compute_drawdown_penalty call. SP15's `sp15_dd_penalty` in + * compute_sp15_final_reward_kernel.cu reads λ_dd / DD threshold + * directly from ISV slots 420/421 — the kernel-arg path is + * obsolete. The Rust caller arg pass at gpu_experience_collector.rs + * is also dropped. The `dd_threshold` Rust config field SURVIVES + * because it is reused as `dd_budget` by `launch_sp15_dd_state`. */ int mirror_active, /* #10 mirror universe: invert direction actions */ float micro_reward_scale, /* v8: dense directional micro-reward scale */ /* #33 Per-episode saboteur params [N, 3]: (spread_mult, fill_prob, slippage_mult). @@ -3791,35 +3797,23 @@ extern "C" __global__ void experience_env_step( reward = r_weighted; /* the post-composition modifiers below operate on r_weighted */ } - /* ---- Drawdown penalty (from trade_physics.cuh) ────────────────────── - * Phase 3: gated by shaping_scale. The drawdown penalty is *behavioral* - * (steers the policy away from large equity drops); the *physics* of - * drawdown impact is already in the per-bar P&L. At shaping_scale=0 - * (validation mode) we measure pure equity-change reward only. + /* ---- Drawdown penalty: DELETED (Class A audit-fix Batch 4-A, 2026-05-08) ── * - * Class A P0-A-downstream (2026-05-08): the scale magnitude (was the - * hardcoded `-5.0f` inside compute_drawdown_penalty, sized to match - * the pre-P0-A REWARD_POS_CAP=5.0f) now scales proportionally with - * the adaptive REWARD_POS_CAP via ISV[REWARD_POS_CAP_ADAPTIVE_INDEX= - * 452]. The 1:1 ratio (penalty saturates at −1×POS_CAP × w_dd) - * preserves the original tuning relative to the cap. Cold-start - * fallback: sentinel ISV (5.0f within EPS) OR ISV outside - * [REWARD_POS_CAP_MIN_BOUND=1.0, REWARD_POS_CAP_MAX_BOUND=50.0] → - * fall back to the original 5.0f magnitude (bit-identical pre-P0-A - * cold-start behavior). Defensive guard mirrors `sp15_apply_sp12_cap` - * and the segment-complete asymmetric cap. Per - * feedback_isv_for_adaptive_bounds + feedback_no_partial_refactor. */ - float dd_penalty_scale = REWARD_POS_CAP; /* fallback = 5.0f from state_layout.cuh */ - if (isv_signals_ptr != NULL) { - const float isv_pos = isv_signals_ptr[REWARD_POS_CAP_ADAPTIVE_INDEX]; - if (isv_pos >= REWARD_POS_CAP_MIN_BOUND - && isv_pos <= REWARD_POS_CAP_MAX_BOUND - && fabsf(isv_pos - SENTINEL_REWARD_POS_CAP) > 1e-6f) { - dd_penalty_scale = isv_pos; - } - } - float equity_now = cash + position * raw_close; - reward += shaping_scale * compute_drawdown_penalty(equity_now, peak_equity, dd_threshold, w_dd, dd_penalty_scale); + * The legacy `compute_drawdown_penalty` call was deleted atomically + * with the device function in `trade_physics.cuh`, the `dd_threshold` + * and `w_dd` kernel arguments, the `w_dd` Rust config field, and the + * arg pass at `gpu_experience_collector.rs:5569-5570`. + * + * Rationale (Item 2 Case A): SP15's quadratic asymmetric DD penalty + * in `compute_sp15_final_reward_kernel.cu:154` (the `sp15_dd_penalty` + * helper applying `λ_dd × max(0, dd_current − dd_threshold)²`) is + * the production-grade replacement. It runs unconditionally as a + * post-modifier on the SP11-composed reward (gated only on + * `isv_signals_dev_ptr != 0 && sp15_alpha_warm_count_dev_ptr != 0`, + * both populated in production). Layering the legacy linear-ramp + * penalty here on top of the SP15 quadratic creates double-counting + * of DD shaping. Per `feedback_no_legacy_aliases.md` and + * `feedback_no_partial_refactor.md`. */ /* ---- Portfolio value for floor check ---- */ float new_portfolio_value_pre_floor = cash + position * raw_close; diff --git a/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs b/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs index 27de97a95..dd062f75d 100644 --- a/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs +++ b/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs @@ -52,8 +52,8 @@ use crate::MLError; use super::gpu_dqn_trainer::{ AUX_HORIZON_UPDATE_CUBIN, AUX_TRUNK_BACKWARD_CUBIN, AUX_TRUNK_FORWARD_CUBIN, - AVG_WIN_HOLD_TIME_UPDATE_CUBIN, H_S2_AUX_RMS_EMA_CUBIN, KELLY_BAYESIAN_PRIORS_UPDATE_CUBIN, - REWARD_CAP_UPDATE_CUBIN, + AVG_WIN_HOLD_TIME_UPDATE_CUBIN, DD_SATURATION_FLOOR_UPDATE_CUBIN, H_S2_AUX_RMS_EMA_CUBIN, + KELLY_BAYESIAN_PRIORS_UPDATE_CUBIN, REWARD_CAP_UPDATE_CUBIN, }; /// Hidden width of the aux trunk's first internal layer (Linear_1 → ELU @@ -857,6 +857,114 @@ impl KellyBayesianPriorsUpdateOps { } } +/// Class A audit-fix Batch 4-A (2026-05-08): adaptive DD saturation floor producer. +/// +/// Sweeps the per-env `dd_state_per_env[n_envs * 6]` tile (populated by +/// `dd_state_kernel` at the same boundary), aggregates DD_MAX (offset 1 +/// within each env's tile) across envs, computes a Welford `mean + +/// Z_75 × sigma` p75 estimator with a `max(p75, mean)` robustness guard, +/// applies a 1.5× safety factor, clamps to dimensional bounds [0.10, 0.50], +/// and writes `ISV[DD_SATURATION_FLOOR_ADAPTIVE_INDEX=458]`. Replaces the +/// hardcoded `0.25f` saturation floor in `trade_physics.cuh::apply_margin_cap` +/// (the upper end of the linear position-size scaling ramp). +/// +/// Single-block kernel; block-tree-reduce over the per-env tile (no +/// atomicAdd per `feedback_no_atomicadd.md`). Pearl-A first-observation +/// bootstrap (sentinel 0.25 matches pre-fix hardcoded constant for +/// bit-identical cold-start); α=0.01 slow EMA thereafter (DD distribution +/// is a slow-moving fold-volatility property — same cadence as the P0-A +/// REWARD_POS_CAP producer). +/// +/// Per-epoch boundary launch — DD distribution shouldn't move fast. +/// +/// # Pearls applied +/// - `feedback_no_atomicadd.md` — single block, block-tree-reduce in shmem. +/// - `pearl_first_observation_bootstrap.md` — sentinel 0.25 → REPLACE. +/// - `pearl_no_host_branches_in_captured_graph.md` — `CudaFunction` +/// pre-loaded at construction; on-device guards. +/// - `pearl_symmetric_clamp_audit.md` — bilateral `fmaxf(lo, fminf(x, hi))`. +/// - `feedback_isv_for_adaptive_bounds.md` — bounds [0.10, 0.50] are +/// Category-1 dimensional safety, NOT tuning. +#[allow(missing_debug_implementations)] +pub(crate) struct DdSaturationFloorUpdateOps { + update_kernel: CudaFunction, +} + +impl DdSaturationFloorUpdateOps { + /// Block dim used by the producer kernel — must match `BLK_DIM` in + /// `dd_saturation_floor_update_kernel.cu`. + const BLK_DIM: u32 = 256; + + pub(crate) fn new(stream: &Arc) -> Result { + let context = stream.context(); + let module = context + .load_cubin(DD_SATURATION_FLOOR_UPDATE_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("dd_saturation_floor_update cubin load: {e}")))?; + let update_kernel = module + .load_function("dd_saturation_floor_update") + .map_err(|e| MLError::ModelError(format!("dd_saturation_floor_update load: {e}")))?; + Ok(Self { update_kernel }) + } + + /// Launch the adaptive DD saturation floor producer. + /// + /// Args: + /// - `dd_state_per_env_ptr`: f32 device ptr `[n_envs * dd_state_stride]` + /// — the same buffer `compute_sp15_final_reward_kernel` reads from. + /// We use the per-env DD_MAX field at offset `dd_max_off`. + /// - `n_envs`: number of envs (alloc_episodes from the collector). + /// - `dd_state_stride`: 6 (mirrors the per-env tile layout in + /// `dd_state_kernel`). + /// - `dd_max_off`: 1 (DD_MAX offset within each env's tile). + /// - `isv_ptr`: ISV[ISV_TOTAL_DIM] device pointer (mapped-pinned). + /// - `floor_idx`: DD_SATURATION_FLOOR_ADAPTIVE_INDEX (458). + /// - `sentinel_floor`: SENTINEL_DD_SATURATION_FLOOR (0.25). + /// - `safety_factor`: DD_SATURATION_FLOOR_SAFETY_FACTOR (1.5). + /// - `floor_min`/`floor_max`: DD_SATURATION_FLOOR_MIN/MAX (0.10, 0.50). + /// - `alpha`: DD_SATURATION_FLOOR_EMA_ALPHA (0.01, slow per-epoch). + #[allow(clippy::too_many_arguments)] + pub(crate) fn launch( + &self, + stream: &Arc, + dd_state_per_env_ptr: u64, + n_envs: i32, + dd_state_stride: i32, + dd_max_off: i32, + isv_ptr: u64, + floor_idx: i32, + sentinel_floor: f32, + safety_factor: f32, + floor_min: f32, + floor_max: f32, + alpha: f32, + ) -> Result<(), MLError> { + // 4 arrays × BLK_DIM × sizeof(f32) — see the kernel comment block. + let smem_bytes = 4 * Self::BLK_DIM * std::mem::size_of::() as u32; + unsafe { + stream + .launch_builder(&self.update_kernel) + .arg(&dd_state_per_env_ptr) + .arg(&n_envs) + .arg(&dd_state_stride) + .arg(&dd_max_off) + .arg(&isv_ptr) + .arg(&floor_idx) + .arg(&sentinel_floor) + .arg(&safety_factor) + .arg(&floor_min) + .arg(&floor_max) + .arg(&alpha) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (Self::BLK_DIM, 1, 1), + shared_mem_bytes: smem_bytes, + }) + .map_err(|e| MLError::ModelError(format!("dd_saturation_floor_update: {e}")))?; + } + Ok(()) + } +} + /// SP14 Layer C Phase C.6 (2026-05-08): h_s2_aux RMS EMA producer. /// /// Computes `RMS(h_s2_aux) = sqrt(mean(h_s2_aux²))` over the diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 4e65a42b6..fabe64a3d 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -2124,6 +2124,22 @@ pub(crate) static REWARD_CAP_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!(" /// priors). pub(crate) static KELLY_BAYESIAN_PRIORS_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/kelly_bayesian_priors_update_kernel.cubin")); +/// Class A audit-fix Batch 4-A (2026-05-08): adaptive DD saturation floor +/// producer cubin. Single-block 256-thread kernel that sweeps the +/// `dd_state_per_env[n_envs * 6]` tile (same source as +/// `compute_sp15_final_reward_kernel` reads from at the same boundary, +/// populated by `dd_state_kernel`), aggregates DD_MAX across envs (offset +/// 1 in each per-env tile), and slow-EMA-blends p75 estimator × 1.5 into +/// `ISV[DD_SATURATION_FLOOR_ADAPTIVE_INDEX=458]`. Replaces the hardcoded +/// `0.25f` saturation floor in `trade_physics.cuh::apply_margin_cap` +/// (the upper end of the linear position-size scaling ramp). Pearl-A +/// first-observation bootstrap (sentinel 0.25 matches pre-fix hardcoded +/// value for bit-identical cold-start) + α=0.01 slow EMA thereafter +/// (DD distribution is a slow-moving fold-volatility property — same +/// cadence as the P0-A REWARD_POS_CAP producer). Loaded by +/// `gpu_aux_trunk::DdSaturationFloorUpdateOps`. Per-epoch boundary launch. +pub(crate) static DD_SATURATION_FLOOR_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/dd_saturation_floor_update_kernel.cubin")); + /// SP14 Layer C Phase C.6 (2026-05-08): h_s2_aux RMS EMA producer cubin. /// Single-block 256-thread kernel computing `sqrt(mean(h_s2_aux²))` over /// `h_s2_aux [B, SH2]` (aux trunk final output) and EMA-blending into @@ -2452,7 +2468,7 @@ const ISV_NETWORK_DIM: usize = 23; /// (shifted 112→116 in Plan 4 Task 6 Commit A). /// Written by the constructor; checked at checkpoint load. Fail-fast only — no migration /// path exists. See spec §4.A.2 and `LAYOUT_FINGERPRINT_CURRENT` for structural-hash rationale. -pub(crate) const ISV_TOTAL_DIM: usize = 458; // SP5 + Layer D D1+D2+D3 + SP7 + SP8 + SP9 + SP10 + SP11 + SP13 + SP13 v3 + SP14 (post-C.1) + SP15 + SP14-C aux trunk control plane + SP14-C Phase C.4b aux horizon + Class A P0-A adaptive reward caps + Class A P1-Producer adaptive Bayesian Kelly priors. Bumped 454 → 458 by Class A P1-Producer (2026-05-08): added 4 adaptive-Bayesian-prior slots [454..458) — KELLY_PRIOR_WINS (454) + KELLY_PRIOR_LOSSES (455) + KELLY_PRIOR_SUM_WINS (456) + KELLY_PRIOR_SUM_LOSSES (457). Per Class A audit, hardcoded `prior_wins=2.0f / prior_losses=2.0f / prior_sum_wins=0.01f / prior_sum_losses=0.01f` (kelly_cap_update_kernel.cu:39-42 + trade_physics.cuh::kelly_position_cap:304-307) was a static "I don't know" Bayesian prior frozen at sprint-1 defaults — the cold-start of every new fold ignored what prior folds had learned about the realized Kelly distribution. Producer kernel `kelly_bayesian_priors_update_kernel` aggregates realized Kelly stats from the existing `portfolio_state[n_envs, PS_STRIDE]` buffer (same source as kelly_cap_update consumes) and slow-EMA blends into the four slots. Pearl-A bootstrap (sentinels 2.0/2.0/0.01/0.01 match pre-P1-Producer hardcoded values for bit-identical cold-start) + α=0.005 slow EMA (per-fold cadence; the Bayesian prior is a *prior belief* and should change slowly across folds). Bounds counts∈[0.5, 100], sums∈[0.001, 1.0] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`. Bumped 452 → 454 by Class A P0-A (2026-05-08): added 2 adaptive-reward-cap slots [452..454) — REWARD_POS_CAP_ADAPTIVE (452) + REWARD_NEG_CAP_ADAPTIVE (453). Per Class A audit, hardcoded `REWARD_POS_CAP=+5.0f` / `REWARD_NEG_CAP=-10.0f` (state_layout.cuh:266-267) was structurally clipping the upper tail of realized alpha — controller signals (sharpe EMA, var_q, q_gap) all derive from the CAPPED buffer, so the controller cannot select for trades it cannot see. The state_layout comment explicitly deferred Phase 2 (ISV-driven) "IF Phase 1 validation reveals adaptive need"; 11 superprojects of WR-stuck-at-46-48% revealed adaptive need. Producer kernel `reward_cap_update_kernel` computes p99(winning_returns) × safety_factor=1.5 → POS cap and NEG = -2 × POS (preserves Kahneman 2:1 loss-aversion asymmetry per `pearl_audit_unboundedness_for_implicit_asymmetry` — moved from hardcoded scalar to producer-time multiplier, single source of truth). Pearl-A bootstrap (sentinels 5.0/-10.0) + Wiener-α=0.01 slow EMA. Bounds POS∈[1, 50], NEG∈[-100, -2] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`, NOT tuning. Bumped 450 → 452 by SP14 Layer C Phase C.4b (2026-05-08): added 2 aux-prediction-horizon slots [450..452) — AUX_PRED_HORIZON_BARS (450) + AVG_WIN_HOLD_TIME_BARS (451). Aux's original label was (p_{t+1} > p_t), pure HFT-scale microstructure noise — pivoted to (p_{t+H} > p_t) with H read from ISV[450], adaptively driven from observed avg winning hold time via Pearl-A bootstrap + Wiener-α EMA. Case B: existing `hold_at_exit_per_sample` + `trade_profitable_per_sample` per-sample buffers, no aggregate slot — added new aggregate slot 451 to expose the EMA target. Bumped 444 → 450 by SP14 Layer C Phase C.1 (2026-05-08): atomically deleted 9 α-machinery slots in [385..396) AND added 6 aux-trunk control-plane slots [444..450) — AUX_TRUNK_LR (444) + AUX_TRUNK_BETA1 (445) + AUX_TRUNK_BETA2 (446) + AUX_TRUNK_EPS (447) + AUX_TRUNK_GRAD_CLIP (448) + H_S2_AUX_RMS_EMA (449). Per `feedback_no_legacy_aliases` and `feedback_no_partial_refactor`. Deleted slots left as RESERVED gap (NOT compacted) to preserve checkpoint layout fingerprint compatibility — the surviving SP14 q_disagreement_* diagnostic slots (383, 384, 389) keep their original indices. Pre-C.1 base: 173 + 210 (164 SP5 slots @ 174..278/280..340, with 2-slot gap before cross-fold-persistent Kelly block; Layer D D1 PnL outputs at [286..290); Layer D D2 health composition outputs at [290..294); Layer D D3 training metrics EMA at [294..297); SP7 T1 loss-balance Wiener stats at [297..313); SP7 activation-flag fix per-(head,branch) flags at [313..321); SP8 Fix 36 train_active_frac canary @ [321..322) + LB_MAX_BUDGET per-(head,branch) at [322..330); SP9 Fix 37 Kelly warmup floor at [330..331) + Q_VAR_MAG_EMA at [331..332) + INTENT_EVAL_DIVERGENCE at [332..333) + 3 EMA targets at [333..336) + 3 eval_dist mag bins at [336..339); SP10 Fix 38 EVAL_THOMPSON_TEMP at [339..340); SP11 Fix 39 reward-subsystem controller at [340..367) — 6 component weights [340..346) + 4 controller scalar outputs [346..350) + 2 val-sharpe canaries [350..352) + 6 mag-ratio canaries [352..358) + saboteur engagement [358] + PnL magnitude EMA [359] + popart-component magnitude EMA [360, B1b fix-up] + per-component variance EMAs [361..367, B1b smoke-recovery z-score normalization]; SP13 directional-skill instrumentation at [372..380) — TARGET_DIR_ACC [372] + AUX_DIR_ACC_SHORT/LONG_EMA [373..375) + AUX_DIR_PREDICTION [375] + DIR_SKILL_BONUS_ALPHA/BETA [376..378) + LUCK_WIN_DISCOUNT [378] + SKILL_BONUS_CAP_RATIO [379]; SP13 v3 P0a.T3 Hold-pricing controller at [380..383) — HOLD_COST [380] + HOLD_RATE_TARGET [381] + HOLD_RATE_OBSERVED_EMA [382]; SP14 q_disagreement diagnostic at [383..385), [389] — Q_DISAGREEMENT_SHORT/LONG_EMA + Q_DISAGREEMENT_VARIANCE_EMA (slots [385..389) and [390..396) RESERVED gap from C.1 deletion); SP15 [397..444) per Phase 1; SP14-C [444..450) per Phase C.1. Intentional 5-slot boundary gap at [367..372)) +pub(crate) const ISV_TOTAL_DIM: usize = 459; // SP5 + Layer D D1+D2+D3 + SP7 + SP8 + SP9 + SP10 + SP11 + SP13 + SP13 v3 + SP14 (post-C.1) + SP15 + SP14-C aux trunk control plane + SP14-C Phase C.4b aux horizon + Class A P0-A adaptive reward caps + Class A P1-Producer adaptive Bayesian Kelly priors + Class A audit-fix Batch 4-A adaptive DD saturation floor. Bumped 458 → 459 by Class A audit-fix Batch 4-A (2026-05-08): added 1 adaptive-DD-saturation-floor slot [458..459) — DD_SATURATION_FLOOR_ADAPTIVE (458). Replaces the hardcoded `0.25f` saturation floor in `trade_physics.cuh::apply_margin_cap` (the upper end of the linear position-size scaling ramp dd_scale = max(0.05, 1.0 − dd_frac / saturation_floor)). Distinct semantic role from `SP15_DD_THRESHOLD_INDEX=421` (the SP15 quadratic DD-penalty *trigger* threshold, a *lower* bound). Producer kernel `dd_saturation_floor_update_kernel` aggregates per-env DD_MAX from the existing `dd_state_per_env[env*6 + 1]` tile (same source as the SP15 dd_state_kernel writes to), computes p75 estimator via Welford rule (mean + Z_75 × sigma) × safety_factor=1.5, and slow-EMA blends. Pearl-A bootstrap (sentinel 0.25 matches pre-fix hardcoded value for bit-identical cold-start) + α=0.01 slow EMA (DD distribution is a slow-moving fold-volatility property, shouldn't track per-epoch noise). Bounds [0.10, 0.50] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`. Atomically deleted alongside this slot addition: legacy `compute_drawdown_penalty` device function in trade_physics.cuh + its single caller at experience_kernels.cu:3822 + the `w_dd` and `dd_threshold` kernel args (Item 2 Case A — SP15's quadratic asymmetric DD penalty in `compute_sp15_final_reward_kernel.cu:154` is the production-grade replacement; layering both creates double-counting per `feedback_no_partial_refactor`). Bumped 454 → 458 by Class A P1-Producer (2026-05-08): added 4 adaptive-Bayesian-prior slots [454..458) — KELLY_PRIOR_WINS (454) + KELLY_PRIOR_LOSSES (455) + KELLY_PRIOR_SUM_WINS (456) + KELLY_PRIOR_SUM_LOSSES (457). Per Class A audit, hardcoded `prior_wins=2.0f / prior_losses=2.0f / prior_sum_wins=0.01f / prior_sum_losses=0.01f` (kelly_cap_update_kernel.cu:39-42 + trade_physics.cuh::kelly_position_cap:304-307) was a static "I don't know" Bayesian prior frozen at sprint-1 defaults — the cold-start of every new fold ignored what prior folds had learned about the realized Kelly distribution. Producer kernel `kelly_bayesian_priors_update_kernel` aggregates realized Kelly stats from the existing `portfolio_state[n_envs, PS_STRIDE]` buffer (same source as kelly_cap_update consumes) and slow-EMA blends into the four slots. Pearl-A bootstrap (sentinels 2.0/2.0/0.01/0.01 match pre-P1-Producer hardcoded values for bit-identical cold-start) + α=0.005 slow EMA (per-fold cadence; the Bayesian prior is a *prior belief* and should change slowly across folds). Bounds counts∈[0.5, 100], sums∈[0.001, 1.0] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`. Bumped 452 → 454 by Class A P0-A (2026-05-08): added 2 adaptive-reward-cap slots [452..454) — REWARD_POS_CAP_ADAPTIVE (452) + REWARD_NEG_CAP_ADAPTIVE (453). Per Class A audit, hardcoded `REWARD_POS_CAP=+5.0f` / `REWARD_NEG_CAP=-10.0f` (state_layout.cuh:266-267) was structurally clipping the upper tail of realized alpha — controller signals (sharpe EMA, var_q, q_gap) all derive from the CAPPED buffer, so the controller cannot select for trades it cannot see. The state_layout comment explicitly deferred Phase 2 (ISV-driven) "IF Phase 1 validation reveals adaptive need"; 11 superprojects of WR-stuck-at-46-48% revealed adaptive need. Producer kernel `reward_cap_update_kernel` computes p99(winning_returns) × safety_factor=1.5 → POS cap and NEG = -2 × POS (preserves Kahneman 2:1 loss-aversion asymmetry per `pearl_audit_unboundedness_for_implicit_asymmetry` — moved from hardcoded scalar to producer-time multiplier, single source of truth). Pearl-A bootstrap (sentinels 5.0/-10.0) + Wiener-α=0.01 slow EMA. Bounds POS∈[1, 50], NEG∈[-100, -2] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`, NOT tuning. Bumped 450 → 452 by SP14 Layer C Phase C.4b (2026-05-08): added 2 aux-prediction-horizon slots [450..452) — AUX_PRED_HORIZON_BARS (450) + AVG_WIN_HOLD_TIME_BARS (451). Aux's original label was (p_{t+1} > p_t), pure HFT-scale microstructure noise — pivoted to (p_{t+H} > p_t) with H read from ISV[450], adaptively driven from observed avg winning hold time via Pearl-A bootstrap + Wiener-α EMA. Case B: existing `hold_at_exit_per_sample` + `trade_profitable_per_sample` per-sample buffers, no aggregate slot — added new aggregate slot 451 to expose the EMA target. Bumped 444 → 450 by SP14 Layer C Phase C.1 (2026-05-08): atomically deleted 9 α-machinery slots in [385..396) AND added 6 aux-trunk control-plane slots [444..450) — AUX_TRUNK_LR (444) + AUX_TRUNK_BETA1 (445) + AUX_TRUNK_BETA2 (446) + AUX_TRUNK_EPS (447) + AUX_TRUNK_GRAD_CLIP (448) + H_S2_AUX_RMS_EMA (449). Per `feedback_no_legacy_aliases` and `feedback_no_partial_refactor`. Deleted slots left as RESERVED gap (NOT compacted) to preserve checkpoint layout fingerprint compatibility — the surviving SP14 q_disagreement_* diagnostic slots (383, 384, 389) keep their original indices. Pre-C.1 base: 173 + 210 (164 SP5 slots @ 174..278/280..340, with 2-slot gap before cross-fold-persistent Kelly block; Layer D D1 PnL outputs at [286..290); Layer D D2 health composition outputs at [290..294); Layer D D3 training metrics EMA at [294..297); SP7 T1 loss-balance Wiener stats at [297..313); SP7 activation-flag fix per-(head,branch) flags at [313..321); SP8 Fix 36 train_active_frac canary @ [321..322) + LB_MAX_BUDGET per-(head,branch) at [322..330); SP9 Fix 37 Kelly warmup floor at [330..331) + Q_VAR_MAG_EMA at [331..332) + INTENT_EVAL_DIVERGENCE at [332..333) + 3 EMA targets at [333..336) + 3 eval_dist mag bins at [336..339); SP10 Fix 38 EVAL_THOMPSON_TEMP at [339..340); SP11 Fix 39 reward-subsystem controller at [340..367) — 6 component weights [340..346) + 4 controller scalar outputs [346..350) + 2 val-sharpe canaries [350..352) + 6 mag-ratio canaries [352..358) + saboteur engagement [358] + PnL magnitude EMA [359] + popart-component magnitude EMA [360, B1b fix-up] + per-component variance EMAs [361..367, B1b smoke-recovery z-score normalization]; SP13 directional-skill instrumentation at [372..380) — TARGET_DIR_ACC [372] + AUX_DIR_ACC_SHORT/LONG_EMA [373..375) + AUX_DIR_PREDICTION [375] + DIR_SKILL_BONUS_ALPHA/BETA [376..378) + LUCK_WIN_DISCOUNT [378] + SKILL_BONUS_CAP_RATIO [379]; SP13 v3 P0a.T3 Hold-pricing controller at [380..383) — HOLD_COST [380] + HOLD_RATE_TARGET [381] + HOLD_RATE_OBSERVED_EMA [382]; SP14 q_disagreement diagnostic at [383..385), [389] — Q_DISAGREEMENT_SHORT/LONG_EMA + Q_DISAGREEMENT_VARIANCE_EMA (slots [385..389) and [390..396) RESERVED gap from C.1 deletion); SP15 [397..444) per Phase 1; SP14-C [444..450) per Phase C.1. Intentional 5-slot boundary gap at [367..372)) /// Legacy alias preserved for call sites that haven't been audited for the /// network-vs-total split. New code should pick `ISV_NETWORK_DIM` (for weight /// tensor sizing) or `ISV_TOTAL_DIM` (for the broadcast bus buffer). @@ -3527,7 +3543,8 @@ const fn layout_fingerprint_seed() -> &'static [u8] { AUX_PRED_HORIZON_BARS=450;AVG_WIN_HOLD_TIME_BARS=451;\ REWARD_POS_CAP_ADAPTIVE=452;REWARD_NEG_CAP_ADAPTIVE=453;\ KELLY_PRIOR_WINS=454;KELLY_PRIOR_LOSSES=455;KELLY_PRIOR_SUM_WINS=456;KELLY_PRIOR_SUM_LOSSES=457;\ - ISV_TOTAL_DIM=458;\ + DD_SATURATION_FLOOR_ADAPTIVE=458;\ + ISV_TOTAL_DIM=459;\ SP14_C_AUX_TRUNK_CONTROL_PLANE=sp14_c_phase_1;\ ALPHA_MACHINERY_DELETED=sp14_c_phase_1;\ TRUNK_INPUT_DD_PCT=sp15_wave_4_1a;\ @@ -7554,6 +7571,16 @@ pub struct GpuDqnTrainer { /// freshly-blended priors). kelly_bayesian_priors_update_ops: super::gpu_aux_trunk::KellyBayesianPriorsUpdateOps, + /// Class A audit-fix Batch 4-A (2026-05-08): adaptive DD saturation + /// floor producer Ops handle. Sweeps `dd_state_per_env[n_envs * 6]`, + /// aggregates per-env DD_MAX (offset 1), computes Welford `mean + + /// Z_75 × sigma` p75 estimator × 1.5 safety factor, slow-EMA blends + /// into `ISV[DD_SATURATION_FLOOR_ADAPTIVE_INDEX=458]`. Replaces the + /// hardcoded `0.25f` saturation floor in + /// `trade_physics.cuh::apply_margin_cap`. Per-epoch boundary launch + /// (mirrors P0-A REWARD_POS_CAP cadence). + dd_saturation_floor_update_ops: super::gpu_aux_trunk::DdSaturationFloorUpdateOps, + // ── Speculative inference cache ── /// Cached trunk output [B, SH2] from between-bar speculative forward. speculative_h_s2: CudaSlice, @@ -14738,6 +14765,64 @@ impl GpuDqnTrainer { Ok(()) } + /// Class A audit-fix Batch 4-A (2026-05-08): launch the adaptive DD + /// saturation floor producer. Per-epoch boundary launch — DD + /// distribution is a slow-moving fold-volatility property and + /// shouldn't track per-epoch noise (α=0.01 EMA mirrors P0-A + /// REWARD_POS_CAP cadence). + /// + /// Aggregates per-env DD_MAX (offset 1 in the 6-wide + /// `dd_state_per_env` tile populated by `dd_state_kernel`) across + /// envs, computes Welford `mean + Z_75 × sigma` p75 estimator with + /// `max(p75, mean)` robustness guard × 1.5 safety factor, and + /// slow-EMA blends into ISV[458]. + /// + /// Pearls applied: + /// - `pearl_first_observation_bootstrap.md` (sentinel = 0.25) + /// - `pearl_no_host_branches_in_captured_graph.md` + /// - `feedback_no_atomicadd.md` (block-tree-reduce in shmem only) + /// - `pearl_symmetric_clamp_audit.md` (bilateral clamp on target) + /// - `feedback_isv_for_adaptive_bounds.md` (bounds [0.10, 0.50] + /// are dimensional safety floors, not tuning) + pub(crate) fn launch_dd_saturation_floor_update( + &self, + dd_state_per_env_ptr: u64, + n_envs: i32, + ) -> Result<(), MLError> { + use crate::cuda_pipeline::sp14_isv_slots::{ + DD_SATURATION_FLOOR_ADAPTIVE_INDEX, DD_SATURATION_FLOOR_EMA_ALPHA, + DD_SATURATION_FLOOR_MAX, DD_SATURATION_FLOOR_MIN, + DD_SATURATION_FLOOR_SAFETY_FACTOR, SENTINEL_DD_SATURATION_FLOOR, + }; + // Per-env tile is 6 floats wide (mirrors dd_state_kernel layout). + // DD_MAX sits at offset 1 (DD_CURRENT=0, DD_MAX=1, DD_RECOVERY_BARS=2, + // DD_PERSISTENCE=3, CALMAR=4, DD_PCT=5). + const DD_STATE_STRIDE: i32 = 6; + const DD_MAX_OFF: i32 = 1; + + debug_assert!( + self.isv_signals_dev_ptr != 0, + "launch_dd_saturation_floor_update: isv_signals_dev_ptr must be allocated" + ); + + self.dd_saturation_floor_update_ops.launch( + &self.stream, + dd_state_per_env_ptr, + n_envs, + DD_STATE_STRIDE, + DD_MAX_OFF, + self.isv_signals_dev_ptr, + DD_SATURATION_FLOOR_ADAPTIVE_INDEX as i32, + SENTINEL_DD_SATURATION_FLOOR, + DD_SATURATION_FLOOR_SAFETY_FACTOR, + DD_SATURATION_FLOOR_MIN, + DD_SATURATION_FLOOR_MAX, + DD_SATURATION_FLOOR_EMA_ALPHA, + )?; + + Ok(()) + } + /// SP8 (Fix 36, 2026-05-03): launch GPU-only `train_active_frac` canary /// producer. /// @@ -23160,6 +23245,14 @@ impl GpuDqnTrainer { let kelly_bayesian_priors_update_ops = super::gpu_aux_trunk::KellyBayesianPriorsUpdateOps::new(&stream)?; + // Class A audit-fix Batch 4-A (2026-05-08): adaptive DD saturation + // floor producer. Pre-loads the `CudaFunction` handle at + // construction per `pearl_no_host_branches_in_captured_graph.md`. + // Per-epoch boundary launch (slow-moving — DD distribution is a + // fold-volatility property; α=0.01 EMA mirrors P0-A REWARD_POS_CAP). + let dd_saturation_floor_update_ops = + super::gpu_aux_trunk::DdSaturationFloorUpdateOps::new(&stream)?; + // ── Speculative inference cache ── let speculative_h_s2 = stream.alloc_zeros::(b * sh2) .map_err(|e| MLError::ModelError(format!("speculative_h_s2 alloc: {e}")))?; @@ -23995,6 +24088,8 @@ impl GpuDqnTrainer { reward_cap_update_ops, // Class A P1-Producer (2026-05-08): adaptive Bayesian Kelly priors producer. kelly_bayesian_priors_update_ops, + // Class A audit-fix Batch 4-A (2026-05-08): adaptive DD saturation floor producer. + dd_saturation_floor_update_ops, speculative_h_s2, speculative_features, speculative_valid: false, diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 381177958..6746e0d51 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -178,10 +178,12 @@ pub struct ExperienceCollectorConfig { pub curiosity_scale: f32, /// Risk penalty weight (e.g. 0.1) pub risk_weight: f32, - /// Drawdown fraction before penalty starts (0.02 = 2%) + /// Drawdown fraction before penalty starts (0.02 = 2%). + /// Class A audit-fix Batch 4-A (2026-05-08): no longer used by the + /// legacy `compute_drawdown_penalty` path (deleted), but reused as + /// the `dd_budget` argument to `launch_sp15_dd_state` (line ~5750) + /// for the SP15 DD_PCT computation in dd_state_kernel. pub dd_threshold: f32, - /// Drawdown penalty weight (1.0 = full penalty) - pub w_dd: f32, /// Transaction cost multiplier (hyperopt-tunable, scales base 0.01% rate) pub tx_cost_multiplier: f32, /// UCB count-bonus coefficient for GPU action selection (0.0 = disabled) @@ -341,7 +343,10 @@ impl Default for ExperienceCollectorConfig { curiosity_scale: 1.0, risk_weight: 0.1, dd_threshold: 0.02, - w_dd: 1.0, + // Class A audit-fix Batch 4-A (2026-05-08): `w_dd` field + // DELETED from the config struct atomically with the legacy + // compute_drawdown_penalty deletion. SP15 path uses ISV + // slot 420 (LAMBDA_DD) instead. tx_cost_multiplier: 1.0, count_bonus_coefficient: 0.0, q_clip_min: -200.0, // Reward v6: tighter than old -500 but covers v_range + safety margin @@ -5566,8 +5571,15 @@ impl GpuExperienceCollector { .arg(&config.spread_cost) // bid-ask spread cost (matches backtest) .arg(&config.contract_multiplier) // futures contract multiplier (e.g. 50 for ES) .arg(&config.margin_pct) // initial margin fraction (e.g. 0.06 = 6%) - .arg(&config.dd_threshold) // drawdown threshold before penalty (0.02 = 2%) - .arg(&config.w_dd) // drawdown penalty weight + // Class A audit-fix Batch 4-A (2026-05-08): legacy + // `dd_threshold` and `w_dd` arg passes DELETED + // atomically with the legacy compute_drawdown_penalty + // path. SP15's compute_sp15_final_reward_kernel reads + // λ_dd / DD threshold directly from ISV slots 420/421. + // `config.dd_threshold` survives — still consumed by + // launch_sp15_dd_state at line 5750 as the DD_PCT + // budget. `config.w_dd` field is also dropped from + // the config struct (was consumed only here). .arg(&mirror_i32) // #10 mirror universe .arg(&config.micro_reward_scale) // v8: dense directional micro-reward (Gem 1) // #33 Per-episode saboteur params (0 = NULL = disabled) @@ -6104,6 +6116,19 @@ impl GpuExperienceCollector { ptr } + /// Class A audit-fix Batch 4-A (2026-05-08): raw device pointer to + /// the per-env DD state tile `sp15_dd_state_per_env[n_envs * 6]` + /// (populated by `dd_state_kernel` at the same boundary). Consumed + /// by `dd_saturation_floor_update_kernel` to aggregate per-env + /// DD_MAX (offset 1) into the adaptive saturation floor in + /// `ISV[DD_SATURATION_FLOOR_ADAPTIVE_INDEX=458]`. The buffer is + /// owned by this collector (mirrors the + /// `step_ret_per_sample_dev_ptr` / `trade_close_per_sample_dev_ptr` + /// pattern from P0-A). + pub fn dd_state_per_env_dev_ptr(&self) -> u64 { + self.sp15_dd_state_per_env.dev_ptr + } + /// Last experience count from the most recent `collect_experiences_gpu` call. pub fn last_experience_count(&self) -> usize { self.last_experience_count } diff --git a/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs index 029bdb04d..24c3740f9 100644 --- a/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs @@ -223,6 +223,66 @@ pub const KELLY_PRIOR_EMA_ALPHA: f32 = 0.005; pub const SP14_P1_SLOT_BASE: usize = 454; pub const SP14_P1_SLOT_END: usize = 458; +// ── Class A audit-fix Batch 4-A (2026-05-08): adaptive DD saturation floor ── +// Replaces the hardcoded `0.25f` saturation floor in +// `trade_physics.cuh::apply_margin_cap` (line ~154) with an ISV-driven +// adaptive bound derived from the realized per-env DD_MAX distribution. +// Producer kernel: `dd_saturation_floor_update_kernel.cu` (per-epoch +// boundary cadence, mirroring the P0-A REWARD_POS_CAP producer pattern). +// +// Distinct semantic role from `SP15_DD_THRESHOLD_INDEX=421` (the SP15 +// quadratic DD-penalty trigger threshold, a *lower* bound). This slot is +// the *upper* end of the linear position-size scaling ramp in +// `apply_margin_cap`: +// +// dd_scale = max(0.05, 1.0 − dd_frac / saturation_floor) +// +// With the static 0.25f, the ramp is calibrated for one specific DD +// distribution. Production training across 11 SPs has shown that DDs +// vary widely (median fold DD_MAX=8% on calm folds, 35% on volatile +// folds), and the static 25% breakpoint either over-throttles (calm +// folds where small DDs unnecessarily clamp positions) or under- +// throttles (volatile folds where 35% DDs hit the 5% floor on most +// envs). Adaptive floor derived from p75(DD_MAX) × 1.5 safety factor +// keeps the breakpoint above 25% of the realized distribution +// regardless of fold volatility. +// +// Pearl-A first-observation bootstrap: SENTINEL_DD_SATURATION_FLOOR=0.25 +// matches the pre-fix hardcoded value for bit-identical cold-start +// behavior before the first valid observation lands. +pub const DD_SATURATION_FLOOR_ADAPTIVE_INDEX: usize = 458; + +// Sentinel — matches the pre-fix hardcoded value (`0.25f` saturation +// floor in `apply_margin_cap`) for bit-identical cold-start behavior +// before the first valid observation lands. Pearl-A bootstrap. +pub const SENTINEL_DD_SATURATION_FLOOR: f32 = 0.25; + +// Bounds — Category-1 dimensional safety floors per +// `feedback_isv_for_adaptive_bounds.md`. Floor in [0.10, 0.50]: below +// 10% the ramp would clip positions on minor noise (over-throttling +// on calm folds); above 50% the ramp barely activates before the +// capital-floor circuit-breaker triggers at 25% PDT. The full ramp's +// lower endpoint (`0.05f` floor on dd_scale) is unaffected — only the +// upper endpoint (where ramp saturates) adapts. +pub const DD_SATURATION_FLOOR_MIN: f32 = 0.10; +pub const DD_SATURATION_FLOOR_MAX: f32 = 0.50; + +// Safety factor on the p75 DD_MAX target — keeps the breakpoint above +// the realized 75th percentile so 25% of envs are still well within the +// linear ramp regardless of fold volatility. Mirrors the P0-A +// REWARD_CAP_SAFETY_FACTOR=1.5 design. +pub const DD_SATURATION_FLOOR_SAFETY_FACTOR: f32 = 1.5; + +// EMA blend rate. Slow — DD distribution is a slow-moving property of +// the fold's volatility regime and shouldn't track per-epoch noise. +// Pearl-A bootstrap replaces sentinel directly on first valid +// observation; α=0.01 thereafter. Mirrors the P0-A +// REWARD_CAP_EMA_ALPHA=0.01 design. +pub const DD_SATURATION_FLOOR_EMA_ALPHA: f32 = 0.01; + +pub const SP14_AUDIT_4A_SLOT_BASE: usize = 458; +pub const SP14_AUDIT_4A_SLOT_END: usize = 459; + #[cfg(test)] mod tests { use super::*; @@ -370,4 +430,36 @@ mod tests { SP14_P1_SLOT_END, ISV_TOTAL_DIM, ); } + + /// Lock Class A audit-fix Batch 4-A (2026-05-08) adaptive DD + /// saturation floor slot. Single contiguous slot [458..459) + /// replacing the hardcoded `0.25f` in + /// `trade_physics.cuh::apply_margin_cap`. Sentinel matches the + /// pre-fix hardcoded value (0.25) for bit-identical cold-start. + #[test] + fn sp14_audit_4a_dd_saturation_floor_slot_layout_locked() { + assert_eq!(SP14_AUDIT_4A_SLOT_BASE, 458); + assert_eq!(SP14_AUDIT_4A_SLOT_END, 459); + assert_eq!(DD_SATURATION_FLOOR_ADAPTIVE_INDEX, 458); + // Sentinel matches pre-fix hardcoded value for bit-identical cold-start. + assert_eq!(SENTINEL_DD_SATURATION_FLOOR, 0.25); + // Category-1 dimensional safety: floor in [0.10, 0.50]. + assert_eq!(DD_SATURATION_FLOOR_MIN, 0.10); + assert_eq!(DD_SATURATION_FLOOR_MAX, 0.50); + // Safety factor mirrors P0-A reward-cap producer (1.5×). + assert_eq!(DD_SATURATION_FLOOR_SAFETY_FACTOR, 1.5); + // EMA α matches P0-A reward-cap (slow per-epoch). + assert_eq!(DD_SATURATION_FLOOR_EMA_ALPHA, 0.01); + } + + #[test] + fn all_sp14_audit_4a_slots_fit_within_isv_total_dim() { + use crate::cuda_pipeline::gpu_dqn_trainer::ISV_TOTAL_DIM; + assert!( + SP14_AUDIT_4A_SLOT_END <= ISV_TOTAL_DIM, + "SP14_AUDIT_4A_SLOT_END={} exceeds ISV_TOTAL_DIM={} — bus too small for Class A audit-fix Batch 4-A DD saturation floor slot; \ + bump ISV_TOTAL_DIM in gpu_dqn_trainer.rs (and update layout_fingerprint_seed()).", + SP14_AUDIT_4A_SLOT_END, ISV_TOTAL_DIM, + ); + } } diff --git a/crates/ml/src/cuda_pipeline/state_layout.cuh b/crates/ml/src/cuda_pipeline/state_layout.cuh index 3ee15d7e3..a130c5d51 100644 --- a/crates/ml/src/cuda_pipeline/state_layout.cuh +++ b/crates/ml/src/cuda_pipeline/state_layout.cuh @@ -336,6 +336,34 @@ #define KELLY_PRIOR_SUM_WINS_DEFAULT 0.01f #define KELLY_PRIOR_SUM_LOSSES_DEFAULT 0.01f +// Class A audit-fix Batch 4-A (2026-05-08) — adaptive DD saturation floor. +// Replaces the hardcoded `0.25f` in `trade_physics.cuh::apply_margin_cap` +// with an ISV-driven adaptive bound derived from the realized per-env +// DD_MAX distribution. Producer kernel: `dd_saturation_floor_update_kernel` +// (per-epoch boundary cadence). Sentinel matches the pre-fix hardcoded +// value (0.25) for bit-identical cold-start. Consumer at the call site +// uses the standard sentinel-detect + range-guard pattern (mirrors the +// SP12-cap and Kelly-prior consumer patterns) to fall back to +// DD_SATURATION_FLOOR_DEFAULT when the slot is at sentinel or +// out-of-range. Mirrors `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` +// constants of the same names (locked by +// `sp14_audit_4a_dd_saturation_floor_slot_layout_locked` test). +#define DD_SATURATION_FLOOR_ADAPTIVE_INDEX 458 +#define SENTINEL_DD_SATURATION_FLOOR 0.25f +// Defensive consumer-side bounds (mirror Rust constants +// DD_SATURATION_FLOOR_MIN/MAX in sp14_isv_slots.rs). Values outside this +// window fall back to the macro on the consumer side as a defense +// against malformed prior state. Producer kernel clamps to the same +// range, so a healthy ISV slot will always pass the guard. +#define DD_SATURATION_FLOOR_MIN_BOUND 0.10f +#define DD_SATURATION_FLOOR_MAX_BOUND 0.50f +// Cold-start fallback for the consumer-side ISV-NULL guard (mirrors the +// SP12-cap fall-back pattern in `sp15_apply_sp12_cap`). When +// `isv_signals_ptr == NULL` OR the slot value is at sentinel within EPS +// OR out of [MIN_BOUND, MAX_BOUND], the consumer falls back to this +// static default. Bit-identical pre-Batch-4A behavior. +#define DD_SATURATION_FLOOR_DEFAULT 0.25f + // ── Compile-time checks ── static_assert(SL_PADDING_START + SL_PADDING_DIM == SL_STATE_DIM, "State layout dimensions must sum to SL_STATE_DIM"); diff --git a/crates/ml/src/cuda_pipeline/trade_physics.cuh b/crates/ml/src/cuda_pipeline/trade_physics.cuh index e693ee266..ff3cc1340 100644 --- a/crates/ml/src/cuda_pipeline/trade_physics.cuh +++ b/crates/ml/src/cuda_pipeline/trade_physics.cuh @@ -141,7 +141,16 @@ __device__ __forceinline__ float apply_margin_cap( float target_position, float equity, float margin_per_contract, - float peak_equity + float peak_equity, + /* Class A audit-fix Batch 4-A (2026-05-08): adaptive DD saturation + * floor (the upper end of the linear position-size scaling ramp). + * NULL-tolerant: when NULL or the slot is at sentinel or out of + * [DD_SATURATION_FLOOR_MIN_BOUND, DD_SATURATION_FLOOR_MAX_BOUND], + * falls back to DD_SATURATION_FLOOR_DEFAULT (0.25f) for + * bit-identical pre-Batch-4A behavior. Mirrors the SP12-cap + * fallback pattern in `sp15_apply_sp12_cap` and the Kelly-prior + * NULL-tolerant pattern in `kelly_position_cap`. */ + const float* __restrict__ isv_signals_ptr /* [ISV_TOTAL_DIM]; NULL = static default */ ) { if (margin_per_contract <= 0.0f || equity <= 0.0f) return target_position; @@ -150,8 +159,24 @@ __device__ __forceinline__ float apply_margin_cap( if (peak_equity > 1.0f) { float dd_frac = (peak_equity - equity) / peak_equity; dd_frac = fmaxf(0.0f, fminf(dd_frac, 1.0f)); - /* Linear ramp: 1.0 at 0% DD → 0.0 at 25% DD (floor threshold) */ - dd_scale = fmaxf(0.05f, 1.0f - dd_frac / 0.25f); /* floor at 5% to avoid zero */ + /* Linear ramp: 1.0 at 0% DD → 0.0 at saturation_floor (then 5% floor). + * Class A audit-fix Batch 4-A (2026-05-08): saturation_floor is now + * ISV-driven (slot 458) — derived from p75(per-env DD_MAX) × 1.5 + * by the dd_saturation_floor_update_kernel producer at per-epoch + * boundary. Cold-start fallback to the original 0.25f when the + * slot is at sentinel (within EPS) or out-of-range. Defensive + * guard mirrors `sp15_apply_sp12_cap`. Per + * feedback_isv_for_adaptive_bounds + feedback_no_partial_refactor. */ + float saturation_floor = DD_SATURATION_FLOOR_DEFAULT; /* fallback = 0.25f */ + if (isv_signals_ptr != NULL) { + const float isv_floor = isv_signals_ptr[DD_SATURATION_FLOOR_ADAPTIVE_INDEX]; + if (isv_floor >= DD_SATURATION_FLOOR_MIN_BOUND + && isv_floor <= DD_SATURATION_FLOOR_MAX_BOUND + && fabsf(isv_floor - SENTINEL_DD_SATURATION_FLOOR) > 1e-6f) { + saturation_floor = isv_floor; + } + } + dd_scale = fmaxf(0.05f, 1.0f - dd_frac / saturation_floor); /* floor at 5% to avoid zero */ } float max_contracts = (equity / margin_per_contract) * dd_scale; @@ -574,30 +599,27 @@ __device__ __forceinline__ float compute_lump_sum_opp_cost( return -(shaping_scale * holding_cost_rate * fabsf(position) * hold_time); } -/* ── Drawdown penalty: smooth ramp from threshold to floor ───────────── */ -/* Returns a penalty in [-dd_penalty_scale * w_dd, 0]. Applied every step */ -/* so the model learns to reduce position size DURING drawdown. */ -/* */ -/* Class A P0-A-downstream (2026-05-08): `dd_penalty_scale` lifted from */ -/* the hardcoded `-5.0f` (= 1× the pre-P0-A REWARD_POS_CAP=5.0f) to a */ -/* caller-supplied scalar so it can scale proportionally with the now- */ -/* adaptive REWARD_POS_CAP (ISV[REWARD_POS_CAP_ADAPTIVE_INDEX=452]). */ -/* Cold-start fallback (sentinel ISV) lands at the original 5.0f at the */ -/* call site, preserving bit-identical pre-P0-A behavior. Per */ -/* feedback_isv_for_adaptive_bounds + feedback_no_partial_refactor. */ - -__device__ __forceinline__ float compute_drawdown_penalty( - float equity, float peak_equity, - float dd_threshold, float w_dd, - float dd_penalty_scale /* magnitude — caller passes pos_cap_eff (production: ISV[452] with sentinel-fallback) */ -) { - if (w_dd <= 0.0f) return 0.0f; - float dd = compute_drawdown(equity, peak_equity); - if (dd <= dd_threshold) return 0.0f; - float floor_dd = 0.25f; // capital floor = 25% DD - float dd_excess = fminf((dd - dd_threshold) / (floor_dd - dd_threshold), 1.0f); - return -dd_penalty_scale * dd_excess * w_dd; -} +/* ── Drawdown penalty: DELETED (Class A audit-fix Batch 4-A, 2026-05-08) ── + * + * The legacy `compute_drawdown_penalty` device function was deleted + * atomically with the `dd_threshold` and `w_dd` kernel arguments and + * the single call site at `experience_kernels.cu:3822`. + * + * Rationale (Item 2 Case A): SP15's quadratic asymmetric DD penalty in + * `compute_sp15_final_reward_kernel.cu:154` (`sp15_dd_penalty` from + * `sp15_reward_axis_helpers.cuh:169`) is the production-grade + * replacement. It runs unconditionally as a post-modifier on the + * SP11-composed reward, with ISV-driven λ_dd (slot 420) and DD + * threshold (slot 421). Layering the legacy linear-ramp penalty inside + * the SP11 composer on top of the SP15 quadratic penalty creates + * double-counting of DD shaping — the exact code-smell the Class A + * audit was designed to eliminate. + * + * Per `feedback_no_legacy_aliases.md` and `feedback_no_partial_refactor.md`, + * when functionality is superseded by SP15, the legacy path must be + * deleted atomically across all consumers in the same commit (function + * + call + caller args + Rust config field). + */ /* ── Capital floor circuit breaker (PDT $25K rule) ───────────────────── */ /* Triggers when equity drops 25% from peak (floor = 75% of peak). */ @@ -821,8 +843,14 @@ __device__ __forceinline__ void unified_env_step_core( /* ── Step 4: Margin cap ── */ if (!is_hold_action) { float margin_per_contract = close * contract_multiplier * margin_pct; + /* Class A audit-fix Batch 4-A (2026-05-08): thread isv_signals_ptr + * through to apply_margin_cap so the DD saturation floor (slot + * 458) is ISV-driven. NULL-tolerant: when isv_signals_ptr is + * NULL (validation envs) or slot is at sentinel, falls back to + * the original 0.25f. */ target_position = apply_margin_cap( - target_position, prev_equity, margin_per_contract, *max_equity + target_position, prev_equity, margin_per_contract, *max_equity, + isv_signals_ptr ); } diff --git a/crates/ml/src/hyperopt/adapters/dqn.rs b/crates/ml/src/hyperopt/adapters/dqn.rs index 942cd4b47..d297a0e4f 100644 --- a/crates/ml/src/hyperopt/adapters/dqn.rs +++ b/crates/ml/src/hyperopt/adapters/dqn.rs @@ -190,7 +190,9 @@ pub struct DQNParams { pub replay_intensity: f64, /// Architecture family: scales hidden_dim_base, dueling_hidden_dim, num_atoms pub architecture_intensity: f64, - /// Risk family: scales max_position, huber_delta, min_profit_factor, w_dd, dd_threshold + /// Risk family: scales max_position, huber_delta, min_profit_factor, dd_threshold. + /// (Pre-Class-A-audit-fix-Batch-4A this also scaled `w_dd`; that field was deleted + /// 2026-05-08 atomically with the legacy compute_drawdown_penalty path.) pub risk_intensity: f64, // === 6 generalization family intensity scalars [0.0, 2.0] === diff --git a/crates/ml/src/trainers/dqn/config.rs b/crates/ml/src/trainers/dqn/config.rs index df5d8a700..512fc3067 100644 --- a/crates/ml/src/trainers/dqn/config.rs +++ b/crates/ml/src/trainers/dqn/config.rs @@ -437,8 +437,9 @@ pub struct DQNHyperparameters { // GPU composite reward weights (6-component composite in CUDA kernel) /// Normalized PnL weight for composite reward pub w_pnl: f32, - /// Drawdown penalty weight for composite reward - pub w_dd: f32, + // `w_dd` field DELETED (Class A audit-fix Batch 4-A, 2026-05-08) + // atomically with the legacy compute_drawdown_penalty path. SP15 + // path uses ISV slot 420 (LAMBDA_DD) — see compute_sp15_final_reward_kernel.cu. /// Idle penalty weight for composite reward pub w_idle: f32, /// Drawdown tolerance before penalty kicks in (fraction, e.g. 0.02 = 2%) @@ -677,7 +678,9 @@ pub struct DQNHyperparameters { pub replay_intensity: f64, /// Architecture family: scales hidden_dim_base, dueling_hidden_dim, num_atoms. pub architecture_intensity: f64, - /// Risk family: scales max_position, huber_delta, min_profit_factor, w_dd, dd_threshold. + /// Risk family: scales max_position, huber_delta, min_profit_factor, dd_threshold. + /// (Pre-Class-A-audit-fix-Batch-4A this also scaled `w_dd`; that field was deleted + /// 2026-05-08 atomically with the legacy compute_drawdown_penalty path.) pub risk_intensity: f64, // ── Generalization family intensity scalars for hyperopt (1.0 = defaults, range [0.0, 2.0]) ── @@ -1254,13 +1257,16 @@ impl DQNHyperparameters { self.num_atoms = (scaled_atoms | 1).max(11); // ensure odd for C51 symmetry // Risk family: scales max_leverage (not max_position_absolute directly), - // huber_delta, minimum_profit_factor, w_dd, dd_threshold. + // huber_delta, minimum_profit_factor, dd_threshold. // max_position_absolute is derived later via compute_max_position(reference_price). + // Class A audit-fix Batch 4-A (2026-05-08): `w_dd *= rki` DELETED + // atomically with the field. SP15's λ_dd (ISV slot 420) is the + // production-grade replacement and is producer-driven (not + // hyperopt-tunable via this RKI multiplier). let rki = self.risk_intensity as f32; self.max_leverage *= rki; self.huber_delta *= rki; self.minimum_profit_factor = 1.0 + (self.minimum_profit_factor - 1.0) * rki; - self.w_dd *= rki; self.dd_threshold *= rki; // ── Generalization families ── @@ -1333,7 +1339,8 @@ impl DQNHyperparameters { min_epochs_before_stopping: 50, hold_penalty: -0.001, // deprecated: kept for hyperopt adapter compat w_pnl: 0.3, - w_dd: 2.0, // was 1.0 — stronger penalty now that threshold actually fires + // `w_dd` field DELETED (Class A audit-fix Batch 4-A, 2026-05-08) + // — SP15 path uses ISV slot 420 (LAMBDA_DD) instead. w_idle: 0.01, dd_threshold: 0.005, // 0.5% — active for bar-level trades (was 2%, effectively dead) loss_aversion: 1.0, diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 88f14ab57..cd791daa2 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -1088,6 +1088,33 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "ISV[KELLY_PRIOR_SUM_LOSSES_INDEX=457] — Class A P1-Producer adaptive Bayesian prior_sum_losses (replaces hardcoded prior_sum_losses=0.01f from kelly_cap_update_kernel.cu:42 + trade_physics.cuh::kelly_position_cap:307). Produced by `kelly_bayesian_priors_update_kernel` from the realized PS_KELLY_SUM_LOSSES field aggregated across envs in the per-epoch `portfolio_state` buffer. FoldReset sentinel SENTINEL_KELLY_PRIOR_SUM_LOSSES=0.01 — Pearl-A first-observation bootstrap (matches pre-P1-Producer hardcoded value for bit-identical cold-start). α=0.005 slow EMA thereafter. Bounds [0.001, 1.0] same as prior_sum_wins. Consumed at the same 2 sites as prior_wins.", }, + // ── Class A audit-fix Batch 4-A (2026-05-08): adaptive DD saturation floor ── + // Single slot [458..459) replacing the hardcoded `0.25f` + // saturation floor in `trade_physics.cuh::apply_margin_cap` + // (the upper end of the linear position-size scaling ramp). + // Producer kernel `dd_saturation_floor_update_kernel` + // aggregates per-env DD_MAX (offset 1) from the existing + // `sp15_dd_state_per_env[n_envs * 6]` tile (same source as + // `compute_sp15_final_reward_kernel` reads from at the same + // per-epoch boundary), then slow-EMA-blends into the slot. + // Pearl-A bootstrap sentinel (0.25) matches the pre-fix + // hardcoded value for bit-identical cold-start behavior — + // consumer falls back to DD_SATURATION_FLOOR_DEFAULT (0.25) + // via the sentinel-detect + range-guard pattern (mirrors + // the SP12-cap fallback in `sp15_apply_sp12_cap`). Once the + // first valid observation lands the adaptive path takes over + // via slow EMA (α=0.01; DD distribution is a slow-moving + // fold-volatility property — same cadence as P0-A + // REWARD_POS_CAP). Without this dispatch arm, the C.10 + // lesson recurs: the FoldReset registry entry exists, but + // no actual reset fires — slot drifts across folds and the + // layout-fingerprint smoke test eventually catches it as a + // runtime crash. + RegistryEntry { + name: "sp14_audit_4a_dd_saturation_floor_adaptive", + category: ResetCategory::FoldReset, + description: "ISV[DD_SATURATION_FLOOR_ADAPTIVE_INDEX=458] — Class A audit-fix Batch 4-A adaptive DD saturation floor (replaces hardcoded `0.25f` in `trade_physics.cuh::apply_margin_cap`, the upper end of the linear position-size scaling ramp `dd_scale = max(0.05, 1.0 − dd_frac/saturation_floor)`). Produced by `dd_saturation_floor_update_kernel` from p75(per-env DD_MAX) × safety_factor=1.5 over the `dd_state_per_env[n_envs * 6]` tile (offset 1 = DD_MAX). FoldReset sentinel SENTINEL_DD_SATURATION_FLOOR=0.25 — Pearl-A first-observation bootstrap (matches pre-fix hardcoded value for bit-identical cold-start). Welford α=0.01 slow EMA thereafter. Bounds [0.10, 0.50] — Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`. Distinct semantic role from SP15_DD_THRESHOLD_INDEX=421 (the SP15 quadratic DD-penalty trigger threshold, a lower bound). Consumed by `apply_margin_cap` at line ~154 (threaded `isv_signals_ptr`, NULL-tolerant cold-start fallback to DD_SATURATION_FLOOR_DEFAULT). Atomically wired with the legacy `compute_drawdown_penalty` deletion (Item 2 Case A) since SP15's quadratic asymmetric DD penalty in `compute_sp15_final_reward_kernel.cu:154` is the production-grade replacement; layering both creates double-counting.", + }, // ── SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots ──────────── // Two ISV slots [407, 408]: // - OFI_IMPACT_LAMBDA_INDEX=407: Invariant-1 anchor (NOT a diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index deb7ca54e..4a7756317 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -443,6 +443,29 @@ impl DQNTrainer { ) { tracing::warn!(epoch, "launch_reward_cap_update failed (non-fatal): {e}"); } + // Class A audit-fix Batch 4-A (2026-05-08): adaptive DD + // saturation floor — per-epoch boundary launch. Sweeps + // the per-env `dd_state_per_env[n_envs * 6]` tile + // (populated by `dd_state_kernel` at the same boundary + // — same source as `compute_sp15_final_reward_kernel` + // reads from), aggregates DD_MAX (offset 1) across envs + // via Welford `mean + Z_75 × sigma` p75 estimator with + // `max(p75, mean)` robustness guard × 1.5 safety factor, + // and slow-EMA blends into ISV[458]. Pearl-A first- + // observation bootstrap (sentinel 0.25 matches pre-fix + // hardcoded `0.25f` in `apply_margin_cap` for bit- + // identical cold-start) + α=0.01 slow EMA. Replaces the + // hardcoded saturation floor in + // `trade_physics.cuh::apply_margin_cap` (the upper end + // of the linear position-size scaling ramp). + let dd_state_per_env_ptr = collector.dd_state_per_env_dev_ptr(); + let n_envs_i32 = collector.alloc_episodes() as i32; + if let Err(e) = fused.trainer().launch_dd_saturation_floor_update( + dd_state_per_env_ptr, + n_envs_i32, + ) { + tracing::warn!(epoch, "launch_dd_saturation_floor_update failed (non-fatal): {e}"); + } } // D.8 Plan 2 Task 6C: update TLOB regime focus EMA in ISV at epoch boundary. @@ -2107,7 +2130,8 @@ impl DQNTrainer { contract_multiplier: self.hyperparams.contract_multiplier as f32, margin_pct: self.hyperparams.margin_pct as f32, dd_threshold: self.hyperparams.dd_threshold as f32, - w_dd: self.hyperparams.w_dd as f32, + // `w_dd` field DELETED (Class A audit-fix Batch 4-A, 2026-05-08) + // atomically with the legacy compute_drawdown_penalty path. time_reversal_mod: self.hyperparams.time_reversal_mod as i32, mirror_active: self.current_epoch % 2 == 1, // Re-enabled: now negates ALL 17 directional features (was only 4) feature_noise_scale: self.hyperparams.feature_noise_scale as f32, @@ -8246,6 +8270,24 @@ impl DQNTrainer { ); } } + // Class A audit-fix Batch 4-A (2026-05-08): adaptive DD saturation + // floor. Single FoldReset slot [458..459) — sentinel matches + // pre-fix hardcoded value (0.25) so the kernel's "first + // observation" check fires cleanly on the new fold's first + // launch (avoids cross-fold EMA contamination). Cold-start + // is bit-identical to pre-Batch-4A until the first valid + // observation lands. + "sp14_audit_4a_dd_saturation_floor_adaptive" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp14_isv_slots::{ + DD_SATURATION_FLOOR_ADAPTIVE_INDEX, SENTINEL_DD_SATURATION_FLOOR, + }; + fused.trainer().write_isv_signal_at( + DD_SATURATION_FLOOR_ADAPTIVE_INDEX, + SENTINEL_DD_SATURATION_FLOOR, + ); + } + } // SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots. // OFI_IMPACT_LAMBDA_INDEX=407 is an Invariant-1 anchor (NOT // a stateful EMA) — rewrite the constructor's value at fold diff --git a/crates/ml/src/training_profile.rs b/crates/ml/src/training_profile.rs index df3ce5d75..b5987afd0 100644 --- a/crates/ml/src/training_profile.rs +++ b/crates/ml/src/training_profile.rs @@ -260,8 +260,10 @@ pub struct RewardSection { pub w_dsr: Option, /// Normalized PnL weight pub w_pnl: Option, - /// Drawdown penalty weight - pub w_dd: Option, + // `w_dd` field DELETED (Class A audit-fix Batch 4-A, 2026-05-08) + // atomically with the legacy compute_drawdown_penalty path. SP15 + // path uses ISV slot 420 (LAMBDA_DD) which is producer-driven — + // not exposed as a profile-tunable hyperparameter. /// Idle penalty weight pub w_idle: Option, /// Drawdown tolerance before penalty (fraction) @@ -418,7 +420,7 @@ pub struct SearchSpaceSection { // Composite reward weights pub w_dsr: Option<[f64; 2]>, pub w_pnl: Option<[f64; 2]>, - pub w_dd: Option<[f64; 2]>, + // `w_dd` field DELETED (Class A audit-fix Batch 4-A, 2026-05-08). pub w_idle: Option<[f64; 2]>, pub dd_threshold: Option<[f64; 2]>, pub loss_aversion: Option<[f64; 2]>, @@ -450,7 +452,7 @@ pub struct PhaseFastSection { // Composite reward weights (fixed in phase fast) pub w_dsr: Option, pub w_pnl: Option, - pub w_dd: Option, + // `w_dd` field DELETED (Class A audit-fix Batch 4-A, 2026-05-08). pub w_idle: Option, pub dd_threshold: Option, pub loss_aversion: Option, @@ -470,9 +472,12 @@ pub enum HyperoptPhase { /// Phase 3: fix dynamics + architecture from Phase 2 JSON, search only reward weights (7D). /// Fastest phase (~10s/trial) — only experience collection changes, not network or training. Reward(Vec), - /// Phase 4: fix dynamics + architecture from Phase 2, search risk params (8D). + /// Phase 4: fix dynamics + architecture from Phase 2, search risk params (7D + /// post-Class-A-audit-fix-Batch-4A; was 8D pre-Batch-4A). /// Risk dims: kelly_fractional, dd_threshold, loss_aversion, time_decay_rate, - /// q_gap_threshold, w_dsr, w_pnl, w_dd. + /// q_gap_threshold, w_dsr, w_pnl. (Pre-Class-A-audit-fix-Batch-4A this also + /// included `w_dd`; field deleted 2026-05-08 atomically with legacy + /// compute_drawdown_penalty.) Risk(Vec), } @@ -582,7 +587,9 @@ impl HyperoptProfile { "her_ratio" => ss.her_ratio, "w_dsr" => ss.w_dsr, "w_pnl" => ss.w_pnl, - "w_dd" => ss.w_dd, + // "w_dd" dispatch arm DELETED (Class A audit-fix Batch 4-A, + // 2026-05-08) atomically with the field. SP15 path drives + // λ_dd from ISV slot 420 (producer-side), not from profile. "w_idle" => ss.w_idle, "dd_threshold" => ss.dd_threshold, "loss_aversion" => ss.loss_aversion, @@ -1052,9 +1059,9 @@ impl DqnTrainingProfile { if let Some(v) = rw.w_pnl { hp.w_pnl = v as f32; } - if let Some(v) = rw.w_dd { - hp.w_dd = v as f32; - } + // `w_dd` profile→hyperparam mapping DELETED (Class A audit-fix + // Batch 4-A, 2026-05-08) atomically with the field. SP15 + // path uses ISV slot 420 (LAMBDA_DD) which is producer-driven. if let Some(v) = rw.w_idle { hp.w_idle = v as f32; } @@ -1481,7 +1488,8 @@ mod tests { // [reward] section — all composite weights assert!((hp.w_pnl - 0.3).abs() < 0.01); - assert!((hp.w_dd - 1.0).abs() < 0.01); + // `w_dd` assertion DELETED (Class A audit-fix Batch 4-A, 2026-05-08) + // atomically with the field — SP15 path uses ISV slot 420. assert!((hp.w_idle - 0.01).abs() < 0.001); assert!((hp.dd_threshold - 0.02).abs() < 0.001); assert!((hp.time_decay_rate - 0.0005).abs() < 0.0001); diff --git a/crates/ml/tests/sp14_oracle_tests.rs b/crates/ml/tests/sp14_oracle_tests.rs index 31d170493..8921485c0 100644 --- a/crates/ml/tests/sp14_oracle_tests.rs +++ b/crates/ml/tests/sp14_oracle_tests.rs @@ -1225,3 +1225,296 @@ fn layout_fingerprint_bumps_after_sp14_wire() { LAYOUT_FINGERPRINT_CURRENT, pre_b8_fp ); } + +// ═══════════════════════════════════════════════════════════════════════════ +// Class A audit-fix Batch 4-A (2026-05-08) — adaptive DD saturation floor +// producer tests. +// +// Verifies the `dd_saturation_floor_update_kernel` producer: +// 1. Cold-start Pearl-A bootstrap: ISV stays at sentinel until first valid +// observation, then REPLACES (no blend). +// 2. No-DD guard: when all envs have DD_MAX=0, slot is preserved bit-exactly. +// 3. Bounds enforced: floor in [0.10, 0.50]. +// 4. Welford slow EMA after bootstrap (α=0.01). +// +// All tests are #[ignore = "requires GPU"]; gated under #[cfg(feature = "cuda")]. +// ═══════════════════════════════════════════════════════════════════════════ + +#[cfg(feature = "cuda")] +#[allow(unsafe_code)] // CUDA kernel launch + mapped-pinned memory. +mod sp14_audit_4a_dd_saturation_floor_gpu { + use std::sync::Arc; + + use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg}; + use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer; + use ml::cuda_pipeline::sp14_isv_slots::{ + DD_SATURATION_FLOOR_ADAPTIVE_INDEX, DD_SATURATION_FLOOR_EMA_ALPHA, + DD_SATURATION_FLOOR_MAX, DD_SATURATION_FLOOR_MIN, + DD_SATURATION_FLOOR_SAFETY_FACTOR, SENTINEL_DD_SATURATION_FLOOR, + }; + + const DD_SATURATION_FLOOR_UPDATE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/dd_saturation_floor_update_kernel.cubin")); + + /// Per-env DD state tile width (mirrors `dd_state_kernel` layout in + /// `dd_state_kernel.cu`: DD_CURRENT, DD_MAX, DD_RECOVERY_BARS, + /// DD_PERSISTENCE, CALMAR, DD_PCT). + const DD_STATE_STRIDE: i32 = 6; + /// DD_MAX offset within each env's tile. + const DD_MAX_OFF: i32 = 1; + + fn make_stream() -> Arc { + let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?"); + ctx.default_stream() + } + + fn load_kernel(stream: &Arc) -> CudaFunction { + let module = stream + .context() + .load_cubin(DD_SATURATION_FLOOR_UPDATE_CUBIN.to_vec()) + .expect("load dd_saturation_floor_update_kernel cubin"); + module + .load_function("dd_saturation_floor_update") + .expect("load dd_saturation_floor_update function") + } + + /// 4 arrays × 256 × 4 bytes = 4096 bytes shmem. + const BLK_DIM: u32 = 256; + const SMEM_BYTES: u32 = 4 * BLK_DIM * std::mem::size_of::() as u32; + + /// Helper: launch with a fixed config. + #[allow(clippy::too_many_arguments)] + fn launch_dd_floor( + stream: &Arc, + kernel: &CudaFunction, + dd_state_ptr: u64, + n_envs: i32, + dd_state_stride: i32, + dd_max_off: i32, + isv_ptr: u64, + floor_idx: i32, + sentinel_floor: f32, + safety_factor: f32, + floor_min: f32, + floor_max: f32, + alpha: f32, + ) { + unsafe { + stream + .launch_builder(kernel) + .arg(&dd_state_ptr) + .arg(&n_envs) + .arg(&dd_state_stride) + .arg(&dd_max_off) + .arg(&isv_ptr) + .arg(&floor_idx) + .arg(&sentinel_floor) + .arg(&safety_factor) + .arg(&floor_min) + .arg(&floor_max) + .arg(&alpha) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (BLK_DIM, 1, 1), + shared_mem_bytes: SMEM_BYTES, + }) + .expect("launch dd_saturation_floor_update"); + } + stream.synchronize().expect("sync after dd_saturation_floor_update"); + } + + /// Build a per-env DD state tile [n_envs * 6] with the given DD_MAX + /// values at offset 1; other offsets zeroed. + fn build_dd_state_tile(dd_max_per_env: &[f32]) -> Vec { + let n_envs = dd_max_per_env.len(); + let mut tile = vec![0.0_f32; n_envs * DD_STATE_STRIDE as usize]; + for (env, dd_max) in dd_max_per_env.iter().enumerate() { + tile[env * DD_STATE_STRIDE as usize + DD_MAX_OFF as usize] = *dd_max; + } + tile + } + + /// Test 1 — Pearl-A first-observation bootstrap. + /// + /// 4 envs with DD_MAX=[0.10, 0.15, 0.20, 0.25]. mean=0.175, var=0.003125, + /// sigma≈0.0559. p75 ≈ 0.175 + 0.6745 × 0.0559 ≈ 0.213. Robust target + /// = max(0.213, 0.175) = 0.213. × safety_factor=1.5 = 0.3194. Clamped to + /// [0.10, 0.50] — stays at 0.3194. + /// + /// Cold-start (current = sentinel 0.25): Pearl-A REPLACES → 0.3194. + #[test] + #[ignore = "requires GPU"] + fn dd_floor_pearl_a_bootstrap() { + let stream = make_stream(); + let kernel = load_kernel(&stream); + + const ISV_DIM: usize = 1024; + let mut isv = vec![0.0_f32; ISV_DIM]; + // Cold-start: sentinel matches pre-fix hardcoded value. + isv[DD_SATURATION_FLOOR_ADAPTIVE_INDEX] = SENTINEL_DD_SATURATION_FLOOR; + + let dd_max: [f32; 4] = [0.10, 0.15, 0.20, 0.25]; + let n_envs = dd_max.len() as i32; + let tile = build_dd_state_tile(&dd_max); + + let tile_buf = unsafe { MappedF32Buffer::new(tile.len()) }.expect("alloc tile"); + tile_buf.write_from_slice(&tile); + let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); + isv_buf.write_from_slice(&isv); + + launch_dd_floor( + &stream, &kernel, + tile_buf.dev_ptr, n_envs, DD_STATE_STRIDE, DD_MAX_OFF, + isv_buf.dev_ptr, DD_SATURATION_FLOOR_ADAPTIVE_INDEX as i32, + SENTINEL_DD_SATURATION_FLOOR, DD_SATURATION_FLOOR_SAFETY_FACTOR, + DD_SATURATION_FLOOR_MIN, DD_SATURATION_FLOOR_MAX, + DD_SATURATION_FLOOR_EMA_ALPHA, + ); + + let result = isv_buf.read_all(); + let floor = result[DD_SATURATION_FLOOR_ADAPTIVE_INDEX]; + + // Pearl-A: sentinel → REPLACE with target. Target = max(p75, mean) × 1.5. + // mean=0.175, var=(0.01+0.0225+0.04+0.0625)/4 - 0.175² = 0.0337/4-... + // sumsq = 0.01+0.0225+0.04+0.0625 = 0.135, sumsq/n = 0.03375 + // variance = 0.03375 - 0.030625 = 0.003125, sigma ≈ 0.0559 + // p75 ≈ 0.175 + 0.6745 × 0.0559 ≈ 0.2127. max(p75, mean)=0.2127. + // × 1.5 = 0.319. Clamped to [0.10, 0.50] → stays at ~0.319. + assert!( + (floor - 0.319).abs() < 0.01, + "Pearl-A bootstrap: expected floor≈0.319, got {floor}" + ); + } + + /// Test 2 — No DD observed → ISV slot preserved bit-exactly. + /// + /// All envs have DD_MAX=0 (no DD observed yet). Kernel must skip the + /// EMA update (count==0 guard) and leave the slot at its seeded value. + #[test] + #[ignore = "requires GPU"] + fn dd_floor_no_dd_preserves_isv() { + let stream = make_stream(); + let kernel = load_kernel(&stream); + + const ISV_DIM: usize = 1024; + const SEED: f32 = 0.37; + let mut isv = vec![0.0_f32; ISV_DIM]; + isv[DD_SATURATION_FLOOR_ADAPTIVE_INDEX] = SEED; + + // All envs at DD_MAX=0.0 — no DD observed. + let dd_max: [f32; 8] = [0.0; 8]; + let n_envs = dd_max.len() as i32; + let tile = build_dd_state_tile(&dd_max); + + let tile_buf = unsafe { MappedF32Buffer::new(tile.len()) }.expect("alloc tile"); + tile_buf.write_from_slice(&tile); + let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); + isv_buf.write_from_slice(&isv); + + launch_dd_floor( + &stream, &kernel, + tile_buf.dev_ptr, n_envs, DD_STATE_STRIDE, DD_MAX_OFF, + isv_buf.dev_ptr, DD_SATURATION_FLOOR_ADAPTIVE_INDEX as i32, + SENTINEL_DD_SATURATION_FLOOR, DD_SATURATION_FLOOR_SAFETY_FACTOR, + DD_SATURATION_FLOOR_MIN, DD_SATURATION_FLOOR_MAX, + DD_SATURATION_FLOOR_EMA_ALPHA, + ); + + let result = isv_buf.read_all(); + assert_eq!( + result[DD_SATURATION_FLOOR_ADAPTIVE_INDEX], SEED, + "DD floor slot must be preserved bit-exactly when no DD observed", + ); + } + + /// Test 3 — Bounds enforcement: huge DD_MAX clamps floor to MAX=0.50. + /// + /// 2 envs with DD_MAX=[0.80, 0.90] — both well above the saturation + /// floor MAX. Target = max(p75, mean) × 1.5 ≈ 0.85 × 1.5 = 1.275 → clamps + /// to 0.50. + #[test] + #[ignore = "requires GPU"] + fn dd_floor_bounds_clamp_extreme_dd() { + let stream = make_stream(); + let kernel = load_kernel(&stream); + + const ISV_DIM: usize = 1024; + let mut isv = vec![0.0_f32; ISV_DIM]; + isv[DD_SATURATION_FLOOR_ADAPTIVE_INDEX] = SENTINEL_DD_SATURATION_FLOOR; + + let dd_max: [f32; 2] = [0.80, 0.90]; + let n_envs = dd_max.len() as i32; + let tile = build_dd_state_tile(&dd_max); + + let tile_buf = unsafe { MappedF32Buffer::new(tile.len()) }.expect("alloc tile"); + tile_buf.write_from_slice(&tile); + let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); + isv_buf.write_from_slice(&isv); + + launch_dd_floor( + &stream, &kernel, + tile_buf.dev_ptr, n_envs, DD_STATE_STRIDE, DD_MAX_OFF, + isv_buf.dev_ptr, DD_SATURATION_FLOOR_ADAPTIVE_INDEX as i32, + SENTINEL_DD_SATURATION_FLOOR, DD_SATURATION_FLOOR_SAFETY_FACTOR, + DD_SATURATION_FLOOR_MIN, DD_SATURATION_FLOOR_MAX, + DD_SATURATION_FLOOR_EMA_ALPHA, + ); + + let result = isv_buf.read_all(); + let floor = result[DD_SATURATION_FLOOR_ADAPTIVE_INDEX]; + + // Pearl-A: sentinel → REPLACE with clamped target. Should hit MAX=0.50. + assert!( + (floor - DD_SATURATION_FLOOR_MAX).abs() < 1e-5, + "DD floor must clamp to MAX=0.50 on extreme DD; got {floor}" + ); + } + + /// Test 4 — Welford slow EMA after bootstrap. + /// + /// Seed slot at 0.30 (NOT sentinel — past the bootstrap). Run with + /// 4 envs at DD_MAX=[0.10, 0.10, 0.10, 0.10] (mean=0.10, sigma=0). + /// Robust target = max(0.10+0×Z_75, 0.10) × 1.5 = 0.15. Clamped to + /// [0.10, 0.50] → stays at 0.15. + /// EMA blend: blended = (1-0.01)×0.30 + 0.01×0.15 = 0.297 + 0.0015 = 0.2985. + #[test] + #[ignore = "requires GPU"] + fn dd_floor_welford_ema_after_bootstrap() { + let stream = make_stream(); + let kernel = load_kernel(&stream); + + const ISV_DIM: usize = 1024; + const SEED: f32 = 0.30; // NOT sentinel — Pearl-A guard fires α-blend. + let mut isv = vec![0.0_f32; ISV_DIM]; + isv[DD_SATURATION_FLOOR_ADAPTIVE_INDEX] = SEED; + + let dd_max: [f32; 4] = [0.10, 0.10, 0.10, 0.10]; + let n_envs = dd_max.len() as i32; + let tile = build_dd_state_tile(&dd_max); + + let tile_buf = unsafe { MappedF32Buffer::new(tile.len()) }.expect("alloc tile"); + tile_buf.write_from_slice(&tile); + let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); + isv_buf.write_from_slice(&isv); + + launch_dd_floor( + &stream, &kernel, + tile_buf.dev_ptr, n_envs, DD_STATE_STRIDE, DD_MAX_OFF, + isv_buf.dev_ptr, DD_SATURATION_FLOOR_ADAPTIVE_INDEX as i32, + SENTINEL_DD_SATURATION_FLOOR, DD_SATURATION_FLOOR_SAFETY_FACTOR, + DD_SATURATION_FLOOR_MIN, DD_SATURATION_FLOOR_MAX, + DD_SATURATION_FLOOR_EMA_ALPHA, + ); + + let result = isv_buf.read_all(); + let floor = result[DD_SATURATION_FLOOR_ADAPTIVE_INDEX]; + + // EMA blend: (1-0.01) × 0.30 + 0.01 × 0.15 = 0.2985. + let expected = (1.0_f32 - DD_SATURATION_FLOOR_EMA_ALPHA) * SEED + + DD_SATURATION_FLOOR_EMA_ALPHA * 0.15; + assert!( + (floor - expected).abs() < 1e-4, + "Welford EMA blend: expected ≈ {expected:.5}, got {floor:.5}" + ); + } +} diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index c26435324..c364f5bd1 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,71 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +## 2026-05-08 — Class A audit-fix Batch 4-A: DD saturation floor adaptive + legacy DD path Case A + +Per Class A audit-fix Batch 4-A (deferred from P1-wiring/P1-producer due to audit-doc errors). Fixes 2 of 4 deferred items; Batch 4-B handles plan_threshold floor + MIN_HOLD_TEMPERATURE in a separate commit. + +### Item 1 — DD saturation floor adaptive (NEW slot 458) + +**Audit-doc error flagged:** the original audit pointed at `trade_physics.cuh:548` for the saturation floor, but inspection shows that line is the magnitude action constant (`magnitude = (mag_idx == MAG_QUARTER) ? 0.25f : ...`). The actual saturation floor lives at `trade_physics.cuh:154` inside `apply_margin_cap`: + +```c +float dd_scale = fmaxf(0.05f, 1.0f - dd_frac / 0.25f); /* floor at 5% to avoid zero */ +``` + +The hardcoded `0.25f` is the upper end of the linear position-size scaling ramp — the DD level at which `dd_scale` floors out at 5%. With a static 25% saturation floor, the ramp is calibrated for one specific DD distribution. Production training across 11 SPs has revealed wide DD variance (median fold DD_MAX=8% on calm folds, 35% on volatile folds), and the static breakpoint either over-throttles (calm folds where small DDs unnecessarily clamp positions) or under-throttles (volatile folds where 35% DDs hit the 5% floor on most envs). + +**Fix:** +- NEW slot `DD_SATURATION_FLOOR_ADAPTIVE_INDEX=458` (post-existing P1 block at [454..458)). +- Producer kernel `dd_saturation_floor_update_kernel.cu` aggregates per-env DD_MAX from the existing `sp15_dd_state_per_env[n_envs * 6]` tile (offset 1, populated by `dd_state_kernel`), computes Welford `mean + Z_75 × sigma` p75 estimator with `max(p75, mean)` robustness guard × 1.5 safety factor. +- Pearl-A first-observation bootstrap (sentinel 0.25 matches pre-fix hardcoded value for bit-identical cold-start) + α=0.01 slow EMA (DD distribution is a slow-moving fold-volatility property — same cadence as P0-A REWARD_POS_CAP). +- Bounds [0.10, 0.50] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`. Below 10% the ramp would clip on noise; above 50% the ramp barely activates before the capital-floor circuit-breaker triggers at 25% PDT. +- Threaded `isv_signals_ptr` into `apply_margin_cap` with NULL-tolerant cold-start fallback to `DD_SATURATION_FLOOR_DEFAULT=0.25` (mirrors the SP12-cap fallback pattern in `sp15_apply_sp12_cap`). +- Distinct semantic role from `SP15_DD_THRESHOLD_INDEX=421` (the SP15 quadratic DD-penalty *trigger* threshold, a *lower* bound). +- Per-epoch boundary launch in `training_loop.rs` right after `launch_reward_cap_update`. +- FoldReset registry entry + dispatch arm wired (per the C.10 lesson). +- 4 oracle tests cover Pearl-A bootstrap, no-DD guard, bounds clamp, Welford EMA after bootstrap. + +### Item 2 — Legacy compute_drawdown_penalty path: Case A (DELETED) + +**Verification before editing:** searched for all consumers of `compute_drawdown_penalty`. Single live call at `experience_kernels.cu:3822` inside the SP11-composer reward path. SP15's `sp15_dd_penalty` in `compute_sp15_final_reward_kernel.cu:154` runs unconditionally as a post-modifier on the SP11-composed reward (gated only on `isv_signals_dev_ptr != 0 && sp15_alpha_warm_count_dev_ptr != 0`, both populated in production), with ISV-driven λ_dd (slot 420) and DD threshold (slot 421). + +**Decision:** Case A — DELETE. Layering the legacy linear-ramp penalty inside the SP11 composer on top of the SP15 quadratic creates double-counting of DD shaping — exactly the code-smell the Class A audit was designed to eliminate. Per `feedback_no_legacy_aliases.md` and `feedback_no_partial_refactor.md`, when functionality is superseded by SP15, the legacy path must be deleted atomically across all consumers in the same commit. + +**Atomic deletion across:** +- `compute_drawdown_penalty` device function (`trade_physics.cuh`). +- Single call site at `experience_kernels.cu:3822` (the call + the dd_penalty_scale lookup that fed it). +- `dd_threshold` and `w_dd` kernel arguments (`experience_kernels.cu:1906-1907`). +- `w_dd` Rust config field (`gpu_experience_collector.rs:184`, `trainers/dqn/config.rs:441` DQNHyperparameters). +- `w_dd` profile section + dispatch (`training_profile.rs` RewardSection field, OptimizableParameterRanges field, FixedRewardParameters field, ParamLookup `"w_dd"` arm, profile→hyperparam mapping `if let Some(v) = rw.w_dd`, test assertion). +- `w_dd *= rki` risk-intensity multiplier (`trainers/dqn/config.rs:1263`). +- `w_dd` TOML keys (`config/training/dqn-hyperopt.toml` × 2 — search range + fixed phase, `dqn-localdev.toml`, `dqn-production.toml`, `dqn-smoketest.toml`). +- Stale doc comments on `hyperopt/adapters/dqn.rs` risk_intensity field + `trainers/dqn/config.rs` risk_intensity field. + +**Survives intentionally:** `config.dd_threshold` Rust field — still consumed by `launch_sp15_dd_state` at `gpu_experience_collector.rs:5750` as the `dd_budget` for DD_PCT scaling in `dd_state_kernel`. Documented in field comment. + +### ISV_TOTAL_DIM bump + +458 → 459 (Item 1 adds 1 slot; Item 2 is pure deletion). Layout fingerprint seed updated atomically. + +### Verification + +- 16 sp14_oracle_tests pass (incl. 4 new Batch-4-A tests). +- 36 sp15_phase1_oracle_tests pass (SP15 DD path unaffected). +- 12 sp14_isv_slots layout tests pass (incl. new layout-locked + fits-in-ISV tests for slot 458). +- 4 state_reset_registry tests pass — `every_fold_and_soft_reset_entry_has_dispatch_arm` confirms the new `sp14_audit_4a_dd_saturation_floor_adaptive` registry arm is wired in `training_loop.rs::reset_named_state`. +- Workspace cargo check clean (only pre-existing unused-import warnings). + +### Cumulative WR-plateau fix series (this is commit 7) + +- Class C bug 1 + P0-B (8f218cab2) +- P0-C (316db416b) +- P0-A (394de7d43) +- P1 wiring (c4b6d6ef2) — 1 of 4 wireable +- P0-A downstream (657972a4b) +- P1 producer (87d597d5d) +- audit-fix 4-A (this commit) + ## 2026-05-08 — Q-side WR-plateau root cause: replay-buffer intent/realized mismatch + Kelly floor wiring gap Two independent bugs identified by Class C (frame-shift) + Class A (hardcoded bounds) audits, both implicated in the **WR-stuck-at-46-48%** plateau across 11 superprojects over months. diff --git a/docs/isv-slots.md b/docs/isv-slots.md index 999dc6f68..a7ff30089 100644 --- a/docs/isv-slots.md +++ b/docs/isv-slots.md @@ -464,3 +464,13 @@ producers and consumers independently without index collisions. **Greenfield sub-worktrees:** `.worktrees/sp15-phase0-egf-retune`, `.worktrees/sp15-phase1-honest-numbers`, `.worktrees/sp15-phase2a-test-scaffold` rebase onto this commit so each phase's producers/consumers see the slot map. **Producers/consumers status:** P.2 lands the slot reservation only; all 46 slots are zero-initialized until per-phase tasks land producers (Phase 0/1/3) and atomic consumer migration (Phase 3.X). Mirrors the SP4/SP5/SP11 pre-allocation pattern. + +## Class A audit-fix Batch 4-A — adaptive DD saturation floor (2026-05-08) + +| Range | Phase | Purpose | +| --- | --- | --- | +| `[458..459)` | Batch 4-A | DD_SATURATION_FLOOR_ADAPTIVE — replaces hardcoded `0.25f` in `trade_physics.cuh::apply_margin_cap` (the upper end of the linear position-size scaling ramp `dd_scale = max(0.05, 1.0 − dd_frac/saturation_floor)`). Producer `dd_saturation_floor_update_kernel.cu` aggregates per-env DD_MAX from `sp15_dd_state_per_env[n_envs * 6]` (offset 1) via Welford `mean + Z_75 × sigma` p75 estimator × 1.5 safety factor. Pearl-A bootstrap (sentinel 0.25 = pre-fix value) + α=0.01 slow EMA (per-epoch boundary). Bounds [0.10, 0.50] = Category-1 dimensional safety. Distinct from `SP15_DD_THRESHOLD_INDEX=421` (the SP15 quadratic DD-penalty *trigger* threshold, a *lower* bound). | + +**Atomic deletion alongside this addition:** legacy `compute_drawdown_penalty` device function in trade_physics.cuh + its single caller at experience_kernels.cu:3822 + the `w_dd` Rust config field + the `dd_threshold` and `w_dd` kernel args + all `w_dd` profile/TOML references. Per `feedback_no_legacy_aliases` + `feedback_no_partial_refactor` — SP15's quadratic asymmetric DD penalty in `compute_sp15_final_reward_kernel.cu:154` is the production-grade replacement; layering both creates double-counting. + +**ISV_TOTAL_DIM:** 458 → 459 (Item 1 adds 1 slot; Item 2 is pure deletion).