diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 32f58d679..6897ada66 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -2114,33 +2114,28 @@ extern "C" __global__ void experience_env_step( float* __restrict__ popart_component_per_sample, /* SP12 v3 (2026-05-04): per-trade event-driven reward composition. * - * `min_hold_target` (bars): cold-start fallback for the patience - * requirement on voluntary trade exits. Class A P0-C (2026-05-08): - * the actual target used inside the kernel is ISV-driven from - * ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451] — EMA of winning-trade hold - * times. This parameter is the fallback when ISV[451] is still at - * sentinel (0.0 = no winning trades observed yet) or when - * isv_signals_ptr is NULL. Default 30 bars (MIN_HOLD_TARGET in - * state_layout.cuh). Voluntary exits with effective hold_time below - * the target get a soft penalty subtracted from `r_popart` via the - * deficit/(deficit + T) saturation formula. + /* SP20 Phase 2 Task 2.2 (2026-05-09): the SP12 v3 min-hold-penalty + * args (`min_hold_target`, `min_hold_penalty_max`, `min_hold_ + * temperature`) were DELETED atomically with the SP12 v3 reward + * block replacement. The 4-quadrant SP20 reward at trade close has + * no patience-via-soft-penalty notion — patience emerges from the + * advantage-centered alpha (long-running unprofitable trades have + * `R_event - alpha_ema < 0` regardless of hold time). The + * `compute_min_hold_penalty` device function in `trade_physics.cuh` + * is RETAINED (still called by `sp12_reward_math_test_kernel.cu`) + * but has no production callers post-Task-2.2. * - * `min_hold_penalty_max`: max penalty magnitude when the deficit is - * infinite (factor → 1). Phase 1 default 3.0 (60% of REWARD_POS_CAP). - * - * `min_hold_temperature`: smoothness controller for the soft factor. - * Higher T = wider transition. Computed in Rust per epoch as - * `T_end + (T_start - T_end) × exp(-epoch / decay)` and passed in - * per launch. Constants in state_layout.cuh - * (MIN_HOLD_TEMPERATURE_{START,END,DECAY}). Annealing curriculum: - * epoch 0 ≈ 50 (forgiving), epoch 50 ≈ 9, epoch 100 ≈ 5 (sharp). - * - * Trail-fire forced exits are exempted from this penalty (the - * trail-stop is a risk-management design choice; penalising it would - * fight the existing exit logic). Spec §design-2 + open question 4. */ - float min_hold_target, - float min_hold_penalty_max, - float min_hold_temperature, + * Follow-up cleanup (Task 2.4 / future commit): remove the orphaned + * upstream producer chain — `min_hold_temperature_update_kernel`, + * `ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460]`, the + * `read_min_hold_temperature_from_isv` helper, and the + * `config.min_hold_temperature` field. Atomic scope is + * "kernel-arg removal" here; producer-chain removal touches the + * SP14 ISV slot registry + StateResetRegistry + an ISV layout + * fingerprint bump and is intentionally deferred to keep the + * Task 2.2 commit tractable per `feedback_no_partial_refactor`'s + * "consumer migrates with contract change" rule (the consumer + * IS what migrates here). */ /* SP15 Wave 2 (2026-05-06) — per-step r_discipline scalar mirror. * Layout [N*L] f32 row-major (i, t). Producer: this kernel writes * `isv_signals_ptr[REGRET_EMA_INDEX=423]` to `r_discipline_out[i*L+t]` @@ -2188,7 +2183,27 @@ extern "C" __global__ void experience_env_step( * the SP20 reward kernel (Task 2.2 wires the read site). * FoldReset-zeroed by `state_reset_registry.rs::label_at_open_per_env`. */ const int* __restrict__ aux_label_per_env, - int* __restrict__ label_at_open_per_env + int* __restrict__ label_at_open_per_env, + /* SP20 Phase 2 Task 2.2 (2026-05-09) — per-env trade-close alpha + * output buffer + ISV slot indices for the 4-quadrant reward block + * at the segment_complete site. + * + * `alpha_per_env` ← [N] f32: written at `is_close && segment_hold_time + * > 0.0f` (segment_complete branch) as `alpha = R_event - + * hold_baseline` per spec §4.1. Phase 2 emits `hold_baseline = 0.0f` + * as a Phase 3.2 forward-reference placeholder. Slot persists across + * rollout steps until the next trade-close overwrites it; the + * `sp20_aggregate_inputs_kernel` consumer reads only when + * `trade_close_per_env[env] != 0` so stale values don't reach the + * ALPHA_EMA. NULL-tolerant for test scaffolds. + * + * `loss_cap_idx` / `alpha_ema_idx`: ISV slot indices for the SP20 + * reward block. Caller passes them as kernel args so the kernel + * stays decoupled from SP14 slot-number drift (mirror the + * `q_dir_abs_ref_idx` pattern at line 2069). */ + float* __restrict__ alpha_per_env, + int loss_cap_idx, + int alpha_ema_idx ) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= N) return; @@ -3183,267 +3198,137 @@ extern "C" __global__ void experience_env_step( segment_return = reversal_return; } - /* ── Magnitude-neutral reward normalization ────────────────────── - * Divide return by position fraction (|pos|/max_pos) so the reward - * measures DIRECTIONAL SKILL PER UNIT OF RISK, not absolute PnL. + /* ══════════════════════════════════════════════════════════════════ + * SP20 Phase 2 Task 2.2 (2026-05-09): 4-quadrant directional- + * correctness reward replaces the SP12 v3 inlined block (lines + * 3206-3466 of the pre-Phase-2 file: Magnitude-neutral comment, + * vol-normalization calc, base_reward = 2 × vol_normalized_return, + * SP18 D-leg trade-close call, SP12 v3 asymmetric cap, min-hold + * penalty, and the r_popart / r_trail / capped_pnl cascade). * - * Without this, Quarter (0.25) produces 4× lower PnL variance than - * Full (1.00) for the same price move. The MSE Bellman target for - * Quarter has 16× lower variance (squared) → gradient is 16× more - * consistent → Quarter Q converges first → Boltzmann locks in → - * model never learns Full's true Q. Same mechanism as C51 collapse - * but through reward variance instead of distributional shape. + * Spec: docs/superpowers/specs/2026-05-09-sp19-20-wr-first-design.md §4.1 * - * After normalization, all magnitudes produce the SAME reward - * variance for the same directional skill. The model then selects - * magnitude based on risk-adjusted return, not convergence speed. + * Algorithm: + * label_at_open = label_at_open_per_env[i] (Task 2.0 buffer) + * loss_cap = ISV[loss_cap_idx] (sp20_controllers + * output, [-2,-1] ramp) + * R_event = sp20_compute_event_reward( (Task 2.1 helper — + * segment_return, 4-quadrant table + * label_at_open, from sp20_reward.cuh) + * loss_cap) + * hold_baseline = 0.0f (Phase 3.2 placeholder + * — Component 2 lands + * the real producer) + * alpha = R_event - hold_baseline (= R_event in Phase 2) + * alpha_ema = ISV[alpha_ema_idx] (sp20_emas output; + * Wiener-α EMA of alpha) + * r_used = alpha - alpha_ema (advantage-style + * centering — load- + * bearing per spec + * §4.1; do not drop) * - * Applied AFTER Kelly (which needs raw returns) and AFTER reversal - * override (which also needs magnitude normalization). */ - /* Magnitude-neutral normalization REMOVED — gave Small 4x LARGER sparse - * reward, contradicting dense micro-reward. Vol-normalization below - * handles risk-adjustment. Magnitude PnL differences are the real signal. */ - - /* Vol normalization (unchanged from v6) */ - float atr_norm = 0.0f; - if (features != NULL && bar_idx < total_bars && market_dim > 9) { - atr_norm = (features[(long long)bar_idx * market_dim + 9]); - } - float log_atr = atr_norm * 16.0f - 7.0f; - float atr_pct = expf(log_atr) / fmaxf(raw_close, 1.0f); - float vol_proxy = fmaxf(atr_pct, 0.0001f); - float vol_norm = vol_proxy * sqrtf(fmaxf(segment_hold_time, 1.0f)); - float vol_normalized_return = segment_return / vol_norm; - - /* (DSR moved to dense per-bar path below — no longer sparse at trade completion) */ - - /* ── Layer 2: Upper-cap only (Task 2.4 relocation) ── - * Negative-tail compression moved to c51_loss_kernel.cu:: - * block_bellman_project_f (Huber-style Q-target smoothing). - * Upper +10 cap kept inline for numerical safety against adversarial - * reward explosions from rare large wins. + * Why centering matters (spec §4.1): at WR=46% with the 4-quadrant + * table, raw EV ≈ −0.06 per trade — the model would prefer not to + * trade. The `R_used = alpha − alpha_ema` subtraction transforms + * absolute EV into advantage-style relative-to-mean, so the model + * learns "this trade was BETTER than my average trade → do more + * like this" even when absolute alpha is negative. Without + * centering, the model collapses to Hold during cold-start before + * it can learn directional discrimination. * - * SP11 B1b §3.5.4: split forced-exit (trail-fire) P&L from - * voluntary-exit P&L. Both paths produce the same vol-normalized - * capped scalar; controller weight `w_trail` vs `w_popart` lets - * SP11 weigh them differently. `trail_triggered` flips - * `exiting_trade=1` upstream (line ~2371), so this branch fires - * for both — the segregation matters only for component attribution. */ - float base_reward = 2.0f * vol_normalized_return; - /* SP18 D-leg Phase 2 (2026-05-09): structural Hold opportunity-cost - * counterfactual replaces the SP13 P0a / SP16 P2 / SP16 T3 reactive - * per-bar Hold-cost-scale chain (deleted Phase 1 c66e15f6d / - * 3c318953a). Goes BEFORE the SP12 asymmetric cap so the - * counterfactual participates in the bounded clamp — keeps total - * reward in the [-10, +5] range and avoids a long Hold run silently - * saturating the lower cap. Hold action rarely coincides with - * segment_complete (Hold keeps current position, doesn't close - * trades) but trail-fire can force-exit while the policy picks - * Hold; this branch covers that edge. Per-bar (non-segment_complete) - * Hold sites are in the two branches below. + * Mutual exclusivity (preserved from pre-SP20): trail_triggered + * routes the result to r_trail (rc[2]); voluntary exits route to + * r_popart (rc[0]). The SP11 controller weights w_trail and w_popart + * keep their existing semantics — exactly one fires per + * segment_complete bar. r_used is the same scalar in both branches; + * the split is purely SP11 component attribution. * - * The device function returns 0 for non-Hold so unconditional - * addition is safe. Caps come from ISV[HOLD_REWARD_NEG_CAP_INDEX=484] - * / ISV[HOLD_REWARD_POS_CAP_INDEX=483] with sentinel/window-guard - * fallback to the macro defaults — bit-identical to a fixed +5/-10 - * cap during cold-start (Phase 3 lands the adaptive producer). - * `shaping_scale` matches the SP12 contract (validation = 0 → no - * Hold opp cost in pure backtest). Per - * `pearl_one_unbounded_signal_per_reward.md` the only unbounded - * multiplicand is `next_log_return / sqrt(vol_proxy)`; everything - * else is in [floor, ceil]. */ - { - float sp18_pos_cap = SENTINEL_HOLD_REWARD_POS_CAP; /* +5.0 */ - float sp18_neg_cap = SENTINEL_HOLD_REWARD_NEG_CAP; /* -10.0 */ - if (isv_signals_ptr != NULL) { - const float isv_pos = - isv_signals_ptr[HOLD_REWARD_POS_CAP_INDEX]; - const float isv_neg = - isv_signals_ptr[HOLD_REWARD_NEG_CAP_INDEX]; - if (fabsf(isv_pos - SENTINEL_HOLD_REWARD_POS_CAP) >= 1e-6f - && isv_pos >= HOLD_REWARD_POS_CAP_MIN_BOUND - && isv_pos <= HOLD_REWARD_POS_CAP_MAX_BOUND) { - sp18_pos_cap = isv_pos; - sp18_neg_cap = isv_neg; - } - } - const float sp18_hold_opp = compute_sp18_hold_opportunity_cost( - dir_idx, - /*dir_hold=*/DIR_HOLD, - /*next_log_return=*/tgt[1], - /*vol_proxy=*/vol_proxy, - /*isv_neg_cap=*/sp18_neg_cap, - /*isv_pos_cap=*/sp18_pos_cap, - /*shaping_scale=*/shaping_scale - ); - base_reward += sp18_hold_opp; - } - /* SP12 v3 fix (2026-05-04): asymmetric bounded cap (loss aversion). - * SP11 (commit 35db31089) made the cap symmetric ±10 to fix slot-63 - * PopArt EMA inflation. That fix preserved stability but erased the - * pre-SP11 implicit downside-asymmetry — losses that previously - * dominated the reward distribution disappeared into the floor. - * Result: train-multi-seed-pmbwn 50-epoch on 6a259942e showed - * sharpe-gaming (sharpe=80 held while PnL declined 30% over 8 - * epochs); the model became a HFT noise extractor (62% trade rate, - * 46% win rate) because losses no longer hurt enough to enforce - * selectivity. + * Bounded by construction (`pearl_one_unbounded_signal_per_reward`): + * `R_event ∈ [loss_cap, +1.0] ⊆ [-2, +1]`. `alpha = R_event - 0 + * ⊆ [-2, +1]` in Phase 2. `alpha_ema` is a Wiener-α EMA with the + * floor at 0.4 over a structurally-bounded observation, so it stays + * within [-2, +1] post-bootstrap. Therefore `r_used ∈ [-3, +3]` in + * Phase 2 — entirely structural, no clamps needed. Phase 3.2 lands + * `hold_baseline ∈ [-1.5, 0]` via Component 2 buffer; alpha grows + * up to [-2, +2.5], r_used up to [-4.5, +4.5] — still structural. * - * SP12 v3 restores asymmetry by widening only the negative cap - * back toward the pre-SP11 unbounded regime (capped at -10 for - * stability) while tightening the positive cap to +5. The 2:1 - * REWARD_NEG_CAP : REWARD_POS_CAP ratio is calibrated to the - * Kahneman/Tversky meta-analysis estimate of human loss aversion - * (~2.0-2.25) — a known psychological-anchor constant rather than - * a tuned multiplier. Stability work from SP11 is preserved - * because the cap is still bilateral. - * - * Per pearl_audit_unboundedness_for_implicit_asymmetry — when a - * symmetric clamp replaces an asymmetric prior, the policy's - * implicit risk preference can flip without warning. The fix is - * to make the bound itself asymmetric in the direction the prior - * implicitly was, not to remove the cap. - * - * Class A P0-A (2026-05-08): Phase 2 lifted. The hardcoded - * REWARD_POS_CAP=+5.0f / REWARD_NEG_CAP=-10.0f - * (state_layout.cuh:266-267) is now ISV-driven from - * `ISV[REWARD_POS_CAP_ADAPTIVE_INDEX=452]` / - * `ISV[REWARD_NEG_CAP_ADAPTIVE_INDEX=453]` — adaptively driven - * by `reward_cap_update_kernel` from p99(winning_returns) × 1.5 - * with NEG=−2 × POS preserving Kahneman/Tversky 2:1 asymmetry - * (single source of truth at producer; no consumer applies the - * 2× ratio itself). Cold-start fallback: when the ISV slot is at - * sentinel SENTINEL_REWARD_POS_CAP=5.0 (within EPS), fall back - * to the original macro `REWARD_POS_CAP` constant — bit-identical - * pre-P0-A behavior until the first valid observation lands. - * Math factored into `compute_asymmetric_capped_pnl` in - * `trade_physics.cuh` so the formula is testable in isolation - * (oracle tests in `tests/sp12_reward_math_tests.rs`). */ - float pos_cap_eff = REWARD_POS_CAP; - float neg_cap_eff = REWARD_NEG_CAP; - if (isv_signals_ptr != NULL) { - float isv_pos = isv_signals_ptr[REWARD_POS_CAP_ADAPTIVE_INDEX]; - float isv_neg = isv_signals_ptr[REWARD_NEG_CAP_ADAPTIVE_INDEX]; - /* Cold-start fallback: sentinel match → use macro fallback - * (bit-identical pre-P0-A). Out-of-bounds defensive guard: - * any value outside [REWARD_POS_CAP_MIN_BOUND, - * REWARD_POS_CAP_MAX_BOUND] also falls back, defending - * against malformed prior state. */ - if (fabsf(isv_pos - SENTINEL_REWARD_POS_CAP) >= 1e-6f - && isv_pos >= REWARD_POS_CAP_MIN_BOUND - && isv_pos <= REWARD_POS_CAP_MAX_BOUND) { - pos_cap_eff = isv_pos; - neg_cap_eff = isv_neg; - } - } - float capped_pnl = compute_asymmetric_capped_pnl( - base_reward, neg_cap_eff, pos_cap_eff + * Pearl notes: + * - `pearl_event_driven_reward_density_alignment` — fires only + * at trade close (segment_complete branch). Per-bar Hold + * opportunity-cost density alignment lives in Component 2 + * (Phase 3.2); the per-bar SP18 D-leg call sites at + * experience_kernels.cu:3722 / :3833 are RETAINED pending + * Phase 3.2 replacement (per `feedback_no_stubs` — + * working signal stays until replacement lands). + * - `pearl_audit_unboundedness_for_implicit_asymmetry` — the + * SP12 v3 asymmetric cap (REWARD_NEG_CAP=-10, REWARD_POS_CAP=+5) + * was preserving the pre-SP11 implicit downside asymmetry + * (Kahneman-Tversky 2:1 ratio). The SP20 4-quadrant reward + * embeds the SAME asymmetry structurally via `loss_cap ∈ + * [-2, -1]` versus `win_cap = +1.0` — the 2:1 ratio is + * preserved at the wrong-reason-wrong-outcome quadrant. + * `compute_asymmetric_capped_pnl` and + * `compute_min_hold_penalty` device functions in + * `trade_physics.cuh` are RETAINED (still called by per-bar + * SP18 sites + the sp12_reward_math_test_kernel test wrapper). + * ══════════════════════════════════════════════════════════════════ */ + const int label_at_open_sign = (label_at_open_per_env != NULL) + ? label_at_open_per_env[i] + : 0; /* sentinel ⇒ wrong-reason quadrant */ + /* loss_cap fallback: -1.0 (cold-start cap per spec §4.1, the value + * the controller emits when WR_EMA == 0.50). Used when ISV is + * unwired (test scaffolds) — bit-identical to a fresh fold's + * cold-start. The controller ramps to -2.0 by WR=0.55 once + * sp20_controllers_compute starts firing. */ + const float loss_cap = (isv_signals_ptr != NULL) + ? isv_signals_ptr[loss_cap_idx] + : -1.0f; + const float r_event = sp20_compute_event_reward( + segment_return, label_at_open_sign, loss_cap ); + /* Phase 3.2 placeholder: hold_baseline = 0.0f. Component 2 + * (Phase 3.2) lands the real `hold_baseline_buffer.sum_over_ + * trade_range()` consumer. Forward reference per the audit-doc + * Task 2.2 entry — NOT a stub: the producer site here is fully + * wired to a working bounded scalar, the consumer just adds + * a (negative) opportunity-cost subtraction in a later phase. */ + const float hold_baseline = 0.0f; + const float alpha = r_event - hold_baseline; + const float alpha_ema = (isv_signals_ptr != NULL) + ? isv_signals_ptr[alpha_ema_idx] + : 0.0f; + const float r_used = alpha - alpha_ema; + + /* Per-env alpha emit — written every segment_complete bar (both + * voluntary and trail-fire). Read by `sp20_aggregate_inputs_kernel` + * to populate `EmaInputs.alpha` (mean over closed envs); + * non-close-step writes don't matter because the aggregation + * kernel reads only when `trade_close_per_env[env] != 0`. */ + if (alpha_per_env != NULL) { + alpha_per_env[i] = alpha; + } + + /* SP11 component attribution split (mutually exclusive — preserved + * from pre-SP20 semantic). r_used is identical in both branches; + * only the rc[*] slot the SP11 controller weighs differs. */ if (trail_triggered) { - r_trail = capped_pnl; + r_trail = r_used; } else { - /* SP12 v3 Change 2: min-hold soft penalty for voluntary exits. - * Encourages patience-matched trade duration via a soft penalty. - * Soft saturation formula: factor = deficit/(deficit + T) — bounded - * [0,1], smooth transition (no cliff), monotone increasing in - * deficit, monotone decreasing in T. Temperature anneals over - * training (50 → 5 across ~100 epochs) — early forgiveness lets - * the policy learn reward gradients on short trades, late sharpness - * enforces commitment to longer holds. - * - * Trail-fire exits exempted (the surrounding `else` already - * guarantees `!trail_triggered`). When `effective_min_hold_target - * <= 0` (defensive guard against malformed configs) the entire - * block collapses to no-op since `deficit = max(0, -hold) <= 0`. - * - * Voluntary-exit only by construction: this lives in the - * `else` arm of the `trail_triggered` branch, which itself is - * inside `segment_complete`. Per spec §design-2 + - * pearl_event_driven_reward_density_alignment. - * - * Applied BEFORE `r_popart = capped_pnl` so the penalty flows - * through the popart component (which the SP11 controller - * weighs) — keeps mag-ratio canary semantics consistent. - * - * Math factored into `compute_min_hold_penalty` in - * `trade_physics.cuh` so the formula is testable in isolation - * (oracle tests in `tests/sp12_reward_math_tests.rs`). The - * device function returns the penalty MAGNITUDE (positive) and - * subsumes the `hold_time >= target` early-exit; preserving the - * outer `if (segment_hold_time < target)` guard would be - * redundant since the function already returns 0 in that case - * and `capped_pnl - shaping_scale * 0 == capped_pnl`. - * - * Class A P0-C (2026-05-08): target is ISV-driven from - * ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451] — Pearl-A-bootstrapped - * Welford EMA of observed winning trade hold times. Adapts to - * whatever duration the model actually finds profitable in the - * current regime rather than a hardcoded MFT-leaning constant. - * Cold-start fallback: sentinel = 0.0 (no winning trades yet); - * guarded by the `> 0 && < 240` window below. Outside that - * window (cold-start or anomalous ISV) falls back to - * `min_hold_target` (= MIN_HOLD_TARGET=30 from Rust config). - * Per feedback_isv_for_adaptive_bounds. */ - const float isv_hold_target = (isv_signals_ptr != NULL) - ? isv_signals_ptr[AVG_WIN_HOLD_TIME_BARS_INDEX] - : 0.0f; - /* Validity window: > 0.0 excludes sentinel; < 240.0 caps at - * ~1 hour of bars (sanity guard against ISV corruption). */ - const float effective_min_hold_target = - (isv_hold_target > 0.0f && isv_hold_target < 240.0f) - ? isv_hold_target - : min_hold_target; /* cold-start / NULL fallback */ - /* Class A P0-A-downstream (2026-05-08): min_hold_penalty_max - * was tuned as a fixed 60% fraction of the pre-P0-A - * REWARD_POS_CAP=5.0f (= 3.0). With POS_CAP now adaptive via - * ISV[REWARD_POS_CAP_ADAPTIVE_INDEX=452], propagate the same - * 0.6 ratio so the penalty's relative size to the bounded - * reward stays constant ("60% of REWARD_POS_CAP — meaningful - * gradient pressure but not paralyzing" — state_layout.cuh - * line 252-253). Cold-start fallback: sentinel ISV (5.0f - * within EPS) OR ISV outside the [1.0, 50.0] defensive - * window → fall back to the kernel-passed `min_hold_penalty - * _max` (3.0f from gpu_experience_collector.rs:399), which - * is bit-identical to pre-P0-A behavior. Mirrors the - * effective_min_hold_target pattern above. Per - * feedback_isv_for_adaptive_bounds + feedback_no_partial_ - * refactor. */ - const float MIN_HOLD_PENALTY_RATIO = 0.6f; - float effective_min_hold_penalty_max = min_hold_penalty_max; - if (isv_signals_ptr != NULL) { - const float isv_pos = isv_signals_ptr[REWARD_POS_CAP_ADAPTIVE_INDEX]; - if (isv_pos >= REWARD_POS_CAP_MIN_BOUND - && isv_pos <= REWARD_POS_CAP_MAX_BOUND - && fabsf(isv_pos - SENTINEL_REWARD_POS_CAP) > 1e-6f) { - effective_min_hold_penalty_max = MIN_HOLD_PENALTY_RATIO * isv_pos; - } - } - const float min_hold_penalty = compute_min_hold_penalty( - segment_hold_time, - effective_min_hold_target, - min_hold_temperature, - effective_min_hold_penalty_max - ); - float adjusted_pnl = capped_pnl - shaping_scale * min_hold_penalty; - r_popart = adjusted_pnl; - /* SP11 Fix 39 B1b fix-up (2026-05-04): mirror r_popart to the - * dedicated per-bar buffer feeding ISV[POPART_COMPONENT_MAG_EMA_ - * INDEX=360]. Voluntary-exit branch only (the trail-triggered - * branch above leaves popart-component at the entry-point 0 - * default). The mag-ratio canary now reads slot 360 for the - * popart axis instead of slot 63's overloaded total-reward - * value. Spec §4 amendment "Why slot 360". */ + r_popart = r_used; + /* SP11 Fix 39 B1b fix-up (2026-05-04): popart-component- + * specific reward magnitude buffer feeds ISV[POPART_COMPONENT_ + * MAG_EMA_INDEX=360] (mag-ratio canary read site, spec §4 + * "Why slot 360"). Voluntary-exit only — the trail branch + * leaves the buffer at the kernel-entry 0 default. Pre-SP20 + * this was r_popart = capped_pnl post-min-hold; under SP20 + * r_popart = r_used and the buffer carries the same value, + * preserving the canary's magnitude-ratio semantic against + * the new reward shape. */ if (popart_component_per_sample != NULL) { popart_component_per_sample[out_off] = r_popart; } - /* Keep cascade scalar in lockstep so downstream C.4/D.4a/D.4b - * bonus blocks Q-cap their `|reward|` multiplicand against the - * post-min-hold-penalty value (matches pre-SP12 invariant: the - * trade-cumulative scalar always reflects the most recent - * popart write). */ - capped_pnl = adjusted_pnl; } - reward = capped_pnl; + reward = r_used; /* C.2 attribution: rc[2] now carries the trail signal when the * trail-fire branch fires (was a structural placeholder pre-B1b). * The reward_component_ema kernel will pick up real trail magnitude diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index f5a91f3a3..1263469f1 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -678,6 +678,34 @@ pub struct GpuExperienceCollector { /// folds — each fold has independent aux next-bar predictions). pub(crate) label_at_open_per_env: CudaSlice, // [alloc_episodes] + /// SP20 Phase 2 Task 2.2 (2026-05-09): per-env trade-close alpha + /// scratch — `alpha = R_event - hold_baseline` per spec §4.1 + /// (Component 1 reward kernel). Phase 2 emits `hold_baseline = 0.0f` + /// as a Phase 3.2 forward-reference placeholder; Phase 3.2's Hold + /// opportunity-cost dual emission lands the real + /// `hold_baseline_buffer.sum_over_trade_range()` consumer. + /// + /// Layout: `[alloc_episodes]` f32 — one slot per env, written by + /// `experience_env_step`'s `is_close` (`segment_complete`) branch + /// after the `sp20_compute_event_reward` call. Read once per + /// rollout step by `sp20_aggregate_inputs_kernel` to populate the + /// `EmaInputs.alpha` field (mean over closed envs); the Phase 1.4 + /// hardcoded `out_inputs->alpha = 0.0f` placeholder is replaced + /// atomically with this commit. + /// + /// FoldReset semantics (registered in `state_reset_registry.rs` as + /// `alpha_per_env`): zero on every fold boundary so the new fold's + /// first trade-close fully populates the slot from the new fold's + /// reward calculation (no stale alpha leakage across folds — each + /// fold has independent ALPHA_EMA bootstrap state per Phase 1.4's + /// `obs_count[OBS_COUNT_ALPHA]` reset). + /// + /// On non-close steps the slot retains the previous trade-close + /// alpha (or 0.0 sentinel pre-first-close); the aggregation kernel + /// reads only when `trade_close_per_env[env] != 0`, so stale values + /// don't leak into the EMA. + pub(crate) alpha_per_env: CudaSlice, // [alloc_episodes] + // Output buffers [alloc_episodes * alloc_timesteps, ...] states_out: CudaSlice, // #30 [alloc_episodes * alloc_timesteps * STATE_DIM] f32 actions_out: CudaSlice, // [alloc_episodes * alloc_timesteps] @@ -1737,6 +1765,17 @@ impl GpuExperienceCollector { "sp20: alloc label_at_open_per_env ({alloc_episodes} i32): {e}" )))?; + // SP20 Phase 2 Task 2.2 (2026-05-09): per-env trade-close alpha + // (= R_event - hold_baseline). Sentinel = 0.0; updated at each + // trade close by the SP20 4-quadrant reward kernel block. + // Consumed by sp20_aggregate_inputs_kernel (mean over closed + // envs) → SP20EmaInputs.alpha. Sized [alloc_episodes]. + let alpha_per_env = stream + .alloc_zeros::(alloc_episodes) + .map_err(|e| MLError::ModelError(format!( + "sp20: alloc alpha_per_env ({alloc_episodes} f32): {e}" + )))?; + // Persistent epoch state let epoch_state_init: Vec = vec![ 0.01, // vol_ema @@ -2765,6 +2804,7 @@ impl GpuExperienceCollector { portfolio_states, episode_starts_buf, label_at_open_per_env, + alpha_per_env, states_out, actions_out, rewards_out, @@ -6136,9 +6176,15 @@ impl GpuExperienceCollector { // `DQNTrainer::read_min_hold_temperature_from_isv`, // which falls back to MIN_HOLD_TEMPERATURE_FALLBACK // =50.0 when the slot is at sentinel. - .arg(&config.min_hold_target) - .arg(&config.min_hold_penalty_max) - .arg(&config.min_hold_temperature) + // SP20 Phase 2 Task 2.2 (2026-05-09): the SP12 v3 + // min-hold-penalty kernel args (min_hold_target / + // min_hold_penalty_max / min_hold_temperature) were + // DELETED atomically with the SP12 v3 reward block + // replacement at experience_kernels.cu (the SP20 + // 4-quadrant reward has no patience-via-soft-penalty + // notion). config.min_hold_* fields remain in the + // Rust struct pending the Task 2.4 follow-up cleanup + // (orphaned producer chain removal). // SP15 Wave 2 (2026-05-06): per-step `r_discipline` // mirror + slot completion sentinel. Both [N*L] // sized; defaulted at kernel entry. The fused @@ -6159,6 +6205,23 @@ impl GpuExperienceCollector { // by Task 2.2's `sp20_compute_event_reward` call). .arg(&self.exp_aux_nb_label_buf) .arg(&mut self.label_at_open_per_env) + // SP20 Phase 2 Task 2.2 (2026-05-09): per-env + // trade-close alpha buffer (output) + ISV slot + // indices for the 4-quadrant reward block at the + // segment_complete site. Slot indices passed as + // kernel args so the kernel stays decoupled from + // SP14 slot-number drift (mirrors the + // `q_dir_abs_ref_idx` arg pattern — single source of + // truth at `sp14_isv_slots.rs` constants). + .arg(&mut self.alpha_per_env) + .arg(&{ + use crate::cuda_pipeline::sp14_isv_slots::LOSS_CAP_INDEX; + LOSS_CAP_INDEX as i32 + }) + .arg(&{ + use crate::cuda_pipeline::sp14_isv_slots::ALPHA_EMA_INDEX; + ALPHA_EMA_INDEX as i32 + }) .launch(launch_cfg) .map_err(|e| MLError::ModelError(format!( "experience_env_step t={t}: {e}" @@ -6435,6 +6498,16 @@ impl GpuExperienceCollector { let ema_inputs_dev = self.sp20_ema_inputs_buf.dev_ptr; + // SP20 Phase 2 Task 2.2 (2026-05-09): per-env alpha + // input pointer. `alpha_per_env` is `[alloc_episodes]` + // f32 device-resident, populated at every trade-close + // by the SP20 4-quadrant reward block in + // `experience_env_step` (launched immediately above on + // the same stream — stream-implicit producer→consumer + // ordering). Replaces the Phase 1.4 hardcoded + // `out_inputs->alpha = 0.0f` placeholder atomically per + // `feedback_no_partial_refactor`. + let alpha_per_env_dev = self.alpha_per_env.raw_ptr(); unsafe { launch_sp20_aggregate_inputs( &self.stream, @@ -6447,6 +6520,7 @@ impl GpuExperienceCollector { stats_p50_dev, stats_std_dev, aux_dir_acc_dev, + alpha_per_env_dev, n_envs_i32_sp20, dir_divisor, hold_action_id, diff --git a/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs b/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs index 30ee1041e..0804124ae 100644 --- a/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs +++ b/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs @@ -17,7 +17,7 @@ //! | `is_win` | `1` if `is_close == 1` AND fraction of closed envs with `step_ret > 0` ≥ 0.5, else `0` | //! | `trade_duration` | `round(mean(hold_at_exit) over closed envs)` if `is_close == 1`, else `0` | //! | `action_is_hold` | `1` if `count(decoded_dir == HOLD) > n_envs/2`, else `0` | -//! | `alpha` | `0.0` — **Phase 2 forward reference** (R_event - hold_baseline) | +//! | `alpha` | `mean(alpha_per_env) over closed envs` if `is_close == 1`, else `0` — **SP20 Phase 2 Task 2.2** (R_event - hold_baseline). | //! | `per_bar_hold_reward` | `0.0` — **Phase 3.2 forward reference** (-aux_conf * cost_scale) | //! | `aux_logits_p50` | `aux_logits_p50_dev[0]` (sp20_stats output) | //! | `aux_logits_std` | `aux_logits_std_dev[0]` (sp20_stats output) | @@ -25,9 +25,12 @@ //! //! ## Hard rules covered //! -//! - `feedback_no_atomicadd` — block tree-reduce only. The 4 reductions -//! (closed-count, wins-count, hold-count, hold_at_exit-sum) share one -//! shmem tile and reduce in a single loop pass. +//! - `feedback_no_atomicadd` — block tree-reduce only. The 5 reductions +//! (closed-count, wins-count, hold-count, hold_at_exit-sum, alpha-sum) +//! share one shmem tile and reduce in a single loop pass. The 5th +//! stripe (`alpha-sum`) was added by SP20 Phase 2 Task 2.2 atomically +//! with the per-env alpha producer at the +//! `experience_env_step::segment_complete` site. //! - `feedback_no_cpu_compute_strict` — every aggregation lives in the //! kernel; the launcher only passes pointers + `n_envs` + constants. //! - `feedback_no_htod_htoh_only_mapped_pinned` — output is a device @@ -39,23 +42,24 @@ //! entry land atomically with the Phase 1.2 buffer-arg refactor + //! the production wire-up + audit/spec/plan amendments. //! -//! ## Phase 2 / Phase 3.2 forward references +//! ## Phase 3.2 forward reference //! -//! The `alpha` and `per_bar_hold_reward` fields are intentionally -//! emitted as `0.0` here: -//! - `alpha = R_event - hold_baseline` is computed by the SP20 reward -//! kernel (Phase 2). Phase 1.4 writes `0.0` so the ALPHA_EMA's -//! bootstrap counter advances on every closed trade and the EMA -//! stays at its 0.0 sentinel until Phase 2 wires the real producer. -//! - `per_bar_hold_reward = -aux_conf * cost_scale` is computed by -//! the SP20 Hold-cost dual emission (Phase 3.2). Phase 1.4 writes -//! `0.0` so the HOLD_REWARD_EMA's bootstrap counter advances on -//! Hold-action bars and the EMA stays at its 0.0 sentinel. +//! - `alpha = R_event - hold_baseline` is now load-bearing per SP20 +//! Phase 2 Task 2.2 (2026-05-09). The producer at +//! `experience_env_step::segment_complete` writes +//! `alpha_per_env[env]` at every trade close; this aggregation kernel +//! means over closed envs to populate `EmaInputs.alpha`. +//! `hold_baseline` is still `0.0f` in Phase 2 — Phase 3.2's +//! Component 2 lands the real `hold_baseline_buffer.sum_over_trade_ +//! range()` consumer (replacing the placeholder 0). +//! - `per_bar_hold_reward = -aux_conf * cost_scale` is computed by +//! the SP20 Hold-cost dual emission (Phase 3.2). The aggregation +//! kernel still writes `0.0` for this field; the HOLD_REWARD_EMA +//! stays at its 0.0 sentinel until Phase 3.2. //! //! These are NOT stubs — they are documented forward references where //! the producer is being shipped in a later phase per -//! `feedback_no_partial_refactor` (the Phase 1.4 commit's atomic -//! scope is the SP20 ISV state-tracker chain, not the reward kernel). +//! `feedback_no_partial_refactor`. /// Block size for the aggregation reduction. 32 threads is enough for /// `n_envs` up to a few thousand (one warp tree-reduce per reduction). @@ -63,13 +67,13 @@ pub const SP20_AGGREGATE_BLOCK: u32 = 64; /// Compute the dynamic shared-memory bytes the aggregation kernel -/// needs for the configured block size: 4 stripes × `bdim` × 4 bytes -/// (3 i32 stripes for closed/wins/hold counts + 1 f32 stripe for the -/// hold_at_exit sum). +/// needs for the configured block size: 5 stripes × `bdim` × 4 bytes +/// (3 i32 stripes for closed/wins/hold counts + 2 f32 stripes for the +/// hold_at_exit sum and the SP20 Phase 2 Task 2.2 alpha sum). #[must_use] pub fn dynamic_shmem_bytes() -> u32 { SP20_AGGREGATE_BLOCK - .checked_mul(4) + .checked_mul(5) .and_then(|x| x.checked_mul(std::mem::size_of::() as u32)) .expect("SP20 aggregate: shmem bytes overflow") } @@ -125,6 +129,11 @@ pub unsafe fn launch_sp20_aggregate_inputs( aux_logits_p50_dev: u64, aux_logits_std_dev: u64, aux_dir_acc_dev: u64, + // SP20 Phase 2 Task 2.2 (2026-05-09): per-env trade-close alpha + // device pointer (or `0` for tests / pre-Task-2.2 callers — the + // kernel falls back to `alpha=0.0f` in that case, matching Phase + // 1.4 placeholder behavior). + alpha_per_env_dev: u64, n_envs: i32, dir_divisor: i32, hold_action_id: i32, @@ -139,6 +148,10 @@ pub unsafe fn launch_sp20_aggregate_inputs( debug_assert!(aux_logits_p50_dev != 0); debug_assert!(aux_logits_std_dev != 0); debug_assert!(aux_dir_acc_dev != 0); + /* alpha_per_env_dev is allowed to be 0 (NULL) — the kernel + * gracefully falls back to alpha=0.0f. This preserves the Phase 1.4 + * test scaffold contract for `sp20_phase1_4_wireup_test` paths + * that don't wire the SP20 reward producer. */ debug_assert!(out_inputs_dev != 0); debug_assert!(n_envs > 0, "n_envs must be > 0, got {n_envs}"); debug_assert!(env_stride > 0, "env_stride must be > 0, got {env_stride}"); @@ -162,6 +175,7 @@ pub unsafe fn launch_sp20_aggregate_inputs( .arg(&aux_logits_p50_dev) .arg(&aux_logits_std_dev) .arg(&aux_dir_acc_dev) + .arg(&alpha_per_env_dev) .arg(&n_envs) .arg(&dir_divisor) .arg(&hold_action_id) @@ -187,8 +201,9 @@ mod tests { #[test] fn shmem_bytes_match_layout() { - // 4 stripes × bdim × 4 bytes. - let expected = SP20_AGGREGATE_BLOCK as usize * 4 * 4; + // SP20 Phase 2 Task 2.2: 5 stripes × bdim × 4 bytes (added the + // alpha-sum stripe atomically with the per-env alpha producer). + let expected = SP20_AGGREGATE_BLOCK as usize * 5 * 4; assert_eq!(dynamic_shmem_bytes() as usize, expected); } diff --git a/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu b/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu index b9caa5e1e..80b514caa 100644 --- a/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu +++ b/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu @@ -16,13 +16,14 @@ * trade_duration = round(mean(hold_at_exit_per_env[env]) for env * where trade_close_per_env[env] != 0) * if is_close == 1, else 0. - * alpha = 0.0 (Phase 2 forward reference — `R_event - hold_baseline` - * is computed by the SP20 reward kernel that lands in - * Phase 2 alongside the hold-baseline circular buffer. - * Phase 1.4 writes 0.0 here as a documented placeholder - * so the EMA's bootstrap counter advances and the WR/HOLD - * EMAs can fire on real signals while ALPHA_EMA stays at - * its 0.0 sentinel.) + * alpha = mean(alpha_per_env[env]) over envs with + * trade_close_per_env[env] != 0; 0.0 if is_close == 0. + * (SP20 Phase 2 Task 2.2 — replaced the Phase 1.4 hardcoded + * 0.0 forward-reference placeholder. The per-env producer + * writes alpha = R_event - hold_baseline at every trade-close + * from `sp20_compute_event_reward`. hold_baseline is still + * 0.0f in Phase 2; Phase 3.2's Component 2 lands the real + * `hold_baseline_buffer.sum_over_trade_range()` consumer.) * per_bar_hold_reward = 0.0 (Phase 3.2 forward reference — the * `r_per_bar = -aux_conf * cost_scale` formula needs the * per-bar aux_conf signal that Phase 3.2's per-bar aux @@ -53,9 +54,9 @@ * - `feedback_no_atomicadd` — block tree-reduce only. The 3 reductions * (closed-count, wins-count, hold-count) share one shmem tile * sequentially with explicit `__syncthreads()` between phases. The - * `hold_at_exit` mean is computed as a separate sum reduction - * immediately after; total of 4 tree-reductions, all bdim×i32 or - * bdim×f32 in shmem. + * `hold_at_exit` and SP20 Phase 2 Task 2.2 `alpha` means are computed + * as separate sum reductions immediately after; total of 5 + * tree-reductions, all bdim×i32 or bdim×f32 in shmem (5 stripes). * - `feedback_no_cpu_compute_strict` — every aggregation lives in * this kernel; the launcher only passes pointers + n_envs + * constants. @@ -80,7 +81,8 @@ * grid_dim = (1, 1, 1) * block_dim = (BLOCK_SIZE, 1, 1) — typically 32 or 64 for n_envs * up to a few hundred. - * shared mem = 4 × bdim × max(sizeof(int), sizeof(float)) = 4 × bdim × 4 + * shared mem = 5 × bdim × max(sizeof(int), sizeof(float)) = 5 × bdim × 4 + * (4 stripes pre-Task-2.2 + 1 stripe for the alpha sum) * ══════════════════════════════════════════════════════════════════════════ */ #include @@ -132,6 +134,13 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel( /* Existing aux_dir_acc_reduce_kernel output [0] — the directional * accuracy of the aux head this step. */ const float* __restrict__ aux_dir_acc_dev, /* [1] */ + /* SP20 Phase 2 Task 2.2 (2026-05-09) — per-env trade-close alpha + * input. `[n_envs]` f32 device buffer; the env's slot is read only + * when `trade_close_per_env[env * env_stride] != 0` so non-close- + * step writes don't reach the EMA. NULL-tolerant — when NULL the + * kernel emits `out_inputs->alpha = 0.0f` (matches Phase 1.4 + * behavior, used by tests that haven't wired the alpha producer). */ + const float* __restrict__ alpha_per_env, /* [n_envs] or NULL */ int n_envs, /* Action decoding parameters (production packed factored action * encoding). `dir_divisor = NUM_MAGNITUDES * NUM_ORD * NUM_URG`, @@ -149,22 +158,26 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel( const int tid = threadIdx.x; const int bdim = blockDim.x; - /* Dynamic shared memory: 4 stripes of bdim × 4 bytes each. We use - * `int` storage for the 3 count reductions and `float` storage for - * the hold-at-exit sum (alias-compatible because both are 4 bytes - * wide; we cast between views as needed). */ + /* Dynamic shared memory: 5 stripes of bdim × 4 bytes each (4 stripes + * pre-Task-2.2 + 1 stripe for the SP20 Phase 2 Task 2.2 alpha sum). + * We use `int` storage for the 3 count reductions and `float` storage + * for the hold-at-exit and alpha sums (alias-compatible because both + * are 4 bytes wide; we cast between views as needed). */ extern __shared__ int sh_int[]; - /* Layout: [closed_count | wins_count | hold_count | hold_at_exit_sum_f32] */ - int* sh_closed_count = &sh_int[0]; - int* sh_wins_count = &sh_int[bdim]; - int* sh_hold_count = &sh_int[2 * bdim]; + /* Layout: [closed_count | wins_count | hold_count | + * hold_at_exit_sum_f32 | alpha_sum_f32] */ + int* sh_closed_count = &sh_int[0]; + int* sh_wins_count = &sh_int[bdim]; + int* sh_hold_count = &sh_int[2 * bdim]; float* sh_hold_at_exit_sum = (float*)&sh_int[3 * bdim]; + float* sh_alpha_sum = (float*)&sh_int[4 * bdim]; /* Per-thread strided accumulation over [0, n_envs). */ - int local_closed_count = 0; - int local_wins_count = 0; - int local_hold_count = 0; + int local_closed_count = 0; + int local_wins_count = 0; + int local_hold_count = 0; float local_hold_at_exit_sum = 0.0f; + float local_alpha_sum = 0.0f; for (int env = tid; env < n_envs; env += bdim) { const int tc = trade_close_per_env[env * env_stride]; @@ -179,6 +192,17 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel( local_closed_count += 1; if (sr > 0.0f) local_wins_count += 1; local_hold_at_exit_sum += he; + /* SP20 Phase 2 Task 2.2: gated alpha read on closed-env + * slots only. Non-close-step alpha_per_env values are + * stale (or sentinel 0 pre-first-close) and MUST NOT + * reach the ALPHA_EMA. NULL-tolerant — when alpha_per_env + * is NULL, the read is skipped and `local_alpha_sum` + * stays at 0 → `out_inputs->alpha = 0.0f`, matching the + * Phase 1.4 placeholder behavior used by tests without + * the alpha producer wired. */ + if (alpha_per_env != NULL) { + local_alpha_sum += alpha_per_env[env]; + } } if (dir == hold_action_id) local_hold_count += 1; } @@ -187,10 +211,11 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel( sh_wins_count[tid] = local_wins_count; sh_hold_count[tid] = local_hold_count; sh_hold_at_exit_sum[tid] = local_hold_at_exit_sum; + sh_alpha_sum[tid] = local_alpha_sum; __syncthreads(); /* Tree reduction across `bdim`. Standard log2(bdim) pattern. We - * reduce all four stripes in one loop body — single barrier per + * reduce all five stripes in one loop body — single barrier per * pass keeps the cost equal to a single reduction. */ for (int s = bdim / 2; s > 0; s >>= 1) { if (tid < s) { @@ -198,6 +223,7 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel( sh_wins_count[tid] += sh_wins_count[tid + s]; sh_hold_count[tid] += sh_hold_count[tid + s]; sh_hold_at_exit_sum[tid] += sh_hold_at_exit_sum[tid + s]; + sh_alpha_sum[tid] += sh_alpha_sum[tid + s]; } __syncthreads(); } @@ -208,6 +234,7 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel( const int wins_count = sh_wins_count[0]; const int hold_count = sh_hold_count[0]; const float hold_at_exit_sum_f = sh_hold_at_exit_sum[0]; + const float alpha_sum_f = sh_alpha_sum[0]; const int is_close_out = (closed_count > 0) ? 1 : 0; @@ -226,6 +253,19 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel( trade_duration_out = (int)__float2int_rn(mean_he); } + /* SP20 Phase 2 Task 2.2: alpha = mean(alpha_per_env) over closed + * envs. Replaces the Phase 1.4 hardcoded `0.0f` placeholder + * (audit doc Task 2.2 forward-reference closure). When + * `is_close_out == 0` (no env closed this step) the EMA gate + * is closed at the consumer, so this slot's value is moot — but + * we still emit 0.0 for cleanliness so downstream debug reads + * don't see stale sums from the reduction. NULL-tolerant via + * the per-thread accumulation gate above. */ + float alpha_out = 0.0f; + if (is_close_out == 1) { + alpha_out = alpha_sum_f / (float)closed_count; + } + /* action_is_hold: majority-vote over envs. > n/2 threshold * matches the spec (strict majority, ties ⇒ 0). */ const int action_is_hold_out = (hold_count * 2 > n_envs) ? 1 : 0; @@ -239,17 +279,17 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel( const float std_in = aux_logits_std_dev[0]; const float dir_acc_in = aux_dir_acc_dev[0]; - /* Write the 9 fields. The float fields `alpha` and - * `per_bar_hold_reward` are intentionally zero — Phase 2 / - * Phase 3.2 wire the real producers; until then the EMAs that - * blend those fields stay at sentinel 0.0 (per - * `pearl_first_observation_bootstrap`). */ + /* Write the 9 fields. `alpha` is now the load-bearing SP20 + * Phase 2 Task 2.2 producer output (mean alpha over closed + * envs); `per_bar_hold_reward` remains a Phase 3.2 forward + * reference until Component 2 lands the per-bar Hold opp-cost + * dual emission. */ out_inputs->is_close = is_close_out; out_inputs->is_win = is_win_out; out_inputs->trade_duration = trade_duration_out; out_inputs->action_is_hold = action_is_hold_out; - out_inputs->alpha = 0.0f; /* Phase 2 forward ref */ - out_inputs->per_bar_hold_reward = 0.0f; /* Phase 3.2 fwd ref */ + out_inputs->alpha = alpha_out; /* Task 2.2 — was 0.0f */ + out_inputs->per_bar_hold_reward = 0.0f; /* Phase 3.2 fwd ref */ out_inputs->aux_logits_p50 = p50_in; out_inputs->aux_logits_std = std_in; out_inputs->aux_dir_acc = dir_acc_in; diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index ebcc18868..847dc4530 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -2122,6 +2122,16 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "GpuExperienceCollector.label_at_open_per_env [alloc_episodes] i32 device-resident scratch — SP20 Phase 2 Task 2.0 (2026-05-09) per-env aux next-bar label sign captured at trade open. Layout: tile[env] in {-1, 0, +1} where +1 = aux predicted up at trade open, -1 = aux predicted not-up, 0 = sentinel (no position currently open or trade-open landed in aux skip-window). Producer: `experience_env_step` `entering_trade` branch — reads `exp_aux_nb_label_buf[env]` (kernel emits `0`/`1`/`-1` per `aux_sign_label_per_step_kernel.cu:77`) and maps `0`→`-1`, `1`→`+1`, `-1`→`0` to match the SP20 spec sign convention (spec §4.1: `label_at_open = sign(SP19_blended_label[trade_open_bar])`). Consumer: `experience_env_step` `is_close` branch — passes `label_at_open_per_env[env]` to `sp20_compute_event_reward(close_pnl, label_at_open_sign, isv[LOSS_CAP_INDEX])`. FoldReset sentinel 0 across all `alloc_episodes` slots — leftover label from the previous fold's open trade would corrupt the new fold's first trade-close `dir_match` check (the new fold's aux head has independent calibration and would not have produced that label). Reset path: `cudarc::driver::CudaStream::memset_zeros` on the `CudaSlice` (device-resident, no host buffer to fill). Per `pearl_first_observation_bootstrap` the sentinel-0 input also flows through `sp20_compute_event_reward` to the wrong-reason quadrant on the corner case where a trade closes WITHOUT a prior `entering_trade` having fired this fold (trade carry-over from before the reset is impossible because portfolio_states is also FoldReset).", }, + // ── SP20 Phase 2 Task 2.2 (2026-05-09): per-env trade-close + // alpha scratch — written at `is_close`, read by + // `sp20_aggregate_inputs_kernel` to populate the + // `EmaInputs.alpha` field (replacing the Phase 1.4 0.0f + // forward-reference placeholder). + RegistryEntry { + name: "alpha_per_env", + category: ResetCategory::FoldReset, + description: "GpuExperienceCollector.alpha_per_env [alloc_episodes] f32 device-resident scratch — SP20 Phase 2 Task 2.2 (2026-05-09) per-env trade-close alpha (= R_event - hold_baseline) per spec §4.1. Phase 2 emits `hold_baseline = 0.0f` as a Phase 3.2 forward-reference placeholder; Phase 3.2's Hold opportunity-cost dual emission lands the real `hold_baseline_buffer.sum_over_trade_range()` consumer. Producer: `experience_env_step` `is_close` (segment_complete) branch — writes `alpha = sp20_compute_event_reward(segment_return, label_at_open_per_env[env], ISV[LOSS_CAP_INDEX])`. Consumer: `sp20_aggregate_inputs_kernel` (Path C aggregation) — sums `alpha_per_env[env]` over closed envs and divides by closed-count to populate `EmaInputs.alpha`. FoldReset sentinel 0.0 across all `alloc_episodes` slots — leftover alpha from the previous fold's last trade-close would corrupt the new fold's first ALPHA_EMA observation. The sentinel-0 reads as `alpha == 0` for non-close steps; the aggregation kernel reads only when `trade_close_per_env[env] != 0` so the FoldReset sentinel never reaches the EMA. Reset path: `cudarc::driver::CudaStream::memset_zeros` on the `CudaSlice` (device-resident, no host buffer to fill). Replaces the Phase 1.4 hardcoded `out_inputs->alpha = 0.0f` placeholder atomically per `feedback_no_partial_refactor` (Task 2.2 lands buffer + kernel-arg + aggregation-kernel-update + 2 placeholder pin tests in one commit).", + }, ]; Self { entries } } @@ -2413,4 +2423,42 @@ mod sp20_registry_tests { entry.description ); } + + /// SP20 Phase 2 Task 2.2: `alpha_per_env` MUST be registered as + /// FoldReset so the new fold's first trade-close fully re-populates + /// from the new fold's reward calculation (preventing stale alpha + /// leakage into the new fold's ALPHA_EMA bootstrap). Description + /// MUST mention SP20 Phase 2 + Task 2.2 + `is_close` (producer) + + /// `sp20_aggregate_inputs_kernel` (consumer) so wire-up audits + /// have a stable search anchor. + #[test] + fn sp20_alpha_per_env_registered_fold_reset() { + let registry = StateResetRegistry::new(); + let entry = registry + .entries + .iter() + .find(|e| e.name == "alpha_per_env") + .expect("alpha_per_env must be registered in StateResetRegistry"); + assert_eq!( + entry.category, + ResetCategory::FoldReset, + "alpha_per_env must be FoldReset (got {:?})", + entry.category + ); + assert!( + entry.description.contains("SP20 Phase 2"), + "description must reference 'SP20 Phase 2'; got: {}", + entry.description + ); + assert!( + entry.description.contains("is_close"), + "description must reference 'is_close' producer site; got: {}", + entry.description + ); + assert!( + entry.description.contains("sp20_aggregate_inputs_kernel"), + "description must reference the aggregation-kernel consumer; got: {}", + entry.description + ); + } } diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 298b0d7bb..4bd5f7d9b 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -9554,6 +9554,18 @@ impl DQNTrainer { )))?; } } + // SP20 Phase 2 Task 2.2 (2026-05-09): per-env trade-close + // alpha scratch. Same device-resident `CudaSlice` + // reset pattern as `label_at_open_per_env` above. + "alpha_per_env" => { + if let Some(ref mut collector) = self.gpu_experience_collector { + let stream = collector.stream().clone(); + stream.memset_zeros(&mut collector.alpha_per_env) + .map_err(|e| crate::MLError::ModelError(format!( + "alpha_per_env memset_zeros: {e}" + )))?; + } + } // SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots. // OFI_IMPACT_LAMBDA_INDEX=407 is an Invariant-1 anchor (NOT // a stateful EMA) — rewrite the constructor's value at fold diff --git a/crates/ml/tests/sp20_aggregate_inputs_test.rs b/crates/ml/tests/sp20_aggregate_inputs_test.rs index d384cbb14..70b111bcd 100644 --- a/crates/ml/tests/sp20_aggregate_inputs_test.rs +++ b/crates/ml/tests/sp20_aggregate_inputs_test.rs @@ -144,6 +144,75 @@ mod gpu { p50_buf.dev_ptr, std_buf.dev_ptr, dir_buf.dev_ptr, + /* alpha_per_env_dev = */ 0, // NULL — test scaffold without SP20 reward producer + /* n_envs = */ n as i32, + DIR_DIVISOR, + HOLD_ACTION_ID, + out_buf.dev_ptr, + ).expect("launch aggregate kernel"); + } + stream.synchronize().expect("sync after aggregate"); + read_ema_inputs(&out_buf) + } + + /// Variant of `run_kernel` that wires the SP20 Phase 2 Task 2.2 + /// per-env alpha producer (`alpha_per_env: &[f32]`). Used by + /// `alpha_mean_over_closed_envs_phase_2_contract` to validate the + /// new alpha aggregation rule (mean over closed envs). + fn run_kernel_with_alpha( + trade_close: &[i32], + step_ret: &[f32], + hold_at_exit: &[f32], + actions: &[i32], + alpha_per_env: &[f32], + aux_logits_p50: f32, + aux_logits_std: f32, + aux_dir_acc: f32, + ) -> (i32, i32, i32, i32, f32, f32, f32, f32, f32) { + let stream = make_test_stream(); + let kernel = load_aggregation_kernel(&stream); + let n = trade_close.len(); + assert!(n > 0, "test n_envs must be > 0"); + assert_eq!(step_ret.len(), n); + assert_eq!(hold_at_exit.len(), n); + assert_eq!(actions.len(), n); + assert_eq!(alpha_per_env.len(), n); + + let tc_buf = unsafe { MappedI32Buffer::new(n) }.expect("alloc tc"); + tc_buf.write_from_slice(trade_close); + let sr_buf = unsafe { MappedF32Buffer::new(n) }.expect("alloc sr"); + sr_buf.write_from_slice(step_ret); + let he_buf = unsafe { MappedF32Buffer::new(n) }.expect("alloc he"); + he_buf.write_from_slice(hold_at_exit); + let act_buf = unsafe { MappedI32Buffer::new(n) }.expect("alloc act"); + act_buf.write_from_slice(actions); + let alpha_buf = unsafe { MappedF32Buffer::new(n) }.expect("alloc alpha"); + alpha_buf.write_from_slice(alpha_per_env); + + let p50_buf = unsafe { MappedF32Buffer::new(1) }.expect("alloc p50"); + p50_buf.write_from_slice(&[aux_logits_p50]); + let std_buf = unsafe { MappedF32Buffer::new(1) }.expect("alloc std"); + std_buf.write_from_slice(&[aux_logits_std]); + let dir_buf = unsafe { MappedF32Buffer::new(1) }.expect("alloc dir"); + dir_buf.write_from_slice(&[aux_dir_acc]); + + let out_buf = unsafe { MappedF32Buffer::new(SP20_EMA_INPUTS_F32_LEN) } + .expect("alloc out"); + out_buf.write_from_slice(&[f32::NAN; SP20_EMA_INPUTS_F32_LEN]); + + unsafe { + launch_sp20_aggregate_inputs( + &stream, + &kernel, + tc_buf.dev_ptr, + sr_buf.dev_ptr, + he_buf.dev_ptr, + act_buf.dev_ptr, + /* env_stride = */ 1, + p50_buf.dev_ptr, + std_buf.dev_ptr, + dir_buf.dev_ptr, + alpha_buf.dev_ptr, // SP20 Phase 2 Task 2.2 alpha producer /* n_envs = */ n as i32, DIR_DIVISOR, HOLD_ACTION_ID, @@ -260,14 +329,18 @@ mod gpu { assert_eq!(td, 3, "trade_duration = round(3.0) = 3"); } - /// Phase 2 / Phase 3.2 forward-reference placeholders are emitted - /// as 0.0 regardless of any input combinations. This test - /// **verifies the documented contract** — the kernel does NOT - /// surface any per-env signal into `alpha` / `per_bar_hold_reward` - /// in Phase 1.4; that wiring lands in Phase 2 / 3.2. + /// SP20 Phase 2 Task 2.2 (2026-05-09) — alpha is now load-bearing + /// (mean over closed envs of `alpha_per_env`). `per_bar_hold_reward` + /// remains a Phase 3.2 forward-reference placeholder. + /// + /// **NULL pre-Task-2.2 contract preserved**: when `alpha_per_env_dev + /// = 0` (NULL), the kernel emits `alpha = 0.0` matching the Phase + /// 1.4 placeholder behavior. Tests that don't exercise the SP20 + /// reward producer (run_kernel) still see the old `alpha = 0.0` + /// invariant. This test validates the NULL-tolerance contract. #[test] #[ignore = "requires GPU"] - fn alpha_and_per_bar_hold_reward_are_phase_2_and_3_2_placeholders() { + fn null_alpha_producer_keeps_phase_1_4_placeholder_contract() { let trade_close = vec![1, 1, 1]; let step_ret = vec![0.5, 0.4, 0.3]; let hold_at_exit = vec![10.0, 10.0, 10.0]; @@ -277,8 +350,8 @@ mod gpu { 0.5, 0.5, 0.99); assert_eq!( alpha, 0.0, - "alpha is Phase 2 placeholder; aggregation kernel must NOT \ - surface a per-env signal here regardless of inputs", + "alpha is 0.0 when alpha_per_env producer is NULL (test scaffold); \ + matches Phase 1.4 placeholder behavior for backward compat", ); assert_eq!( hr, 0.0, @@ -286,4 +359,59 @@ mod gpu { kernel must NOT surface a per-env signal here", ); } + + /// SP20 Phase 2 Task 2.2 (2026-05-09) — alpha aggregation contract. + /// When `alpha_per_env` is wired (non-NULL), `EmaInputs.alpha` is + /// the **mean** of `alpha_per_env[env]` over envs with + /// `trade_close_per_env[env] != 0`. Non-close envs do NOT contribute + /// to the mean (gated read at the producer site). + /// + /// Test setup: 4 envs, 3 close (1 with alpha=+1.0, 2 with + /// alpha=-0.5), 1 doesn't close (alpha=+99 — should be ignored). + /// Expected: alpha_out = mean(1.0, -0.5, -0.5) = 0.0. + #[test] + #[ignore = "requires GPU"] + fn alpha_mean_over_closed_envs_phase_2_contract() { + let trade_close = vec![1, 1, 1, 0]; // env 3 doesn't close + let step_ret = vec![0.5, -0.3, -0.2, 0.0]; + let hold_at_exit = vec![5.0, 6.0, 7.0, 0.0]; + let actions = vec![encode_dir(0); 4]; + let alpha_per_env = vec![1.0, -0.5, -0.5, 99.0]; // env 3 ignored + let (ic, _, _, _, alpha, hr, _, _, _) = + run_kernel_with_alpha(&trade_close, &step_ret, &hold_at_exit, + &actions, &alpha_per_env, 0.0, 0.0, 0.0); + assert_eq!(ic, 1, "is_close = 1 (3 envs closed)"); + // mean(1.0 + (-0.5) + (-0.5)) / 3 = 0.0 — env 3's alpha=99 must + // not leak in (gated by trade_close_per_env[3] == 0 at producer). + assert!( + alpha.abs() < 1e-6, + "expected alpha = mean(1.0, -0.5, -0.5) over closed envs = 0.0; got {alpha}", + ); + assert_eq!(hr, 0.0, "per_bar_hold_reward stays Phase 3.2 placeholder"); + } + + /// SP20 Phase 2 Task 2.2 — alpha is forced to `0.0` when + /// `is_close == 0` (no env closed this step), regardless of any + /// stale `alpha_per_env` content. Validates the close-gated emit + /// at the kernel's `tid == 0` write block. + #[test] + #[ignore = "requires GPU"] + fn alpha_zero_when_no_close_phase_2_contract() { + let trade_close = vec![0_i32; 4]; // no env closes + let step_ret = vec![0.0_f32; 4]; + let hold_at_exit = vec![0.0_f32; 4]; + let actions = vec![encode_dir(0); 4]; + // Stale alpha values from prior trade closes — kernel MUST gate + // these out via the close-count check. + let alpha_per_env = vec![0.5, -0.7, 2.3, -1.4]; + let (ic, _, _, _, alpha, _, _, _, _) = + run_kernel_with_alpha(&trade_close, &step_ret, &hold_at_exit, + &actions, &alpha_per_env, 0.0, 0.0, 0.0); + assert_eq!(ic, 0, "is_close = 0 (no envs closed)"); + assert_eq!( + alpha, 0.0, + "alpha must be 0.0 when no env closes — stale alpha_per_env \ + values must NOT leak into the EMA on non-close steps", + ); + } } diff --git a/crates/ml/tests/sp20_phase1_4_wireup_test.rs b/crates/ml/tests/sp20_phase1_4_wireup_test.rs index d8574116c..699ddf2f4 100644 --- a/crates/ml/tests/sp20_phase1_4_wireup_test.rs +++ b/crates/ml/tests/sp20_phase1_4_wireup_test.rs @@ -163,6 +163,13 @@ mod gpu { stats_p50, stats_std, dir_acc_buf.dev_ptr, + /* alpha_per_env_dev = */ 0, // SP20 Phase 2 Task 2.2 — Phase 1.4 + // wireup test stays NULL; the alpha + // producer is exercised at the + // segment_complete site in the full + // collector path. Tests in this file + // assert the placeholder contract + // (alpha_ema stays at sentinel 0.0). n_envs as i32, DIR_DIVISOR, HOLD_ACTION_ID, @@ -226,12 +233,18 @@ mod gpu { // ── ISV invariants ── - // ALPHA_EMA: aggregation kernel writes alpha=0.0 (Phase 2 - // forward ref); first-observation REPLACE on a closed trade ⇒ - // ALPHA_EMA stays at 0.0. **This validates the documented - // forward reference contract.** + // ALPHA_EMA: aggregation kernel sees `alpha_per_env_dev = 0` + // (NULL in this test scaffold — the full SP20 reward producer + // is exercised at the `experience_env_step::segment_complete` + // site, not in this Phase 1.4 wireup test). The kernel falls + // back to `alpha = 0.0`, first-observation REPLACE on a closed + // trade ⇒ ALPHA_EMA stays at 0.0. **This validates the SP20 + // Phase 2 Task 2.2 NULL-tolerance contract** — the + // aggregation kernel still emits 0.0 when the alpha producer + // is unwired, preserving the Phase 1.4 placeholder behavior + // for tests that don't exercise the reward producer. assert_eq!(isv[ALPHA_EMA_INDEX], 0.0, - "ALPHA_EMA stays at 0.0 (Phase 2 forward ref) on first close"); + "ALPHA_EMA stays at 0.0 with NULL alpha_per_env producer (Task 2.2 NULL contract)"); // WR_EMA: 1 win out of 1 close ⇒ is_win = 1, first-obs replace // ⇒ WR_EMA = 1.0. @@ -295,14 +308,26 @@ mod gpu { "HOLD_COST_SCALE in [0.01, 0.5]; got {hcs}"); } - /// Locks the Phase 2 / Phase 3.2 forward-reference contract: the - /// aggregation kernel writes 0.0 to `alpha` and `per_bar_hold_reward` - /// regardless of inputs, so the corresponding EMAs (ALPHA_EMA, - /// HOLD_REWARD_EMA) stay at 0.0 sentinel under Phase 1.4 even after - /// many fired observations. + /// SP20 Phase 2 Task 2.2 contract — when the aggregation kernel is + /// invoked with `alpha_per_env_dev = 0` (NULL, this test scaffold's + /// configuration) AND the per-bar-hold-reward producer is still in + /// Phase 3.2 forward-reference, both EMAs stay at sentinel 0.0 + /// regardless of any input combinations. Validates the + /// NULL-tolerance contract of the aggregation kernel post-Task-2.2: + /// downstream tests that don't wire the SP20 reward producer (e.g. + /// this Phase 1.4 wireup smoke) preserve the Phase 1.4 ALPHA_EMA + /// placeholder behavior. + /// + /// HOLD_REWARD_EMA stays at 0.0 because Phase 3.2 lands the per-bar + /// `r_per_bar = -aux_conf * cost_scale` producer; the aggregation + /// kernel writes 0.0 for `per_bar_hold_reward` until then. + /// + /// **Pinned forward refs locked here**: + /// - ALPHA_EMA bound to NULL-alpha-producer scaffold contract. + /// - HOLD_REWARD_EMA bound to Phase 3.2 placeholder. #[test] #[ignore = "requires GPU"] - fn alpha_and_hold_reward_emas_stay_at_zero_phase_2_3_2_placeholders() { + fn alpha_and_hold_reward_emas_stay_at_zero_with_null_producers() { let n_envs = 4; // 4 closing wins per call, 4 Hold actions per call let trade_close = vec![1; n_envs]; @@ -311,20 +336,23 @@ mod gpu { let actions = vec![encode_dir(HOLD_ACTION_ID); n_envs]; let aux_logits = vec![0.5_f32; n_envs * AUX_K_CLASSES]; - // Single chain run is enough — first-obs REPLACE on placeholder - // 0.0 lands the EMA at 0.0; subsequent observations would only - // re-blend toward 0.0 from 0.0. + // Single chain run is enough — `run_chain_once` passes + // alpha_per_env_dev=0 (NULL), so the aggregation kernel falls + // back to `alpha = 0.0`; first-obs REPLACE on 0.0 lands the + // EMA at 0.0. let isv = run_chain_once( n_envs, &aux_logits, &trade_close, &step_ret, &hold_at_exit, &actions, 0.99, ); assert_eq!(isv[ALPHA_EMA_INDEX], 0.0, - "ALPHA_EMA must be 0.0 — Phase 2 forward ref placeholder. \ - Wire-up must NOT surface a per-env signal in Phase 1.4"); + "ALPHA_EMA stays at 0.0 with NULL alpha_per_env producer \ + (Task 2.2 NULL-tolerance contract). The full SP20 reward \ + producer at experience_env_step::segment_complete is NOT \ + exercised in this Phase 1.4 wireup test."); assert_eq!(isv[HOLD_REWARD_EMA_INDEX], 0.0, "HOLD_REWARD_EMA must be 0.0 — Phase 3.2 forward ref \ - placeholder. Wire-up must NOT surface a per-env signal \ - in Phase 1.4"); + placeholder. Aggregation kernel writes 0.0 until \ + Component 2 lands the per-bar producer."); } } diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 9819c87d1..67c92e12e 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,181 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +## 2026-05-09 — SP20 Phase 2 Task 2.2: SP12 v3 reward block replacement + per-env alpha plumbing (atomic) + +Lands the SP20 4-quadrant reward at the `experience_env_step:: +segment_complete` site, atomically with the per-env alpha plumbing +(`alpha_per_env` → `sp20_aggregate_inputs_kernel` → `EmaInputs.alpha` → +`sp20_emas_compute_kernel` → `ISV[ALPHA_EMA_INDEX]`) per +`feedback_no_partial_refactor`. Replaces the SP12 v3 inlined block +(experience_kernels.cu pre-Task-2.2 lines 3206-3466) with the SP20 +4-quadrant reward via `sp20_compute_event_reward` (Task 2.1 helper). + +### Replaced block scope + +The trade-close-site SP12 v3 reward inline block, comprising: + - vol-normalization (lines ~3208-3216): `atr_norm`, `log_atr`, + `atr_pct`, `vol_proxy`, `vol_norm`, `vol_normalized_return` + - base reward (line ~3232): `base_reward = 2 × vol_normalized_return` + - SP18 D-leg trade-close call (lines ~3255-3279): adds + `compute_sp18_hold_opportunity_cost(...)` to `base_reward` (DELETED + at trade-close site only — per-bar SP18 sites at + experience_kernels.cu:3722 / :3833 RETAINED pending Phase 3.2 + Component 2 replacement per `feedback_no_stubs`) + - SP12 v3 asymmetric cap (lines ~3282-3341): `pos_cap_eff` / + `neg_cap_eff` ISV reads, `compute_asymmetric_capped_pnl` call + - min-hold penalty (lines ~3345-3445): `compute_min_hold_penalty` + + min-hold-target ISV-driven fallback chain + - r_popart / r_trail / capped_pnl assignment cascade + - `reward = capped_pnl;` (line ~3466) + +Replaced with the SP20 4-quadrant reward block (~70 LoC) computing: + +``` +label_at_open = label_at_open_per_env[i] (Task 2.0 buffer) +loss_cap = ISV[loss_cap_idx] (sp20_controllers output) +R_event = sp20_compute_event_reward( (Task 2.1 helper) + segment_return, + label_at_open, + loss_cap) +hold_baseline = 0.0f (Phase 3.2 placeholder) +alpha = R_event - hold_baseline +alpha_ema = ISV[alpha_ema_idx] (sp20_emas output) +r_used = alpha - alpha_ema (advantage centering) +alpha_per_env[i] = alpha (sp20_aggregate input) +r_popart / r_trail = r_used (SP11 component split) +reward = r_used +``` + +### Per-env alpha plumbing (replaces Phase 1.4 forward-ref placeholder) + +`sp20_aggregate_inputs_kernel.cu`: + - New kernel arg: `const float* __restrict__ alpha_per_env`. + - New shmem stripe: `sh_alpha_sum` (5th f32 stripe; 5 stripes × + bdim × 4 bytes total, was 4). + - New per-thread accumulation gated on `trade_close_per_env[env] + != 0` — non-close-step writes do NOT contribute to the mean. + - New tree-reduce loop body: `sh_alpha_sum[tid] += + sh_alpha_sum[tid + s]` (single-pass with the existing 4 + stripes). + - Output write: `out_inputs->alpha = (closed_count > 0) ? alpha_sum + / (float)closed_count : 0.0f` — replaces the hardcoded `= 0.0f` + placeholder. + +NULL-tolerant: when `alpha_per_env_dev = 0` (test scaffolds that +don't wire the SP20 reward producer, e.g. the `sp20_phase1_4_wireup_ +test` integration test), the per-thread accumulation is skipped and +the kernel emits `alpha = 0.0f` matching the Phase 1.4 placeholder +behavior. This preserves backward compat for the 8-test Phase 1.4 +GPU integration suite. + +### Buffer + reset infrastructure + +New `alpha_per_env: CudaSlice` field on `GpuExperienceCollector` +(allocated `[alloc_episodes]` next to `label_at_open_per_env`). +Registered in `StateResetRegistry` with FoldReset category; reset path +mirrors `label_at_open_per_env`'s `stream.memset_zeros` async GPU +memset (no host buffer to fill). + +### Pin-test updates + +The Phase 1.4 placeholder pin tests are RENAMED + UPDATED to assert +the new contract: + +| Old test | New test | +|----------------------------------------------------------------------------|-------------------------------------------------------| +| `sp20_aggregate_inputs_test::alpha_and_per_bar_hold_reward_are_phase_2_and_3_2_placeholders` | `null_alpha_producer_keeps_phase_1_4_placeholder_contract` | +| `sp20_phase1_4_wireup_test::alpha_and_hold_reward_emas_stay_at_zero_phase_2_3_2_placeholders` | `alpha_and_hold_reward_emas_stay_at_zero_with_null_producers` | + +The renamed tests assert: with `alpha_per_env_dev = 0` (NULL), the +kernel falls back to `alpha = 0.0f`. These pin tests document and +lock the NULL-tolerance semantic. + +Two NEW tests in `sp20_aggregate_inputs_test.rs` validate the +load-bearing contract: + + 1. `alpha_mean_over_closed_envs_phase_2_contract` — 4 envs (3 close, + 1 doesn't); `alpha_per_env = [+1.0, -0.5, -0.5, +99.0]` with + env 3 not closing. Asserts `EmaInputs.alpha = mean(+1.0, -0.5, + -0.5) = 0.0` (env 3's alpha=99 must NOT leak into the mean). + 2. `alpha_zero_when_no_close_phase_2_contract` — 4 envs, none + close. Asserts `EmaInputs.alpha = 0.0` regardless of stale + `alpha_per_env` content. Validates the close-gated emit at the + kernel's `tid == 0` write block. + +### Files modified + +| File | Status | Purpose | +|------|--------|---------| +| `crates/ml/src/cuda_pipeline/experience_kernels.cu` | ~70 LoC delete + ~80 LoC insert | SP12 v3 block → SP20 4-quadrant reward; +3 kernel args (`alpha_per_env`, `loss_cap_idx`, `alpha_ema_idx`) | +| `crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu` | +1 kernel arg, +1 shmem stripe | per-env alpha consumer (replaces Phase 1.4 placeholder) | +| `crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs` | +1 launcher arg | thread `alpha_per_env_dev` through | +| `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` | +field, +alloc, +launch args (×2) | per-env alpha buffer + production wire-up | +| `crates/ml/src/trainers/dqn/state_reset_registry.rs` | +entry + invariant test | FoldReset for `alpha_per_env` | +| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | +match arm | `reset_named_state` dispatch | +| `crates/ml/tests/sp20_aggregate_inputs_test.rs` | +2 tests, +1 helper, 1 rename | new contract + null contract pins | +| `crates/ml/tests/sp20_phase1_4_wireup_test.rs` | 1 rename, 2 docstring updates | null-producer contract pin | +| `docs/dqn-wire-up-audit.md` | This entry | Audit log | + +### Retained (not deleted) per `feedback_no_stubs.md` + +The following continue to compile / be called by other sites: + - `compute_sp18_hold_opportunity_cost` (`trade_physics.cuh:655`) — + called by per-bar SP18 D-leg sites at `experience_kernels.cu:3722` + and `:3833`. Phase 3.2 Component 2 replaces these per-bar sites + with the SP20 Hold-cost dual emission. + - `compute_asymmetric_capped_pnl` (`trade_physics.cuh:545`) — called + by `compute_sp18_hold_opportunity_cost` (line 680 of trade_physics.cuh) + and by `sp12_reward_math_test_kernel.cu`. + - `compute_min_hold_penalty` (`trade_physics.cuh:567`) — called by + `sp12_reward_math_test_kernel.cu` only post-Task-2.2; retained + against `feedback_no_quickfixes` (the test wrapper preserves the + SP12 v3 oracle test surface even after the production caller is + removed). + +### Deleted: the `min_hold_*` kernel-arg trio + +Per `feedback_no_hiding` "wire up or delete; never suppress with `_` +or `#[allow]`", the now-unused `min_hold_target`, `min_hold_penalty_ +max`, `min_hold_temperature` kernel args of `experience_env_step` are +DELETED in this commit (kernel signature + launch-site `.arg(...)` +calls). The corresponding upstream producer chain +(`min_hold_temperature_update_kernel`, `ISV[MIN_HOLD_TEMPERATURE_ +ADAPTIVE_INDEX=460]`, the `read_min_hold_temperature_from_isv` +helper, and `config.min_hold_temperature` Rust field) is INTENTIONALLY +DEFERRED to a follow-up cleanup commit ("Task 2.4: remove orphaned +SP12 v3 min_hold_* producer chain"). Atomic scope rationale: the +kernel-arg removal is mechanical and tied directly to the kernel-body +change at the same site; the producer-chain removal touches the +SP14 ISV slot registry + StateResetRegistry + an ISV layout +fingerprint bump and is independently testable. Per +`feedback_no_partial_refactor`'s "consumer migrates with contract +change" rule, the consumer IS what migrates here — the producer can +keep emitting into ISV[460] in the meantime as a documented orphan +without breaking any contract. Task 2.4 cleanup commit is tracked +explicitly in this audit-doc entry as a known follow-up. + +### Verification + +``` +SQLX_OFFLINE=true cargo build -p ml # cubin compile +SQLX_OFFLINE=true cargo check -p ml --tests --features cuda # tests compile +SQLX_OFFLINE=true cargo test -p ml --lib sp20 # 20/20 lib tests +``` + +### Phase 2 → Phase 3.2 forward references + +Phase 3.2 Component 2 lands: + - `hold_cost_scale_compute_kernel` (separate fine-grained controller + sibling of `sp20_controllers_compute`). + - Per-bar Hold opp-cost dual emission at the per-bar SP18 D-leg + sites (replacing `compute_sp18_hold_opportunity_cost` calls at + `experience_kernels.cu:3722` and `:3833`). + - `hold_baseline_buffer` (circular [N, 30] f32) populated by the + per-bar emission, consumed by the Task 2.2 trade-close site to + replace the `hold_baseline = 0.0f` placeholder with + `sum_hold_baseline_buffer(buffer, trade_open_bar, trade_close_bar)`. + ## 2026-05-09 — SP20 Phase 2 Task 2.1: sp20_compute_event_reward device function + GPU oracle tests (additive) Lands the SP20 Component 1 reward math as a header-only `__device__