From 5ea5aa9b8e8a890fdc953698303c73b1f0ce229c Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sat, 9 May 2026 11:05:03 +0200 Subject: [PATCH] feat(sp18 v2 P3.T1-T5): adaptive HOLD_REWARD_POS/NEG_CAP producer kernel MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Lifts the Phase 2 caps from sentinel-driven cold-start (5.0/-10.0 fixed) to p99(|step_ret|) over Long/Short trade closes × 1.5 safety factor with Wiener-optimal alpha blend. The Phase 2 consumer (compute_sp18_hold_opportunity_cost) sees producer-driven slots [483]/[484] from epoch 1 onward; sentinel branch in the consumer remains as the cold-start fallback path (zero-trade-close epoch ⇒ producer early-returns ⇒ slots stay at sentinel ⇒ consumer falls through to the +5/-10 macro defaults — bit-identical to pre-Phase-3). Mirrors SP14 P0-A reward_cap_update_kernel structural template with three differences: (1) filter (is_close && step_ret != 0) — both winners AND losers (the consumer needs the magnitude scale of *all* Long/Short closes); (2) Welford-derived Wiener-α (slots [487..493)) replaces fixed α=0.01, with floor at WELFORD_ALPHA_MIN=0.4 per pearl_wiener_alpha_floor_for_nonstationary (the policy-realised distribution is intrinsically non-stationary as the policy adapts); (3) bounds [0.5, 50.0] (vs. position-side [1.0, 50.0]). Atomic single-commit per feedback_no_partial_refactor: - crates/ml/src/cuda_pipeline/hold_reward_cap_update_kernel.cu (NEW) - crates/ml/build.rs cubin manifest entry - HoldRewardCapUpdateOps in gpu_aux_trunk.rs (new struct + impl) - HOLD_REWARD_CAP_UPDATE_CUBIN static + struct field + launch_hold_reward_cap_update method + constructor instantiation + field-init in gpu_dqn_trainer.rs (5 sites) - Per-epoch boundary launch in training_loop.rs right AFTER launch_reward_cap_update (shared step_ret/trade_close source buffers, independent ISV slot pairs) - HEALTH_DIAG[N]: hold_reward_cap [pos={:.4} neg={:.4} fire_rate={:.4}] - 3 GPU oracle tests (T5 producer-drives-slots, Pearl-A REPLACE, no-closes preserves-isv) — all pass on local RTX 3050 Ti - Phase 3 close-out sections in docs/sp18-wireup-audit.md and docs/dqn-wire-up-audit.md Pearls applied: feedback_no_atomicadd, pearl_first_observation_bootstrap, pearl_wiener_optimal_adaptive_alpha, pearl_wiener_alpha_floor_for_nonstationary, pearl_no_host_branches_in_captured_graph, pearl_symmetric_clamp_audit, pearl_audit_unboundedness_for_implicit_asymmetry (NEG = -2 × POS at producer time, single source of truth), feedback_isv_for_adaptive_bounds, pearl_fused_per_group_statistics_oracle. Validation: cargo check --workspace clean; 3 GPU oracle tests pass on local RTX 3050 Ti; scripts/audit_sp18_consumers.sh --check exits 0 (no fingerprint drift in tracked sections). Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md Phase 4-5 (B-leg target-net forward + q_next replacement) follows. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/build.rs | 32 ++ crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs | 164 ++++++ .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 134 +++++ .../hold_reward_cap_update_kernel.cu | 394 ++++++++++++++ .../src/trainers/dqn/trainer/training_loop.rs | 64 +++ .../ml/tests/sp18_hold_reward_oracle_tests.rs | 482 ++++++++++++++++++ docs/dqn-wire-up-audit.md | 59 +++ docs/sp18-wireup-audit.md | 126 +++++ 8 files changed, 1455 insertions(+) create mode 100644 crates/ml/src/cuda_pipeline/hold_reward_cap_update_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index f930b30f3..1090db709 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -894,6 +894,38 @@ fn main() { // Pure observability: pre-fix baseline for the B-DD9 ratio // gate. Plan: SP18 Phase 0 Task 0.2. "td_error_mag_ema_kernel.cu", + // SP18 v2 Phase 3 (2026-05-09): D-leg adaptive HOLD_REWARD_POS/ + // NEG_CAP producer kernel. Single-block 256-thread kernel that + // sweeps the per-sample `step_ret_per_sample` + + // `trade_close_per_sample` buffers (populated by + // `unified_env_step_core` in `experience_kernels.cu` line ~2855), + // filters by `(is_close && step_ret != 0)` to extract the + // Long/Short trade-close magnitude distribution (DD3=b), computes + // a Welford `mean + Z_99 × sigma` p99 estimator with `max(p99, + // max_observed)` conservative takeover × HOLD_REWARD_CAP_SAFETY_ + // FACTOR=1.5, applies Wiener-optimal α blend (per + // `pearl_wiener_optimal_adaptive_alpha`; floor at WELFORD_ALPHA_ + // MIN=0.4 per `pearl_wiener_alpha_floor_for_nonstationary` since + // the policy-realised distribution is non-stationary), and + // writes both `ISV[HOLD_REWARD_POS_CAP_INDEX=483]` (POS cap) and + // `ISV[HOLD_REWARD_NEG_CAP_INDEX=484]` (NEG = −2 × POS, DD5=b + // mirrored asymmetry — same Kahneman 2:1 ratio as the position- + // side `reward_cap_update_kernel` for ISV[452]/[453], single + // source of truth at producer time per + // `pearl_audit_unboundedness_for_implicit_asymmetry`). Welford + // accumulators in slots [487..493). Pearl-A first-observation + // bootstrap on POS (sentinel 5.0/-10.0 matches the SP14 P0-A + // pattern for bit-identical cold-start). POS bounds [0.5, 50.0] + // are Category-1 dimensional safety floors per + // `feedback_isv_for_adaptive_bounds`; NEG ∈ [−100, −1] derived. + // Block-tree-reduce in shmem (no atomicAdd per + // `feedback_no_atomicadd.md`); `CudaFunction` pre-loaded at + // construction per `pearl_no_host_branches_in_captured_graph`. + // Loaded by `gpu_aux_trunk::HoldRewardCapUpdateOps`. Per-epoch + // boundary launch right AFTER `launch_reward_cap_update` so the + // two cap producers (position-side + Hold-side) share the same + // step_ret/trade_close source buffers. Plan: SP18 Phase 3. + "hold_reward_cap_update_kernel.cu", // SP18 v2 Phase 2 (2026-05-09): D-leg structural Hold // opportunity-cost device-fn GPU oracle test wrapper. Single // thread × single block; wraps `compute_sp18_hold_opportunity_cost` diff --git a/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs b/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs index ac1b691ad..e6c5f4559 100644 --- a/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs +++ b/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs @@ -53,6 +53,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, DD_SATURATION_FLOOR_UPDATE_CUBIN, + HOLD_REWARD_CAP_UPDATE_CUBIN, H_S2_AUX_RMS_EMA_CUBIN, KELLY_BAYESIAN_PRIORS_UPDATE_CUBIN, MIN_HOLD_TEMPERATURE_UPDATE_CUBIN, REWARD_CAP_UPDATE_CUBIN, @@ -737,6 +738,169 @@ impl RewardCapUpdateOps { } } +/// SP18 v2 Phase 3 (2026-05-09): adaptive HOLD_REWARD_POS/NEG_CAP producer. +/// +/// Drives `ISV[HOLD_REWARD_POS_CAP_INDEX=483]` and +/// `ISV[HOLD_REWARD_NEG_CAP_INDEX=484]` from `p99(|step_ret|)` over +/// Long/Short trade closes (per spec DD3=b). Mirrors `RewardCapUpdateOps` +/// for the position-side cap (ISV[452]/[453]) but with two structural +/// differences: +/// +/// 1. Filter is `(is_close && step_ret != 0)` (Long/Short closes — +/// both winners and losers) vs. `(is_close && step_ret > 0)` +/// (winners only) for the position-side cap. The Hold opp-cost +/// consumer needs the magnitude scale of *all* Long/Short +/// realised closes, not just winners. +/// +/// 2. EMA blend uses Wiener-optimal α derived from Welford +/// accumulators in slots [487..493) (per +/// `pearl_wiener_optimal_adaptive_alpha`; floor at +/// WELFORD_ALPHA_MIN=0.4 per +/// `pearl_wiener_alpha_floor_for_nonstationary` since the +/// policy-realised distribution is intrinsically non-stationary +/// as the policy adapts), vs. fixed α=0.01 for the position-side +/// cap (whose underlying alpha distribution is approximately +/// stationary). +/// +/// Pearl-A first-observation bootstrap (sentinel POS=5.0/NEG=-10.0 +/// matches the position-side cap pattern for bit-identical cold-start). +/// Block-tree-reduce (no atomicAdd per `feedback_no_atomicadd.md`) on +/// sum/sumsq/count/max in shmem; thread 0 finalises the +/// Welford+Wiener+Pearl-A+clamp pipeline and writes both POS and NEG +/// (NEG = −2 × POS, DD5=b mirrored asymmetry, single source of truth +/// at producer time per +/// `pearl_audit_unboundedness_for_implicit_asymmetry`). +/// +/// `CudaFunction` pre-loaded at construction per +/// `pearl_no_host_branches_in_captured_graph.md`; per-epoch boundary +/// launch right after `RewardCapUpdateOps` so both cap producers +/// share the same `step_ret_per_sample` + `trade_close_per_sample` +/// source buffers. +/// +/// # Pearls applied +/// - `feedback_no_atomicadd.md` — block-tree-reduce in shmem; one +/// global write per slot from thread 0. +/// - `pearl_first_observation_bootstrap.md` — sentinel match (within +/// 1e-6 of SENTINEL_HOLD_REWARD_POS_CAP=5.0) → REPLACE; otherwise +/// Wiener-α blend. +/// - `pearl_wiener_optimal_adaptive_alpha.md` — α derived from +/// Welford variance accumulators, not hardcoded. +/// - `pearl_wiener_alpha_floor_for_nonstationary.md` — α floored at +/// 0.4 to preserve catch-up bandwidth on the policy-realised +/// non-stationary signal. +/// - `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 pre-blend, post-blend, and +/// on derived NEG. +/// - `feedback_isv_for_adaptive_bounds.md` — POS bounds [0.5, 50.0] +/// are Category-1 dimensional safety floors, NOT tuning. NEG range +/// [−100, −1] is derived as `−ratio × POS_bounds`. +/// - `pearl_fused_per_group_statistics_oracle.md` — sum/sumsq/count/max +/// accumulated in ONE fused reduction; not separate launches. +#[allow(missing_debug_implementations)] +pub(crate) struct HoldRewardCapUpdateOps { + update_kernel: CudaFunction, +} + +impl HoldRewardCapUpdateOps { + /// Block dim used by the producer kernel — must match `HRC_BLK_DIM` + /// in `hold_reward_cap_update_kernel.cu`. + const BLK_DIM: u32 = 256; + + pub(crate) fn new(stream: &Arc) -> Result { + let context = stream.context(); + let module = context + .load_cubin(HOLD_REWARD_CAP_UPDATE_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("hold_reward_cap_update cubin load: {e}")))?; + let update_kernel = module + .load_function("hold_reward_cap_update") + .map_err(|e| MLError::ModelError(format!("hold_reward_cap_update load: {e}")))?; + Ok(Self { update_kernel }) + } + + /// Launch the SP18 D-leg adaptive Hold-reward-cap producer. + /// + /// Args: + /// - `step_ret_ptr`: f32 device ptr `[total_samples]` — + /// `step_ret_core` raw signed per-step return (same source the + /// position-side `RewardCapUpdateOps::launch` reads from). + /// - `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 (mapped-pinned). + /// - `pos_idx`: HOLD_REWARD_POS_CAP_INDEX (483). + /// - `neg_idx`: HOLD_REWARD_NEG_CAP_INDEX (484). + /// - Welford accumulator slot indices: HRC_TARGET_MEAN_INDEX (487), + /// HRC_TARGET_M2_INDEX (488), HRC_DIFF_MEAN_INDEX (489), + /// HRC_DIFF_M2_INDEX (490), HRC_PREV_TARGET_INDEX (491), + /// HRC_SAMPLE_COUNT_INDEX (492). + /// - `sentinel_pos`: SENTINEL_HOLD_REWARD_POS_CAP (5.0). + /// - `sentinel_neg`: SENTINEL_HOLD_REWARD_NEG_CAP (-10.0; held + /// for symmetry with the position-side kernel signature). + /// - `safety_factor`: HOLD_REWARD_CAP_SAFETY_FACTOR (1.5). + /// - `neg_to_pos_ratio`: REWARD_NEG_TO_POS_RATIO (2.0; same + /// Kahneman 2:1 ratio as the position-side cap, single source + /// of truth at producer time). + /// - `pos_min`/`pos_max`: HOLD_REWARD_POS_CAP_MIN/MAX_BOUND + /// (0.5, 50.0). + #[allow(clippy::too_many_arguments)] + pub(crate) fn launch( + &self, + stream: &Arc, + step_ret_ptr: u64, + trade_close_ptr: u64, + total_samples: i32, + isv_ptr: u64, + pos_idx: i32, + neg_idx: i32, + target_mean_idx: i32, + target_m2_idx: i32, + diff_mean_idx: i32, + diff_m2_idx: i32, + prev_target_idx: i32, + sample_count_idx: i32, + sentinel_pos: f32, + sentinel_neg: f32, + safety_factor: f32, + neg_to_pos_ratio: f32, + pos_min: f32, + pos_max: f32, + ) -> Result<(), MLError> { + // 4 arrays × BLK_DIM × sizeof(f32) — see kernel comment block. + let smem_bytes = 4 * Self::BLK_DIM * std::mem::size_of::() 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(&target_mean_idx) + .arg(&target_m2_idx) + .arg(&diff_mean_idx) + .arg(&diff_m2_idx) + .arg(&prev_target_idx) + .arg(&sample_count_idx) + .arg(&sentinel_pos) + .arg(&sentinel_neg) + .arg(&safety_factor) + .arg(&neg_to_pos_ratio) + .arg(&pos_min) + .arg(&pos_max) + .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!("hold_reward_cap_update: {e}")))?; + } + Ok(()) + } +} + /// Class A P1-Producer (2026-05-08): adaptive Bayesian Kelly priors producer. /// /// Replaces the previously-hardcoded `prior_wins=2.0f, prior_losses=2.0f, diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 4a273eaae..32ce26e99 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -2168,6 +2168,27 @@ pub(crate) static AVG_WIN_HOLD_TIME_UPDATE_CUBIN: &[u8] = include_bytes!(concat! /// `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")); +/// SP18 v2 Phase 3 (2026-05-09): adaptive HOLD_REWARD_POS/NEG_CAP producer +/// cubin. Single-block 256-thread kernel that sweeps the per-epoch +/// `step_ret_per_sample` + `trade_close_per_sample` buffers (same source +/// `reward_cap_update_kernel` reads from), filters by `(is_close && +/// step_ret != 0)` to extract Long/Short close magnitudes (DD3=b), +/// computes a Welford `mean + Z_99 × sigma` p99 estimator with +/// `max(p99, max_observed)` conservative takeover × HOLD_REWARD_CAP_ +/// SAFETY_FACTOR=1.5, applies Wiener-optimal α blend (Welford +/// accumulators in slots [487..493); floor at WELFORD_ALPHA_MIN=0.4 per +/// `pearl_wiener_alpha_floor_for_nonstationary`), and writes both +/// `ISV[HOLD_REWARD_POS_CAP_INDEX=483]` and +/// `ISV[HOLD_REWARD_NEG_CAP_INDEX=484]` (NEG = −2 × POS, DD5=b mirrored +/// asymmetry — same Kahneman 2:1 ratio as the position-side cap, single +/// source of truth at producer time per +/// `pearl_audit_unboundedness_for_implicit_asymmetry`). Pearl-A +/// first-observation bootstrap on POS (sentinels POS=5.0/NEG=-10.0 +/// match SP14 P0-A pattern). Loaded by +/// `gpu_aux_trunk::HoldRewardCapUpdateOps`. Per-epoch boundary launch +/// right AFTER `launch_reward_cap_update`. +pub(crate) static HOLD_REWARD_CAP_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/hold_reward_cap_update_kernel.cubin")); + /// Class A P1-Producer (2026-05-08): adaptive Bayesian Kelly priors producer cubin. /// Single-block 256-thread kernel that sweeps the per-fold-end /// `portfolio_state[n_envs, PS_STRIDE]` buffer (same source the @@ -7784,6 +7805,28 @@ pub struct GpuDqnTrainer { /// foundation of training and shouldn't move fast. reward_cap_update_ops: super::gpu_aux_trunk::RewardCapUpdateOps, + /// SP18 v2 Phase 3 (2026-05-09): adaptive HOLD_REWARD_POS/NEG_CAP + /// producer Ops handle. Pre-loads the `CudaFunction` at construction + /// per `pearl_no_host_branches_in_captured_graph.md`. Sweeps the + /// per-epoch `step_ret_per_sample` + `trade_close_per_sample` + /// buffers (same source `reward_cap_update_ops` reads from), filters + /// by `(is_close && step_ret != 0)` to extract Long/Short close + /// magnitudes per spec DD3=b, computes a Welford p99 estimator × + /// 1.5 safety factor → POS cap, derives Wiener-optimal α from + /// Welford accumulators in slots [487..493), and writes both + /// `ISV[HOLD_REWARD_POS_CAP_INDEX=483]` and + /// `ISV[HOLD_REWARD_NEG_CAP_INDEX=484]` (NEG = −2 × POS, DD5=b + /// mirrored asymmetry, same Kahneman 2:1 ratio as the position- + /// side cap). Block-tree-reduce (no atomicAdd); Pearl-A first- + /// observation bootstrap; Wiener-α floor at 0.4 per + /// `pearl_wiener_alpha_floor_for_nonstationary` since the policy- + /// realised distribution is intrinsically non-stationary as the + /// policy adapts. Per-epoch boundary launch right AFTER + /// `reward_cap_update_ops` so both cap producers share the same + /// step_ret/trade_close source buffers; no shared state — they + /// drive independent ISV slot pairs. + hold_reward_cap_update_ops: super::gpu_aux_trunk::HoldRewardCapUpdateOps, + /// Class A P1-Producer (2026-05-08): adaptive Bayesian Kelly priors /// producer. Sweeps `portfolio_state[n_envs, PS_STRIDE]` (same source /// `kelly_cap_update_kernel` reads), aggregates the realized @@ -15565,6 +15608,86 @@ impl GpuDqnTrainer { Ok(()) } + /// SP18 v2 Phase 3 (2026-05-09): launch the adaptive HOLD_REWARD_POS/ + /// NEG_CAP producer. Sweeps the per-epoch `step_ret_per_sample` + + /// `trade_close_per_sample` buffers (same source `launch_reward_cap_update` + /// reads from), filters by `(is_close && step_ret != 0)` to extract + /// Long/Short close magnitudes (per spec DD3=b — Hold/Flat closes + /// don't generate non-zero step_ret per `experience_kernels.cu:2855`), + /// computes a Welford `mean + Z_99 × sigma` p99 estimator × + /// `HOLD_REWARD_CAP_SAFETY_FACTOR=1.5`, applies Wiener-optimal α + /// blend (Welford accumulators in slots [487..493); floor at + /// `WELFORD_ALPHA_MIN=0.4` per + /// `pearl_wiener_alpha_floor_for_nonstationary`), and writes both + /// `ISV[HOLD_REWARD_POS_CAP_INDEX=483]` (POS cap) and + /// `ISV[HOLD_REWARD_NEG_CAP_INDEX=484]` (NEG = −2 × POS, DD5=b + /// mirrored asymmetry — same Kahneman 2:1 ratio as the position- + /// side `launch_reward_cap_update`, single source of truth at + /// producer time per + /// `pearl_audit_unboundedness_for_implicit_asymmetry`). + /// + /// Per-epoch boundary launch — call AFTER `launch_reward_cap_update` + /// so both cap producers share the same source buffers (no shared + /// state; they drive independent ISV slot pairs at [452]/[453] + /// and [483]/[484]). + /// + /// Pearls applied: + /// - `pearl_first_observation_bootstrap.md` (sentinel POS=5.0/ + /// NEG=-10.0 match SP14 P0-A pattern for bit-identical cold-start) + /// - `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 [0.5, 50.0] + /// are dimensional safety floors, not tuning) + /// - `pearl_wiener_optimal_adaptive_alpha.md` (Welford-derived α) + /// - `pearl_wiener_alpha_floor_for_nonstationary.md` (α ≥ 0.4) + /// - `pearl_fused_per_group_statistics_oracle.md` (sum/sumsq/count/ + /// max accumulated in ONE fused reduction) + pub(crate) fn launch_hold_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::{ + HOLD_REWARD_CAP_SAFETY_FACTOR, HOLD_REWARD_NEG_CAP_INDEX, + HOLD_REWARD_POS_CAP_INDEX, HOLD_REWARD_POS_CAP_MAX_BOUND, + HOLD_REWARD_POS_CAP_MIN_BOUND, HRC_DIFF_M2_INDEX, HRC_DIFF_MEAN_INDEX, + HRC_PREV_TARGET_INDEX, HRC_SAMPLE_COUNT_INDEX, HRC_TARGET_M2_INDEX, + HRC_TARGET_MEAN_INDEX, REWARD_NEG_TO_POS_RATIO, SENTINEL_HOLD_REWARD_NEG_CAP, + SENTINEL_HOLD_REWARD_POS_CAP, + }; + + debug_assert!( + self.isv_signals_dev_ptr != 0, + "launch_hold_reward_cap_update: isv_signals_dev_ptr must be allocated" + ); + + self.hold_reward_cap_update_ops.launch( + &self.stream, + step_ret_dev_ptr, + trade_close_dev_ptr, + total_samples, + self.isv_signals_dev_ptr, + HOLD_REWARD_POS_CAP_INDEX as i32, + HOLD_REWARD_NEG_CAP_INDEX as i32, + HRC_TARGET_MEAN_INDEX as i32, + HRC_TARGET_M2_INDEX as i32, + HRC_DIFF_MEAN_INDEX as i32, + HRC_DIFF_M2_INDEX as i32, + HRC_PREV_TARGET_INDEX as i32, + HRC_SAMPLE_COUNT_INDEX as i32, + SENTINEL_HOLD_REWARD_POS_CAP, + SENTINEL_HOLD_REWARD_NEG_CAP, + HOLD_REWARD_CAP_SAFETY_FACTOR, + REWARD_NEG_TO_POS_RATIO, + HOLD_REWARD_POS_CAP_MIN_BOUND, + HOLD_REWARD_POS_CAP_MAX_BOUND, + )?; + + Ok(()) + } + /// Class A P1-Producer (2026-05-08): launch the adaptive Bayesian Kelly /// priors producer. Per-fold-end (per-epoch boundary in current /// scheduling) launch — Bayesian prior is a *prior belief* and should @@ -24277,6 +24400,15 @@ impl GpuDqnTrainer { let reward_cap_update_ops = super::gpu_aux_trunk::RewardCapUpdateOps::new(&stream)?; + // SP18 v2 Phase 3 (2026-05-09): adaptive HOLD_REWARD_POS/NEG_CAP + // producer. Pre-loads the `CudaFunction` handle at construction + // per `pearl_no_host_branches_in_captured_graph.md`. Per-epoch + // boundary launch (Wiener-α blend with floor at 0.4 — non- + // stationary control loop tracks the policy-realised Long/Short + // close magnitude distribution as the policy adapts). + let hold_reward_cap_update_ops = + super::gpu_aux_trunk::HoldRewardCapUpdateOps::new(&stream)?; + // Class A P1-Producer (2026-05-08): adaptive Bayesian Kelly priors // producer. Pre-loads the `CudaFunction` handle at construction per // `pearl_no_host_branches_in_captured_graph.md`. Per-fold-end (per-epoch @@ -25157,6 +25289,8 @@ impl GpuDqnTrainer { avg_win_hold_time_update_ops, // Class A P0-A (2026-05-08): adaptive reward-cap producer. reward_cap_update_ops, + // SP18 v2 Phase 3 (2026-05-09): adaptive Hold-reward-cap producer. + hold_reward_cap_update_ops, // Class A P1-Producer (2026-05-08): adaptive Bayesian Kelly priors producer. kelly_bayesian_priors_update_ops, // Class A audit-fix Batch 4-A (2026-05-08): adaptive DD saturation floor producer. diff --git a/crates/ml/src/cuda_pipeline/hold_reward_cap_update_kernel.cu b/crates/ml/src/cuda_pipeline/hold_reward_cap_update_kernel.cu new file mode 100644 index 000000000..cf3d004ee --- /dev/null +++ b/crates/ml/src/cuda_pipeline/hold_reward_cap_update_kernel.cu @@ -0,0 +1,394 @@ +/* ══════════════════════════════════════════════════════════════════════════ + * SP18 v2 Phase 3 (2026-05-09) — D-leg adaptive HOLD_REWARD_POS/NEG_CAP + * producer. + * + * Drives `ISV[HOLD_REWARD_POS_CAP_INDEX=483]` and + * `ISV[HOLD_REWARD_NEG_CAP_INDEX=484]` from the realised Long/Short + * trade-close `|step_ret|` distribution per spec DD3=b. The Hold + * opportunity-cost reward (`compute_sp18_hold_opportunity_cost`, + * Phase 2) reads slots [483]/[484] as its asymmetric cap pair on + * every per-bar Hold subtraction site (`experience_kernels.cu` + * lines 3170/3172, 3687/3689, 3798/3800). Without this kernel the + * caps stay at the cold-start sentinels (5.0/−10.0), unable to track + * the policy's own realised distribution as it adapts. + * + * Why it mirrors `reward_cap_update_kernel` (the position-side cap + * producer for ISV[452]/[453]): both producers serve the same role — + * "p99-of-realised-trade-magnitude × safety_factor → adaptive cap" — + * and both must satisfy `feedback_no_atomicadd`, + * `pearl_no_host_branches_in_captured_graph`, and + * `feedback_isv_for_adaptive_bounds`. The difference is the input + * filter: + * + * - `reward_cap_update` : `is_close && step_ret > 0` (winners only) + * → drives realised-WIN-distribution cap + * - `hold_reward_cap_update`: `is_close && step_ret != 0` (winners + * AND losers) → drives Long/Short + * realised-trade |step_ret| distribution, + * which is the magnitude scale the Hold + * opp-cost reward must be capped to. + * + * step_ret_per_sample[i] is non-zero only on close events from + * positions that had non-zero pre_sign — which can only be Long or + * Short (Hold/Flat positions don't generate close events with + * non-zero step_ret). Filtering by `(is_close && step_ret != 0)` + * is the cheapest correct extraction of the Long/Short close + * distribution; no extra dir_idx_per_sample buffer is needed. + * + * Why Wiener-α (vs. fixed α=0.01 in the position-side kernel): + * + * The position-side cap drifts slowly because the underlying alpha + * distribution is approximately stationary over training. The + * Hold-side cap distribution is intrinsically NON-stationary: as + * the policy adapts, Long/Short close magnitudes sharpen (better + * sizing, tighter trail). The Wiener-optimal α formula + * `α = diff_var / (diff_var + sample_var + ε)` lets the EMA bandwidth + * track that drift; the floor at WELFORD_ALPHA_MIN=0.4 prevents the + * standard "α decays to 0 once the policy converges" trap (per + * `pearl_wiener_alpha_floor_for_nonstationary` and the SP16 T3 + * canonical regression at train-multi-seed-b5gmp sha 641aa0dfd). + * + * Algorithm (single-block, BLK_DIM=256): + * + * Phase 1a (per-thread): sweep `step_ret_per_sample` strided by + * BLK_DIM, accumulating sum/sumsq/count/max over `|step_ret[i]|` + * when `(trade_close[i] != 0 && step_ret[i] != 0.0f)` — the + * Long/Short close filter. Local registers, no shared mem yet. + * + * Phase 1b (block-tree-reduce): standard halving tree on + * sum/sumsq/count/max via shmem + __syncthreads(). No atomicAdd. + * + * Phase 2 (thread 0): compute target_cap. + * - If count == 0 (no closes this epoch): bail without writing — + * cold-start sentinel persists, consumers fall back to 5.0/−10.0 + * per the structural cold-start contract. + * - mean = total_sum / count + * - var = max(0, total_sumsq / count − mean²) (clamp under-flow) + * - sigma = sqrt(var) + * - p99_est = mean + Z_99 × sigma (Welford-rule) + * - target_cap = max(p99_est, max_observed) × safety_factor + * + * Phase 3 (thread 0): Welford-derived Wiener-optimal α + Pearl-A. + * Read 6 Welford accumulator slots [487..493). Update them with the + * newly-computed target_cap, then derive α from the variance + * ratio: + * sample_count += 1 + * update Welford on target_cap: + * delta1 = target_cap − target_mean + * target_mean += delta1 / sample_count + * delta2 = target_cap − target_mean + * target_M2 += delta1 × delta2 + * if sample_count ≥ 2: update Welford on (target_cap − prev_target). + * If sample_count < WELFORD_MIN_SAMPLE_COUNT_F=3: + * α = WELFORD_COLD_START_ALPHA_F=1.0 (REPLACE — variance not + * yet defined) + * else: + * sample_var = target_M2 / (sample_count − 1) + * diff_var = (diff_count ≥ 2) ? diff_M2 / (diff_count − 1) + * : diff_M2 + * α = clamp(diff_var / (diff_var + sample_var + ε), + * WELFORD_ALPHA_MIN_F=0.4, WELFORD_ALPHA_MAX_F=0.95) + * + * Phase 4 (thread 0): bilateral clamp + Pearl-A bootstrap + write. + * pos_target = clamp(target_cap, pos_min=0.5, pos_max=50.0) + * current = isv[pos_idx] + * if |current − sentinel_pos| < EPS_F: + * blended = pos_target (Pearl-A REPLACE on first observation) + * else: + * blended = (1 − α) × current + α × pos_target + * blended = clamp(blended, pos_min, pos_max) (defensive re-clamp) + * + * NEG_CAP = −neg_to_pos_ratio × POS_CAP (DD5=b mirrored + * asymmetry; same + * 2:1 Kahneman ratio + * as the position-side + * cap, single source + * of truth at producer + * time per + * `pearl_audit_unboundedness_for_implicit_asymmetry`) + * blended_neg = clamp(blended_neg, neg_lo=−ratio×pos_max=−100, + * neg_hi=−ratio×pos_min=−1) + * + * Phase 5 (thread 0): write back POS, NEG, all 6 Welford slots, + * and prev_target = target_cap (for the next launch's diff). + * + * Pearls + invariants: + * + * - `feedback_no_atomicadd.md` — single-block kernel; block-tree- + * reduce in shmem; no atomics across blocks. + * + * - `pearl_fused_per_group_statistics_oracle.md` — sum, sumsq, + * count, max in ONE fused reduction; not separate launches. + * + * - `pearl_no_host_branches_in_captured_graph.md` — `CudaFunction` + * pre-loaded at construction; on-device guards for "no closes" + * and Pearl-A sentinel detection. + * + * - `pearl_first_observation_bootstrap.md` — first valid sample + * replaces sentinel directly; no blend. + * + * - `pearl_wiener_optimal_adaptive_alpha.md` — α derived from + * Welford variance accumulators, not hardcoded. + * + * - `pearl_wiener_alpha_floor_for_nonstationary.md` — α floored at + * 0.4 to preserve catch-up bandwidth on the non-stationary + * Hold-cap signal. + * + * - `pearl_symmetric_clamp_audit.md` — bilateral + * `fmaxf(lo, fminf(x, hi))` clamp on POS and NEG before write. + * + * - `feedback_no_stubs.md` — full body, no return-zero placeholders. + * + * - `feedback_isv_for_adaptive_bounds.md` — POS bounds [0.5, 50.0] + * and NEG range [−100, −1] are dimensional-safety floors, not + * tuning. Adaptive cap is signal-driven; the bounds catch + * pathological p99 outliers. + * + * 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 ~2855). + * trade_close_per_sample — `[total_samples]` device i32; 1 iff + * exiting_trade || reversing_trade + * (`experience_kernels.cu` line ~2856). + * total_samples — N*L (B*T total samples this epoch). + * isv — `[ISV_TOTAL_DIM]` device f32, modified + * in place at indices listed below. + * pos_idx — HOLD_REWARD_POS_CAP_INDEX (483). + * neg_idx — HOLD_REWARD_NEG_CAP_INDEX (484). + * target_mean_idx — HRC_TARGET_MEAN_INDEX (487). + * target_M2_idx — HRC_TARGET_M2_INDEX (488). + * diff_mean_idx — HRC_DIFF_MEAN_INDEX (489). + * diff_M2_idx — HRC_DIFF_M2_INDEX (490). + * prev_target_idx — HRC_PREV_TARGET_INDEX (491). + * sample_count_idx — HRC_SAMPLE_COUNT_INDEX (492). + * sentinel_pos — SENTINEL_HOLD_REWARD_POS_CAP (5.0). + * sentinel_neg — SENTINEL_HOLD_REWARD_NEG_CAP (-10.0; + * unused arg held for symmetry with + * `reward_cap_update`). + * safety_factor — HOLD_REWARD_CAP_SAFETY_FACTOR (1.5). + * neg_to_pos_ratio — REWARD_NEG_TO_POS_RATIO (2.0; + * matches the position-side cap so the + * Hold and Long/Short caps share the + * same Kahneman 2:1 asymmetry). + * pos_min — HOLD_REWARD_POS_CAP_MIN_BOUND (0.5). + * pos_max — HOLD_REWARD_POS_CAP_MAX_BOUND (50.0). + * + * Launch: grid=(1, 1, 1), block=(BLK_DIM=256, 1, 1). + * Shared memory: 4 × BLK_DIM × sizeof(float) = 4096 bytes (BLK_DIM=256). + * sh_sum — sum of |step_ret| + * sh_sumsq — sum of step_ret² + * sh_count — count of qualifying close events (f32 for shmem unification; + * total_samples ≤ 2^16 so f32 representation is exact for + * counts up to 2^24). + * sh_max — max of |step_ret| + * ══════════════════════════════════════════════════════════════════════════ */ + +#include + +#define HRC_BLK_DIM 256 +#define HRC_EPS_F 1e-6f +/* Z-score for ~p99 of a normal distribution (one-sided). Used as a + * lightweight estimator: p99 ≈ mean + Z_99 × sigma. Same constant as + * the position-side `reward_cap_update_kernel`. */ +#define HRC_Z_99 2.326f + +/* Wiener-optimal-α floor for non-stationary control loops. Same value + * as `min_hold_temperature_update_kernel` and `td_error_mag_ema_kernel` + * — consistent across all SP18-related Welford producers per + * `pearl_wiener_alpha_floor_for_nonstationary`. */ +#define HRC_WELFORD_ALPHA_MIN_F 0.4f +#define HRC_WELFORD_ALPHA_MAX_F 0.95f +#define HRC_WELFORD_MIN_SAMPLE_COUNT_F 3.0f +#define HRC_WELFORD_COLD_START_ALPHA_F 1.0f + +extern "C" __global__ +void hold_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, + int target_mean_idx, + int target_M2_idx, + int diff_mean_idx, + int diff_M2_idx, + int prev_target_idx, + int sample_count_idx, + float sentinel_pos, + float sentinel_neg, + float safety_factor, + float neg_to_pos_ratio, + float pos_min, + float pos_max) +{ + /* 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 + HRC_BLK_DIM; /* [BLK_DIM] */ + float* sh_count = shmem + 2 * HRC_BLK_DIM; /* [BLK_DIM] */ + float* sh_max = shmem + 3 * HRC_BLK_DIM; /* [BLK_DIM] */ + + int tid = threadIdx.x; + + /* Phase 1a: stride-based sweep — accumulate stats over + * Long/Short trade closes. The filter `(is_close && ret != 0)` + * implicitly selects positions that had a non-zero pre_sign at + * close time (i.e. Long or Short — Hold/Flat closes don't + * generate non-zero step_ret per `experience_kernels.cu:2855`). */ + 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 += HRC_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) { + float abs_ret = fabsf(ret); + local_sum += abs_ret; + local_sumsq += ret * ret; /* = abs_ret² */ + local_count += 1.0f; + if (abs_ret > local_max) local_max = abs_ret; + } + } + 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). Sum/sumsq/count fold + * via addition; max folds via fmaxf. */ + for (int s = HRC_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-5: thread 0 finalises. */ + 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 Long/Short trade closes observed this epoch. Keep + * ALL slots (POS, NEG, every Welford accumulator) unchanged so + * the Welford state machine starts cleanly on the next epoch's + * first valid observation. Consumers fall back to the cold-start + * sentinel behavior (5.0/−10.0) per the structural cold-start + * contract — which is exactly what `compute_sp18_hold_opportunity_cost` + * does on its `is_sentinel` branch. */ + if (total_count <= 0.0f) { + return; + } + + /* Welford-rule p99 estimator: mean + Z_99 × sigma. 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 + HRC_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 Long/Short close magnitude. */ + float target_cap = fmaxf(p99_estimate, total_max) * safety_factor; + + /* Phase 3: Welford accumulator update + Wiener-optimal α. */ + float sample_count = isv[sample_count_idx]; + float target_mean = isv[target_mean_idx]; + float target_M2 = isv[target_M2_idx]; + float diff_mean = isv[diff_mean_idx]; + float diff_M2 = isv[diff_M2_idx]; + const float prev_target = isv[prev_target_idx]; + + /* Update Welford on target_cap. */ + sample_count += 1.0f; + const float t_delta = target_cap - target_mean; + target_mean += t_delta / sample_count; + const float t_delta2 = target_cap - target_mean; + target_M2 += t_delta * t_delta2; + + /* Update Welford on consecutive diff (need ≥ 2 samples). */ + const float diff_count = sample_count - 1.0f; + if (sample_count >= 2.0f) { + const float diff = target_cap - prev_target; + const float d_delta = diff - diff_mean; + diff_mean += d_delta / diff_count; + const float d_delta2 = diff - diff_mean; + diff_M2 += d_delta * d_delta2; + } + + /* Wiener-optimal α with cold-start REPLACE + non-stationary floor. */ + float alpha; + if (sample_count < HRC_WELFORD_MIN_SAMPLE_COUNT_F) { + alpha = HRC_WELFORD_COLD_START_ALPHA_F; + } else { + const float sample_var = target_M2 / (sample_count - 1.0f); + const float diff_var = (diff_count >= 2.0f) + ? diff_M2 / (diff_count - 1.0f) + : diff_M2; + alpha = diff_var / (diff_var + sample_var + HRC_EPS_F); + alpha = fmaxf(HRC_WELFORD_ALPHA_MIN_F, fminf(alpha, HRC_WELFORD_ALPHA_MAX_F)); + } + + /* Phase 4: bilateral clamp on POS, Pearl-A bootstrap, then blend. + * Bilateral pre-clamp per `pearl_symmetric_clamp_audit.md`. */ + float pos_target = fmaxf(pos_min, fminf(target_cap, pos_max)); + + float current_pos = isv[pos_idx]; + float blended_pos; + if (fabsf(current_pos - sentinel_pos) < HRC_EPS_F) { + /* Pearl-A first-observation bootstrap: REPLACE. */ + blended_pos = pos_target; + } else { + blended_pos = (1.0f - alpha) * current_pos + alpha * pos_target; + } + /* Re-clamp post-blend (defensive — handles malformed prior state). */ + blended_pos = fmaxf(pos_min, fminf(blended_pos, pos_max)); + + /* NEG cap = −neg_to_pos_ratio × POS cap (DD5=b mirrored + * asymmetry, single source of truth at producer time per + * `pearl_audit_unboundedness_for_implicit_asymmetry`). NEG range + * derived from POS bounds × ratio: NEG ∈ [−ratio × pos_max, + * −ratio × pos_min] = [−100, −1] for the SP18 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; /* −1 */ + blended_neg = fmaxf(neg_lo, fminf(blended_neg, neg_hi)); + + /* Suppress unused-parameter warning for sentinel_neg — held in + * the signature for symmetry with `reward_cap_update_kernel` and + * to allow a future Pearl-A variant on the NEG slot if the 2:1 + * ratio spec ever evolves. */ + (void)sentinel_neg; + + /* Phase 5: write back POS, NEG, all 6 Welford accumulators, and + * prev_target = target_cap (for the next launch's diff). The + * __threadfence_system() ensures PCIe-visible to mapped-pinned + * host_ptr before the stream sync returns. */ + isv[pos_idx] = blended_pos; + isv[neg_idx] = blended_neg; + isv[sample_count_idx] = sample_count; + isv[target_mean_idx] = target_mean; + isv[target_M2_idx] = target_M2; + isv[diff_mean_idx] = diff_mean; + isv[diff_M2_idx] = diff_M2; + isv[prev_target_idx] = target_cap; + __threadfence_system(); +} diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 133202c73..8fb20959b 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -484,6 +484,37 @@ impl DQNTrainer { ) { tracing::warn!(epoch, "launch_reward_cap_update failed (non-fatal): {e}"); } + // SP18 v2 Phase 3 (2026-05-09): adaptive HOLD_REWARD_POS/ + // NEG_CAP — per-epoch boundary launch RIGHT AFTER the + // position-side `launch_reward_cap_update` so both cap + // producers share the same `step_ret_per_sample` + + // `trade_close_per_sample` source buffers (read-only; + // independent ISV slot pairs at [452]/[453] and + // [483]/[484]). Filters by `(is_close && step_ret != 0)` + // to extract Long/Short close magnitudes (DD3=b — Hold/ + // Flat closes don't generate non-zero step_ret per + // `experience_kernels.cu:2855`), computes Welford p99 × + // 1.5 → POS cap, derives Wiener-optimal α from Welford + // accumulators in slots [487..493) (floor at + // WELFORD_ALPHA_MIN=0.4 per + // `pearl_wiener_alpha_floor_for_nonstationary` — the + // policy-realised distribution is intrinsically non- + // stationary as the policy adapts), Pearl-A REPLACE on + // sentinel detection, and writes ISV[483]=POS and + // ISV[484]=−2×POS (DD5=b mirrored asymmetry, single + // source of truth at producer time per + // `pearl_audit_unboundedness_for_implicit_asymmetry`). + // The Hold opportunity-cost reward + // (`compute_sp18_hold_opportunity_cost`, Phase 2) reads + // these slots as its asymmetric cap pair on every + // per-bar Hold subtraction site. + if let Err(e) = fused.trainer().launch_hold_reward_cap_update( + step_ret_ptr, + trade_close_ptr, + total_samples, + ) { + tracing::warn!(epoch, "launch_hold_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 @@ -5507,6 +5538,39 @@ impl DQNTrainer { "HEALTH_DIAG[{}]: td_error_pre [magnitude_ema={:.6}]", epoch, td_mag, ); + + // SP18 v2 Phase 3 (2026-05-09): D-leg adaptive Hold- + // reward-cap producer observability. Surfaces the + // current ISV[483]/[484] (POS / NEG caps) plus + // ISV[486] (HOLD_OPP_COST_FIRE_RATE_EMA — the + // engagement-canary diag for the Hold opportunity-cost + // reward, populated by `compute_sp18_hold_opportunity_cost` + // consumers in `experience_kernels.cu`). Reads ISV + // directly via `read_isv_signal_at` — the producer + // ran at the per-epoch boundary above; the values + // here are the post-blend caps the next epoch's + // `experience_kernels.cu:3170/3172/3687/3689/3798/3800` + // consumer sites will use. NEG = −2 × POS at producer + // time (DD5=b mirrored asymmetry, single source of + // truth). + let (hold_pos_cap, hold_neg_cap, hold_fire_rate) = + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp14_isv_slots::{ + HOLD_OPP_COST_FIRE_RATE_EMA_INDEX, HOLD_REWARD_NEG_CAP_INDEX, + HOLD_REWARD_POS_CAP_INDEX, + }; + ( + fused.trainer().read_isv_signal_at(HOLD_REWARD_POS_CAP_INDEX), + fused.trainer().read_isv_signal_at(HOLD_REWARD_NEG_CAP_INDEX), + fused.trainer().read_isv_signal_at(HOLD_OPP_COST_FIRE_RATE_EMA_INDEX), + ) + } else { + (f32::NAN, f32::NAN, f32::NAN) + }; + tracing::info!( + "HEALTH_DIAG[{}]: hold_reward_cap [pos={:.4} neg={:.4} fire_rate={:.4}]", + epoch, hold_pos_cap, hold_neg_cap, hold_fire_rate, + ); } // SP7 observability (grad_decomp_pinned): surface the exact contents of diff --git a/crates/ml/tests/sp18_hold_reward_oracle_tests.rs b/crates/ml/tests/sp18_hold_reward_oracle_tests.rs index 30adc362d..9012d8f2f 100644 --- a/crates/ml/tests/sp18_hold_reward_oracle_tests.rs +++ b/crates/ml/tests/sp18_hold_reward_oracle_tests.rs @@ -769,4 +769,486 @@ mod gpu { assert_eq!(gpu_f, 0.0, "Case F (shaping_scale=0): expected 0.0 exactly, got {gpu_f}"); } + + // ── SP18 Phase 3 — D-leg adaptive HOLD_REWARD_POS/NEG_CAP producer ── + // + // Tests for `hold_reward_cap_update_kernel.cu`. Mirrors the SP14 P0-A + // `reward_cap_update_kernel` test pattern (sp14_oracle_tests.rs) with + // three structural differences: + // 1. Filter: (is_close && step_ret != 0) — both winners AND losers + // contribute, since the consumer needs the magnitude scale of + // *all* Long/Short closes. + // 2. Welford accumulators in slots [487..493) feed Wiener-optimal α + // — sample_count < 3 → α=1.0 (REPLACE), else + // α ∈ [WELFORD_ALPHA_MIN=0.4, WELFORD_ALPHA_MAX=0.95]. + // 3. Bounds [HOLD_REWARD_POS_CAP_MIN_BOUND=0.5, + // HOLD_REWARD_POS_CAP_MAX_BOUND=50.0] (vs. [1.0, 50.0] for the + // position-side cap). + // + // All tests are #[ignore = "requires GPU"]; gated under + // `#[cfg(feature = "cuda")]` (we are already inside that gate). + + /// Cubin handle for the SP18 D-leg adaptive Hold-reward-cap producer. + const SP18_HOLD_REWARD_CAP_UPDATE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/hold_reward_cap_update_kernel.cubin")); + + fn load_hold_reward_cap_update(stream: &Arc) -> CudaFunction { + let module = stream + .context() + .load_cubin(SP18_HOLD_REWARD_CAP_UPDATE_CUBIN.to_vec()) + .expect("load hold_reward_cap_update_kernel cubin"); + module + .load_function("hold_reward_cap_update") + .expect("load hold_reward_cap_update function") + } + + /// 4 arrays × 256 × 4 bytes = 4096 bytes shmem (matches kernel). + const HRC_BLK_DIM: u32 = 256; + const HRC_SMEM_BYTES: u32 = 4 * HRC_BLK_DIM * std::mem::size_of::() as u32; + + /// Helper: launch the SP18 hold-reward-cap producer with a fixed + /// config. Mirrors the position-side `launch_reward_cap` helper but + /// passes the 6 Welford slot indices instead of a single fixed α. + #[allow(clippy::too_many_arguments)] + fn launch_hold_reward_cap( + stream: &Arc, + kernel: &CudaFunction, + step_ret_ptr: u64, + trade_close_ptr: u64, + total_samples: i32, + isv_ptr: u64, + pos_idx: i32, + neg_idx: i32, + target_mean_idx: i32, + target_m2_idx: i32, + diff_mean_idx: i32, + diff_m2_idx: i32, + prev_target_idx: i32, + sample_count_idx: i32, + sentinel_pos: f32, + sentinel_neg: f32, + safety_factor: f32, + neg_to_pos_ratio: f32, + pos_min: f32, + pos_max: 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(&target_mean_idx) + .arg(&target_m2_idx) + .arg(&diff_mean_idx) + .arg(&diff_m2_idx) + .arg(&prev_target_idx) + .arg(&sample_count_idx) + .arg(&sentinel_pos) + .arg(&sentinel_neg) + .arg(&safety_factor) + .arg(&neg_to_pos_ratio) + .arg(&pos_min) + .arg(&pos_max) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (HRC_BLK_DIM, 1, 1), + shared_mem_bytes: HRC_SMEM_BYTES, + }) + .expect("launch hold_reward_cap_update"); + } + stream.synchronize().expect("sync after hold_reward_cap_update"); + } + + /// SP18 Phase 3 T5 — adaptive Hold-reward-cap producer drives slots + /// [483]/[484] from Long/Short trade-close p99. + /// + /// Construct 100 simulated Long/Short close events with `step_ret` + /// drawn from a centred Gaussian-ish distribution scaled to ~0.01 + /// (per the spec's T5 anchor: `step_ret ∈ N(0, 0.01)`). Mix winners + /// (positive) and losers (negative) so the kernel's + /// `(is_close && step_ret != 0)` filter takes both. Also include + /// non-close samples (trade_close=0) with non-zero step_ret to + /// verify they're FILTERED OUT (close events only contribute). + /// + /// Per the spec's T5 acceptance: ISV[483] ≈ p99(|step_ret|) × 1.5 + /// (small, ~0.04 for σ ≈ 0.01); ISV[484] = −2 × ISV[483] + /// (mirrored); both within [HOLD_REWARD_POS_CAP_MIN_BOUND=0.5, + /// HOLD_REWARD_POS_CAP_MAX_BOUND=50.0]. The lower-bound clamp + /// at 0.5 is the dominating constraint at this magnitude scale — + /// verifying the bilateral clamp activation per + /// `pearl_symmetric_clamp_audit.md`. + /// + /// Jitter: per `pearl_sp4_histogram_warp_tile_undercount`, synthetic + /// uniform test data can trip warp-tile undercount in + /// `sp4_histogram_p99`. This producer doesn't use `sp4_histogram_p99` + /// (it uses Welford `mean + Z_99 × sigma` like the position-side + /// kernel), but the input is jittered anyway as defensive practice. + #[test] + #[ignore = "requires GPU"] + fn hold_reward_cap_producer_drives_slots() { + use ml::cuda_pipeline::sp14_isv_slots::{ + HOLD_REWARD_CAP_SAFETY_FACTOR, HOLD_REWARD_NEG_CAP_INDEX, HOLD_REWARD_POS_CAP_INDEX, + HOLD_REWARD_POS_CAP_MAX_BOUND, HOLD_REWARD_POS_CAP_MIN_BOUND, HRC_DIFF_M2_INDEX, + HRC_DIFF_MEAN_INDEX, HRC_PREV_TARGET_INDEX, HRC_SAMPLE_COUNT_INDEX, + HRC_TARGET_M2_INDEX, HRC_TARGET_MEAN_INDEX, REWARD_NEG_TO_POS_RATIO, + SENTINEL_HOLD_REWARD_NEG_CAP, SENTINEL_HOLD_REWARD_POS_CAP, SENTINEL_WELFORD_ZERO, + }; + + let stream = make_test_stream(); + let kernel = load_hold_reward_cap_update(&stream); + + const ISV_DIM: usize = 1024; + + // Build 100 Long/Short close events with jittered ~N(0, 0.01) + // step_ret. Mix winners and losers; sprinkle in non-close + // samples that should be filtered out. + let n_close = 100usize; + let n_filler = 50usize; + let total = n_close + n_filler; + let mut step_ret = vec![0.0_f32; total]; + let mut trade_close = vec![0_i32; total]; + + // Deterministic pseudo-random: alternating signs + slight phase + // jitter so the histogram (had we used one) wouldn't undercount. + for i in 0..n_close { + let sign = if i % 2 == 0 { 1.0_f32 } else { -1.0_f32 }; + // Stretch the magnitude across [0.001, 0.025] so p99 ≈ 0.025. + let mag = 0.001 + (i as f32 / n_close as f32) * 0.024 + (i as f32 * 0.0001); + step_ret[i] = sign * mag; + trade_close[i] = 1; + } + // Filler: non-close samples with non-zero step_ret. The kernel + // MUST filter these out via `is_close == 0`. + for i in 0..n_filler { + step_ret[n_close + i] = 1000.0; // outlier magnitude + trade_close[n_close + i] = 0; // not a close + } + + let mut isv = vec![0.0_f32; ISV_DIM]; + // Cold-start: sentinels match SP18 PP.2 sentinel allocations. + isv[HOLD_REWARD_POS_CAP_INDEX] = SENTINEL_HOLD_REWARD_POS_CAP; + isv[HOLD_REWARD_NEG_CAP_INDEX] = SENTINEL_HOLD_REWARD_NEG_CAP; + isv[HRC_TARGET_MEAN_INDEX] = SENTINEL_WELFORD_ZERO; + isv[HRC_TARGET_M2_INDEX] = SENTINEL_WELFORD_ZERO; + isv[HRC_DIFF_MEAN_INDEX] = SENTINEL_WELFORD_ZERO; + isv[HRC_DIFF_M2_INDEX] = SENTINEL_WELFORD_ZERO; + isv[HRC_PREV_TARGET_INDEX] = SENTINEL_WELFORD_ZERO; + isv[HRC_SAMPLE_COUNT_INDEX] = SENTINEL_WELFORD_ZERO; + + let n = total 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_hold_reward_cap( + &stream, &kernel, + ret_buf.dev_ptr, close_buf.dev_ptr, n, isv_buf.dev_ptr, + HOLD_REWARD_POS_CAP_INDEX as i32, + HOLD_REWARD_NEG_CAP_INDEX as i32, + HRC_TARGET_MEAN_INDEX as i32, + HRC_TARGET_M2_INDEX as i32, + HRC_DIFF_MEAN_INDEX as i32, + HRC_DIFF_M2_INDEX as i32, + HRC_PREV_TARGET_INDEX as i32, + HRC_SAMPLE_COUNT_INDEX as i32, + SENTINEL_HOLD_REWARD_POS_CAP, + SENTINEL_HOLD_REWARD_NEG_CAP, + HOLD_REWARD_CAP_SAFETY_FACTOR, + REWARD_NEG_TO_POS_RATIO, + HOLD_REWARD_POS_CAP_MIN_BOUND, + HOLD_REWARD_POS_CAP_MAX_BOUND, + ); + + let result = isv_buf.read_all(); + let pos = result[HOLD_REWARD_POS_CAP_INDEX]; + let neg = result[HOLD_REWARD_NEG_CAP_INDEX]; + + // The unclamped target_cap at this scale is ~p99 × 1.5 ≈ + // 0.025 × 1.5 ≈ 0.04, well below HOLD_REWARD_POS_CAP_MIN_BOUND + // = 0.5. The bilateral clamp pre-blend pulls it up to the + // floor. Pearl-A then REPLACES the sentinel with the clamped + // value → POS = 0.5 exactly. + assert!( + (pos - HOLD_REWARD_POS_CAP_MIN_BOUND).abs() < 1e-5, + "Pearl-A bootstrap: small p99 → bilateral clamp pulls POS to MIN_BOUND. \ + Expected POS={HOLD_REWARD_POS_CAP_MIN_BOUND}, got {pos}", + ); + + // POS must be within bounds. + assert!( + pos >= HOLD_REWARD_POS_CAP_MIN_BOUND && pos <= HOLD_REWARD_POS_CAP_MAX_BOUND, + "POS out of bounds: {pos} not in [{HOLD_REWARD_POS_CAP_MIN_BOUND}, {HOLD_REWARD_POS_CAP_MAX_BOUND}]", + ); + + // NEG = −2 × POS (DD5=b mirrored asymmetry, single source of + // truth at producer time). + assert!( + (neg - (-REWARD_NEG_TO_POS_RATIO * pos)).abs() < 1e-4, + "NEG must equal −{REWARD_NEG_TO_POS_RATIO} × POS at producer-time \ + (Kahneman 2:1 single source of truth); got POS={pos}, NEG={neg}", + ); + + // NEG within derived range [−ratio × pos_max, −ratio × pos_min] = [−100, −1]. + let neg_lo = -REWARD_NEG_TO_POS_RATIO * HOLD_REWARD_POS_CAP_MAX_BOUND; + let neg_hi = -REWARD_NEG_TO_POS_RATIO * HOLD_REWARD_POS_CAP_MIN_BOUND; + assert!( + neg >= neg_lo && neg <= neg_hi, + "NEG out of bounds: {neg} not in [{neg_lo}, {neg_hi}]", + ); + + // Welford state updated: sample_count == 1 (first observation). + let count = result[HRC_SAMPLE_COUNT_INDEX]; + assert!((count - 1.0).abs() < 1e-5, + "First-observation Welford: sample_count expected 1.0, got {count}"); + + // Filler outliers MUST NOT have contributed (they were + // is_close=0). We can't read max directly from ISV, but we + // can verify the cap stayed clamped to 0.5 — if the filler + // outliers had leaked, target_cap would be ≈ 1000 × 1.5 = 1500, + // clamped to MAX_BOUND=50.0. The fact POS == 0.5 proves the + // filter worked. + assert!( + pos < 1.0, + "Filter regression: filler outliers (close=0) leaked into accumulator; \ + cap should be clamped to MIN_BOUND=0.5, got {pos}", + ); + } + + /// SP18 Phase 3 — Pearl-A first-observation bootstrap test. + /// + /// Sentinel state → first observation REPLACES (no blend). Verify + /// ISV[483] equals `target_cap × 1.5` post-clamp exactly, NOT + /// blended with the sentinel 5.0. Per + /// `pearl_first_observation_bootstrap.md` the kernel must detect + /// `|isv[pos_idx] − sentinel_pos| < EPS_F` and route through + /// REPLACE (α=1.0), not the standard EMA blend. + /// + /// Construction: 4 closes with `step_ret = [+10, +20, +30, +40]` + /// (large magnitudes — ensures the unclamped target_cap is well + /// above MIN_BOUND so Pearl-A's REPLACE is observable). Welford + /// rule on |step_ret|: mean=25, var = (100+400+900+1600)/4 − 625 + /// = 750 − 625 = 125, sigma ≈ 11.18, p99 ≈ 25 + 2.326 × 11.18 ≈ + /// 51.0. max=40. Estimator picks `max(p99, max) = 51.0`. ×1.5 = + /// 76.5. Clamped to MAX_BOUND=50.0 → POS = 50.0. + /// + /// With sample_count < WELFORD_MIN_SAMPLE_COUNT_F=3, α=1.0 + /// (cold-start REPLACE per `pearl_wiener_alpha_floor_for_nonstationary` + /// — Welford variance not yet defined). Combined with Pearl-A + /// sentinel detection, the result is bit-identical to a direct + /// REPLACE: blended = 50.0, NOT (1−α)×5.0 + α×50.0 if α had been + /// some EMA value. + #[test] + #[ignore = "requires GPU"] + fn hold_reward_cap_pearl_a_first_observation_replace() { + use ml::cuda_pipeline::sp14_isv_slots::{ + HOLD_REWARD_CAP_SAFETY_FACTOR, HOLD_REWARD_NEG_CAP_INDEX, HOLD_REWARD_POS_CAP_INDEX, + HOLD_REWARD_POS_CAP_MAX_BOUND, HOLD_REWARD_POS_CAP_MIN_BOUND, HRC_DIFF_M2_INDEX, + HRC_DIFF_MEAN_INDEX, HRC_PREV_TARGET_INDEX, HRC_SAMPLE_COUNT_INDEX, + HRC_TARGET_M2_INDEX, HRC_TARGET_MEAN_INDEX, REWARD_NEG_TO_POS_RATIO, + SENTINEL_HOLD_REWARD_NEG_CAP, SENTINEL_HOLD_REWARD_POS_CAP, SENTINEL_WELFORD_ZERO, + }; + + let stream = make_test_stream(); + let kernel = load_hold_reward_cap_update(&stream); + + const ISV_DIM: usize = 1024; + + // 4 closes with mixed-sign large step_ret values. The kernel's + // |step_ret| filter takes both winners and losers. + let step_ret: Vec = vec![10.0, -20.0, 30.0, -40.0]; + let trade_close: Vec = vec![1, 1, 1, 1]; + + let mut isv = vec![0.0_f32; ISV_DIM]; + isv[HOLD_REWARD_POS_CAP_INDEX] = SENTINEL_HOLD_REWARD_POS_CAP; // 5.0 + isv[HOLD_REWARD_NEG_CAP_INDEX] = SENTINEL_HOLD_REWARD_NEG_CAP; // -10.0 + isv[HRC_TARGET_MEAN_INDEX] = SENTINEL_WELFORD_ZERO; + isv[HRC_TARGET_M2_INDEX] = SENTINEL_WELFORD_ZERO; + isv[HRC_DIFF_MEAN_INDEX] = SENTINEL_WELFORD_ZERO; + isv[HRC_DIFF_M2_INDEX] = SENTINEL_WELFORD_ZERO; + isv[HRC_PREV_TARGET_INDEX] = SENTINEL_WELFORD_ZERO; + isv[HRC_SAMPLE_COUNT_INDEX] = SENTINEL_WELFORD_ZERO; + + 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_hold_reward_cap( + &stream, &kernel, + ret_buf.dev_ptr, close_buf.dev_ptr, n, isv_buf.dev_ptr, + HOLD_REWARD_POS_CAP_INDEX as i32, + HOLD_REWARD_NEG_CAP_INDEX as i32, + HRC_TARGET_MEAN_INDEX as i32, + HRC_TARGET_M2_INDEX as i32, + HRC_DIFF_MEAN_INDEX as i32, + HRC_DIFF_M2_INDEX as i32, + HRC_PREV_TARGET_INDEX as i32, + HRC_SAMPLE_COUNT_INDEX as i32, + SENTINEL_HOLD_REWARD_POS_CAP, + SENTINEL_HOLD_REWARD_NEG_CAP, + HOLD_REWARD_CAP_SAFETY_FACTOR, + REWARD_NEG_TO_POS_RATIO, + HOLD_REWARD_POS_CAP_MIN_BOUND, + HOLD_REWARD_POS_CAP_MAX_BOUND, + ); + + let result = isv_buf.read_all(); + let pos = result[HOLD_REWARD_POS_CAP_INDEX]; + let neg = result[HOLD_REWARD_NEG_CAP_INDEX]; + + // Pearl-A REPLACE: target_cap = max(p99, max) × 1.5 ≈ max(51, 40) + // × 1.5 = 76.5. Bilateral clamp to [0.5, 50.0] → 50.0. The + // sentinel-detection branch in the kernel writes blended = 50.0 + // directly (not (1−α)×5.0 + α×50.0). With α=1.0 these would + // collapse to the same value, but the test pins behaviour: + // the kernel MUST take the `if (|current − sentinel| < EPS)` + // branch and write target_cap unblended. + assert!( + (pos - HOLD_REWARD_POS_CAP_MAX_BOUND).abs() < 1e-4, + "Pearl-A REPLACE on sentinel: expected POS exactly at MAX_BOUND={HOLD_REWARD_POS_CAP_MAX_BOUND}, \ + got {pos}", + ); + + // Confirm we did NOT blend with sentinel: a blended result with + // any α < 1 would land between 5.0 and 50.0, never exactly at + // the clamp boundary. + assert!(pos > 49.99 && pos <= 50.01, + "Pearl-A signature: POS ≈ MAX_BOUND exactly (no sentinel blend), got {pos}"); + + // NEG mirrors at producer time. + assert!( + (neg - (-REWARD_NEG_TO_POS_RATIO * pos)).abs() < 1e-4, + "NEG must equal −{REWARD_NEG_TO_POS_RATIO} × POS; got POS={pos}, NEG={neg}", + ); + assert!( + (neg - (-100.0)).abs() < 1e-3, + "NEG = −2 × 50.0 = −100; got {neg}", + ); + + // sample_count == 1 (first observation). + let count = result[HRC_SAMPLE_COUNT_INDEX]; + assert!((count - 1.0).abs() < 1e-5, + "Welford sample_count after 1 observation: expected 1.0, got {count}"); + + // prev_target stores the unclamped target_cap (= max(p99, max) + // × safety_factor) for the next launch's diff. Useful sanity: + // proves the Welford state is being updated even though the + // CAP write was REPLACED. + let prev = result[HRC_PREV_TARGET_INDEX]; + assert!(prev > 1.0, + "prev_target should be the unclamped target_cap (~76.5), got {prev}"); + } + + /// SP18 Phase 3 — no Long/Short closes ⇒ ISV slots preserved bit-exactly. + /// + /// All trade_close=0 (no closes) AND all step_ret=0. The kernel + /// must skip the EMA update (count==0 guard) and leave POS, NEG, + /// AND every Welford accumulator at their seeded values. + /// + /// Models the cold-start "Hold-heavy phase with no real trade + /// activity" regime — exactly the case where the producer SHOULD + /// keep the cap at sentinel and let the consumer fall back to the + /// structural cold-start sentinel value (5.0/-10.0). + #[test] + #[ignore = "requires GPU"] + fn hold_reward_cap_no_closes_preserves_isv() { + use ml::cuda_pipeline::sp14_isv_slots::{ + HOLD_REWARD_CAP_SAFETY_FACTOR, HOLD_REWARD_NEG_CAP_INDEX, HOLD_REWARD_POS_CAP_INDEX, + HOLD_REWARD_POS_CAP_MAX_BOUND, HOLD_REWARD_POS_CAP_MIN_BOUND, HRC_DIFF_M2_INDEX, + HRC_DIFF_MEAN_INDEX, HRC_PREV_TARGET_INDEX, HRC_SAMPLE_COUNT_INDEX, + HRC_TARGET_M2_INDEX, HRC_TARGET_MEAN_INDEX, REWARD_NEG_TO_POS_RATIO, + SENTINEL_HOLD_REWARD_NEG_CAP, SENTINEL_HOLD_REWARD_POS_CAP, + }; + + let stream = make_test_stream(); + let kernel = load_hold_reward_cap_update(&stream); + + const ISV_DIM: usize = 1024; + // Seed with deliberately non-sentinel values to verify + // bit-exact preservation (NOT just "sentinel was preserved"). + const SEED_POS: f32 = 12.5; + const SEED_NEG: f32 = -25.0; + const SEED_TARGET_MEAN: f32 = 7.7; + const SEED_TARGET_M2: f32 = 0.123; + const SEED_DIFF_MEAN: f32 = 0.05; + const SEED_DIFF_M2: f32 = 0.0001; + const SEED_PREV_TARGET: f32 = 7.6; + const SEED_SAMPLE_COUNT: f32 = 5.0; + + let mut isv = vec![0.0_f32; ISV_DIM]; + isv[HOLD_REWARD_POS_CAP_INDEX] = SEED_POS; + isv[HOLD_REWARD_NEG_CAP_INDEX] = SEED_NEG; + isv[HRC_TARGET_MEAN_INDEX] = SEED_TARGET_MEAN; + isv[HRC_TARGET_M2_INDEX] = SEED_TARGET_M2; + isv[HRC_DIFF_MEAN_INDEX] = SEED_DIFF_MEAN; + isv[HRC_DIFF_M2_INDEX] = SEED_DIFF_M2; + isv[HRC_PREV_TARGET_INDEX] = SEED_PREV_TARGET; + isv[HRC_SAMPLE_COUNT_INDEX] = SEED_SAMPLE_COUNT; + + // Step returns present but no closes — none should contribute. + let step_ret: Vec = vec![3.0, 5.0, -2.0]; + let trade_close: Vec = 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_hold_reward_cap( + &stream, &kernel, + ret_buf.dev_ptr, close_buf.dev_ptr, n, isv_buf.dev_ptr, + HOLD_REWARD_POS_CAP_INDEX as i32, + HOLD_REWARD_NEG_CAP_INDEX as i32, + HRC_TARGET_MEAN_INDEX as i32, + HRC_TARGET_M2_INDEX as i32, + HRC_DIFF_MEAN_INDEX as i32, + HRC_DIFF_M2_INDEX as i32, + HRC_PREV_TARGET_INDEX as i32, + HRC_SAMPLE_COUNT_INDEX as i32, + SENTINEL_HOLD_REWARD_POS_CAP, + SENTINEL_HOLD_REWARD_NEG_CAP, + HOLD_REWARD_CAP_SAFETY_FACTOR, + REWARD_NEG_TO_POS_RATIO, + HOLD_REWARD_POS_CAP_MIN_BOUND, + HOLD_REWARD_POS_CAP_MAX_BOUND, + ); + + let result = isv_buf.read_all(); + // Every slot must be preserved bit-exactly. + assert_eq!(result[HOLD_REWARD_POS_CAP_INDEX], SEED_POS, + "POS preserved bit-exactly when no closes"); + assert_eq!(result[HOLD_REWARD_NEG_CAP_INDEX], SEED_NEG, + "NEG preserved bit-exactly when no closes"); + assert_eq!(result[HRC_TARGET_MEAN_INDEX], SEED_TARGET_MEAN, + "target_mean preserved"); + assert_eq!(result[HRC_TARGET_M2_INDEX], SEED_TARGET_M2, + "target_M2 preserved"); + assert_eq!(result[HRC_DIFF_MEAN_INDEX], SEED_DIFF_MEAN, + "diff_mean preserved"); + assert_eq!(result[HRC_DIFF_M2_INDEX], SEED_DIFF_M2, + "diff_M2 preserved"); + assert_eq!(result[HRC_PREV_TARGET_INDEX], SEED_PREV_TARGET, + "prev_target preserved"); + assert_eq!(result[HRC_SAMPLE_COUNT_INDEX], SEED_SAMPLE_COUNT, + "sample_count preserved"); + } } diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index fe433e658..8f616ec57 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,65 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +## 2026-05-09 — SP18 v2 Phase 3 Tasks 3.1-3.5: adaptive HOLD_REWARD_POS/NEG_CAP producer + +D-leg adaptive cap producer kernel + launcher + per-epoch wiring + GPU +oracle tests. Lifts Phase 2's sentinel-driven cold-start caps (5.0/-10.0 +fixed) to `p99(|step_ret|)` over Long/Short trade closes × 1.5 safety +factor with Wiener-optimal α blend. The Phase 2 consumer +(`compute_sp18_hold_opportunity_cost`) sees producer-driven slots +[483]/[484] from epoch 1 onward; the consumer's sentinel branch +remains as the cold-start fallback path (zero-trade-close epoch ⇒ +producer early-returns ⇒ slots stay at sentinel ⇒ consumer falls +through to the +5/-10 macro defaults — bit-identical to pre-Phase-3). + +Mirrors the SP14 P0-A `reward_cap_update_kernel` structural template +with three differences: (1) filter `(is_close && step_ret != 0)` — +both winners AND losers (the consumer needs the magnitude scale of +*all* Long/Short closes); (2) Welford-derived Wiener-α (slots +[487..493)) replaces fixed α=0.01, with floor at WELFORD_ALPHA_MIN=0.4 +per `pearl_wiener_alpha_floor_for_nonstationary` (the policy-realised +distribution is intrinsically non-stationary as the policy adapts); +(3) bounds [0.5, 50.0] (vs. position-side [1.0, 50.0]). + +Atomic commit per `feedback_no_partial_refactor`: +- `crates/ml/src/cuda_pipeline/hold_reward_cap_update_kernel.cu` (NEW) +- `crates/ml/build.rs` cubin manifest entry +- `HoldRewardCapUpdateOps` in `gpu_aux_trunk.rs` (new struct + impl) +- `HOLD_REWARD_CAP_UPDATE_CUBIN` static + struct field + + `launch_hold_reward_cap_update` method + constructor instantiation + + field-init in `gpu_dqn_trainer.rs` +- Per-epoch boundary launch in `training_loop.rs` right AFTER + `launch_reward_cap_update` (shared step_ret/trade_close source + buffers, independent ISV slot pairs) +- `HEALTH_DIAG[N]: hold_reward_cap [pos={:.4} neg={:.4} fire_rate={:.4}]` + emit (reads ISV[483]/[484]/[486]) +- 3 GPU oracle tests in `tests/sp18_hold_reward_oracle_tests.rs` + (T5 producer-drives-slots, Pearl-A REPLACE, no-closes preserves-isv) — + all pass on local RTX 3050 Ti +- Phase 3 close-out section in `docs/sp18-wireup-audit.md` + +Wire-up audit (per Task 3.4): 1 kernel definition + 1 cubin static + +1 ops struct + 1 launcher + 1 per-epoch call site + ISV slot reads +in 3 Phase-2 consumer sites (`experience_kernels.cu` lines 3170/3172, +3687/3689, 3798/3800) + GPU oracle tests. No orphans. + +Pearls applied: `feedback_no_atomicadd`, `pearl_first_observation_bootstrap`, +`pearl_wiener_optimal_adaptive_alpha`, `pearl_wiener_alpha_floor_for_nonstationary`, +`pearl_no_host_branches_in_captured_graph`, `pearl_symmetric_clamp_audit`, +`pearl_audit_unboundedness_for_implicit_asymmetry` (NEG = -2 × POS at +producer time, single source of truth), `feedback_isv_for_adaptive_bounds`, +`pearl_fused_per_group_statistics_oracle`. + +Validation: `cargo check --workspace` clean (1m 12s); 3 GPU oracle +tests pass on local RTX 3050 Ti (1.93s wallclock); `bash scripts/audit_sp18_consumers.sh --check` +exit 0 (no fingerprint drift in tracked sections — the audit tracks +the DELETED SP13/SP16 legacy chain; the new Phase 3 producer is in +its own symbol-namespace by design). + +Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md +Phase 4-5 (B-leg target-net forward + q_next replacement) follows. + ## 2026-05-09 — SP18 archaeology cleanup pass (post-Phase 2) Pure-comment cleanup follow-up to the SP18 Phase 1 atomic deletion (3c318953a) diff --git a/docs/sp18-wireup-audit.md b/docs/sp18-wireup-audit.md index 48e099218..414c929c8 100644 --- a/docs/sp18-wireup-audit.md +++ b/docs/sp18-wireup-audit.md @@ -175,6 +175,132 @@ all OK).** fn returns 0.0 → no Hold opp cost in pure backtest, matching the SP12 lump-sum `compute_lump_sum_opp_cost` contract. +### Phase 3 sites introduced (atomic commit, 2026-05-09) + +D-leg adaptive HOLD_REWARD_POS/NEG_CAP producer kernel + launcher + +per-epoch wiring + GPU oracle tests. Lifts the Phase 2 caps from +"sentinel-driven cold-start" (5.0/-10.0 fixed) to "p99(|step_ret|) over +Long/Short trade closes × 1.5 safety factor" with Wiener-optimal α blend +per `pearl_wiener_optimal_adaptive_alpha`. The Phase 2 consumer +(`compute_sp18_hold_opportunity_cost`) sees the producer-driven slots +[483]/[484] from epoch 1 onward; the sentinel branch in the consumer +remains as the cold-start fallback path (zero-trade-close epoch ⇒ +producer early-returns ⇒ slots stay at sentinel ⇒ consumer falls +through to the +5/-10 macro defaults — bit-identical to pre-Phase-3 +behavior). + +Sites introduced (single atomic commit per `feedback_no_partial_refactor`): + +- **Kernel**: `crates/ml/src/cuda_pipeline/hold_reward_cap_update_kernel.cu` + (NEW, ~310 lines). Single-block 256-thread producer; mirrors the SP14 + P0-A `reward_cap_update_kernel` structural template with three + differences: + 1. Filter: `(is_close && step_ret != 0)` — both winners AND losers + (the consumer needs the magnitude scale of *all* Long/Short + closes, not just winners). Position-side cap filters + `(is_close && step_ret > 0)`. + 2. Welford-derived Wiener-α (slots [487..493)) replaces fixed α=0.01. + Floor at WELFORD_ALPHA_MIN=0.4 per + `pearl_wiener_alpha_floor_for_nonstationary` — the policy-realised + Long/Short close magnitude distribution is intrinsically + non-stationary as the policy adapts, sharpening tails. + 3. Bounds [HOLD_REWARD_POS_CAP_MIN_BOUND=0.5, + HOLD_REWARD_POS_CAP_MAX_BOUND=50.0] (vs. position-side [1.0, 50.0]). + +- **Cubin manifest**: `crates/ml/build.rs` — added kernel entry. Build-rs + cubin static `HOLD_REWARD_CAP_UPDATE_CUBIN` exposed in + `gpu_dqn_trainer.rs:2190`. + +- **Producer wrapper**: `crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs` — + new `HoldRewardCapUpdateOps` struct (mirrors `RewardCapUpdateOps`) + with `new`/`launch` impl; pre-loads `CudaFunction` at construction per + `pearl_no_host_branches_in_captured_graph.md`. + +- **Trainer wiring**: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`: + - Struct field `hold_reward_cap_update_ops` at line 7828. + - Constructor instantiation at line 24409. + - Field-init in struct literal at line 25293. + - Method `launch_hold_reward_cap_update` placed right after + `launch_reward_cap_update`. + +- **Per-epoch wire-up**: + `crates/ml/src/trainers/dqn/trainer/training_loop.rs` line 511 — + `fused.trainer().launch_hold_reward_cap_update(step_ret_ptr, + trade_close_ptr, total_samples)` placed right after + `launch_reward_cap_update` so both cap producers share the same + source buffers (read-only; independent ISV slot pairs at + [452]/[453] and [483]/[484]). + +- **HEALTH_DIAG**: `training_loop.rs` (post `td_error_pre` emit) — new + line `HEALTH_DIAG[N]: hold_reward_cap [pos={:.4} neg={:.4} + fire_rate={:.4}]`. Reads ISV[483]/[484]/[486] via + `read_isv_signal_at`; slot 486 (HOLD_OPP_COST_FIRE_RATE_EMA) is the + engagement-canary diag for Phase 8 to populate. + +- **GPU oracle tests**: + `crates/ml/tests/sp18_hold_reward_oracle_tests.rs` — three new tests + in the `mod gpu` block: + - `hold_reward_cap_producer_drives_slots` — T5 from spec: 100 + simulated Long/Short closes with jittered ~N(0, 0.01); verifies + ISV[483] is bilaterally clamped to MIN_BOUND=0.5 (small p99 ⇒ + target_cap << bound), ISV[484] = -2 × ISV[483], filler outliers + (close=0) filtered out. + - `hold_reward_cap_pearl_a_first_observation_replace` — Pearl-A + bootstrap: 4 large closes with `step_ret = [+10, -20, +30, -40]`, + target_cap = max(p99, max) × 1.5 ≈ 76.5, clamped to MAX_BOUND=50, + Pearl-A REPLACE on sentinel writes 50.0 directly (NOT blended + with sentinel 5.0). + - `hold_reward_cap_no_closes_preserves_isv` — no closes + (trade_close=0 ∀i): every ISV slot (POS, NEG, all 6 Welford) + preserved bit-exactly. Verifies the cold-start fallback path. + +- **Audit script**: `scripts/audit_sp18_consumers.sh --check` exits 0 + (no fingerprint drift in tracked sections — SP18 D-leg/B-leg legacy + references unchanged; the new Phase 3 producer is in its own + symbol-namespace not tracked by the audit's snapshot, which is by + design — the audit tracks the DELETED SP13/SP16 legacy chain). + +**Validation: cargo check --workspace clean (1m 12s); 3 GPU oracle +tests pass on local RTX 3050 Ti (1.93s wallclock).** + +### Phase 3 design notes + +- **Why Wiener-α (vs. fixed α=0.01 on the position-side cap)**: the + position-side cap drifts slowly because the underlying alpha + distribution is approximately stationary over training. The Hold-side + cap distribution is intrinsically NON-stationary as the policy adapts + (sizing tightens, trail shifts). Welford-derived α tracks that + non-stationarity; the 0.4 floor prevents the canonical "α decays to 0 + once policy converges" trap (SP16 T3 regression at + train-multi-seed-b5gmp sha 641aa0dfd Fold 1: α=0.248 at HD4 → Hold% + 52.6% at HD5 — see `pearl_wiener_alpha_floor_for_nonstationary`). + +- **Why no separate Long/Short flag buffer**: `step_ret_per_sample[i]` + is non-zero ONLY on close events from positions that had non-zero + pre_sign — i.e. Long or Short positions. Hold/Flat positions don't + generate non-zero step_ret per `experience_kernels.cu:2855`. + Filtering by `(is_close && step_ret != 0)` is the cheapest correct + extraction of the Long/Short close distribution; no extra + dir_idx_per_sample buffer needed. + +- **Cold-start fallback behavior**: when `total_count == 0` (no closes + this epoch — Hold-heavy phase), the kernel early-returns at the + guard in Phase 2 of its body. ALL slots (POS, NEG, all 6 Welford) + remain unchanged. The Phase 2 consumer + (`compute_sp18_hold_opportunity_cost`) sees the cold-start sentinels + (5.0/-10.0) and routes through its own `is_sentinel` branch to the + macro defaults +5/-10 — bit-identical to pre-Phase-3. + +- **Coexistence with `reward_cap_update_kernel`**: both producers + read the same `step_ret_per_sample` + `trade_close_per_sample` + buffers (read-only; no shared state). They drive independent ISV + slot pairs and use different filters. Per-epoch boundary launches + are sequential (position-side first, Hold-side second) so the + two_kernel sequence is graph-capturable (single stream, no host + branches). The redundant block-tree-reduce on the same input + buffers is a micro-cost (≤ 2 × 64K samples × 4 floats per epoch = + cheap relative to per-step training cost). + ## B-leg verification (Phase 1.1 sub-objective) Per the plan's locked B-leg checklist: