feat(class-a-p0a): REWARD_POS/NEG_CAP → ISV-driven adaptive caps from realized return distribution
Per Class A audit ranking, the highest-suspected-impact fix for the months-long WR-stuck-at-46-48% plateau across 11 superprojects. 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 this CAPPED buffer, so the controller cannot select for trades it cannot see. Selectivity gradient evaporates — small wins clip to +5 alongside large wins also clipping to +5. The state_layout comment lines 261-264 explicitly deferred Phase 2 (ISV-driven) "IF Phase 1 validation reveals adaptive need". Phase 1 has been running 11 SPs without budging WR — adaptive need revealed. Architecture: - 2 new ISV slots [452..454): REWARD_POS_CAP_ADAPTIVE, REWARD_NEG_CAP_ADAPTIVE. - Producer kernel reward_cap_update_kernel.cu: block-tree-reduce Welford `mean + Z_99 × sigma` p99 estimator over winning realized returns + conservative `max(p99, max_win)` takeover, × 1.5 safety factor → POS cap. NEG cap = -2 × POS cap (preserves Kahneman 2:1 asymmetry per pearl_audit_unboundedness_for_implicit_asymmetry — asymmetry stays, but moved from hardcoded scalar to producer-time multiplier; single source of truth, no consumer applies the 2× ratio itself). - Pearl-A first-observation bootstrap from sentinel (5.0 / -10.0, matching pre-P0-A hardcoded values for bit-identical cold-start). Welford EMA α=0.01 thereafter (slow blend — reward distribution is the foundation of training and shouldn't move fast). - Bounds: POS in [1, 50], NEG in [-100, -2] (Category-1 dimensional safety per feedback_isv_for_adaptive_bounds, NOT tuning). - 3 consumer sites migrated atomically per feedback_no_partial_refactor: experience_kernels.cu:3112-3114 (segment_complete cap), compute_sp15_final_reward_kernel.cu:163 (Stage 4 helper invocation), sp15_reward_axis_helpers.cuh:211 (sp15_apply_sp12_cap device fn signature change to take isv ptr). - Cold-start fallback: when ISV slot at sentinel OR outside [REWARD_POS_CAP_MIN_BOUND=1, REWARD_POS_CAP_MAX_BOUND=50], consumers fall back to original macros (still defined in state_layout.cuh). - 2 new device-ptr accessors on the experience collector (step_ret_per_sample_dev_ptr, trade_close_per_sample_dev_ptr) — reuses existing per-sample buffers; no new buffer allocated. - Per-epoch boundary launch (cold path) in training_loop.rs alongside launch_aux_horizon_chain. - Reset registry entries + dispatch arms in reset_named_state per the C.10 lesson (missing dispatch causes runtime crash). - Layout fingerprint seed updated: ISV_TOTAL_DIM 452→454 + AUX_PRED_HORIZON_BARS=450 + AVG_WIN_HOLD_TIME_BARS=451 (previously missing from seed) + REWARD_POS_CAP_ADAPTIVE=452 + REWARD_NEG_CAP_ADAPTIVE=453. Per feedback_isv_for_adaptive_bounds: every adaptive bound in ISV. Verification: - cargo check -p ml --tests --all-targets: clean (19 pre-existing warnings, 0 new). - sp14_isv_slots tests: 8/8 pass (4 layout + 4 fits-within). - sp14_oracle_tests with --features cuda --ignored: 8/8 pass (4 existing q_disagreement/dir_concat + 4 new P0-A tests covering Pearl-A bootstrap, no-winning-trades preservation, bounds clamping to [1,50], Welford α=0.01 EMA blend). - sp15_phase1_oracle_tests with --features cuda --ignored: 36/36 pass (no regression from sp15_apply_sp12_cap signature change). Cumulative WR-plateau fix series: - Class C bug 1 (8f218cab2): replay buffer intent→realized. - Class A P0-B (8f218cab2): Kelly warmup floor wiring. - Class A P0-C (316db416b): MIN_HOLD_TARGET adaptive. - Class A P0-A (this commit): REWARD_POS/NEG_CAP adaptive — restores upper-tail alpha to the controller's view. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -786,6 +786,26 @@ fn main() {
|
||||
// (no atomicAdd) per `feedback_no_atomicadd.md`. Pearl-A
|
||||
// first-observation bootstrap; α=0.05 EMA thereafter. Plan: §C.4b.
|
||||
"avg_win_hold_time_update_kernel.cu",
|
||||
// Class A P0-A (2026-05-08): adaptive REWARD_POS/NEG_CAP producer.
|
||||
// Single-block kernel that sweeps the per-epoch
|
||||
// `step_ret_per_sample` + `trade_close_per_sample` buffers
|
||||
// (populated by `unified_env_step_core` in
|
||||
// `experience_kernels.cu`), computes p99(winning_returns) ×
|
||||
// safety_factor=1.5 → POS cap, NEG = −2 × POS (Kahneman 2:1
|
||||
// asymmetry preserved at producer-time per
|
||||
// `pearl_audit_unboundedness_for_implicit_asymmetry` — single
|
||||
// source of truth, no consumer applies the multiplier itself).
|
||||
// Pearl-A first-observation bootstrap (sentinels 5.0/-10.0
|
||||
// match pre-P0-A hardcoded constants for bit-identical
|
||||
// cold-start); α=0.01 slow EMA thereafter. Bounds POS∈[1, 50],
|
||||
// NEG∈[-100, -2] are Category-1 dimensional safety floors per
|
||||
// `feedback_isv_for_adaptive_bounds`. Replaces hardcoded
|
||||
// REWARD_POS_CAP=+5.0f / REWARD_NEG_CAP=-10.0f
|
||||
// (state_layout.cuh:266-267) — Class A audit
|
||||
// highest-suspected-impact fix for the months-long
|
||||
// WR-stuck-at-46-48% plateau across 11 superprojects. Per-epoch
|
||||
// boundary launch.
|
||||
"reward_cap_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)
|
||||
|
||||
@@ -153,8 +153,14 @@ extern "C" __global__ void compute_sp15_final_reward_kernel(
|
||||
const float dd_thr = isv[SP15_DD_THRESHOLD_INDEX];
|
||||
r = sp15_dd_penalty(r, lambda_dd, dd_thr, dd_current);
|
||||
|
||||
/* Stage 4: SP12 v3 asymmetric bilateral cap via macros. */
|
||||
r = sp15_apply_sp12_cap(r);
|
||||
/* Stage 4: SP12 v3 asymmetric bilateral cap. Class A P0-A (2026-05-08):
|
||||
* caps now ISV-driven from `ISV[REWARD_POS_CAP_ADAPTIVE_INDEX=452]` /
|
||||
* `ISV[REWARD_NEG_CAP_ADAPTIVE_INDEX=453]`. Cold-start fallback: when
|
||||
* the ISV slot is at sentinel SENTINEL_REWARD_POS_CAP=5.0 (within
|
||||
* EPS), the helper falls back to `state_layout.cuh` macros for
|
||||
* bit-identical pre-P0-A behavior. The helper handles the fallback
|
||||
* internally — caller only passes the ISV pointer. */
|
||||
r = sp15_apply_sp12_cap(r, isv);
|
||||
|
||||
/* Final NaN guard mirrors experience_kernels.cu's pattern: if a
|
||||
* helper produced NaN/Inf, fall back to the SP11 input value
|
||||
|
||||
@@ -3102,15 +3102,40 @@ extern "C" __global__ void experience_env_step(
|
||||
* to make the bound itself asymmetric in the direction the prior
|
||||
* implicitly was, not to remove the cap.
|
||||
*
|
||||
* Constants live in state_layout.cuh as Invariant-1 numerical
|
||||
* anchors (Phase 1). Phase 2 (deferred) lifts the ratio to an
|
||||
* ISV-driven controller per feedback_isv_for_adaptive_bounds.
|
||||
*
|
||||
* Class A P0-A (2026-05-08): Phase 2 lifted. The hardcoded
|
||||
* REWARD_POS_CAP=+5.0f / REWARD_NEG_CAP=-10.0f
|
||||
* (state_layout.cuh:266-267) is now ISV-driven from
|
||||
* `ISV[REWARD_POS_CAP_ADAPTIVE_INDEX=452]` /
|
||||
* `ISV[REWARD_NEG_CAP_ADAPTIVE_INDEX=453]` — adaptively driven
|
||||
* by `reward_cap_update_kernel` from p99(winning_returns) × 1.5
|
||||
* with NEG=−2 × POS preserving Kahneman/Tversky 2:1 asymmetry
|
||||
* (single source of truth at producer; no consumer applies the
|
||||
* 2× ratio itself). Cold-start fallback: when the ISV slot is at
|
||||
* sentinel SENTINEL_REWARD_POS_CAP=5.0 (within EPS), fall back
|
||||
* to the original macro `REWARD_POS_CAP` constant — bit-identical
|
||||
* pre-P0-A behavior until the first valid observation lands.
|
||||
* Math factored into `compute_asymmetric_capped_pnl` in
|
||||
* `trade_physics.cuh` so the formula is testable in isolation
|
||||
* (oracle tests in `tests/sp12_reward_math_tests.rs`). */
|
||||
float pos_cap_eff = REWARD_POS_CAP;
|
||||
float neg_cap_eff = REWARD_NEG_CAP;
|
||||
if (isv_signals_ptr != NULL) {
|
||||
float isv_pos = isv_signals_ptr[REWARD_POS_CAP_ADAPTIVE_INDEX];
|
||||
float isv_neg = isv_signals_ptr[REWARD_NEG_CAP_ADAPTIVE_INDEX];
|
||||
/* Cold-start fallback: sentinel match → use macro fallback
|
||||
* (bit-identical pre-P0-A). Out-of-bounds defensive guard:
|
||||
* any value outside [REWARD_POS_CAP_MIN_BOUND,
|
||||
* REWARD_POS_CAP_MAX_BOUND] also falls back, defending
|
||||
* against malformed prior state. */
|
||||
if (fabsf(isv_pos - SENTINEL_REWARD_POS_CAP) >= 1e-6f
|
||||
&& isv_pos >= REWARD_POS_CAP_MIN_BOUND
|
||||
&& isv_pos <= REWARD_POS_CAP_MAX_BOUND) {
|
||||
pos_cap_eff = isv_pos;
|
||||
neg_cap_eff = isv_neg;
|
||||
}
|
||||
}
|
||||
float capped_pnl = compute_asymmetric_capped_pnl(
|
||||
base_reward, REWARD_NEG_CAP, REWARD_POS_CAP
|
||||
base_reward, neg_cap_eff, pos_cap_eff
|
||||
);
|
||||
if (trail_triggered) {
|
||||
r_trail = capped_pnl;
|
||||
|
||||
@@ -52,7 +52,7 @@ 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,
|
||||
AVG_WIN_HOLD_TIME_UPDATE_CUBIN, H_S2_AUX_RMS_EMA_CUBIN, REWARD_CAP_UPDATE_CUBIN,
|
||||
};
|
||||
|
||||
/// Hidden width of the aux trunk's first internal layer (Linear_1 → ELU
|
||||
@@ -618,6 +618,122 @@ impl AvgWinHoldTimeUpdateOps {
|
||||
}
|
||||
}
|
||||
|
||||
/// Class A P0-A (2026-05-08): adaptive REWARD_POS/NEG_CAP producer.
|
||||
///
|
||||
/// Sweeps the per-epoch `step_ret_per_sample` + `trade_close_per_sample`
|
||||
/// buffers (populated by `unified_env_step_core` in `experience_kernels.cu`),
|
||||
/// computes a Welford `mean + Z_99 × sigma` p99 estimator over winning
|
||||
/// realized returns plus a `max(winning_returns)` conservative takeover,
|
||||
/// applies a 1.5× safety factor, clamps to dimensional bounds [1, 50],
|
||||
/// and writes both `ISV[REWARD_POS_CAP_ADAPTIVE_INDEX=452]` (POS cap)
|
||||
/// and `ISV[REWARD_NEG_CAP_ADAPTIVE_INDEX=453]` (NEG = −2 × POS,
|
||||
/// preserving Kahneman/Tversky 2:1 loss-aversion asymmetry per
|
||||
/// `pearl_audit_unboundedness_for_implicit_asymmetry` — moved from
|
||||
/// hardcoded scalar to producer-time multiplier; SINGLE source of truth).
|
||||
///
|
||||
/// Single-block kernel; block-tree-reduce over the per-sample buffer
|
||||
/// (no atomicAdd per `feedback_no_atomicadd.md`). Pearl-A
|
||||
/// first-observation bootstrap (sentinel 5.0/-10.0 matches pre-fix
|
||||
/// hardcoded constants for bit-identical cold-start); α=0.01 slow EMA
|
||||
/// thereafter.
|
||||
///
|
||||
/// Per-epoch boundary launch — reward distribution is the foundation
|
||||
/// of training and shouldn't move fast; per-step would track sample noise.
|
||||
///
|
||||
/// # Pearls applied
|
||||
/// - `feedback_no_atomicadd.md` — single block, block-tree-reduce in shmem.
|
||||
/// - `pearl_first_observation_bootstrap.md` — sentinel = 5.0 → 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))`
|
||||
/// on POS, NEG derived clamp.
|
||||
/// - `feedback_isv_for_adaptive_bounds.md` — POS bounds [1, 50] are
|
||||
/// Category-1 dimensional safety, NOT tuning.
|
||||
#[allow(missing_debug_implementations)]
|
||||
pub(crate) struct RewardCapUpdateOps {
|
||||
update_kernel: CudaFunction,
|
||||
}
|
||||
|
||||
impl RewardCapUpdateOps {
|
||||
/// Block dim used by the producer kernel — must match `BLK_DIM` in
|
||||
/// `reward_cap_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(REWARD_CAP_UPDATE_CUBIN.to_vec())
|
||||
.map_err(|e| MLError::ModelError(format!("reward_cap_update cubin load: {e}")))?;
|
||||
let update_kernel = module
|
||||
.load_function("reward_cap_update")
|
||||
.map_err(|e| MLError::ModelError(format!("reward_cap_update load: {e}")))?;
|
||||
Ok(Self { update_kernel })
|
||||
}
|
||||
|
||||
/// Launch the adaptive reward-cap producer.
|
||||
///
|
||||
/// Args:
|
||||
/// - `step_ret_ptr`: f32 device ptr `[total_samples]` —
|
||||
/// `step_ret_core` raw signed per-step return.
|
||||
/// - `trade_close_ptr`: i32 device ptr `[total_samples]` — 1 iff
|
||||
/// exiting_trade || reversing_trade.
|
||||
/// - `total_samples`: N*L (B*T total samples this epoch).
|
||||
/// - `isv_ptr`: ISV[ISV_TOTAL_DIM] device pointer.
|
||||
/// - `pos_idx`: REWARD_POS_CAP_ADAPTIVE_INDEX (452).
|
||||
/// - `neg_idx`: REWARD_NEG_CAP_ADAPTIVE_INDEX (453).
|
||||
/// - `sentinel_pos`: SENTINEL_REWARD_POS_CAP (5.0).
|
||||
/// - `sentinel_neg`: SENTINEL_REWARD_NEG_CAP (-10.0).
|
||||
/// - `safety_factor`: REWARD_CAP_SAFETY_FACTOR (1.5).
|
||||
/// - `neg_to_pos_ratio`: REWARD_NEG_TO_POS_RATIO (2.0).
|
||||
/// - `pos_min`, `pos_max`: REWARD_POS_CAP_MIN/MAX (1.0, 50.0).
|
||||
/// - `alpha`: REWARD_CAP_EMA_ALPHA (0.01).
|
||||
#[allow(clippy::too_many_arguments)]
|
||||
pub(crate) fn launch(
|
||||
&self,
|
||||
stream: &Arc<CudaStream>,
|
||||
step_ret_ptr: u64,
|
||||
trade_close_ptr: u64,
|
||||
total_samples: i32,
|
||||
isv_ptr: u64,
|
||||
pos_idx: i32,
|
||||
neg_idx: i32,
|
||||
sentinel_pos: f32,
|
||||
sentinel_neg: f32,
|
||||
safety_factor: f32,
|
||||
neg_to_pos_ratio: f32,
|
||||
pos_min: f32,
|
||||
pos_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(&step_ret_ptr)
|
||||
.arg(&trade_close_ptr)
|
||||
.arg(&total_samples)
|
||||
.arg(&isv_ptr)
|
||||
.arg(&pos_idx)
|
||||
.arg(&neg_idx)
|
||||
.arg(&sentinel_pos)
|
||||
.arg(&sentinel_neg)
|
||||
.arg(&safety_factor)
|
||||
.arg(&neg_to_pos_ratio)
|
||||
.arg(&pos_min)
|
||||
.arg(&pos_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!("reward_cap_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
|
||||
|
||||
@@ -2099,6 +2099,17 @@ pub(crate) static AUX_HORIZON_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!(
|
||||
/// `gpu_aux_trunk::AvgWinHoldTimeUpdateOps`. Per-epoch boundary launch.
|
||||
pub(crate) static AVG_WIN_HOLD_TIME_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/avg_win_hold_time_update_kernel.cubin"));
|
||||
|
||||
/// Class A P0-A (2026-05-08): adaptive REWARD_POS/NEG_CAP producer cubin.
|
||||
/// Single-block kernel that sweeps the per-epoch `step_ret_per_sample` +
|
||||
/// `trade_close_per_sample` buffers, computes a Welford p99 estimator of
|
||||
/// realized winning returns × 1.5× safety factor → POS cap, and writes
|
||||
/// both `ISV[REWARD_POS_CAP_ADAPTIVE_INDEX=452]` and
|
||||
/// `ISV[REWARD_NEG_CAP_ADAPTIVE_INDEX=453]` (NEG = −2 × POS, Kahneman
|
||||
/// asymmetry preserved at producer-time). Pearl-A first-observation
|
||||
/// bootstrap; α=0.01 slow EMA thereafter. Loaded by
|
||||
/// `gpu_aux_trunk::RewardCapUpdateOps`. Per-epoch boundary launch.
|
||||
pub(crate) static REWARD_CAP_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/reward_cap_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
|
||||
@@ -2427,7 +2438,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 = 452; // 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. 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 = 454; // 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. 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).
|
||||
@@ -3499,7 +3510,9 @@ const fn layout_fingerprint_seed() -> &'static [u8] {
|
||||
DD_PERSISTENCE_MAX=443;\
|
||||
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;\
|
||||
ISV_TOTAL_DIM=450;\
|
||||
AUX_PRED_HORIZON_BARS=450;AVG_WIN_HOLD_TIME_BARS=451;\
|
||||
REWARD_POS_CAP_ADAPTIVE=452;REWARD_NEG_CAP_ADAPTIVE=453;\
|
||||
ISV_TOTAL_DIM=454;\
|
||||
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;\
|
||||
@@ -7502,6 +7515,16 @@ pub struct GpuDqnTrainer {
|
||||
/// (no atomicAdd); per-epoch boundary launch.
|
||||
avg_win_hold_time_update_ops: super::gpu_aux_trunk::AvgWinHoldTimeUpdateOps,
|
||||
|
||||
/// Class A P0-A (2026-05-08): adaptive REWARD_POS/NEG_CAP producer.
|
||||
/// Sweeps the per-epoch `step_ret_per_sample` + `trade_close_per_sample`
|
||||
/// buffers, computes p99(winning_returns) × 1.5 → POS cap, NEG = −2 × POS,
|
||||
/// writes both `ISV[REWARD_POS_CAP_ADAPTIVE_INDEX=452]` and
|
||||
/// `ISV[REWARD_NEG_CAP_ADAPTIVE_INDEX=453]`. Block-tree-reduce
|
||||
/// (no atomicAdd); Pearl-A first-observation bootstrap; α=0.01 slow
|
||||
/// EMA. Per-epoch boundary launch — reward distribution is the
|
||||
/// foundation of training and shouldn't move fast.
|
||||
reward_cap_update_ops: super::gpu_aux_trunk::RewardCapUpdateOps,
|
||||
|
||||
// ── Speculative inference cache ──
|
||||
/// Cached trunk output [B, SH2] from between-bar speculative forward.
|
||||
speculative_h_s2: CudaSlice<f32>,
|
||||
@@ -14558,6 +14581,68 @@ impl GpuDqnTrainer {
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// Class A P0-A (2026-05-08): launch the adaptive reward-cap producer
|
||||
/// at per-epoch boundary.
|
||||
///
|
||||
/// Sweeps the per-epoch `step_ret_per_sample` + `trade_close_per_sample`
|
||||
/// buffers (populated by `unified_env_step_core` in
|
||||
/// `experience_kernels.cu`), computes a Welford `mean + Z_99 × sigma`
|
||||
/// p99 estimator over winning realized returns plus a
|
||||
/// `max(winning_returns)` conservative takeover, applies a 1.5×
|
||||
/// safety factor, clamps to dimensional bounds [1, 50], and writes
|
||||
/// both `ISV[REWARD_POS_CAP_ADAPTIVE_INDEX=452]` (POS cap) and
|
||||
/// `ISV[REWARD_NEG_CAP_ADAPTIVE_INDEX=453]` (NEG = −2 × POS,
|
||||
/// preserving Kahneman/Tversky 2:1 loss-aversion asymmetry per
|
||||
/// `pearl_audit_unboundedness_for_implicit_asymmetry`).
|
||||
///
|
||||
/// Per-epoch boundary launch — reward distribution is the foundation
|
||||
/// of training and shouldn't move fast; per-step would track sample
|
||||
/// noise.
|
||||
///
|
||||
/// Pearls applied:
|
||||
/// - `pearl_first_observation_bootstrap.md` (sentinel = 5.0/-10.0)
|
||||
/// - `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 POS, NEG)
|
||||
/// - `feedback_isv_for_adaptive_bounds.md` (POS bounds [1, 50] are
|
||||
/// dimensional safety floors, not tuning)
|
||||
pub(crate) fn launch_reward_cap_update(
|
||||
&self,
|
||||
step_ret_dev_ptr: u64,
|
||||
trade_close_dev_ptr: u64,
|
||||
total_samples: i32,
|
||||
) -> Result<(), MLError> {
|
||||
use crate::cuda_pipeline::sp14_isv_slots::{
|
||||
REWARD_CAP_EMA_ALPHA, REWARD_CAP_SAFETY_FACTOR, REWARD_NEG_CAP_ADAPTIVE_INDEX,
|
||||
REWARD_NEG_TO_POS_RATIO, REWARD_POS_CAP_ADAPTIVE_INDEX, REWARD_POS_CAP_MAX,
|
||||
REWARD_POS_CAP_MIN, SENTINEL_REWARD_NEG_CAP, SENTINEL_REWARD_POS_CAP,
|
||||
};
|
||||
|
||||
debug_assert!(
|
||||
self.isv_signals_dev_ptr != 0,
|
||||
"launch_reward_cap_update: isv_signals_dev_ptr must be allocated"
|
||||
);
|
||||
|
||||
self.reward_cap_update_ops.launch(
|
||||
&self.stream,
|
||||
step_ret_dev_ptr,
|
||||
trade_close_dev_ptr,
|
||||
total_samples,
|
||||
self.isv_signals_dev_ptr,
|
||||
REWARD_POS_CAP_ADAPTIVE_INDEX as i32,
|
||||
REWARD_NEG_CAP_ADAPTIVE_INDEX as i32,
|
||||
SENTINEL_REWARD_POS_CAP,
|
||||
SENTINEL_REWARD_NEG_CAP,
|
||||
REWARD_CAP_SAFETY_FACTOR,
|
||||
REWARD_NEG_TO_POS_RATIO,
|
||||
REWARD_POS_CAP_MIN,
|
||||
REWARD_POS_CAP_MAX,
|
||||
REWARD_CAP_EMA_ALPHA,
|
||||
)?;
|
||||
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// SP8 (Fix 36, 2026-05-03): launch GPU-only `train_active_frac` canary
|
||||
/// producer.
|
||||
///
|
||||
@@ -22964,6 +23049,14 @@ impl GpuDqnTrainer {
|
||||
let avg_win_hold_time_update_ops =
|
||||
super::gpu_aux_trunk::AvgWinHoldTimeUpdateOps::new(&stream)?;
|
||||
|
||||
// Class A P0-A (2026-05-08): adaptive reward-cap producer. Pre-loads
|
||||
// the `CudaFunction` handle at construction per
|
||||
// `pearl_no_host_branches_in_captured_graph.md`. Per-epoch boundary
|
||||
// launch (slow-moving — reward distribution is the foundation of
|
||||
// training and shouldn't move fast).
|
||||
let reward_cap_update_ops =
|
||||
super::gpu_aux_trunk::RewardCapUpdateOps::new(&stream)?;
|
||||
|
||||
// ── Speculative inference cache ──
|
||||
let speculative_h_s2 = stream.alloc_zeros::<f32>(b * sh2)
|
||||
.map_err(|e| MLError::ModelError(format!("speculative_h_s2 alloc: {e}")))?;
|
||||
@@ -23795,6 +23888,8 @@ impl GpuDqnTrainer {
|
||||
// SP14 Layer C Phase C.4b (2026-05-08): aux horizon + winning-hold-time producers.
|
||||
aux_horizon_update_ops,
|
||||
avg_win_hold_time_update_ops,
|
||||
// Class A P0-A (2026-05-08): adaptive reward-cap producer.
|
||||
reward_cap_update_ops,
|
||||
speculative_h_s2,
|
||||
speculative_features,
|
||||
speculative_valid: false,
|
||||
|
||||
@@ -6081,6 +6081,29 @@ impl GpuExperienceCollector {
|
||||
ptr
|
||||
}
|
||||
|
||||
/// Class A P0-A (2026-05-08): raw device pointer to
|
||||
/// `step_ret_per_sample` `[alloc_episodes * alloc_timesteps]` —
|
||||
/// `step_ret_core` raw signed per-step return (populated by
|
||||
/// `unified_env_step_core` in `experience_kernels.cu`, line ~2779).
|
||||
/// Consumed by `reward_cap_update_kernel` at the per-epoch boundary
|
||||
/// to compute the p99 of winning realized returns for the adaptive
|
||||
/// REWARD_POS/NEG_CAP ISV slots.
|
||||
pub fn step_ret_per_sample_dev_ptr(&self) -> u64 {
|
||||
let (ptr, _guard) = self.step_ret_per_sample.device_ptr(&self.stream);
|
||||
ptr
|
||||
}
|
||||
|
||||
/// Class A P0-A (2026-05-08): raw device pointer to
|
||||
/// `trade_close_per_sample` `[alloc_episodes * alloc_timesteps]` —
|
||||
/// 1 iff (exiting_trade || reversing_trade) (populated by
|
||||
/// `unified_env_step_core` in `experience_kernels.cu`, line ~2780).
|
||||
/// Consumed by `reward_cap_update_kernel` to mask the realized-return
|
||||
/// distribution to actual trade-close events.
|
||||
pub fn trade_close_per_sample_dev_ptr(&self) -> u64 {
|
||||
let (ptr, _guard) = self.trade_close_per_sample.device_ptr(&self.stream);
|
||||
ptr
|
||||
}
|
||||
|
||||
/// Last experience count from the most recent `collect_experiences_gpu` call.
|
||||
pub fn last_experience_count(&self) -> usize { self.last_experience_count }
|
||||
|
||||
|
||||
264
crates/ml/src/cuda_pipeline/reward_cap_update_kernel.cu
Normal file
264
crates/ml/src/cuda_pipeline/reward_cap_update_kernel.cu
Normal file
@@ -0,0 +1,264 @@
|
||||
/* ══════════════════════════════════════════════════════════════════════════
|
||||
* Class A P0-A — adaptive REWARD_POS_CAP / REWARD_NEG_CAP producer (2026-05-08).
|
||||
*
|
||||
* Replaces the hardcoded `REWARD_POS_CAP=+5.0f` / `REWARD_NEG_CAP=-10.0f`
|
||||
* (state_layout.cuh:266-267) with ISV-driven adaptive bounds derived from
|
||||
* the realized winning-trade return distribution.
|
||||
*
|
||||
* Why this matters (per Class A audit, the highest-suspected-impact fix
|
||||
* for the months-long WR-stuck-at-46-48% plateau across 11 superprojects):
|
||||
*
|
||||
* EV(trade) = WR × min(realized_win, +5)
|
||||
* + (1−WR) × max(realized_loss, −10)
|
||||
*
|
||||
* With WR ≈ 0.46, the cap asymmetry is fully active. Wins clip to +5
|
||||
* regardless of magnitude. Controller signals (sharpe EMA, var_q,
|
||||
* q_gap) all derive from this CAPPED buffer, so the controller cannot
|
||||
* select for trades it cannot see. Selectivity gradient evaporates —
|
||||
* small wins clip to +5 alongside large wins also clipping to +5.
|
||||
*
|
||||
* The state_layout.cuh comment block (lines 261-264) explicitly
|
||||
* deferred Phase 2 (ISV-driven) "IF Phase 1 validation reveals
|
||||
* adaptive need". Phase 1 has been running across 11 SPs without
|
||||
* budging WR — adaptive need is revealed.
|
||||
*
|
||||
* Algorithm:
|
||||
*
|
||||
* Phase 1 (block-tree-reduce): each block of `BLK_DIM` threads sweeps
|
||||
* its tile of N=B*T samples, masking by `trade_close_per_sample[i]`
|
||||
* and partitioning by `step_ret_per_sample[i] > 0` (winning) vs
|
||||
* `< 0` (losing). Local block reduce in shared memory accumulates a
|
||||
* per-block sum-of-squares over winning returns AND a count, plus a
|
||||
* max-of-winning-returns. Single-block kernel — the per-epoch sample
|
||||
* buffer is bounded (≤ ~64k slots).
|
||||
*
|
||||
* Phase 2 (Pearl-A + EMA, single thread): if no winning trades
|
||||
* observed, keep ISV slots unchanged (zero contributors → no signal).
|
||||
* The "p99 estimator" without an actual histogram is the
|
||||
* one-step-of-Welford rule of thumb:
|
||||
*
|
||||
* p99(winning_returns) ≈ mean + Z_99 × sigma (Z_99 ≈ 2.326)
|
||||
*
|
||||
* for a roughly-normal positive tail, OR `max(winning_returns)` for
|
||||
* small-N regimes, whichever is larger. We do NOT histogram (no
|
||||
* atomicAdd; would also need full sample-buffer materialisation in
|
||||
* shmem). The mean+sigma estimator with the conservative max
|
||||
* takeover handles both heavy-tail outliers (max wins) and the
|
||||
* well-behaved bulk (Gaussian-ish wins). Per
|
||||
* `feedback_no_atomicadd.md`.
|
||||
*
|
||||
* POS_target = max(p99_estimate, max_win) × safety_factor (1.5×)
|
||||
*
|
||||
* NEG_target = − REWARD_NEG_TO_POS_RATIO × POS_target (= −2 × POS,
|
||||
* preserving Kahneman/Tversky loss-aversion asymmetry per
|
||||
* `pearl_audit_unboundedness_for_implicit_asymmetry`). Asymmetry
|
||||
* stays — moved from hardcoded scalar to producer-time multiplier.
|
||||
* SINGLE SOURCE OF TRUTH: every consumer reads its corresponding
|
||||
* ISV slot, no consumer applies the 2× ratio itself.
|
||||
*
|
||||
* Pearl-A first-observation bootstrap: if `current_pos == sentinel`
|
||||
* (5.0 within 1e-6f), REPLACE with target directly. Otherwise EMA
|
||||
* blend with α=REWARD_CAP_EMA_ALPHA (0.01, slow — reward
|
||||
* distribution is the foundation of training and shouldn't move
|
||||
* fast). Per `pearl_first_observation_bootstrap.md`.
|
||||
*
|
||||
* Bounds: POS clamped to [REWARD_POS_CAP_MIN=1.0, REWARD_POS_CAP_MAX=50.0]
|
||||
* (Category-1 dimensional safety per `feedback_isv_for_adaptive_bounds`).
|
||||
* NEG = −2 × POS by construction; resulting NEG ∈ [−100, −2].
|
||||
*
|
||||
* 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 winning
|
||||
* trades" 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` — POS/NEG ranges
|
||||
* [1, 50] / [−100, −2] are dimensional safety bounds, NOT tuning.
|
||||
*
|
||||
* - `pearl_symmetric_clamp_audit.md` — bilateral `fmaxf(lo, fminf(x, hi))`
|
||||
* clamp on POS_target before writing the ISV slot. NEG ratio
|
||||
* applied AFTER clamp on POS, then NEG bilaterally clamped to its
|
||||
* own [−100, −2] range as a defensive layer.
|
||||
*
|
||||
* Args:
|
||||
* step_ret_per_sample — `[total_samples]` device f32; raw
|
||||
* signed per-step return populated by
|
||||
* `unified_env_step_core` in
|
||||
* `experience_kernels.cu` (line ~2779).
|
||||
* trade_close_per_sample — `[total_samples]` device i32; 1 iff
|
||||
* exiting_trade || reversing_trade
|
||||
* (`experience_kernels.cu` line ~2780).
|
||||
* total_samples — N*L (B*T total samples in the epoch).
|
||||
* isv — `[ISV_TOTAL_DIM]` device f32, modified
|
||||
* in place at indices `pos_idx` and
|
||||
* `neg_idx`.
|
||||
* pos_idx — REWARD_POS_CAP_ADAPTIVE_INDEX (452).
|
||||
* neg_idx — REWARD_NEG_CAP_ADAPTIVE_INDEX (453).
|
||||
* sentinel_pos — SENTINEL_REWARD_POS_CAP (5.0).
|
||||
* sentinel_neg — SENTINEL_REWARD_NEG_CAP (-10.0).
|
||||
* safety_factor — REWARD_CAP_SAFETY_FACTOR (1.5).
|
||||
* neg_to_pos_ratio — REWARD_NEG_TO_POS_RATIO (2.0).
|
||||
* pos_min — REWARD_POS_CAP_MIN (1.0).
|
||||
* pos_max — REWARD_POS_CAP_MAX (50.0).
|
||||
* alpha — REWARD_CAP_EMA_ALPHA (0.01, slow).
|
||||
*
|
||||
* Launch: grid=(1, 1, 1), block=(BLK_DIM=256, 1, 1).
|
||||
* Shared memory: 4 * BLK_DIM floats:
|
||||
* sh_sum_pos — sum of positive (winning) returns
|
||||
* sh_sumsq_pos — sum-of-squares of positive returns
|
||||
* sh_count_pos — count of positive returns
|
||||
* sh_max_pos — max of positive returns
|
||||
* The count is summed as float for unified shmem layout — values are
|
||||
* bounded by total_samples ≤ 2^16 so f32 representation is exact for
|
||||
* counts up to 2^24.
|
||||
* ══════════════════════════════════════════════════════════════════════════ */
|
||||
|
||||
#include <cuda_runtime.h>
|
||||
|
||||
#define BLK_DIM 256
|
||||
#define EPS_F 1e-6f
|
||||
/* Z-score for ~p99 of a normal distribution (one-sided). Used as a
|
||||
* lightweight estimator: p99 ≈ mean + Z_99 × sigma. */
|
||||
#define Z_99 2.326f
|
||||
|
||||
extern "C" __global__
|
||||
void reward_cap_update(
|
||||
const float* __restrict__ step_ret_per_sample,
|
||||
const int* __restrict__ trade_close_per_sample,
|
||||
int total_samples,
|
||||
float* __restrict__ isv,
|
||||
int pos_idx,
|
||||
int neg_idx,
|
||||
float sentinel_pos,
|
||||
float sentinel_neg,
|
||||
float safety_factor,
|
||||
float neg_to_pos_ratio,
|
||||
float pos_min,
|
||||
float pos_max,
|
||||
float alpha)
|
||||
{
|
||||
/* Single-block kernel — only block 0 does anything. */
|
||||
if (blockIdx.x != 0) return;
|
||||
|
||||
extern __shared__ float shmem[];
|
||||
float* sh_sum_pos = shmem; /* [BLK_DIM] */
|
||||
float* sh_sumsq_pos = shmem + BLK_DIM; /* [BLK_DIM] */
|
||||
float* sh_count_pos = shmem + 2 * BLK_DIM; /* [BLK_DIM] */
|
||||
float* sh_max_pos = shmem + 3 * BLK_DIM; /* [BLK_DIM] */
|
||||
|
||||
int tid = threadIdx.x;
|
||||
|
||||
/* Phase 1a: stride-based sweep — accumulate stats on closed-trade
|
||||
* winning returns only. Losing closes are NOT stored (we don't
|
||||
* adapt the negative cap from realized losses; the 2:1 ratio is
|
||||
* spec'd-asymmetric, not data-driven). */
|
||||
float local_sum = 0.0f;
|
||||
float local_sumsq = 0.0f;
|
||||
float local_count = 0.0f;
|
||||
float local_max = 0.0f;
|
||||
for (int i = tid; i < total_samples; i += BLK_DIM) {
|
||||
int is_close = trade_close_per_sample[i]; /* 0 or 1 */
|
||||
float ret = step_ret_per_sample[i]; /* signed */
|
||||
if (is_close != 0 && ret > 0.0f) {
|
||||
local_sum += ret;
|
||||
local_sumsq += ret * ret;
|
||||
local_count += 1.0f;
|
||||
if (ret > local_max) local_max = ret;
|
||||
}
|
||||
}
|
||||
sh_sum_pos[tid] = local_sum;
|
||||
sh_sumsq_pos[tid] = local_sumsq;
|
||||
sh_count_pos[tid] = local_count;
|
||||
sh_max_pos[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_pos[tid] += sh_sum_pos[tid + s];
|
||||
sh_sumsq_pos[tid] += sh_sumsq_pos[tid + s];
|
||||
sh_count_pos[tid] += sh_count_pos[tid + s];
|
||||
sh_max_pos[tid] = fmaxf(sh_max_pos[tid], sh_max_pos[tid + s]);
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
|
||||
/* Phase 2: thread 0 finalises EMA + Pearl-A + clamp + writes BOTH slots. */
|
||||
if (tid != 0) return;
|
||||
|
||||
float total_sum = sh_sum_pos[0];
|
||||
float total_sumsq = sh_sumsq_pos[0];
|
||||
float total_count = sh_count_pos[0];
|
||||
float total_max = sh_max_pos[0];
|
||||
|
||||
/* Guard: no winning trades observed this epoch — keep ISV slots
|
||||
* unchanged (cold-start sentinel persists; consumers fall back to
|
||||
* `state_layout.cuh` constants per the cold-start behavior). */
|
||||
if (total_count <= 0.0f) {
|
||||
return;
|
||||
}
|
||||
|
||||
/* Welford-rule p99 estimator: mean + Z_99 × 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 p99_estimate = mean + Z_99 * sigma;
|
||||
|
||||
/* Conservative takeover by realized max: in small-N or heavy-tail
|
||||
* regimes the parametric estimate underestimates the true tail.
|
||||
* Using `max(p99_estimate, total_max)` ensures the cap does NOT
|
||||
* actively clip any observed winning return. */
|
||||
float pos_target = fmaxf(p99_estimate, total_max) * safety_factor;
|
||||
|
||||
/* Clamp POS to dimensional-safety bounds [pos_min, pos_max].
|
||||
* Bilateral per `pearl_symmetric_clamp_audit.md`. */
|
||||
pos_target = fmaxf(pos_min, fminf(pos_target, pos_max));
|
||||
|
||||
/* Pearl-A first-observation bootstrap on POS slot. The sentinel is
|
||||
* exactly SENTINEL_REWARD_POS_CAP=5.0; any value within EPS_F of
|
||||
* the sentinel is treated as sentinel. (Direct equality on f32
|
||||
* literals would also work, but `|x − sentinel| < eps` is the
|
||||
* established pattern for ISV slot bootstrap detection.) */
|
||||
float current_pos = isv[pos_idx];
|
||||
float blended_pos;
|
||||
if (fabsf(current_pos - sentinel_pos) < EPS_F) {
|
||||
blended_pos = pos_target;
|
||||
} else {
|
||||
blended_pos = (1.0f - alpha) * current_pos + alpha * pos_target;
|
||||
}
|
||||
/* Re-clamp post-blend (defensive — handles malformed prior state
|
||||
* outside [pos_min, pos_max]). */
|
||||
blended_pos = fmaxf(pos_min, fminf(blended_pos, pos_max));
|
||||
|
||||
/* NEG cap = −neg_to_pos_ratio × POS cap. Asymmetry preserved at
|
||||
* producer-time, single source of truth. NEG bounds derived from
|
||||
* POS bounds × ratio: NEG ∈ [−ratio × pos_max, −ratio × pos_min] =
|
||||
* [−100, −2] for the production constants. The bilateral clamp
|
||||
* here is a defensive layer (POS clamp above already guarantees
|
||||
* NEG sits within its derived range). */
|
||||
float blended_neg = -neg_to_pos_ratio * blended_pos;
|
||||
float neg_lo = -neg_to_pos_ratio * pos_max; /* −100 */
|
||||
float neg_hi = -neg_to_pos_ratio * pos_min; /* −2 */
|
||||
blended_neg = fmaxf(neg_lo, fminf(blended_neg, neg_hi));
|
||||
|
||||
/* Suppress unused-parameter warning for sentinel_neg — kept in the
|
||||
* signature for symmetry with the POS sentinel and to allow a
|
||||
* future Pearl-A variant on the NEG slot if the 2:1 ratio
|
||||
* spec ever evolves. */
|
||||
(void)sentinel_neg;
|
||||
|
||||
isv[pos_idx] = blended_pos;
|
||||
isv[neg_idx] = blended_neg;
|
||||
}
|
||||
@@ -105,6 +105,65 @@ pub const AUX_PRED_HORIZON_BARS_MAX: f32 = 240.0; // ceiling (rollout bu
|
||||
pub const SP14_C4B_SLOT_BASE: usize = 450;
|
||||
pub const SP14_C4B_SLOT_END: usize = 452;
|
||||
|
||||
// ── Class A P0-A (2026-05-08): adaptive REWARD_POS/NEG_CAP slots ─────────
|
||||
// Replaces hardcoded `REWARD_POS_CAP=+5.0f` / `REWARD_NEG_CAP=-10.0f`
|
||||
// (state_layout.cuh:266-267) with ISV-driven values fed by the
|
||||
// `reward_cap_update_kernel` producer at per-epoch boundary. Per Class A
|
||||
// audit, these caps were structurally clipping the upper tail of realized
|
||||
// alpha — the controller signals (sharpe EMA, var_q, q_gap) all derive
|
||||
// from the CAPPED buffer, so the controller cannot select for trades it
|
||||
// cannot see. Selectivity gradient evaporates.
|
||||
//
|
||||
// The state_layout comment lines 261-264 explicitly deferred Phase 2
|
||||
// (ISV-driven) "IF Phase 1 validation reveals adaptive need". 11
|
||||
// superprojects of WR-stuck-at-46-48% revealed adaptive need.
|
||||
//
|
||||
// Asymmetry rationale (Kahneman/Tversky 2:1 loss aversion, per
|
||||
// `pearl_audit_unboundedness_for_implicit_asymmetry`) IS preserved. The
|
||||
// 2:1 asymmetry comes from the producer maintaining the
|
||||
// `NEG_CAP = -2 * POS_CAP` ratio internally (single source of truth in the
|
||||
// kernel), NOT from a fixed scalar. The scalar values themselves adapt to
|
||||
// the realized winning-return distribution: POS=p99(winning_returns) ×
|
||||
// safety_factor=1.5, with bounds [1.0, 50.0] clamping POS and [-100.0,
|
||||
// -2.0] clamping NEG (Category-1 dimensional safety floors per
|
||||
// `feedback_isv_for_adaptive_bounds`, NOT tuning).
|
||||
//
|
||||
// Slot layout: [452..454) — appended after SP14 Layer C Phase C.4b's
|
||||
// AVG_WIN_HOLD_TIME_BARS_INDEX=451 to preserve the checkpoint layout
|
||||
// fingerprint compatibility (no slot insertion before existing slots).
|
||||
pub const REWARD_POS_CAP_ADAPTIVE_INDEX: usize = 452; // adaptive +cap (replaces hardcoded +5.0)
|
||||
pub const REWARD_NEG_CAP_ADAPTIVE_INDEX: usize = 453; // adaptive −cap (= −2 × POS, replaces hardcoded −10.0)
|
||||
|
||||
// Sentinels (initial values; reset at fold boundary). Match the
|
||||
// pre-Class-A-P0-A hardcoded constants so cold-start behavior is
|
||||
// bit-identical to the prior Phase 1 anchor before the first valid
|
||||
// observation lands.
|
||||
pub const SENTINEL_REWARD_POS_CAP: f32 = 5.0; // matches pre-Class-A-P0-A REWARD_POS_CAP
|
||||
pub const SENTINEL_REWARD_NEG_CAP: f32 = -10.0; // matches pre-Class-A-P0-A REWARD_NEG_CAP
|
||||
|
||||
// Producer kernel constants (Category-1 dimensional safety bounds).
|
||||
// POS bounds: [1.0, 50.0] — cap below 1 makes the entire reward signal
|
||||
// flat (selectivity collapses); cap above 50 lets a single adversarial
|
||||
// outlier dominate any per-batch EMA. NEG = -2 × POS by construction;
|
||||
// resulting NEG range is therefore [-100.0, -2.0].
|
||||
pub const REWARD_POS_CAP_MIN: f32 = 1.0;
|
||||
pub const REWARD_POS_CAP_MAX: f32 = 50.0;
|
||||
// 2:1 ratio (Kahneman/Tversky meta-analysis ~2.0-2.25) preserved at
|
||||
// producer time inside `reward_cap_update_kernel`. Single source of
|
||||
// truth — every consumer reads its cap directly from the corresponding
|
||||
// ISV slot, no consumer applies the multiplier itself.
|
||||
pub const REWARD_NEG_TO_POS_RATIO: f32 = 2.0;
|
||||
// Safety factor on the p99 win-return target — keeps headroom above the
|
||||
// realized 99th percentile so we don't actively clip the upper tail.
|
||||
pub const REWARD_CAP_SAFETY_FACTOR: f32 = 1.5;
|
||||
// EMA blend rate. Slow — reward distribution is the foundation of
|
||||
// training and shouldn't move fast. Pearl-A bootstrap replaces sentinel
|
||||
// directly on first valid observation; α=0.01 thereafter.
|
||||
pub const REWARD_CAP_EMA_ALPHA: f32 = 0.01;
|
||||
|
||||
pub const SP14_P0A_SLOT_BASE: usize = 452;
|
||||
pub const SP14_P0A_SLOT_END: usize = 454;
|
||||
|
||||
#[cfg(test)]
|
||||
mod tests {
|
||||
use super::*;
|
||||
@@ -181,4 +240,37 @@ mod tests {
|
||||
SP14_C4B_SLOT_END, ISV_TOTAL_DIM,
|
||||
);
|
||||
}
|
||||
|
||||
/// Lock Class A P0-A (2026-05-08) adaptive reward-cap slot layout.
|
||||
/// Two contiguous slots [452..454) — REWARD_POS_CAP_ADAPTIVE +
|
||||
/// REWARD_NEG_CAP_ADAPTIVE. The pre-existing hardcoded constants
|
||||
/// `REWARD_POS_CAP=+5.0f` / `REWARD_NEG_CAP=-10.0f` (state_layout.cuh)
|
||||
/// remain as cold-start fallback when the ISV slot is at sentinel.
|
||||
#[test]
|
||||
fn sp14_p0a_reward_cap_slot_layout_locked() {
|
||||
assert_eq!(SP14_P0A_SLOT_BASE, 452);
|
||||
assert_eq!(SP14_P0A_SLOT_END, 454);
|
||||
assert_eq!(REWARD_POS_CAP_ADAPTIVE_INDEX, 452);
|
||||
assert_eq!(REWARD_NEG_CAP_ADAPTIVE_INDEX, 453);
|
||||
// 2:1 asymmetry preserved at producer-time, not consumer-time.
|
||||
assert_eq!(REWARD_NEG_TO_POS_RATIO, 2.0);
|
||||
// Sentinels match pre-P0-A hardcoded constants for bit-identical
|
||||
// cold-start behavior before first valid observation.
|
||||
assert_eq!(SENTINEL_REWARD_POS_CAP, 5.0);
|
||||
assert_eq!(SENTINEL_REWARD_NEG_CAP, -10.0);
|
||||
// Category-1 dimensional safety: POS in [1, 50], NEG in [-100, -2].
|
||||
assert_eq!(REWARD_POS_CAP_MIN, 1.0);
|
||||
assert_eq!(REWARD_POS_CAP_MAX, 50.0);
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn all_sp14_p0a_slots_fit_within_isv_total_dim() {
|
||||
use crate::cuda_pipeline::gpu_dqn_trainer::ISV_TOTAL_DIM;
|
||||
assert!(
|
||||
SP14_P0A_SLOT_END <= ISV_TOTAL_DIM,
|
||||
"SP14_P0A_SLOT_END={} exceeds ISV_TOTAL_DIM={} — bus too small for Class A P0-A reward-cap slots; \
|
||||
bump ISV_TOTAL_DIM in gpu_dqn_trainer.rs (and update layout_fingerprint_seed()).",
|
||||
SP14_P0A_SLOT_END, ISV_TOTAL_DIM,
|
||||
);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -180,17 +180,51 @@ __device__ inline float sp15_dd_penalty(
|
||||
// --------------------------------------------------------------------
|
||||
// `sp15_apply_sp12_cap` — Stage 4: asymmetric bilateral clamp
|
||||
//
|
||||
// `r_capped = fmaxf(REWARD_NEG_CAP, fminf(REWARD_POS_CAP, r))`
|
||||
// `r_capped = fmaxf(neg_cap_eff, fminf(pos_cap_eff, r))`
|
||||
//
|
||||
// Per Q1 resolution: REWARD_NEG_CAP / REWARD_POS_CAP stay as
|
||||
// `state_layout.cuh` macros (NOT lifted to ISV) — spec'd constants per
|
||||
// the SP12 v3 design, not adaptive bounds. Asymmetric NEG=−10 / POS=+5
|
||||
// is the loss-aversion encoding per
|
||||
// `pearl_audit_unboundedness_for_implicit_asymmetry` and matches the
|
||||
// existing `compute_asymmetric_capped_pnl` helper in trade_physics.cuh.
|
||||
// Class A P0-A (2026-05-08): caps lifted from `state_layout.cuh`
|
||||
// macros to ISV-driven adaptive slots (the Q1 resolution comment
|
||||
// below was the Phase 1 design — Phase 2 is now active per the
|
||||
// audit-revealed adaptive need). The helper reads
|
||||
// `ISV[REWARD_POS_CAP_ADAPTIVE_INDEX=452]` /
|
||||
// `ISV[REWARD_NEG_CAP_ADAPTIVE_INDEX=453]` first and falls back to
|
||||
// the original macros when the POS slot is at sentinel
|
||||
// SENTINEL_REWARD_POS_CAP=5.0 (within EPS) — bit-identical pre-P0-A
|
||||
// cold-start behavior until the first valid observation lands.
|
||||
//
|
||||
// Asymmetric NEG = −2 × POS is the loss-aversion encoding per
|
||||
// `pearl_audit_unboundedness_for_implicit_asymmetry`. The 2:1 ratio
|
||||
// is preserved at producer-time inside `reward_cap_update_kernel`
|
||||
// (single source of truth) — the consumer here applies whatever NEG
|
||||
// value sits in slot 453 without re-multiplying.
|
||||
//
|
||||
// Defensive bounds guard: any POS value outside
|
||||
// [REWARD_POS_CAP_MIN_BOUND=1.0, REWARD_POS_CAP_MAX_BOUND=50.0] also
|
||||
// falls back to the macro, defending against malformed prior state.
|
||||
// Producer kernel clamps to the same range, so a healthy ISV slot
|
||||
// always passes.
|
||||
//
|
||||
// Per `pearl_symmetric_clamp_audit.md`: bilateral
|
||||
// `fmaxf(NEG, fminf(POS, x))` even though the bounds are
|
||||
// intentionally asymmetric (NEG ≈ −2 × POS).
|
||||
// --------------------------------------------------------------------
|
||||
__device__ inline float sp15_apply_sp12_cap(float r_in) {
|
||||
return fmaxf(REWARD_NEG_CAP, fminf(REWARD_POS_CAP, r_in));
|
||||
__device__ inline float sp15_apply_sp12_cap(
|
||||
float r_in,
|
||||
const float* __restrict__ isv
|
||||
) {
|
||||
float pos_cap_eff = REWARD_POS_CAP;
|
||||
float neg_cap_eff = REWARD_NEG_CAP;
|
||||
if (isv != nullptr) {
|
||||
const float isv_pos = isv[REWARD_POS_CAP_ADAPTIVE_INDEX];
|
||||
const float isv_neg = isv[REWARD_NEG_CAP_ADAPTIVE_INDEX];
|
||||
if (fabsf(isv_pos - SENTINEL_REWARD_POS_CAP) >= 1e-6f
|
||||
&& isv_pos >= REWARD_POS_CAP_MIN_BOUND
|
||||
&& isv_pos <= REWARD_POS_CAP_MAX_BOUND) {
|
||||
pos_cap_eff = isv_pos;
|
||||
neg_cap_eff = isv_neg;
|
||||
}
|
||||
}
|
||||
return fmaxf(neg_cap_eff, fminf(pos_cap_eff, r_in));
|
||||
}
|
||||
|
||||
#endif // SP15_REWARD_AXIS_HELPERS_CUH_
|
||||
|
||||
@@ -271,6 +271,34 @@
|
||||
#define MIN_HOLD_TEMPERATURE_END 5.0f
|
||||
#define MIN_HOLD_TEMPERATURE_DECAY 20.0f
|
||||
|
||||
// Class A P0-A (2026-05-08) — adaptive REWARD_POS/NEG_CAP ISV slots.
|
||||
// REWARD_POS_CAP / REWARD_NEG_CAP above remain as Phase-1 cold-start
|
||||
// fallback values; consumers read the ISV slots first and fall back to
|
||||
// the macros when the slot is at sentinel (bit-identical pre-P0-A
|
||||
// behavior until the first valid observation). Slot indices mirror
|
||||
// `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` (locked by
|
||||
// `sp14_p0a_reward_cap_slot_layout_locked` test). Producer kernel:
|
||||
// `reward_cap_update_kernel.cu`; sentinels match the macros above for
|
||||
// bit-identical bootstrap. Per Class A audit, the highest-suspected-
|
||||
// impact fix for the months-long WR-stuck-at-46-48% plateau across
|
||||
// 11 superprojects — controller signals (sharpe EMA, var_q, q_gap)
|
||||
// derive from the CAPPED reward buffer, so the controller cannot
|
||||
// select for trades it cannot see. Per
|
||||
// `pearl_audit_unboundedness_for_implicit_asymmetry`, the 2:1
|
||||
// asymmetry stays — moved from hardcoded scalar to producer-time
|
||||
// multiplier (single source of truth).
|
||||
#define REWARD_POS_CAP_ADAPTIVE_INDEX 452
|
||||
#define REWARD_NEG_CAP_ADAPTIVE_INDEX 453
|
||||
#define SENTINEL_REWARD_POS_CAP 5.0f
|
||||
#define SENTINEL_REWARD_NEG_CAP -10.0f
|
||||
// Defensive consumer-side bounds (mirror Rust constants
|
||||
// REWARD_POS_CAP_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 REWARD_POS_CAP_MIN_BOUND 1.0f
|
||||
#define REWARD_POS_CAP_MAX_BOUND 50.0f
|
||||
|
||||
// ── Compile-time checks ──
|
||||
static_assert(SL_PADDING_START + SL_PADDING_DIM == SL_STATE_DIM,
|
||||
"State layout dimensions must sum to SL_STATE_DIM");
|
||||
|
||||
@@ -1021,6 +1021,31 @@ impl StateResetRegistry {
|
||||
category: ResetCategory::FoldReset,
|
||||
description: "ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451] — SP14 Layer C Phase C.4b avg winning hold time EMA (bars). Aggregated from `hold_at_exit_per_sample` + `trade_profitable_per_sample` per-sample buffers (populated by `unified_env_step_core` in `experience_kernels.cu`) by `avg_win_hold_time_update_kernel` at per-epoch boundary. FoldReset sentinel SENTINEL_AVG_WIN_HOLD_TIME_BARS=0.0 — Pearl-A first-observation bootstrap. Drives the horizon producer for ISV[450]. Plan: §C.4b.",
|
||||
},
|
||||
// ── Class A P0-A (2026-05-08): adaptive REWARD_POS/NEG_CAP ──
|
||||
// Two slots [452, 453] replacing the hardcoded
|
||||
// `REWARD_POS_CAP=+5.0f` / `REWARD_NEG_CAP=-10.0f`
|
||||
// (state_layout.cuh:266-267) per Class A audit. Producer
|
||||
// kernel `reward_cap_update_kernel` computes
|
||||
// p99(winning_returns) × 1.5 → POS cap, NEG = −2 × POS,
|
||||
// writes both at per-epoch boundary. Pearl-A bootstrap
|
||||
// sentinels match pre-P0-A hardcoded constants for
|
||||
// bit-identical cold-start behavior — consumers fall back
|
||||
// to the macro when the ISV slot is at sentinel. Once the
|
||||
// first valid observation lands the adaptive path takes
|
||||
// over via Welford α=0.01 slow EMA (reward distribution
|
||||
// is the foundation of training; shouldn't move fast).
|
||||
// Per `feedback_isv_for_adaptive_bounds.md` and
|
||||
// `pearl_audit_unboundedness_for_implicit_asymmetry.md`.
|
||||
RegistryEntry {
|
||||
name: "sp14_p0a_reward_pos_cap_adaptive",
|
||||
category: ResetCategory::FoldReset,
|
||||
description: "ISV[REWARD_POS_CAP_ADAPTIVE_INDEX=452] — Class A P0-A adaptive positive reward cap (replaces hardcoded REWARD_POS_CAP=+5.0f from state_layout.cuh:266). Produced by `reward_cap_update_kernel` from p99(winning_returns) × safety_factor=1.5 over the per-epoch `step_ret_per_sample` + `trade_close_per_sample` buffers. FoldReset sentinel SENTINEL_REWARD_POS_CAP=5.0 — Pearl-A first-observation bootstrap (matches pre-P0-A hardcoded value for bit-identical cold-start). Welford α=0.01 slow EMA thereafter. Bounds [1, 50] — Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`. Consumed by experience_kernels.cu (segment_complete cap), compute_sp15_final_reward_kernel.cu (Stage 4 cap), and sp15_reward_axis_helpers.cuh (sp15_apply_sp12_cap helper). Per Class A audit (highest-suspected-impact fix for the months-long WR-stuck-at-46-48% plateau across 11 superprojects).",
|
||||
},
|
||||
RegistryEntry {
|
||||
name: "sp14_p0a_reward_neg_cap_adaptive",
|
||||
category: ResetCategory::FoldReset,
|
||||
description: "ISV[REWARD_NEG_CAP_ADAPTIVE_INDEX=453] — Class A P0-A adaptive negative reward cap (replaces hardcoded REWARD_NEG_CAP=-10.0f from state_layout.cuh:267). Produced by `reward_cap_update_kernel` as NEG = −REWARD_NEG_TO_POS_RATIO=2.0 × POS_CAP. The 2:1 ratio (Kahneman/Tversky meta-analysis ~2.0-2.25) is preserved at producer-time inside the kernel (single source of truth — no consumer applies the multiplier itself), per `pearl_audit_unboundedness_for_implicit_asymmetry`. FoldReset sentinel SENTINEL_REWARD_NEG_CAP=-10.0 (matches pre-P0-A hardcoded value for bit-identical cold-start). Resulting bounds [-100, -2] derived from POS bounds × ratio. Consumed alongside POS slot at the same 3 sites.",
|
||||
},
|
||||
// ── 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
|
||||
|
||||
@@ -411,6 +411,26 @@ impl DQNTrainer {
|
||||
) {
|
||||
tracing::warn!(epoch, "launch_aux_horizon_chain failed (non-fatal): {e}");
|
||||
}
|
||||
// Class A P0-A (2026-05-08): adaptive REWARD_POS/NEG_CAP
|
||||
// — per-epoch boundary launch. Sweeps the per-sample
|
||||
// `step_ret_per_sample` + `trade_close_per_sample` buffers
|
||||
// (populated by `unified_env_step_core` in
|
||||
// `experience_kernels.cu`), computes p99(winning_returns)
|
||||
// × 1.5 → POS cap, NEG = −2 × POS, writes both
|
||||
// ISV[452, 453]. Pearl-A first-observation bootstrap +
|
||||
// α=0.01 slow EMA. 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
|
||||
// — months-long WR-stuck-at-46-48% plateau across 11 SPs.
|
||||
let step_ret_ptr = collector.step_ret_per_sample_dev_ptr();
|
||||
let trade_close_ptr = collector.trade_close_per_sample_dev_ptr();
|
||||
if let Err(e) = fused.trainer().launch_reward_cap_update(
|
||||
step_ret_ptr,
|
||||
trade_close_ptr,
|
||||
total_samples,
|
||||
) {
|
||||
tracing::warn!(epoch, "launch_reward_cap_update failed (non-fatal): {e}");
|
||||
}
|
||||
}
|
||||
|
||||
// D.8 Plan 2 Task 6C: update TLOB regime focus EMA in ISV at epoch boundary.
|
||||
@@ -8129,6 +8149,40 @@ impl DQNTrainer {
|
||||
);
|
||||
}
|
||||
}
|
||||
// Class A P0-A (2026-05-08): adaptive REWARD_POS/NEG_CAP
|
||||
// sentinels. Both slots use Pearl-A first-observation
|
||||
// bootstrap — the FoldReset rewrites to sentinel so the
|
||||
// kernel's "current_pos == sentinel" check fires cleanly on
|
||||
// the new fold's first launch (avoids cross-fold EMA
|
||||
// contamination). Sentinel values match pre-P0-A hardcoded
|
||||
// REWARD_POS_CAP=+5.0 / REWARD_NEG_CAP=-10.0 for bit-identical
|
||||
// cold-start until the first valid observation lands.
|
||||
// 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.
|
||||
"sp14_p0a_reward_pos_cap_adaptive" => {
|
||||
if let Some(ref fused) = self.fused_ctx {
|
||||
use crate::cuda_pipeline::sp14_isv_slots::{
|
||||
REWARD_POS_CAP_ADAPTIVE_INDEX, SENTINEL_REWARD_POS_CAP,
|
||||
};
|
||||
fused.trainer().write_isv_signal_at(
|
||||
REWARD_POS_CAP_ADAPTIVE_INDEX,
|
||||
SENTINEL_REWARD_POS_CAP,
|
||||
);
|
||||
}
|
||||
}
|
||||
"sp14_p0a_reward_neg_cap_adaptive" => {
|
||||
if let Some(ref fused) = self.fused_ctx {
|
||||
use crate::cuda_pipeline::sp14_isv_slots::{
|
||||
REWARD_NEG_CAP_ADAPTIVE_INDEX, SENTINEL_REWARD_NEG_CAP,
|
||||
};
|
||||
fused.trainer().write_isv_signal_at(
|
||||
REWARD_NEG_CAP_ADAPTIVE_INDEX,
|
||||
SENTINEL_REWARD_NEG_CAP,
|
||||
);
|
||||
}
|
||||
}
|
||||
// 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
|
||||
|
||||
@@ -448,6 +448,345 @@ mod gpu {
|
||||
}
|
||||
}
|
||||
|
||||
// ═══════════════════════════════════════════════════════════════════════════
|
||||
// Class A P0-A (2026-05-08) — adaptive REWARD_POS/NEG_CAP producer tests.
|
||||
//
|
||||
// Verifies the `reward_cap_update_kernel` producer:
|
||||
// 1. Cold-start Pearl-A bootstrap: ISV stays at sentinel until first valid
|
||||
// observation, then REPLACES (no blend).
|
||||
// 2. 2:1 asymmetry preserved at producer-time: NEG_CAP = -2 × POS_CAP.
|
||||
// 3. Bounds enforced: POS in [1, 50], NEG in [-100, -2].
|
||||
// 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_p0a_reward_cap_gpu {
|
||||
use std::sync::Arc;
|
||||
|
||||
use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg};
|
||||
use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedI32Buffer};
|
||||
use ml::cuda_pipeline::sp14_isv_slots::{
|
||||
REWARD_CAP_EMA_ALPHA, REWARD_CAP_SAFETY_FACTOR, REWARD_NEG_CAP_ADAPTIVE_INDEX,
|
||||
REWARD_NEG_TO_POS_RATIO, REWARD_POS_CAP_ADAPTIVE_INDEX, REWARD_POS_CAP_MAX,
|
||||
REWARD_POS_CAP_MIN, SENTINEL_REWARD_NEG_CAP, SENTINEL_REWARD_POS_CAP,
|
||||
};
|
||||
|
||||
const REWARD_CAP_UPDATE_CUBIN: &[u8] =
|
||||
include_bytes!(concat!(env!("OUT_DIR"), "/reward_cap_update_kernel.cubin"));
|
||||
|
||||
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(REWARD_CAP_UPDATE_CUBIN.to_vec())
|
||||
.expect("load reward_cap_update_kernel cubin");
|
||||
module
|
||||
.load_function("reward_cap_update")
|
||||
.expect("load reward_cap_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_reward_cap(
|
||||
stream: &Arc<CudaStream>,
|
||||
kernel: &CudaFunction,
|
||||
step_ret_ptr: u64,
|
||||
trade_close_ptr: u64,
|
||||
total_samples: i32,
|
||||
isv_ptr: u64,
|
||||
pos_idx: i32,
|
||||
neg_idx: i32,
|
||||
sentinel_pos: f32,
|
||||
sentinel_neg: f32,
|
||||
safety_factor: f32,
|
||||
neg_to_pos_ratio: f32,
|
||||
pos_min: f32,
|
||||
pos_max: f32,
|
||||
alpha: f32,
|
||||
) {
|
||||
unsafe {
|
||||
stream
|
||||
.launch_builder(kernel)
|
||||
.arg(&step_ret_ptr)
|
||||
.arg(&trade_close_ptr)
|
||||
.arg(&total_samples)
|
||||
.arg(&isv_ptr)
|
||||
.arg(&pos_idx)
|
||||
.arg(&neg_idx)
|
||||
.arg(&sentinel_pos)
|
||||
.arg(&sentinel_neg)
|
||||
.arg(&safety_factor)
|
||||
.arg(&neg_to_pos_ratio)
|
||||
.arg(&pos_min)
|
||||
.arg(&pos_max)
|
||||
.arg(&alpha)
|
||||
.launch(LaunchConfig {
|
||||
grid_dim: (1, 1, 1),
|
||||
block_dim: (BLK_DIM, 1, 1),
|
||||
shared_mem_bytes: SMEM_BYTES,
|
||||
})
|
||||
.expect("launch reward_cap_update");
|
||||
}
|
||||
stream.synchronize().expect("sync after reward_cap_update");
|
||||
}
|
||||
|
||||
/// Test 1 — Pearl-A first-observation bootstrap.
|
||||
///
|
||||
/// 4 winning closes (all step_ret>0) with values [1.0, 2.0, 3.0, 4.0],
|
||||
/// all trade_close=1. Mean=2.5, sigma≈1.118, p99≈2.5+2.326×1.118≈5.10,
|
||||
/// max=4.0. Estimator picks `max(p99, max)=5.10`, then × 1.5 = 7.65.
|
||||
/// Clamped to [1, 50] → 7.65 stays.
|
||||
///
|
||||
/// Cold-start (current = sentinel 5.0): Pearl-A REPLACES → POS=7.65,
|
||||
/// NEG=-2 × 7.65 = -15.3.
|
||||
#[test]
|
||||
#[ignore = "requires GPU"]
|
||||
fn reward_cap_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: sentinels (matches pre-P0-A constants).
|
||||
isv[REWARD_POS_CAP_ADAPTIVE_INDEX] = SENTINEL_REWARD_POS_CAP;
|
||||
isv[REWARD_NEG_CAP_ADAPTIVE_INDEX] = SENTINEL_REWARD_NEG_CAP;
|
||||
|
||||
let step_ret: Vec<f32> = vec![1.0, 2.0, 3.0, 4.0];
|
||||
let trade_close: Vec<i32> = vec![1, 1, 1, 1];
|
||||
|
||||
let n = step_ret.len() as i32;
|
||||
let ret_buf = unsafe { MappedF32Buffer::new(step_ret.len()) }
|
||||
.expect("alloc step_ret");
|
||||
ret_buf.write_from_slice(&step_ret);
|
||||
// i32 buffer via MappedF32Buffer — repurpose: the kernel reads
|
||||
// `int*`, so we need a u32-aligned i32 slice. We use a separate
|
||||
// raw allocation for the i32 close mask.
|
||||
let close_buf = unsafe { MappedI32Buffer::new(trade_close.len()) }
|
||||
.expect("alloc trade_close");
|
||||
close_buf.write_from_slice(&trade_close);
|
||||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }
|
||||
.expect("alloc isv");
|
||||
isv_buf.write_from_slice(&isv);
|
||||
|
||||
launch_reward_cap(
|
||||
&stream, &kernel,
|
||||
ret_buf.dev_ptr, close_buf.dev_ptr, n, isv_buf.dev_ptr,
|
||||
REWARD_POS_CAP_ADAPTIVE_INDEX as i32,
|
||||
REWARD_NEG_CAP_ADAPTIVE_INDEX as i32,
|
||||
SENTINEL_REWARD_POS_CAP, SENTINEL_REWARD_NEG_CAP,
|
||||
REWARD_CAP_SAFETY_FACTOR, REWARD_NEG_TO_POS_RATIO,
|
||||
REWARD_POS_CAP_MIN, REWARD_POS_CAP_MAX,
|
||||
REWARD_CAP_EMA_ALPHA,
|
||||
);
|
||||
|
||||
let result = isv_buf.read_all();
|
||||
let pos = result[REWARD_POS_CAP_ADAPTIVE_INDEX];
|
||||
let neg = result[REWARD_NEG_CAP_ADAPTIVE_INDEX];
|
||||
|
||||
// Pearl-A: sentinel → REPLACE with target. Target = max(p99, max) × 1.5.
|
||||
// mean=2.5, var=(1+4+9+16)/4 - 6.25 = 7.5 - 6.25 = 1.25, sigma≈1.118
|
||||
// p99 ≈ 2.5 + 2.326 × 1.118 ≈ 5.10. max=4.0. Picks 5.10 × 1.5 ≈ 7.65.
|
||||
// Clamped to [1, 50] → stays at ~7.65.
|
||||
assert!(
|
||||
(pos - 7.65).abs() < 0.05,
|
||||
"Pearl-A bootstrap: expected POS≈7.65, got {pos}"
|
||||
);
|
||||
// 2:1 asymmetry preserved at producer-time: NEG = -2 × POS.
|
||||
assert!(
|
||||
(neg - (-2.0 * pos)).abs() < 1e-4,
|
||||
"NEG must equal -2 × POS at producer-time (Kahneman 2:1 single source of truth); got POS={pos}, NEG={neg}"
|
||||
);
|
||||
}
|
||||
|
||||
/// Test 2 — No winning trades → ISV slots preserved bit-exactly.
|
||||
///
|
||||
/// All trade_close=0 (no closes). Kernel must skip the EMA update
|
||||
/// (count==0 guard) and leave both slots at their seeded values.
|
||||
#[test]
|
||||
#[ignore = "requires GPU"]
|
||||
fn reward_cap_no_winning_trades_preserves_isv() {
|
||||
let stream = make_stream();
|
||||
let kernel = load_kernel(&stream);
|
||||
|
||||
const ISV_DIM: usize = 1024;
|
||||
const SEED_POS: f32 = 12.5;
|
||||
const SEED_NEG: f32 = -25.0;
|
||||
let mut isv = vec![0.0_f32; ISV_DIM];
|
||||
isv[REWARD_POS_CAP_ADAPTIVE_INDEX] = SEED_POS;
|
||||
isv[REWARD_NEG_CAP_ADAPTIVE_INDEX] = SEED_NEG;
|
||||
|
||||
// Step returns present but no closes — none should contribute.
|
||||
let step_ret: Vec<f32> = vec![3.0, 5.0, -2.0];
|
||||
let trade_close: Vec<i32> = vec![0, 0, 0];
|
||||
|
||||
let n = step_ret.len() as i32;
|
||||
let ret_buf = unsafe { MappedF32Buffer::new(step_ret.len()) }
|
||||
.expect("alloc step_ret");
|
||||
ret_buf.write_from_slice(&step_ret);
|
||||
let close_buf = unsafe { MappedI32Buffer::new(trade_close.len()) }
|
||||
.expect("alloc trade_close");
|
||||
close_buf.write_from_slice(&trade_close);
|
||||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }
|
||||
.expect("alloc isv");
|
||||
isv_buf.write_from_slice(&isv);
|
||||
|
||||
launch_reward_cap(
|
||||
&stream, &kernel,
|
||||
ret_buf.dev_ptr, close_buf.dev_ptr, n, isv_buf.dev_ptr,
|
||||
REWARD_POS_CAP_ADAPTIVE_INDEX as i32,
|
||||
REWARD_NEG_CAP_ADAPTIVE_INDEX as i32,
|
||||
SENTINEL_REWARD_POS_CAP, SENTINEL_REWARD_NEG_CAP,
|
||||
REWARD_CAP_SAFETY_FACTOR, REWARD_NEG_TO_POS_RATIO,
|
||||
REWARD_POS_CAP_MIN, REWARD_POS_CAP_MAX,
|
||||
REWARD_CAP_EMA_ALPHA,
|
||||
);
|
||||
|
||||
let result = isv_buf.read_all();
|
||||
assert_eq!(
|
||||
result[REWARD_POS_CAP_ADAPTIVE_INDEX], SEED_POS,
|
||||
"POS slot must be preserved bit-exactly when no winning closes",
|
||||
);
|
||||
assert_eq!(
|
||||
result[REWARD_NEG_CAP_ADAPTIVE_INDEX], SEED_NEG,
|
||||
"NEG slot must be preserved bit-exactly when no winning closes",
|
||||
);
|
||||
}
|
||||
|
||||
/// Test 3 — Bounds enforcement: huge winning return clamps POS to MAX=50.
|
||||
///
|
||||
/// Single huge winning close at 1000.0. p99 estimate ≈ 1000 + 0 = 1000,
|
||||
/// max=1000, × safety_factor=1.5 = 1500. Must clamp to POS_MAX=50.
|
||||
/// NEG = -2 × 50 = -100 (also clamped at the lower bound).
|
||||
#[test]
|
||||
#[ignore = "requires GPU"]
|
||||
fn reward_cap_bounds_clamp_extreme_outlier() {
|
||||
let stream = make_stream();
|
||||
let kernel = load_kernel(&stream);
|
||||
|
||||
const ISV_DIM: usize = 1024;
|
||||
let mut isv = vec![0.0_f32; ISV_DIM];
|
||||
isv[REWARD_POS_CAP_ADAPTIVE_INDEX] = SENTINEL_REWARD_POS_CAP;
|
||||
isv[REWARD_NEG_CAP_ADAPTIVE_INDEX] = SENTINEL_REWARD_NEG_CAP;
|
||||
|
||||
// Single huge winning close — would explode the cap if unclamped.
|
||||
let step_ret: Vec<f32> = vec![1000.0];
|
||||
let trade_close: Vec<i32> = vec![1];
|
||||
|
||||
let n = step_ret.len() as i32;
|
||||
let ret_buf = unsafe { MappedF32Buffer::new(step_ret.len()) }
|
||||
.expect("alloc step_ret");
|
||||
ret_buf.write_from_slice(&step_ret);
|
||||
let close_buf = unsafe { MappedI32Buffer::new(trade_close.len()) }
|
||||
.expect("alloc trade_close");
|
||||
close_buf.write_from_slice(&trade_close);
|
||||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }
|
||||
.expect("alloc isv");
|
||||
isv_buf.write_from_slice(&isv);
|
||||
|
||||
launch_reward_cap(
|
||||
&stream, &kernel,
|
||||
ret_buf.dev_ptr, close_buf.dev_ptr, n, isv_buf.dev_ptr,
|
||||
REWARD_POS_CAP_ADAPTIVE_INDEX as i32,
|
||||
REWARD_NEG_CAP_ADAPTIVE_INDEX as i32,
|
||||
SENTINEL_REWARD_POS_CAP, SENTINEL_REWARD_NEG_CAP,
|
||||
REWARD_CAP_SAFETY_FACTOR, REWARD_NEG_TO_POS_RATIO,
|
||||
REWARD_POS_CAP_MIN, REWARD_POS_CAP_MAX,
|
||||
REWARD_CAP_EMA_ALPHA,
|
||||
);
|
||||
|
||||
let result = isv_buf.read_all();
|
||||
let pos = result[REWARD_POS_CAP_ADAPTIVE_INDEX];
|
||||
let neg = result[REWARD_NEG_CAP_ADAPTIVE_INDEX];
|
||||
|
||||
// Bounds: POS clamped to MAX=50.
|
||||
assert!(
|
||||
(pos - REWARD_POS_CAP_MAX).abs() < 1e-4,
|
||||
"POS must clamp to MAX=50 on extreme outlier; got {pos}"
|
||||
);
|
||||
// NEG = -2 × 50 = -100 (lower bound by construction).
|
||||
let expected_neg = -REWARD_NEG_TO_POS_RATIO * REWARD_POS_CAP_MAX;
|
||||
assert!(
|
||||
(neg - expected_neg).abs() < 1e-4,
|
||||
"NEG must clamp to -2 × MAX = {expected_neg}; got {neg}"
|
||||
);
|
||||
}
|
||||
|
||||
/// Test 4 — Welford EMA blend after bootstrap.
|
||||
///
|
||||
/// Pre-seed POS=10.0 (NOT sentinel — bootstrap path NOT taken).
|
||||
/// Single winning close at 4.0. mean=4, sigma=0, p99=4, max=4.
|
||||
/// Target = max(p99, max) × 1.5 = 4 × 1.5 = 6. Clamped to [1, 50] → 6.
|
||||
/// EMA blend: (1 - α) × 10 + α × 6 with α=0.01:
|
||||
/// = 0.99 × 10 + 0.01 × 6 = 9.96. NEG = -2 × 9.96 = -19.92.
|
||||
#[test]
|
||||
#[ignore = "requires GPU"]
|
||||
fn reward_cap_welford_ema_after_bootstrap() {
|
||||
let stream = make_stream();
|
||||
let kernel = load_kernel(&stream);
|
||||
|
||||
const ISV_DIM: usize = 1024;
|
||||
const SEEDED_POS: f32 = 10.0;
|
||||
const SEEDED_NEG: f32 = -20.0;
|
||||
let mut isv = vec![0.0_f32; ISV_DIM];
|
||||
// NOT sentinel — Pearl-A skipped, EMA path taken.
|
||||
isv[REWARD_POS_CAP_ADAPTIVE_INDEX] = SEEDED_POS;
|
||||
isv[REWARD_NEG_CAP_ADAPTIVE_INDEX] = SEEDED_NEG;
|
||||
|
||||
let step_ret: Vec<f32> = vec![4.0];
|
||||
let trade_close: Vec<i32> = vec![1];
|
||||
|
||||
let n = step_ret.len() as i32;
|
||||
let ret_buf = unsafe { MappedF32Buffer::new(step_ret.len()) }
|
||||
.expect("alloc step_ret");
|
||||
ret_buf.write_from_slice(&step_ret);
|
||||
let close_buf = unsafe { MappedI32Buffer::new(trade_close.len()) }
|
||||
.expect("alloc trade_close");
|
||||
close_buf.write_from_slice(&trade_close);
|
||||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }
|
||||
.expect("alloc isv");
|
||||
isv_buf.write_from_slice(&isv);
|
||||
|
||||
launch_reward_cap(
|
||||
&stream, &kernel,
|
||||
ret_buf.dev_ptr, close_buf.dev_ptr, n, isv_buf.dev_ptr,
|
||||
REWARD_POS_CAP_ADAPTIVE_INDEX as i32,
|
||||
REWARD_NEG_CAP_ADAPTIVE_INDEX as i32,
|
||||
SENTINEL_REWARD_POS_CAP, SENTINEL_REWARD_NEG_CAP,
|
||||
REWARD_CAP_SAFETY_FACTOR, REWARD_NEG_TO_POS_RATIO,
|
||||
REWARD_POS_CAP_MIN, REWARD_POS_CAP_MAX,
|
||||
REWARD_CAP_EMA_ALPHA,
|
||||
);
|
||||
|
||||
let result = isv_buf.read_all();
|
||||
let pos = result[REWARD_POS_CAP_ADAPTIVE_INDEX];
|
||||
let neg = result[REWARD_NEG_CAP_ADAPTIVE_INDEX];
|
||||
|
||||
// EMA blend: 0.99 × 10 + 0.01 × 6 = 9.96.
|
||||
let expected_pos = (1.0 - REWARD_CAP_EMA_ALPHA) * SEEDED_POS
|
||||
+ REWARD_CAP_EMA_ALPHA * (4.0 * REWARD_CAP_SAFETY_FACTOR);
|
||||
assert!(
|
||||
(pos - expected_pos).abs() < 1e-3,
|
||||
"POS Welford EMA: expected {expected_pos}, got {pos}"
|
||||
);
|
||||
// 2:1 ratio preserved.
|
||||
assert!(
|
||||
(neg - (-REWARD_NEG_TO_POS_RATIO * pos)).abs() < 1e-4,
|
||||
"NEG must equal -2 × POS at producer; got POS={pos}, NEG={neg}"
|
||||
);
|
||||
}
|
||||
}
|
||||
|
||||
// ═══════════════════════════════════════════════════════════════════════════
|
||||
// Test B.8: Layout-fingerprint regression for direction Q-head input bump.
|
||||
//
|
||||
|
||||
@@ -7810,3 +7810,103 @@ Wired `ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451]` — Pearl-A-bootstrapped Welford E
|
||||
### Consumer audit
|
||||
|
||||
Only one consumer of `MIN_HOLD_TARGET` was found in the CUDA call path (the `compute_min_hold_penalty` call in `experience_kernels.cu`). No additional consumers discovered. `trade_physics.cuh` takes `min_hold_target` as a parameter — no constant reads — no change needed there. Per `feedback_no_partial_refactor`: complete.
|
||||
|
||||
---
|
||||
|
||||
## 2026-05-08 — Class A P0-A: REWARD_POS/NEG_CAP ISV-driven from realized return p99/p1 EMA
|
||||
|
||||
### Problem
|
||||
|
||||
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. Per Class A audit:
|
||||
|
||||
```
|
||||
EV(trade) = WR × min(realized_win, +5)
|
||||
+ (1−WR) × max(realized_loss, −10)
|
||||
```
|
||||
|
||||
With WR ≈ 0.46 (the months-long plateau), the cap asymmetry is fully active. Wins clip to +5 regardless of magnitude. Controller signals (sharpe EMA, var_q, q_gap) all derive from this CAPPED buffer — the controller cannot select for trades it cannot see. Selectivity gradient evaporates: small wins clip to +5 alongside large wins also clipping to +5.
|
||||
|
||||
The state_layout comment lines 261-264 explicitly deferred Phase 2 (ISV-driven) "IF Phase 1 validation reveals adaptive need". 11 superprojects of WR-stuck-at-46-48% revealed adaptive need.
|
||||
|
||||
### Fix
|
||||
|
||||
Two new ISV slots [452, 453] driven by a new producer kernel `reward_cap_update_kernel.cu`:
|
||||
|
||||
- **POS cap producer**: block-tree-reduce `mean + Z_99 × sigma` p99 estimator over winning realized returns from the per-epoch `step_ret_per_sample` × `trade_close_per_sample` mask, with conservative `max(p99_estimate, max_winning_return)` takeover, × 1.5 safety factor, clamped to [1, 50].
|
||||
- **NEG cap = −2 × POS cap**: 2:1 Kahneman/Tversky asymmetry (per `pearl_audit_unboundedness_for_implicit_asymmetry`) preserved at producer-time — single source of truth, no consumer applies the multiplier itself. Resulting NEG ∈ [−100, −2] derived from POS bounds × ratio.
|
||||
|
||||
### Cold-start behavior
|
||||
|
||||
- Sentinels (5.0 / −10.0) match pre-P0-A hardcoded constants for bit-identical cold-start until the first valid observation.
|
||||
- Pearl-A bootstrap: when `|isv_pos − sentinel| < EPS`, REPLACE with target directly (no blend).
|
||||
- After bootstrap: Welford α=0.01 slow EMA (reward distribution is the foundation of training; shouldn't move fast).
|
||||
- Consumer-side fallback: when ISV slot at sentinel OR outside [REWARD_POS_CAP_MIN_BOUND=1, REWARD_POS_CAP_MAX_BOUND=50], fall back to `state_layout.cuh` macro (defensive — defends against malformed prior state).
|
||||
|
||||
### Cadence
|
||||
|
||||
Per-epoch boundary launch (cold path). Reward distribution moves on policy-evolution timescales, not per-step; per-step would track sample noise.
|
||||
|
||||
### Source of trade-close return data
|
||||
|
||||
Existing per-sample buffers populated by `unified_env_step_core` in `experience_kernels.cu`:
|
||||
|
||||
- `step_ret_per_sample[N*L]` (line ~2779) — `step_ret_core` raw signed per-step return.
|
||||
- `trade_close_per_sample[N*L]` (line ~2780) — 1 iff `exiting_trade || reversing_trade`.
|
||||
|
||||
Mask `step_ret > 0 && trade_close == 1` selects realized winning closes. No new buffer added.
|
||||
|
||||
### Sites modified
|
||||
|
||||
| File | LOC delta | Change |
|
||||
|------|-----------|--------|
|
||||
| `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` | +63 | 2 slot constants + 2 sentinels + 7 producer constants + locked-layout test |
|
||||
| `crates/ml/src/cuda_pipeline/state_layout.cuh` | +27 | 2 slot defines + 2 sentinels + 2 bound defines (consumer-side fallback) |
|
||||
| `crates/ml/src/cuda_pipeline/reward_cap_update_kernel.cu` | +233 (new) | Producer kernel: block-tree-reduce p99 + Pearl-A + EMA + bounds |
|
||||
| `crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs` | +110 | `RewardCapUpdateOps` struct + cubin loader + launcher |
|
||||
| `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` | +83 | Cubin static + struct field + constructor wire + `launch_reward_cap_update` fn + ISV_TOTAL_DIM 452→454 + fingerprint seed update |
|
||||
| `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` | +22 | 2 device-ptr accessors (`step_ret_per_sample_dev_ptr`, `trade_close_per_sample_dev_ptr`) |
|
||||
| `crates/ml/src/cuda_pipeline/experience_kernels.cu` | +27 / −3 | Consumer 1: ISV read + sentinel-fallback at segment_complete cap |
|
||||
| `crates/ml/src/cuda_pipeline/compute_sp15_final_reward_kernel.cu` | +6 / −1 | Consumer 2: pass `isv` ptr to `sp15_apply_sp12_cap` helper |
|
||||
| `crates/ml/src/cuda_pipeline/sp15_reward_axis_helpers.cuh` | +35 / −12 | Consumer 3: device fn signature `(r_in, isv*)` + ISV read + sentinel-fallback |
|
||||
| `crates/ml/src/trainers/dqn/state_reset_registry.rs` | +24 | 2 FoldReset registry entries (`sp14_p0a_reward_pos_cap_adaptive`, `sp14_p0a_reward_neg_cap_adaptive`) |
|
||||
| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | +37 | 2 dispatch arms in `reset_named_state` (per C.10 lesson) + 1 launch call at per-epoch boundary |
|
||||
| `crates/ml/build.rs` | +21 | cubin manifest entry for `reward_cap_update_kernel.cu` |
|
||||
| `crates/ml/tests/sp14_oracle_tests.rs` | +280 | 4 GPU oracle tests: Pearl-A bootstrap, no-winning-trades preservation, bounds clamping, Welford EMA |
|
||||
|
||||
### Architectural decisions
|
||||
|
||||
1. **Asymmetry at producer, not consumer**: `NEG = −2 × POS` enforced inside `reward_cap_update_kernel`. Single source of truth — every consumer reads its corresponding ISV slot directly, no consumer applies the 2× ratio itself. Per `pearl_audit_unboundedness_for_implicit_asymmetry`.
|
||||
|
||||
2. **Slot ordering for layout fingerprint compat**: Appended at 452-453 (immediately after SP14 Layer C Phase C.4b's slot 451). No insertion in earlier slots; checkpoint compatibility preserved beyond ISV_TOTAL_DIM.
|
||||
|
||||
3. **Pearl-A bootstrap mandatory**: Cold-start sentinel REPLACES on first valid observation; no blend. Per `pearl_first_observation_bootstrap.md`.
|
||||
|
||||
4. **Safety bounds Category-1**: POS ∈ [1, 50] are dimensional safety floors per `feedback_isv_for_adaptive_bounds`, NOT tuning. Cap below 1 makes the entire reward signal flat (selectivity collapses); cap above 50 lets a single adversarial outlier dominate any per-batch EMA. Producer kernel and consumer-side fallback both enforce.
|
||||
|
||||
5. **No new buffer**: Reused existing `step_ret_per_sample` + `trade_close_per_sample` — no widening of the experience-collector tile.
|
||||
|
||||
6. **Dispatch arms for FoldReset**: Per the C.10 lesson (missing `reset_named_state` arm for `sp14_q_disagreement_variance_ema` caused a runtime crash), both `sp14_p0a_reward_pos_cap_adaptive` and `sp14_p0a_reward_neg_cap_adaptive` arms added in `training_loop.rs::reset_named_state`.
|
||||
|
||||
### Verification
|
||||
|
||||
- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --tests --all-targets` — clean (19 pre-existing warnings, 0 new).
|
||||
- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --lib --release sp14_isv_slots` — 8/8 pass (4 layout + 4 fits-within tests).
|
||||
- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests --release --features cuda -- --ignored --nocapture` — 8/8 pass (4 existing + 4 new P0-A oracle tests).
|
||||
- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests --release --features cuda -- --ignored --nocapture` — 36/36 pass (no regression from `sp15_apply_sp12_cap` signature change).
|
||||
|
||||
### Consumer audit
|
||||
|
||||
3 consumer sites identified by Class A audit, all migrated atomically per `feedback_no_partial_refactor`:
|
||||
|
||||
1. `experience_kernels.cu:3112-3114` — segment-complete asymmetric cap. Reads ISV[452, 453]; falls back to macros at sentinel.
|
||||
2. `compute_sp15_final_reward_kernel.cu:163` — Stage 4 SP12 cap. Now passes `isv` ptr into the helper.
|
||||
3. `sp15_reward_axis_helpers.cuh:211` — `sp15_apply_sp12_cap` device fn. Refactored signature `(r_in, isv*)`; reads ISV[452, 453] with sentinel-fallback to macros internally.
|
||||
|
||||
No additional consumers found via `grep -rnE "REWARD_POS_CAP|REWARD_NEG_CAP" crates/ml/src/cuda_pipeline/`. Remaining matches are doc comments (gpu_dqn_trainer.rs:1462-63, gpu_experience_collector.rs:5697) and `trade_physics.cuh` parameter-name comments (no constant reads).
|
||||
|
||||
### Cumulative WR-plateau fix series
|
||||
|
||||
- Class C bug 1 (`8f218cab2`): replay buffer intent→realized — corrects the buffer's ground truth.
|
||||
- Class A P0-B (`8f218cab2`): Kelly warmup floor wiring — keeps Kelly cap from going to 0 cold-start.
|
||||
- Class A P0-C (`316db416b`): MIN_HOLD_TARGET adaptive from AVG_WIN_HOLD_TIME — patience-matched holds.
|
||||
- Class A P0-A (this commit): REWARD_POS/NEG_CAP adaptive from realized return distribution — restores upper-tail alpha to the controller's view.
|
||||
|
||||
Reference in New Issue
Block a user