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

feat(sp18 v2 P3.T1-T5): adaptive HOLD_REWARD_POS/NEG_CAP producer kernel

Browse Source 
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-09 11:05:03 +02:00
parent 1f4cc0f207
commit 5ea5aa9b8e
8 changed files with 1455 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

32
crates/ml/build.rs
View File

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

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

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

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

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

394
crates/ml/src/cuda_pipeline/hold_reward_cap_update_kernel.cu Normal file
View File

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

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

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

482
crates/ml/tests/sp18_hold_reward_oracle_tests.rs
View File

@@ -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<CudaStream>) -> 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::<f32>() 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<CudaStream>,
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<f32> = vec![10.0, -20.0, 30.0, -40.0];
let trade_close: Vec<i32> = 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<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_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");
}
}

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

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

126
docs/sp18-wireup-audit.md
View File

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