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

feat(sp20): Phase 2 Task 2.2 — SP12 v3 reward block → 4-quadrant + alpha plumbing (atomic)

Browse Source 
Lands the SP20 4-quadrant reward at the experience_env_step::
segment_complete site, atomically with per-env alpha plumbing through
the SP20 aggregation kernel per `feedback_no_partial_refactor`:

* `experience_env_step.cu` — replace SP12 v3 inlined block (vol-norm
  calc + base_reward + SP18 D-leg trade-close call + asymmetric cap +
  min-hold penalty + r_popart/r_trail cascade, ~260 LoC) with the SP20
  4-quadrant reward (~80 LoC) computing:
    R_event = sp20_compute_event_reward(segment_return,
                                         label_at_open_per_env[i],
                                         ISV[LOSS_CAP_INDEX])
    alpha   = R_event - 0.0  (Phase 3.2 hold_baseline placeholder)
    r_used  = alpha - ISV[ALPHA_EMA_INDEX]    (advantage centering)
    alpha_per_env[i] = alpha   (consumed by sp20_aggregate)
    r_popart / r_trail = r_used (SP11 component split preserved)

  Adds 3 new kernel args at the end (preserves Task 2.0 args):
    float* alpha_per_env, int loss_cap_idx, int alpha_ema_idx
  Deletes 3 SP12 v3 args (min_hold_target, min_hold_penalty_max,
  min_hold_temperature) per `feedback_no_hiding` "wire up or delete".

* `sp20_aggregate_inputs_kernel.cu` — add `alpha_per_env` input arg
  (NULL-tolerant), 5th shmem stripe (sh_alpha_sum, was 4 stripes),
  per-thread close-gated accumulation, 5-way tree-reduce loop body,
  output write `alpha = mean / closed_count`. Replaces the Phase 1.4
  hardcoded `out_inputs->alpha = 0.0f` placeholder atomically.

* `sp20_aggregate_inputs.rs` launcher — +1 arg `alpha_per_env_dev`
  (passes 0/NULL for tests). Shmem byte-count: 4 stripes → 5 stripes.

* `gpu_experience_collector.rs` — new `alpha_per_env: CudaSlice<f32>`
  field [alloc_episodes], allocated next to `label_at_open_per_env`,
  threaded into env_step launch + sp20_aggregate launch. Removes the
  3 obsolete config.min_hold_* `.arg(...)` calls from env_step launch.

* `state_reset_registry.rs` — `alpha_per_env` FoldReset entry +
  invariant test `sp20_alpha_per_env_registered_fold_reset`.

* `training_loop.rs::reset_named_state` — `alpha_per_env` dispatch arm
  via `stream.memset_zeros` (mirrors `label_at_open_per_env` pattern).

* Pin-test updates per spec (atomic with the contract change):
  - `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` (asserts
    NULL alpha producer ⇒ alpha=0.0 backward-compat).
  - `sp20_phase1_4_wireup_test::alpha_and_hold_reward_emas_stay_at_
    zero_phase_2_3_2_placeholders` →
    `..._with_null_producers` (same NULL-tolerance pin).
  - +2 NEW tests in `sp20_aggregate_inputs_test`:
    * `alpha_mean_over_closed_envs_phase_2_contract` — 4 envs (3
      close, 1 doesn't), asserts close-gated mean (env 3's alpha=99
      ignored).
    * `alpha_zero_when_no_close_phase_2_contract` — no-close steps
      emit alpha=0.0 regardless of stale alpha_per_env content.

* `dqn-wire-up-audit.md` — Task 2.2 entry documents the deleted SP12 v3
  block, the per-env alpha plumbing, the retained device functions
  per `feedback_no_stubs`, and the deferred Task 2.4 cleanup of the
  orphaned min_hold_temperature_update_kernel producer chain.

Per `pearl_event_driven_reward_density_alignment`: per-bar SP18 D-leg
sites at experience_kernels.cu:3722 / :3833 are RETAINED pending
Phase 3.2 Component 2 replacement. Per `feedback_no_stubs`:
compute_sp18_hold_opportunity_cost / compute_asymmetric_capped_pnl /
compute_min_hold_penalty device functions are RETAINED (still called
by per-bar SP18 sites + sp12_reward_math_test_kernel test wrapper).

Per spec §4.1: alpha_ema centering (`R_used = alpha - alpha_ema`) is
load-bearing for cold-start learning at WR=46% with raw EV ≈ -0.06,
NOT just Q-target stability. Documented in the kernel comment block.

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

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-10 00:28:52 +02:00
parent eaaab152fc
commit 3f9b34030c
9 changed files with 762 additions and 357 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

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

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

@@ -678,6 +678,34 @@ pub struct GpuExperienceCollector {
/// folds — each fold has independent aux next-bar predictions).
pub(crate) label_at_open_per_env: CudaSlice<i32>, // [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<f32>, // [alloc_episodes]
// Output buffers [alloc_episodes * alloc_timesteps, ...]
states_out: CudaSlice<f32>, // #30 [alloc_episodes * alloc_timesteps * STATE_DIM] f32
actions_out: CudaSlice<i32>, // [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::<f32>(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<f32> = 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,

61
crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs
View File

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

100
crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu
View File

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

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

@@ -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<i32>` (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<f32>` (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
);
}
}

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

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

144
crates/ml/tests/sp20_aggregate_inputs_test.rs
View File

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

64
crates/ml/tests/sp20_phase1_4_wireup_test.rs
View File

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

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

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