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

fix(class-a-audit-batch-4a): DD saturation floor adaptive + legacy DD path Case A

Browse Source 
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-08 10:41:07 +02:00
parent 87d597d5d7
commit 7e9a8f6ef1
21 changed files with 1191 additions and 93 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
config/training/dqn-hyperopt.toml
View File

@@ -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

2
config/training/dqn-localdev.toml
View File

@@ -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

2
config/training/dqn-production.toml
View File

@@ -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)

2
config/training/dqn-smoketest.toml
View File

@@ -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

29
crates/ml/build.rs
View File

@@ -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)

243
crates/ml/src/cuda_pipeline/dd_saturation_floor_update_kernel.cu Normal file
View File

@@ -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 <cuda_runtime.h>
#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;
}

54
crates/ml/src/cuda_pipeline/experience_kernels.cu
View File

@@ -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;

112
crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs
View File

@@ -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<CudaStream>) -> Result<Self, MLError> {
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<CudaStream>,
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::<f32>() 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

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

File diff suppressed because one or more lines are too long

37
crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
View File

@@ -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 }

92
crates/ml/src/cuda_pipeline/sp14_isv_slots.rs
View File

@@ -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,
);
}
}

28
crates/ml/src/cuda_pipeline/state_layout.cuh
View File

@@ -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");

84
crates/ml/src/cuda_pipeline/trade_physics.cuh
View File

@@ -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
);
}

4
crates/ml/src/hyperopt/adapters/dqn.rs
View File

@@ -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] ===

19
crates/ml/src/trainers/dqn/config.rs
View File

@@ -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,

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

@@ -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

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

@@ -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

30
crates/ml/src/training_profile.rs
View File

@@ -260,8 +260,10 @@ pub struct RewardSection {
pub w_dsr: Option<f64>,
/// Normalized PnL weight
pub w_pnl: Option<f64>,
/// Drawdown penalty weight
pub w_dd: Option<f64>,
// `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<f64>,
/// 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<f64>,
pub w_pnl: Option<f64>,
pub w_dd: Option<f64>,
// `w_dd` field DELETED (Class A audit-fix Batch 4-A, 2026-05-08).
pub w_idle: Option<f64>,
pub dd_threshold: Option<f64>,
pub loss_aversion: Option<f64>,
@@ -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<f64>),
/// 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<f64>),
}
@@ -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);

293
crates/ml/tests/sp14_oracle_tests.rs
View File

@@ -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<CudaStream> {
let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
ctx.default_stream()
}
fn load_kernel(stream: &Arc<CudaStream>) -> 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::<f32>() as u32;
/// Helper: launch with a fixed config.
#[allow(clippy::too_many_arguments)]
fn launch_dd_floor(
stream: &Arc<CudaStream>,
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<f32> {
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}"
);
}
}

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

@@ -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.

10
docs/isv-slots.md
View File

@@ -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).