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

fix(sp12): per-trade event-driven reward composition

Browse Source 
Three architectural changes in one atomic commit per spec ecab584c3:

1. Asymmetric bounded cap (REWARD_NEG_CAP=-10, REWARD_POS_CAP=+5):
   restores prospect-theory loss aversion (2:1 ratio) erased by SP11
   symmetric cap. Per pearl_audit_unboundedness_for_implicit_asymmetry.

2. Min-hold soft penalty with temperature curriculum:
   patience requirement on voluntary exits via deficit/(deficit+T) soft
   factor. Temperature anneals 50->5 across ~100 epochs (curriculum,
   half-life 20 epochs). Trail-fire exits exempted (preserves stop-loss
   design). Penalty flows through r_popart so SP11 controller weighs it
   consistently with other voluntary-exit terms.

3. Zero per-bar shaping (Change 3):
   r_micro = 0 entirely on positioned-non-event bars (per-bar shaping is
   anti-pattern; Q-learning handles credit assignment via TD).
   Per-bar Flat opp_cost zeroed (was the symmetric counterpart to micro).
   r_opp_cost preserved as lump-sum at exit:
     r_opp_cost = -shaping_scale * holding_cost_rate * |position| * hold_time
   on segment_complete (both voluntary and trail-fire). Preserves
   carrying-cost economic concept without per-bar density bias.
   Per pearl_event_driven_reward_density_alignment.

Empirical motivation: train-multi-seed-pmbwn 50-epoch on 6a259942e
showed sharpe-gaming (PnL -30% over 8 epochs while sharpe held at 80).
Three causes: lost loss aversion (Change 1), per-bar gradient (Change
3), no commitment (Change 2). Unified per-trade event-driven design
fixes all three together.

Constants live in state_layout.cuh as Invariant-1 numerical anchors
(REWARD_POS_CAP, REWARD_NEG_CAP, MIN_HOLD_TARGET, MIN_HOLD_PENALTY_MAX,
MIN_HOLD_TEMPERATURE_{START,END,DECAY}). Min-hold temperature is
recomputed in Rust per epoch via min_hold_temperature_for_epoch in
training_loop.rs and passed as a launch scalar. New HEALTH_DIAG line
sp12_event_reward emits the constants per epoch alongside sp11_reward.

Phase 1 = constants only. Phase 2 (ISV adaptive bounds per
feedback_isv_for_adaptive_bounds) deferred until validation results
indicate adaptive need.

LOC: ~344 added (includes spec-required documentation comments)
across experience_kernels.cu, state_layout.cuh,
gpu_experience_collector.rs, training_loop.rs, dqn-wire-up-audit.md.

Build: SQLX_OFFLINE=true cargo check -p ml --lib clean.
Tests: SQLX_OFFLINE=true cargo test -p ml --lib — 938 passed,
13 failed (same 13 failures pre-existing on HEAD ecab584c3, none
related to SP12 changes).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-04 19:04:26 +02:00
parent ecab584c32
commit 17cfbb2503
5 changed files with 360 additions and 187 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -1833,7 +1833,31 @@ extern "C" __global__ void experience_env_step(
* guards every write with the NULL check. Pure GPU buffer (the
* trainer holds it as a `MappedF32Buffer`); reads happen only inside
* `popart_component_ema_kernel` (kernel-to-kernel dataflow). */
float* __restrict__ popart_component_per_sample
float* __restrict__ popart_component_per_sample,
/* SP12 v3 (2026-05-04): per-trade event-driven reward composition.
*
* `min_hold_target` (bars): patience requirement on voluntary trade
* exits. Phase 1 default 30 bars (Invariant-1 numerical anchor in
* state_layout.cuh::MIN_HOLD_TARGET). Voluntary exits with
* `hold_time < min_hold_target` get a soft penalty subtracted from
* `r_popart` via the deficit/(deficit + T) saturation formula.
*
* `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
) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i >= N) return;
@@ -2744,22 +2768,70 @@ extern "C" __global__ void experience_env_step(
* `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;
/* SP11 fix (2026-05-04): symmetric cap. The previous unilateral
* `fminf(base_reward, 10.0f)` capped profits at +10 but left losses
* unbounded — vol_normalized_return has no lower bound, so a single
* large adverse segment_return propagated as a large negative reward
* straight through the modifier chain into reward_components[+0],
* inflating PopArt input EMA (slot 63) and breaking C51/IQN/Bellman
* Q-target normalization. Empirical fingerprint from
* smoke-test-k9drh @ 774d7552a: ep1 r_popart min=-9186, ep2
* min=-79089 (growing), max stuck at +10. Spec semantic is bilateral
* [-10, +10]; enforce bilaterally per pearl_bounded_modifier_outputs_
* require_structural_activation. */
float capped_pnl = fmaxf(-10.0f, fminf(base_reward, 10.0f));
/* 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.
*
* 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.
*
* Constants live in state_layout.cuh as Invariant-1 numerical
* anchors (Phase 1). Phase 2 (deferred) lifts the ratio to an
* ISV-driven controller per feedback_isv_for_adaptive_bounds. */
float capped_pnl = fmaxf(REWARD_NEG_CAP, fminf(base_reward, REWARD_POS_CAP));
if (trail_triggered) {
r_trail = capped_pnl;
} else {
r_popart = capped_pnl;
/* SP12 v3 Change 2: min-hold soft penalty for voluntary exits.
* Encourages MFT-leaning trade duration via patience requirement.
* 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 `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. */
float adjusted_pnl = capped_pnl;
if (segment_hold_time < min_hold_target) {
const float deficit = min_hold_target - segment_hold_time;
const float soft_factor = deficit
/ fmaxf(deficit + min_hold_temperature, 1e-6f);
const float min_hold_penalty = min_hold_penalty_max * soft_factor;
adjusted_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
@@ -2770,6 +2842,12 @@ extern "C" __global__ void experience_env_step(
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;
/* C.2 attribution: rc[2] now carries the trail signal when the
@@ -2778,6 +2856,43 @@ extern "C" __global__ void experience_env_step(
* for ISV-driven component-mag-ratio canaries. */
reward_components_per_sample[out_off * 6 + 2] = r_trail;
/* SP12 v3 Change 3: lump-sum opportunity cost on trade close.
*
* Per spec §design-3 + open question 2 ("lump-sum at exit"): the
* pre-SP12 per-bar Flat opp_cost (Plan 3 B.1) was an exposure-density
* gradient — zeroed in the per-bar branch above. The carrying-cost
* economic concept it modelled is preserved as an integrated
* lump-sum at trade close: a position of size |pre_trade_position|
* held for `segment_hold_time` bars at rate `holding_cost_rate`
* accumulates a total inventory cost of
* k × |pre_trade_position| × hold_time
* regardless of whether the bars were spread across the trade or
* collapsed at exit. Under γ=1 the integrated and per-bar
* formulations are mathematically equivalent; under γ<1 they differ
* slightly due to discount, but the difference is small for short
* trades and the event-driven density alignment dominates per
* pearl_event_driven_reward_density_alignment.
*
* Fires on BOTH voluntary exits and trail-fire (any segment_complete)
* — the carrying-cost economic concept doesn't care WHY the trade
* closed, only that the position was held. Routed into r_opp_cost
* (rc[4]) so the SP11 controller weighs it through w_oc; preserves
* the existing component-attribution semantics.
*
* `pre_trade_position` is the position held during the segment
* (already saved upstream before unified_env_step_core flipped
* `position` for the exit). `segment_hold_time` is the saved
* pre-helper hold count (set to `saved_hold_time` at line ~2493).
* `shaping_scale` gates the term off in pure-validation mode (=0)
* matching the spec's "step_returns are pure per-bar P&L"
* convention for backtest reward shape. */
const float lump_opp_cost = shaping_scale
* holding_cost_rate
* fabsf(pre_trade_position)
* segment_hold_time;
r_opp_cost = -lump_opp_cost;
reward_components_per_sample[out_off * 6 + 4] = r_opp_cost;
/* ── Layer 3: CEA counterfactual loop REMOVED (4-branch: direction(4) replaces exposure(9)).
* Dense micro-reward (Gem 1) below provides directional gradient signal. ── */
@@ -2964,102 +3079,57 @@ extern "C" __global__ void experience_env_step(
* removed. To activate, set micro_reward_scale > 0 in the loaded TOML
* profile (no recompile needed).
* ──────────────────────────────────────────────────────────────────── */
if (!segment_complete && fabsf(position) > 0.001f && micro_reward_scale > 0.0f
&& full_batch_states != NULL && full_state_dim >= SL_OFI_START + SL_OFI_DIM) {
/* ── Dense micro-reward: OFI momentum + MBP-10 price confirmation ──
* Replaces flat holding cost with informed per-bar signal using
* order flow imbalance deltas and mid-price mark-to-market. */
const float* ofi_cur = full_batch_states + (long long)i * full_state_dim + SL_OFI_START;
float sign_pos = (position > 0.0f) ? 1.0f : -1.0f;
/* OFI momentum: delta from previous bar (stored in ps[30..37]) */
float delta_depth = ofi_cur[1] - ps[PS_OFI_PREV_BASE + 1]; /* depth_imbalance_L1 */
float delta_vpin = ofi_cur[6] - ps[PS_OFI_PREV_BASE + 6]; /* VPIN */
float delta_queue = ofi_cur[2] - ps[PS_OFI_PREV_BASE + 2]; /* queue_pressure */
float delta_tarr = ofi_cur[5] - ps[PS_OFI_PREV_BASE + 5]; /* trade_arrival_rate */
float flow_momentum = sign_pos * 0.25f * (delta_depth + delta_vpin + delta_queue + delta_tarr);
/* ATR for normalization (same decode as vol normalization above) */
float atr_abs = 1.0f;
if (features != NULL && bar_idx < total_bars && market_dim > 9) {
float atr_n = features[(long long)bar_idx * market_dim + 9];
float log_a = atr_n * 16.0f - 7.0f;
atr_abs = fmaxf(expf(log_a), 0.01f);
}
float atr_pct = atr_abs / fmaxf(raw_close, 1.0f);
/* Price confirmation: MBP-10 mid-price mark-to-market at decision points */
float raw_open = tgt[4];
float mid_open = tgt[5];
float bar_move = (raw_close - raw_open) / fmaxf(atr_abs, 1e-6f);
float price_confirm = sign_pos * bar_move;
/* Adaptive cost tolerance: profitable models trade more freely */
float spread_at_entry = fabsf(raw_open - mid_open);
float capital_ratio = equity / fmaxf(peak_equity, 1.0f);
float sharpe_ema = ps[DSR_A_SLOT];
float cost_tolerance = fmaxf(capital_ratio * fmaxf(sharpe_ema, 0.0f), 0.1f);
float spread_penalty = spread_at_entry / (fmaxf(atr_abs, 1e-6f) * cost_tolerance);
/* Book aggression from state vector (precomputed from 10-level MBP-10).
* Located at OFI[16] = SL_OFI_START + 16 in canonical layout. */
float book_aggression = (full_state_dim > SL_OFI_START + 16)
? full_batch_states[(long long)i * full_state_dim + SL_OFI_START + 16]
: 0.0f;
/* Informed flow intensity */
float vol_informed = ofi_cur[6] * ofi_cur[7]; /* VPIN x Kyle's_lambda */
/* Retrospective hold quality: was previous bar's hold correct? */
float prev_mid = ps[PREV_MID_SLOT];
float mid_now = tgt[5];
float hold_quality = 0.0f;
if (prev_mid > 0.0f && mid_now > 0.0f) {
float hold_return = (mid_now - prev_mid) / fmaxf(prev_mid, 1.0f);
hold_quality = sign_pos * hold_return / fmaxf(atr_pct, 1e-6f);
}
/* ISV-adaptive reward weights: regime routing adjusts micro-reward composition.
* Volatile regime → weight momentum (price_confirm) higher.
* Stable regime → weight book aggression + hold quality higher.
* Regime transition (low stability) → reduce hold quality (don't hold through shifts). */
float regime_stab = (isv_signals_ptr != NULL) ? isv_signals_ptr[11] : 0.5f;
float volatility = (isv_signals_ptr != NULL) ? isv_signals_ptr[2] : 0.5f;
float w_price = price_confirm_weight * (0.5f + 0.5f * volatility);
float w_book = book_aggression_weight * (0.5f + 0.5f * regime_stab);
float w_hold = hold_quality_weight * regime_stab;
/* Composite quality — regime-adaptive */
float quality = flow_momentum * (1.0f + vol_informed)
+ w_price * price_confirm
+ w_book * sign_pos * book_aggression
+ w_hold * hold_quality
- spread_penalty;
r_micro = shaping_scale * micro_reward_scale * tanhf(quality / micro_reward_temp);
reward = r_micro; /* mutually exclusive with segment_complete; reward starts at 0 here */
/* Task 0.8: captures the micro-reward magnitude BEFORE downstream
* drawdown / conviction / sign-flip transforms. The assignment to
* `r_micro` above is an OVERWRITE (not additive) — the segment_complete
* branch that precedes this `else if` already returned before reaching
* here if it fired, so `r_micro` at this line IS the micro term in
* isolation. Downstream transforms (drawdown penalty, conviction
* scaling, counterfactual sign flip) mutate the post-composition
* `reward` further before `total_reward_per_sample` is written, so
* the ratio
* Σ|micro_pre_transform| / Σ|total_post_transform|
* is a magnitude ratio (diagnostic for relative term scale), not a
* true additive decomposition. See accessor docstring. */
micro_reward_per_sample[out_off] = r_micro;
/* C.2: capture micro component before downstream transforms. */
reward_components_per_sample[out_off * 6 + 3] = r_micro;
} else if (!segment_complete && fabsf(position) > 0.001f) {
/* Fallback: original holding cost when micro_reward_scale=0 or no OFI.
* SP11 B1b: maps to r_micro (positioned-non-completion slot in the
* mutual-exclusivity chain — same semantic role as the OFI branch
* above, just degraded). Captured in micro_reward_per_sample +
* rc[3] for diagnostic continuity with the OFI path. */
r_micro = -shaping_scale * holding_cost_rate * fabsf(position);
/* SP12 v3 Change 3 (2026-05-04): per-bar micro-reward is zero.
*
* Per pearl_event_driven_reward_density_alignment, per-bar shaping for
* credit assignment is the anti-pattern. The 50-epoch validation
* `train-multi-seed-pmbwn` on commit 6a259942e showed sharpe-gaming
* (sharpe=80, PnL -30%, 62% trade rate) — micro_reward + opp_cost firing
* every bar with position injected an exposure-positive density
* gradient that pulled the policy toward HFT noise extraction.
*
* Phase 1 fix: zero `r_micro` entirely. Q-learning handles credit
* assignment via TD targets — that is the right mechanism for "is
* exposure good right now". The OFI/MBP-10 informational signal
* remains available to the encoder via the state vector
* (state[SL_OFI_START..62)); only the reward-layer per-bar density
* bonus is removed. Spec §design-3, open question 3 ("zero entirely").
*
* The branch structure (positioned-non-event vs flat-non-event) is
* preserved because the flat branch still updates the Welford-style
* pre-entry conviction EMA used by D.4c. Both branches now emit
* r_micro = 0 / r_opp_cost = 0 instead of computing per-bar shaping.
* micro_reward_scale, price_confirm_weight, book_aggression_weight,
* hold_quality_weight, micro_reward_temp, holding_cost_rate become
* structurally inert for the per-bar paths (they remain wired for
* the inventory_penalty modifier at the bottom of the kernel which
* is a separate concern; spec carve-out: "ONLY micro_reward and
* opp_cost are the per-bar shaping anti-patterns"). */
if (!segment_complete && fabsf(position) > 0.001f) {
/* SP12 v3: positioned-non-event bar → r_micro = 0.
*
* Pre-SP12 this branch fired EITHER the OFI dense-micro reward (when
* micro_reward_scale > 0 and OFI features wired) OR the
* `-shaping_scale * holding_cost_rate * |position|` fallback. Both
* formulations injected per-bar density-positive (or density-negative)
* reward, contradicting pearl_event_driven_reward_density_alignment.
*
* The OFI informational signal still flows to the encoder via the
* state vector at state[SL_OFI_START..SL_OFI_START + SL_OFI_DIM);
* removing the per-bar reward path only stops the critic-side
* gradient injection. The PREV_MID-based retrospective hold-quality
* signal is the one piece of information NOT in the state vector,
* but per spec §design-3 Q-learning's TD targets handle "was that
* hold correct?" via the eventual segment_complete payoff — the
* dense per-bar version was redundant pressure on the same signal.
*
* `holding_cost_rate`, `micro_reward_scale`, `price_confirm_weight`,
* `book_aggression_weight`, `hold_quality_weight`, `micro_reward_temp`
* all become structurally inert for this branch. They remain wired
* for the inventory_penalty modifier at the bottom of the kernel
* (which is a separate concern; spec carve-out: "ONLY micro_reward
* and opp_cost are the per-bar shaping anti-patterns"). */
r_micro = 0.0f;
reward = r_micro;
micro_reward_per_sample[out_off] = r_micro;
reward_components_per_sample[out_off * 6 + 3] = r_micro;
@@ -3089,86 +3159,33 @@ extern "C" __global__ void experience_env_step(
ps[PS_PRE_ENTRY_CONVICTION_EMA] = new_m;
ps[PS_PRE_ENTRY_CONVICTION_VAR_EMA] = new_v;
}
/* ── Flat opportunity cost — ISV-driven self-adapting via conviction ──
* (val-Flat-collapse fix #3 revision, 2026-04-24)
/* SP12 v3: per-bar Flat opportunity cost is zero.
*
* Prior revision used `0.5 × holding_cost_rate × vol_proxy` — a
* hardcoded 0.5 tuned constant violating
* `feedback_isv_for_adaptive_bounds.md`. The magnitude of
* "missed-opportunity" isn't a number — it's a signal that
* depends on how strong the model's own directional belief is.
* Pre-SP12 this branch fired the ISV-driven conviction × vol × Q-ref
* × KL-amp opportunity-cost penalty on every flat bar — the symmetric
* counterpart to the positioned-bar micro_reward, both formulations
* baking a per-bar density gradient into the reward signal. Per
* pearl_event_driven_reward_density_alignment that gradient pulled
* the policy toward HFT noise extraction in the
* `train-multi-seed-pmbwn` 50-epoch validation on commit 6a259942e.
*
* Semantics this branch covers:
* position ≈ 0 AND not a trade-completion bar. This fires on
* both `dir=Flat` (explicit close-to-zero) and `dir=Hold while
* position=0` (no-op stay-flat). The Hold-while-in-trade case
* (position≠0) hits the positioned-bar holding-cost branch
* above and is NOT penalised here — staying in a position IS
* an active choice that the holding-cost mechanism already
* prices. The distinction is structural: Hold-at-zero and Flat
* produce identical portfolio outcomes so their reward shaping
* must match.
* The carrying-cost economic concept is preserved as a lump-sum
* applied in the segment_complete branch (see "SP12 v3 Change 3:
* lump-sum opportunity cost" block above) — the cost magnitude is
* the same `−shaping_scale × holding_cost_rate × |position| ×
* hold_time` integrated over the trade's duration, just emitted at
* the close event instead of distributed per-bar. Spec §design-3,
* open question 2 ("lump-sum at exit").
*
* Self-adapting via conviction (∈ [0, 1]):
* conviction = direction-branch Q-range normalised by ISV[21]
* (q_dir_abs_ref EMA). Computed per-sample in
* `experience_action_select` and threaded into this kernel via
* `conviction_ptr`. Already clamped to [0, 1] at source.
* - Cold start (ISV[21] ≈ 0): fallback conviction = 1.0 → full
* penalty, encourages early exploration out of the flat
* equilibrium.
* - Low conviction: network is uncertain about direction → the
* penalty scales toward zero → don't force trades on noise.
* - High conviction: network has learned strong directional
* edge → penalty scales up → Flat becomes expensive ONLY
* where the model itself says there's an edge to capture.
*
* This is the right shape for "incentive for taking action":
* it's ISV-bus-driven, self-adapting via EMA state, and
* piggybacks on a signal the network's own output defines.
* No tuned multiplier.
*
* Volatility factor (vol_proxy) retained:
* Scales with realised ATR so quiet markets produce little
* penalty. Capped at 0.01 (1%) as numerical-safety bound
* against single-bar vol spikes. */
if (features != NULL && bar_idx < total_bars && market_dim > 9) {
float atr_norm_flat = features[(long long)bar_idx * market_dim + 9];
float log_atr_flat = atr_norm_flat * 16.0f - 7.0f;
float atr_pct_flat = expf(log_atr_flat) / fmaxf(raw_close, 1.0f);
float vol_proxy_flat = fmaxf(atr_pct_flat, 0.0001f);
if (vol_proxy_flat > 0.01f) vol_proxy_flat = 0.01f;
/* B.1 (Plan 3 Task 2): scale Flat opp-cost by ISV[21] (q_dir_abs_ref
* EMA, max(|Q_mean|) across direction bins). Self-scaling — as Q
* magnitudes drift during training, opp-cost tracks proportionally.
* When Q ≈ ±50, conviction_core keeps the penalty visible relative
* to action Q-values. No tuned multiplier. Populated by
* update_eval_v_range / q_stats_kernel.
* Floor 1e-3 ensures a non-zero penalty during the first few
* epochs before the EMA is warm (cold-start protection). */
const float q_abs_ref = (isv_signals_ptr != NULL)
? fmaxf(1e-3f, isv_signals_ptr[21])
: 1e-3f;
/* C.3 Plan 3 Task 7: amplify Flat-trap escape pressure when the
* train/val state distribution diverges. `kl_amp` ∈ [1, 2] is
* GPU-written by `state_kl_moment_match`. The fmaxf(1.0, …) guard
* is critical: cold-start is 0 if the constructor write was skipped
* (or 1 if it ran), and we never want to scale rewards by < 1. Bound
* is structural (kernel emits target ∈ [1, 2] via clamp), so this
* stacks safely with B.1's q_abs_ref unbounded multiplicand per
* pearl_one_unbounded_signal_per_reward.md (the unbounded slot is
* already there; multiplying by a bounded amp keeps it safe). */
const float kl_amp_b1 = (isv_signals_ptr != NULL)
? fmaxf(1.0f, isv_signals_ptr[ISV_STATE_KL_AMP_IDX])
: 1.0f;
/* conviction_core ∈ [0, 1] — ISV-driven per-sample adaptive
* multiplier. Replaces the prior 0.5 hardcoded constant. */
r_opp_cost = -shaping_scale * holding_cost_rate
* conviction_core * vol_proxy_flat * q_abs_ref * kl_amp_b1;
reward = r_opp_cost; /* mutually exclusive with segment_complete + micro paths */
/* C.2 attribution: record opp_cost component. */
reward_components_per_sample[out_off * 6 + 4] = r_opp_cost;
}
* The Welford-style D.4c conviction EMA above MUST stay because it
* is a state-tracking signal (read by the entry-bonus block to
* detect stable pre-entry conviction), not reward shaping. Removing
* it would break the D.4c novelty mechanism. Spec carve-out
* preserved: only the Flat opp_cost EMITTER is zeroed; the
* conviction Welford observer is independent. */
r_opp_cost = 0.0f;
reward = r_opp_cost;
reward_components_per_sample[out_off * 6 + 4] = r_opp_cost;
}
/* ── B.2 (Plan 3 Task 3): ISV-driven trade-attempt bonus ──

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

@@ -296,6 +296,27 @@ pub struct ExperienceCollectorConfig {
/// Branch order: [direction(0), magnitude(1), order(2), urgency(3)].
/// 0.0 per-branch disables noise for that branch. Default: [0.1; 4].
pub noise_sigma_per_branch: [f32; 4],
/// SP12 v3 (2026-05-04): per-trade event-driven reward composition.
///
/// Patience target on voluntary trade exits (in bars). Voluntary exits
/// with `hold_time < min_hold_target` get a soft penalty subtracted
/// from `r_popart` via the deficit/(deficit + T) saturation formula.
/// Phase 1 default = `MIN_HOLD_TARGET=30` (state_layout.cuh).
/// Trail-fire forced exits exempted — see kernel docstring.
pub min_hold_target: f32,
/// SP12 v3: max penalty magnitude when min-hold deficit is infinite
/// (soft factor → 1). Phase 1 default = `MIN_HOLD_PENALTY_MAX=3.0`
/// (60% of `REWARD_POS_CAP`).
pub min_hold_penalty_max: f32,
/// SP12 v3: temperature controller for the min-hold soft saturation.
/// Computed per-epoch in `training_loop.rs::min_hold_temperature_for_epoch`
/// via the annealing schedule
/// `T_end + (T_start - T_end) × exp(-epoch / decay)` with constants
/// from `state_layout.cuh::MIN_HOLD_TEMPERATURE_{START,END,DECAY}`.
/// Default at construction = `MIN_HOLD_TEMPERATURE_START=50.0`
/// (epoch-0 forgiving regime; replaced per epoch by training loop).
pub min_hold_temperature: f32,
}
impl Default for ExperienceCollectorConfig {
@@ -367,6 +388,13 @@ impl Default for ExperienceCollectorConfig {
hindsight_fraction: 0.0,
hindsight_lookahead: 10,
noise_sigma_per_branch: [0.1; 4],
// SP12 v3 defaults — match Invariant-1 anchors in state_layout.cuh.
// The training loop overrides `min_hold_temperature` per epoch via
// the annealing schedule (50 → 5 across ~100 epochs); the other
// two are static Phase 1 constants.
min_hold_target: 30.0,
min_hold_penalty_max: 3.0,
min_hold_temperature: 50.0,
}
}
}
@@ -4108,6 +4136,18 @@ impl GpuExperienceCollector {
// in training_loop.rs init). Spec §4 amendment "Why
// slot 360".
.arg(&mut self.popart_component_per_sample)
// SP12 v3 (2026-05-04): per-trade event-driven reward
// composition parameters. The kernel applies a soft
// min-hold penalty on voluntary exits (deficit/(deficit+T)
// saturation) using these three scalars; constants live
// in state_layout.cuh (MIN_HOLD_TARGET, MIN_HOLD_PENALTY_MAX,
// MIN_HOLD_TEMPERATURE_{START,END,DECAY}). The temperature
// is recomputed in Rust per epoch via the annealing
// schedule before this launch (see
// `training_loop.rs::min_hold_temperature_for_epoch`).
.arg(&config.min_hold_target)
.arg(&config.min_hold_penalty_max)
.arg(&config.min_hold_temperature)
.launch(launch_cfg)
.map_err(|e| MLError::ModelError(format!(
"experience_env_step t={t}: {e}"

40
crates/ml/src/cuda_pipeline/state_layout.cuh
View File

@@ -177,6 +177,46 @@
#define SL_NUM_FEATURE_GROUPS 6
#define SL_MAX_FEATURE_GROUP_DIM 42 // == SL_MARKET_DIM (largest group)
// ────────────────────────────────────────────────────────────────────────────
// SP12 v3 (2026-05-04): per-trade event-driven reward composition constants.
// Spec: docs/superpowers/specs/2026-05-04-sp12-quality-over-quantity-reward-shaping.md
//
// REWARD_POS_CAP / REWARD_NEG_CAP — asymmetric bounded reward cap (loss aversion).
// 2:1 ratio matches Kahneman/Tversky prospect-theory meta-analysis pegging
// human loss aversion at ~2.0-2.25. Restores selectivity that SP11's
// symmetric ±10 cap (commit 35db31089) erased while preserving the bilateral
// bound that fixed slot-63 PopArt EMA inflation. Site:
// experience_kernels.cu line ~2758 (segment_complete vol-normalized P&L).
// Per pearl_audit_unboundedness_for_implicit_asymmetry.
//
// MIN_HOLD_TARGET — patience requirement on voluntary trade exits, in bars.
// 30 bars = MFT-leaning but not pure MFT (full MFT ≈ 100 bars,
// HFT-tolerant ≈ 10 bars). Matches the user-profile target band
// "between HFT and MFT".
//
// MIN_HOLD_PENALTY_MAX — max penalty magnitude when voluntary exit at hold=0.
// 3.0 = 60% of REWARD_POS_CAP. Meaningful gradient pressure but not
// paralyzing — a profitable scalp can still net positive with the penalty.
//
// MIN_HOLD_TEMPERATURE_{START,END,DECAY} — annealing schedule for the
// soft-saturation temperature in the deficit/(deficit + T) formula. Computed
// in Rust per epoch via T(e) = T_end + (T_start - T_end) × exp(-e/decay)
// and passed to the kernel as a scalar (not a constant lookup). Curriculum:
// early training forgiving (T=50), late training sharp (T=5).
//
// Phase 1 = Invariant-1 numerical anchors (static constants). Phase 2 lifts
// the target / penalty_max / temperature schedule to ISV-driven controllers
// per pearl_kelly_cap_signal_driven_floors and feedback_isv_for_adaptive_bounds
// IF Phase 1 validation reveals adaptive need. Spec §implementation phases.
// ────────────────────────────────────────────────────────────────────────────
#define REWARD_POS_CAP 5.0f
#define REWARD_NEG_CAP -10.0f
#define MIN_HOLD_TARGET 30.0f
#define MIN_HOLD_PENALTY_MAX 3.0f
#define MIN_HOLD_TEMPERATURE_START 50.0f
#define MIN_HOLD_TEMPERATURE_END 5.0f
#define MIN_HOLD_TEMPERATURE_DECAY 20.0f
// ── Compile-time checks ──
static_assert(SL_PADDING_START + SL_PADDING_DIM == SL_STATE_DIM,
"State layout dimensions must sum to SL_STATE_DIM");

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

@@ -42,6 +42,36 @@ use super::super::financials::compute_epoch_financials;
use super::super::monitoring::TrainingMonitor;
use super::DQNTrainer;
/// SP12 v3 (2026-05-04): per-epoch min-hold soft-penalty temperature schedule.
///
/// Computes the temperature T(e) controlling the smoothness of the
/// `deficit/(deficit + T)` saturation factor inside the
/// `experience_env_step` kernel. Annealing curriculum:
/// `T(e) = T_end + (T_start - T_end) × exp(-e / decay)`
///
/// With Phase 1 anchors `T_start=50.0`, `T_end=5.0`, `decay=20.0`
/// (state_layout.cuh::MIN_HOLD_TEMPERATURE_{START,END,DECAY}):
///
/// | epoch | T | regime |
/// |------:|----:|------------------|
/// | 0 | 50 | forgiving |
/// | 20 | 21 | mid-transition |
/// | 50 | 9 | sharp |
/// | 100 | 5 | full-MFT enforce |
///
/// Per spec §design-2 — early training needs forgiveness so the policy
/// can learn reward gradients on short trades; late training needs
/// sharpness so the min-hold penalty bites and enforces commitment.
/// Constants live in CUDA headers; this Rust mirror duplicates them
/// (kernel side keeps T as a launch scalar so we can switch to an
/// ISV-driven formulation in Phase 2 without recompiling cubin).
pub(crate) fn min_hold_temperature_for_epoch(epoch: usize) -> f32 {
const T_START: f32 = 50.0;
const T_END: f32 = 5.0;
const DECAY_RATE: f32 = 20.0;
T_END + (T_START - T_END) * (-(epoch as f32) / DECAY_RATE).exp()
}
/// Normalized Shannon entropy of a discrete probability distribution.
///
/// Returns the entropy divided by log(N) where N is the number of bins, so
@@ -1816,6 +1846,18 @@ impl DQNTrainer {
}
arr
},
// SP12 v3 (2026-05-04): per-trade event-driven reward composition.
// `min_hold_target` and `min_hold_penalty_max` are Phase 1 static
// constants — defaults are read directly from ExperienceCollectorConfig::default()
// (MIN_HOLD_TARGET=30, MIN_HOLD_PENALTY_MAX=3.0). Phase 2 lifts these
// to ISV-driven controllers per feedback_isv_for_adaptive_bounds when/if
// validation results indicate adaptive need.
// `min_hold_temperature` is computed PER EPOCH from the annealing
// schedule. Curriculum: forgiving (T=50) at epoch 0 → sharp (T=5)
// around epoch 100. Half-life ≈ MIN_HOLD_TEMPERATURE_DECAY (=20)
// epochs. The kernel evaluates `deficit/(deficit + T)` per voluntary
// exit; lower T narrows the transition into a sharper penalty cliff.
min_hold_temperature: min_hold_temperature_for_epoch(self.current_epoch),
..Default::default()
};
@@ -4448,6 +4490,24 @@ impl DQNTrainer {
curiosity_p, sab_mult, w_floor, curiosity_b,
improvement_z,
);
// SP12 v3 (2026-05-04): per-trade event-driven reward composition diagnostic.
// Mirrors the constants emitted into the experience_env_step kernel for
// this epoch — lets `multi_seed_validation` / smoke pipelines verify
// (1) the asymmetric cap is in force [pos=+5, neg=-10], (2) the
// min-hold target / penalty_max anchors match the spec defaults, and
// (3) the temperature annealing schedule is on the expected curve
// (forgiving early → sharp late). Same emit cadence as sp11_reward
// above (per-epoch on the metrics path).
tracing::info!(
"HEALTH_DIAG[{}]: sp12_event_reward [pos_cap={:.1} neg_cap={:.1} min_hold_tgt={:.1} min_hold_T={:.2} min_hold_pen={:.1}]",
epoch,
5.0_f32, // REWARD_POS_CAP — Invariant-1 anchor in state_layout.cuh
-10.0_f32, // REWARD_NEG_CAP
30.0_f32, // MIN_HOLD_TARGET (Phase 1 default; Phase 2 lifts to ISV)
min_hold_temperature_for_epoch(epoch),
3.0_f32, // MIN_HOLD_PENALTY_MAX
);
}
// SP6 Pearl 2: per-branch budget HEALTH_DIAG lines — enables Layer C debugging

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

@@ -2,6 +2,22 @@
**Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7.
SP12 v3 — per-trade event-driven reward composition (2026-05-04): three architectural changes in one atomic commit per spec `2026-05-04-sp12-quality-over-quantity-reward-shaping.md` (committed at `ecab584c3`). Empirical motivation: 50-epoch validation `train-multi-seed-pmbwn` on commit `6a259942e` showed sharpe-gaming (sharpe=80 held while PnL declined 30% over 8 epochs; 62% trade rate, 46% win rate stuck). **Three causes** addressed in one commit per `feedback_no_partial_refactor`:
(1) **Asymmetric bounded cap** at `experience_kernels.cu:2758` — replaces SP11's symmetric ±10 cap (commit `35db31089`) with `fmaxf(REWARD_NEG_CAP=-10, fminf(base_reward, REWARD_POS_CAP=+5))`. The 2:1 ratio matches Kahneman/Tversky meta-analysis pegging human loss aversion at ~2.0-2.25 — calibration to a known psychological constant (Invariant-1 numerical anchor in `state_layout.cuh`), not a tuned multiplier. SP11's stability work preserved (cap still bilateral). Per `pearl_audit_unboundedness_for_implicit_asymmetry`: SP11 erased the pre-fix implicit downside-asymmetry by symmetrizing the bound; the model became risk-neutral and locked into HFT noise extraction. SP12 widens only the negative cap back toward the pre-SP11 unbounded regime.
(2) **Min-hold soft penalty** with temperature-annealing curriculum at `experience_kernels.cu` voluntary-exit branch (inside `segment_complete && !trail_triggered`). Computes `soft_factor = deficit / (deficit + T)` where `deficit = max(0, MIN_HOLD_TARGET - segment_hold_time)` and applies `min_hold_penalty_max × soft_factor` (default `MIN_HOLD_PENALTY_MAX=3.0` = 60% of REWARD_POS_CAP) to `r_popart` before the SP11 controller weighting. The penalty flows through the popart component slot so the SP11 mag-ratio canary observes it consistently. Trail-fire forced exits exempted by control-flow structure (the `else` branch of `if (trail_triggered)`); preserves stop-loss design semantics. Temperature is recomputed in Rust per epoch via `min_hold_temperature_for_epoch(epoch) = T_END + (T_START - T_END) × exp(-epoch / DECAY)` with anchors `T_START=50, T_END=5, DECAY=20` from `state_layout.cuh::MIN_HOLD_TEMPERATURE_{START,END,DECAY}`. Curriculum: epoch 0 → T≈50 (forgiving so the policy can learn reward gradients on short trades), epoch 50 → T≈9, epoch 100 → T≈5 (sharp enforcement of MFT commitment). Soft-saturation formulation (no cliff) chosen over hard threshold to avoid gradient discontinuity at the boundary.
(3) **Zero per-bar shaping (event-driven density)** at `experience_kernels.cu` per-bar branches. The pre-SP12 `r_micro` (positioned-non-event) and `r_opp_cost` (flat-non-event) terms fired every bar, injecting an exposure-positive density gradient that signaled "exposure earns reward density independent of trade outcomes" — the structural mechanism that pulled the policy toward HFT. Per `pearl_event_driven_reward_density_alignment`: per-bar shaping for credit assignment is the anti-pattern; Q-learning's TD targets are the correct mechanism for "is exposure good right now?" (a) `r_micro = 0` on positioned-non-event bars (pre-SP12 had OFI dense-micro path OR fallback holding-cost; both are removed; OFI informational signal still flows to encoder via state vector at `state[SL_OFI_START..62)` — only the critic-side per-bar reward injection is gone; `holding_cost_rate`, `micro_reward_scale`, `price_confirm_weight`, `book_aggression_weight`, `hold_quality_weight`, `micro_reward_temp` become structurally inert for this branch). (b) Flat-bar `r_opp_cost = 0` (pre-SP12 was conviction × vol × Q-ref × KL-amp formulation); the Welford-style D.4c conviction EMA is preserved (state-tracking signal, not reward shaping — reads by the entry-bonus block). (c) **Lump-sum carrying cost added at trade close**: in `segment_complete` branch, after the cap and before the cf block, computes `r_opp_cost = -shaping_scale × holding_cost_rate × |pre_trade_position| × segment_hold_time`. Fires on BOTH voluntary exits and trail-fire (the carrying-cost economic concept doesn't care WHY the trade closed). Mathematically equivalent to the deleted per-bar formulation under γ=1; small discount difference under γ<1 dominated by the event-driven density alignment. Routed through `r_opp_cost` (rc[4]) so the SP11 controller weighs it through `w_oc` — preserves component-attribution semantics.
**Files touched** (atomic per `feedback_no_partial_refactor`): `crates/ml/src/cuda_pipeline/state_layout.cuh` (+7 Invariant-1 anchors: REWARD_POS_CAP, REWARD_NEG_CAP, MIN_HOLD_TARGET, MIN_HOLD_PENALTY_MAX, MIN_HOLD_TEMPERATURE_{START,END,DECAY}); `crates/ml/src/cuda_pipeline/experience_kernels.cu` (+3 kernel parameters at end of `experience_env_step` signature; cap site at line ~2758 changed to asymmetric; min-hold soft penalty inserted in voluntary-exit branch; lump-sum opp_cost added in segment_complete; r_micro positioned-bar branch collapsed to zero-write; flat-bar r_opp_cost zeroed; OFI dense-micro computation block removed); `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (+3 fields on `ExperienceCollectorConfig` with default values matching state_layout.cuh anchors; +3 launch args in `experience_env_step` invocation); `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (+1 helper fn `min_hold_temperature_for_epoch`; +1 field write `min_hold_temperature: min_hold_temperature_for_epoch(self.current_epoch)` in collector config builder; +1 HEALTH_DIAG line `sp12_event_reward [pos_cap=+5.0 neg_cap=-10.0 min_hold_tgt=30.0 min_hold_T=… min_hold_pen=3.0]` for per-epoch observability). LOC: ~344 added (heavy with spec-required documentation comments per `feedback_no_quickfixes`).
**Phase split**: Phase 1 = Invariant-1 numerical anchors only (this commit). Phase 2 (deferred — only if Phase 1 validation reveals adaptive need) lifts MIN_HOLD_TARGET / MIN_HOLD_PENALTY_MAX / temperature-schedule to ISV-driven controllers per `feedback_isv_for_adaptive_bounds.md` and `pearl_kelly_cap_signal_driven_floors.md` precedent. The temperature is already structurally lift-ready: kernel reads it as a launch scalar, so Phase 2 only swaps the Rust producer (annealing-schedule fn → ISV-driven controller) without touching the kernel.
**Hard rules upheld**: `feedback_no_atomicadd` (no new producers), `feedback_no_partial_refactor` (single atomic commit covering all 3 changes plus all consumers — `ExperienceCollectorConfig` defaults, kernel signature, kernel launch site, training loop temperature computation, HEALTH_DIAG emit), `feedback_isv_for_adaptive_bounds` (Phase 1 = static Invariant-1 anchors documented as such; Phase 2 path identified), `feedback_no_quickfixes` (the architectural fix per pearls, not a tuning knob), `feedback_no_feature_flags` (behavior change unconditional), `feedback_no_todo_fixme` (no TODO markers), `pearl_event_driven_reward_density_alignment` (drives Change 3), `pearl_audit_unboundedness_for_implicit_asymmetry` (drives Change 1). Build: `SQLX_OFFLINE=true cargo check -p ml --lib` clean (only pre-existing warnings). Tests: `SQLX_OFFLINE=true cargo test -p ml --lib` 938 passed, 13 failed — all 13 failures pre-existing on `ecab584c3` (verified via `git stash` baseline run); none introduced by SP12. Smoke validation pending (5-epoch L40S per spec §validation-smoke).
**Backtest-kernel cap site investigation**: spec §design-1 mentioned `crates/ml/src/cuda_pipeline/backtest_plan_kernel.cu` as the matching cap location. Investigation confirmed `backtest_plan_kernel.cu` lines 167/171 contain `fmaxf(-2.0f, fminf(unrealized / (plan_profit × equity + 1e-6f), 2.0f))` for `pisv[PLAN_ISV_PNL_VS_TARGET]` and `pisv[PLAN_ISV_PNL_VS_STOP]` — these are STATE FEATURE normalizations (the policy reads them as input features), not reward caps. Same pattern mirrored in `experience_kernels.cu` lines 852-856. The state-feature representations need to remain symmetric so the policy's perception of "way over target" vs "way under stop" is balanced. Asymmetrizing those would be a different change with different rationale; not in scope for SP12. The reward cap is exclusively at `experience_kernels.cu:2758`, the only site touched by Change 1.
SP6 clean-compile — wire-or-delete all 13 ml + 2 ml-dqn warnings (2026-05-01): W1 SEMANTIC: Pearl 2's `iqn_branch[b]` (per-branch IQN budget, `[f32; 4]` returned by `compute_adaptive_budgets`) was destructured but never consumed — both parallel and sequential Pearl 5 IQN paths used `iqn_trunk / 4.0` (the mean), silently averaging out Pearl 2's per-branch differentiation. Fixed in both paths by moving `let iqn_budget_per_branch = iqn_branch[branch_idx] / 4.0_f32` inside the `for branch_idx in 0..4` loop so each pass uses its branch-specific budget. `iqn_trunk` is now only referenced in comments → renamed to `_iqn_trunk` in the destructure. W2 IMPORT DELETE: `ATOM_NUM_ATOMS_BASE` and `NOISY_SIGMA_BASE` removed from `fused_training.rs:44` — consumers are in Layer B's `atoms_update_kernel.cu` and `experience_kernels.cu`, not `fused_training.rs`; imports were speculative additions from an earlier draft. W3 DELETE: `v_blocks` at `gpu_iql_trainer.rs:756` — computed `((b+255)/256)` but both downstream kernel launches used `v_blocks2` (the 2× variant); stale dead code. W4 IMPORT DELETE: `OrderRouter` removed from `ml-dqn/src/dqn.rs:25` — import only, never used in the file body. W5 IMPORT DELETE: `get_snapshots_for_timestamp` removed from `data_loading.rs:312` — imported inside an `if !mbp10_data_dir.is_empty()` block but the function was never called; `get_trades_for_bar` and `OFICalculator` ARE used. W6 DELETE: `update_target_networks` method in `ml-dqn/src/dqn.rs:1621` — 42-line orphan with docstring claiming it's "shared implementation used by train_step and apply_accumulated_gradients" but `grep` showed zero call sites; real training path uses the fused CUDA EMA kernel on flat buffers. Cascaded: `convergence_half_life` import (now unused) deleted from line 14. W7-W12 FILE-LEVEL `#![allow(unsafe_code)]`: `cublas_algo_deterministic.rs` and `sp4_wiener_ema.rs` — workspace `Cargo.toml` has `unsafe_code = "warn"` which fires for every `unsafe impl / unsafe fn`. Both files require unsafe for cuBLAS-Lt handle passing and CUDA kernel launches; added file-level attr matching the established pattern in `mapped_pinned.rs`, `gpu_iqn_head.rs`, `gpu_weights.rs`, etc. W13 same treatment in `sp4_wiener_ema.rs` (covered by the same file-level attr). W14-W15 VISIBILITY SCOPE: `ControllerPrevValues` and `ControllerFireCounts` in `trainer/mod.rs:172,197` changed from `pub struct` to `pub(crate) struct` — both are consumed only within the `ml` crate; `pub` was unreachable from outside. Per `feedback_no_hiding`: zero `#[allow(...)]` suppressions added at item level; 2 file-level `#![allow(unsafe_code)]` follow established crate pattern. cargo check (ml + ml-dqn) + cargo build --release + cargo test --lib sp5/sp4/state_reset_registry all clean (0 warnings; 13/13 pass).
SP5 Layer A bug-fix — add reset_named_state dispatch arms for 21 SP5 entries (2026-05-02): L40S smoke at SP5 Layer B HEAD (commit `3ad5e011b`) failed in 9.4s at the first fold-reset boundary with `StateResetRegistry reset dispatch: unknown name 'sp5_atom_v_center'`. Every Layer A task added `RegistryEntry` blocks to `state_reset_registry.rs` but **none added the matching dispatch arms in `reset_named_state`** — the runtime validator catches this on first fold boundary, crashing all 3 folds before training runs (0/3 checkpoints saved). Local unit tests verified slot layouts but never exercised the fold-reset dispatch path, so the gap was masked through Layer A close-out. Added 21 dispatch arms covering all SP5 Layer A FoldReset entries: `sp5_atom_v_center/v_half/headroom/clip_rate` (Pearl 1), `sp5_branch_entropy/q_var_per_branch` (shared signal), `sp5_wiener_state` (no-op — `reset_sp4_wiener_state` already bulk-zeros the full 543-float buffer including the SP5 block at offsets [213..543)), `sp5_noisy_sigma/sigma_fraction` (Pearl 3), `sp5_budget_c51/iqn/cql/ens/flatness` (Pearl 2), `sp5_adam_beta1/beta2/eps` (Pearl 4), `sp5_grad_prev_buf` (Pearl 4 aux — calls a new `reset_sp5_grad_prev_buf` bulk-zero method on `GpuDqnTrainer`, mirrors `reset_sp4_clamp_engage_counters` pattern), `sp5_iqn_tau` (Pearl 5), `sp5_trail_dist_per_dir` (Pearl 8), `sp5_atom_num_atoms` (Pearl 1-ext). **Pearl 6 (slots 280..286) intentionally absent** — those slots are cross-fold-persistent (the whole point of A6 per `project_magnitude_eval_collapse_kelly_capped`). Each ISV-range arm zeros its slot range to fire Pearl A's first-observation sentinel on the new fold's first producer launch. Touched: `cuda_pipeline/gpu_dqn_trainer.rs` (+1 method `reset_sp5_grad_prev_buf`), `trainers/dqn/trainer/training_loop.rs` (+21 dispatch arms before the `_ =>` fallback). cargo check + cargo test --lib (sp4+sp5+state_reset_registry: 13/13 pass) both clean. Smoke re-deployment pending.