Three architectural changes in one unified design to push the model from
HFT noise extraction (62% trade rate, sharpe-gaming) toward MFT alpha-hunting:
1. Asymmetric bounded cap (-10/+5) — restores loss aversion erased by
SP11 symmetric cap. 2:1 ratio matches Kahneman/Tversky prospect theory.
Anchor: pearl_audit_unboundedness_for_implicit_asymmetry.
2. Min-hold soft penalty with temperature curriculum — patience requirement
at exit. Soft factor = deficit/(deficit+T), T anneals 50→5 over 50 epochs.
Forces commitment without paralysis in early training.
3. Zero per-bar shaping (gate micro/opp_cost on events) — eliminates
continuous-reward gradient that pulls toward continuous exposure.
Anchor: pearl_event_driven_reward_density_alignment.
Combined: reward fires only on trade events with prospect-theory loss
aversion + commitment requirement. Pure per-trade event-driven Q-learning
properly aligned with per-trade P&L objective.
~50 LOC across 3-4 files. No new ISV slots in Phase 1 (constants only).
Cost ~€1.30 (€0.30 smoke + €1.00 30-epoch validation).
Empirical motivation: train-multi-seed-pmbwn 50-epoch on commit 6a259942e
showed sharpe-gaming pattern (PnL -30% over 8 epochs while sharpe held).
SP11 cap fix unmasked the per-bar shaping bias plus erased loss-aversion
that the unbounded loss path was implicitly providing.
Continues on sp11-reward-as-controlled-subsystem branch — SP12 is
architectural continuation of SP11, not separate work.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Documents the SP11 sharpe-degradation root cause investigation:
- 35db31089: symmetric reward cap (THE bug)
- 348f6078b: plan_isv symmetric clamp mirrors (4 sites)
- b92dcc3df: diagnostic instrumentation cleanup
- this commit: close-out
New memory pearl: pearl_symmetric_clamp_audit
Distinct from pearl_bounded_modifier_outputs_require_structural_activation:
- That pearl: model OUTPUTS need structural activation (sigmoid/tanh)
- This pearl: intermediate kernel scalars need bilateral fminf/fmaxf
Both share: bounded contracts must be enforced structurally.
Audit doc: cross-references bug, fix, validation smoke, both pearls.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Instrumentation from 774d7552a served its purpose — empirically
identified the asymmetric reward cap at experience_kernels.cu:2788
as the inflater (popart min reaching -210336 by F0 ep3 pre-fix).
Fixed in 35db31089; validated by smoke-test-trk72 (PASSED).
Removed:
- reward_chain_diag_reduce_kernel.cu
- 12 per-sample diagnostic buffers in gpu_experience_collector
- kernel parameter threading in experience_kernels.cu
- launcher + reader + mapped-pinned output in gpu_dqn_trainer
- wire-up site + HEALTH_DIAG emit in training_loop
- audit doc section for the transient instrumentation
This brings the worktree back to its pre-instrumentation state on
the SP11 reward chain. Symmetric reward cap fix (35db31089) and
plan_isv symmetric clamps (Commit A) remain.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Implementer of 35db31089 (symmetric reward cap) flagged 4 additional
asymmetric-clamp sites in policy-input features (plan_isv slots),
mirroring the same bug class but on the feature side rather than
reward side:
experience_kernels.cu:850 plan_isv[PNL_VS_TARGET] capped above only
experience_kernels.cu:853 plan_isv[PNL_VS_STOP] capped above only
backtest_plan_kernel.cu:164 — mirror of 850
backtest_plan_kernel.cu:168 — mirror of 853
unrealized P&L can be negative -> fminf(x, 2.0) leaves an unbounded
lower tail that produces feature jitter at the policy input, hurting
the network's state representation under losing-trade conditions.
Fix: symmetric clamp via fmaxf(-2.0, fminf(x, 2.0)).
Per pearl_symmetric_clamp_audit (added in close-out commit) and
pearl_bounded_modifier_outputs_require_structural_activation: any
spec-bounded scalar requires bilateral enforcement.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
experience_kernels.cu:2788:
float capped_pnl = fminf(base_reward, 10.0f);
^^^^^^^^^^^^^ caps profits, NOT losses
Diagnostic instrumentation in smoke-test-k9drh on commit 774d7552a
captured the asymmetry empirically:
ep1: r_popart min=-9186, max=+10
ep2: r_popart min=-79089, max=+10 (growing)
`base_reward = 2.0f * vol_normalized_return` and
`vol_normalized_return = segment_return / vol_norm` where
`segment_return` has no structural lower bound (signed P&L). The
unilateral `fminf(base_reward, 10.0f)` capped the upper tail only,
so a single large adverse segment_return produced an arbitrarily
negative `capped_pnl` → r_popart → r_weighted →
reward_components[+0] → slot 63 (PopArt input EMA), inflating
C51/IQN/Bellman normalization scale and breaking Q-target
consistency across epochs. Empirical fingerprint matched the
within-fold sharpe degradation observed in smoke-test-gwfn8 on
commit fd24b5383 (10→4 within F0, 9→2.5 within F1).
Spec semantic: reward bounded in [-10, +10]. Fix:
float capped_pnl = fmaxf(-10.0f, fminf(base_reward, 10.0f));
Audit findings (per task §2): all related fminf/fmaxf clamps in
experience_kernels.cu reviewed. Reward modifier chain (3260-3500)
verified bounded once r_popart is bilateral. Other bilateral
clamps already correct (515-517, 2174, 2191, 2227, 3300, 3452,
3729, 3763, 5076, 5210, 5213, 5461, 5535, 6353, 6693, 6694).
Intentional asymmetries verified at 1078, 1082-1085, 2564-2565,
2873, 2949, 4574, 5896 (each documented in the audit doc with
the structural reason the lower side is unbounded).
Latent finding flagged separately (NOT fixed here — feature-side
requires consumer audit per feedback_no_partial_refactor):
plan_isv[PNL_VS_TARGET] at line 850 and plan_isv[PNL_VS_STOP]
at line 853 are upper-clamped at 2.0 but lower-unbounded. These
feed assemble_state as policy features (not reward components),
so out of scope of this reward-chain fix. Mirrored in
backtest_plan_kernel.cu:164,168 (same pattern). Tracking as
follow-up.
Per pearl_bounded_modifier_outputs_require_structural_activation:
spec-bounded values require BILATERAL structural enforcement.
This bug was the asymmetric counterpart to the conviction sigmoid
(which IS correctly bounded structurally).
Diagnostic instrumentation (commit 774d7552a) NOT removed in this
commit — will be removed in a follow-up after the symmetric-cap
smoke validates the fix on L40S.
docs/dqn-wire-up-audit.md updated with Resolution (2026-05-04)
section reflecting root cause, fix, audit follow-through, and the
latent plan_isv finding (per Invariant 7).
Build: SQLX_OFFLINE=true cargo check -p ml --lib — clean.
Test: trainers::dqn::trainer::tests::test_reward_function_price_changes
passes. (PPO test_reward_computation pre-existing failure on
HEAD 774d7552a, unrelated — verified via stash+rerun.)
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Smoke smoke-test-gwfn8 on fd24b5383 showed mean(|reward|) hitting
5054 at F0 ep2 despite all known multiplicative modifiers being
structurally bounded (conviction in (0,1) via sigmoid at line 7579,
cf_flip in +/-1, drawdown in [-5*w_dd,0], shaping_scale in [0,1]). The
inflater is somewhere in the pre-composition or modifier chain that
isn't currently visible in HEALTH_DIAG.
Adds 12 per-sample diagnostic buffers + reduction kernel + HEALTH_DIAG
emit for min/mean/max at every checkpoint in the reward chain:
- per-component (r_popart, r_trail, r_micro, r_opp_cost, r_bonus)
- r_weighted (post-composition, pre-modifier)
- post-modifier sequential (post_dd, post_inv, post_churn, post_conv)
- sanity checks (position_abs, conviction)
Implementation:
- 12 new per-sample CudaSlice<f32> buffers in gpu_experience_collector
(alloc_episodes * alloc_timesteps each); zero-init at every (i,t)
in the kernel entry block before any early-return; written at
each checkpoint with NULL-tolerant guards.
- new reward_chain_diag_reduce_kernel.cu: single-block 256-thread
block-tree-reduce over the 12 buffers, three lockstep reductions
(sum/min/max) per buffer in shared memory; outputs 36 floats
(3 stats x 12 buffers) to a 36-slot mapped-pinned scratch buffer
on the trainer; no atomicAdd per feedback_no_atomicadd, pure GPU
compute per feedback_no_cpu_compute_strict, mapped-pinned host
visibility per feedback_no_htod_htoh_only_mapped_pinned.
- trainer: cubin load, mapped-pinned 36-f32 output, set_sp11_reward_
chain_diag_bufs setter, launch_sp11_reward_chain_diag_reduce
launcher, read_sp11_reward_chain_diag host accessor.
- training_loop.rs: wires the 12 collector buffers post-construction
(mirror of the popart-component wire-up); HEALTH_DIAG `reward_
chain_diag` emit added immediately after `reward_split` — launches
reduction kernel, syncs stream, reads the 36 floats.
- build.rs: adds reward_chain_diag_reduce_kernel.cu to the cubin
manifest.
- docs/dqn-wire-up-audit.md: new section documenting the
instrumentation scope, additions, exclusions (no ISV slots, no
state-reset registry entries), and removal plan.
No state-reset registry entry: this is a transient diagnostic, not
persistent state — buffers reset to 0 every step via the kernel
entry-block default writes (same pattern as the other per-sample
diagnostic buffers like trail_triggered_per_sample). No ISV slots
are added: the host reads the mapped-pinned scratch directly to keep
this lightweight and avoid permanent ISV growth.
Will be removed in a follow-up commit once the inflater is identified
and properly fixed per pearl_bounded_modifier_outputs_require_
structural_activation.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Layer C close-out for SP11 B1b smoke-recovery work. Two atomic commits
land the fix:
- b3b4d0278: z-score normalization for mag-ratio canary (Bug 5)
- fd24b5383: launch-order — reward_component_ema before canary (Bug 6)
Audit doc (docs/dqn-wire-up-audit.md) adds the close-out entry covering
all 6 bugs (slot 63 overload, stale rc[] init, cf_flip ordering,
cf-component feedback loop, magnitude-asymmetric ratios, launch-order),
the 6 new ISV variance EMA slots [361..367) producer/consumer wiring,
and validation evidence from the 3 smoke-test workflows
(smoke-test-6wd2c killed → smoke-test-4rbv9 killed → smoke-test-gwfn8
PASSED 15m11s on commit fd24b5383).
3 memory pearls written to user memory (not in repo):
- pearl_zscore_normalization_for_magnitude_asymmetric_signals
- pearl_canary_input_freshness_launch_order
- pearl_controller_amplifies_dominant_magnitude_trap
MEMORY.md index updated under Controllers and signals.
Open follow-up flagged: residual within-fold sharpe degradation
across epochs is a separate triage scoped to post-T10 validation,
not part of B1b smoke-recovery.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
smoke-test-4rbv9 on b3b4d0278 (z-score implementation) showed bit-
identical w_pop=2.000 at ep1 to pre-z-score B1b smoke, proving the
z-score formula was structurally a no-op:
z[c] = mag[c] / fmaxf(sqrtf(0), EPS_DIV) = mag[c] x 1e6
ratio[c] = (mag[c] x 1e6) / (1e6 x sum(mag)) = mag[c] / sum(mag) -> linear
Root cause: launch_reward_component_ema_inplace at line 3707 ran
AFTER launch_sp11_mag_ratio_compute at line 3465. So ISV[64..68]
and ISV[362..366] held sentinel-0 values at ep1's canary read
(ep0 had no segment_complete fires). z[c]=0 for c=1..5 -> popart
ratio collapsed to 1.0 -> controller saturated.
This was structurally the same bug that motivated adding
launch_sp11_popart_component_ema at line 3441 (B1b follow-up).
That fix-up addressed popart but left cf/trail/micro/opp_cost/
bonus stale.
Moved launch_reward_component_ema_inplace from line 3707 to before
launch_sp11_popart_component_ema. Other launches at the original
site (trade_attempt_rate_ema, plan_threshold_update, etc.) stay
where they were — different consumers, different timing constraints.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Linear magnitude ratios in reward_component_mag_ratio_compute_kernel
amplified popart's intrinsic O(100) magnitude over the other 5
components' O(0.1-2) magnitudes, causing controller to saturate
w_pop toward MAX_WEIGHT regardless of actual signal quality.
Replaced with z-score: z[c] = mag[c] / max(sqrt(var[c]), EPS_DIV).
6 new ISV slots [361..367) for per-component variance EMAs computed
via Welford's online algorithm in extended popart_component_ema_kernel
and reward_component_ema_kernel.
Atomic per feedback_no_partial_refactor: slot allocation +
state-reset registry + 2 producer kernels + canary signature +
launcher Pearls A+D + tests.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
smoke-test-6wd2c on commit 61b2fa962 (B1b + 4 bug fix-ups) revealed a 5th
pathology not covered by the previous fix-ups: the mag-ratio canary's
linear-magnitude-ratio formula amplifies whichever component is
intrinsically largest, regardless of whether that's a useful signal.
Popart (trade P&L on segment_complete) is O(100) per fire while the
other 5 components (cf/trail/micro/opp_cost/bonus) are O(0.1-2). Even
with the slot 360 fix preventing total-reward contamination, popart's
intrinsic magnitude makes popart_mag / Σ ≈ 0.93. The controller blend
'winner_weight = ratio' then amplifies popart further. Smoke trajectory:
w_pop=2.0 → 2.44 → 2.57, curiosity_b=30 → 120 → 199, sharpe_ema=10.7 →
2.4 → 0.75 (cascading collapse).
Resolution: z-score normalization. Each component's magnitude divided
by its own running standard deviation before computing the ratio:
popart_z = popart_mag_ema / max(sqrt(popart_var_ema), EPS_DIV) ≈ O(1)
cf_z = cf_mag_ema / max(sqrt(cf_var_ema), EPS_DIV) ≈ O(1)
...
ratio[c] = component_z / Σ component_z ≈ ~1/6 each when stable
Allocates 6 new ISV slots [361..367) for per-component variance EMAs.
Producers: extend popart_component_ema_kernel + reward_component_ema_kernel
to also emit variance via Welford's online algorithm (single-pass).
SP5_SLOT_END = 367, ISV_TOTAL_DIM = 367.
Carries forward main's slot 360 amendment (commit 52c0b7521 on main)
which the sp11 branch was missing, plus this z-score amendment.
Per-component gradient ratios (the original spec intent) don't fix this
either — in DQN there's no per-component gradient pathway; grad norm
scales with current_weight × magnitude, so it's the same bias. Z-score
normalization is the standard scale-invariant measure of significance
and matches what the SP11 controller is trying to express.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Deep audit on b435d25be found a third post-B1b bug in same class as
the slot 63 overload: experience_env_step writes cf_reward_weighted
(= w_cf × cf_reward, post-controller-weight) to reward_components_
per_sample[+1] at line ~3696. SP4's reward_component_ema_kernel EMAs
this into ISV slot 64 (REWARD_CF_EMA_INDEX) which the SP11 mag-ratio
canary reads as ratio[1] = slot[64] / Σ.
Self-reinforcing loop:
high w_cf → high cf_reward_weighted → high slot 64 →
high ratio[1] → controller raises w_cf → tighter loop
Until mean=1 normalization saturates other components to floor.
The other 5 component slots correctly write RAW pre-weight values:
- rc[+0] = total_reward (intentional, PopArt input — post-composition)
- rc[+2..+5] = r_trail / r_micro / r_opp_cost / r_bonus (all raw, pre-Σ)
- Slot 360 (popart-component) fed by `popart_component_per_sample`
which receives `r_popart` raw
Only rc[+1] was wrong. Fixed: write raw cf_reward to rc[+1] so the
canary tracks intrinsic cf magnitude. The replay-buffer cf-tuple
reward (out_rewards[cf_off]) still uses cf_reward_weighted — that's
correct, loss kernels train on controller-weighted signal.
This was the third bug in a class — pre-SP11 invariants exposed by
post-decomposition semantic. Trio: (1) slot 63 overload (5e16b67ca),
(2) stale rc[] init + cf_flip ordering (b435d25be), (3) cf-component
feedback loop (this commit).
cargo check + build clean; 6/6 SP11 GPU tests + 14/14 contract tests
still pass. After this fix-up, all known SP11 reward-system bugs
identified by deep audit are resolved. L40S smoke validates empirical
sharpe recovery.
Bug-hunt review on 5e16b67ca found two real bugs in experience_env_step
post-B1b structural refactor — same class as the slot 63 overload
(pre-SP11 invariants exposed by post-decomposition semantic).
Bug 1 (Critical): reward_components_per_sample[+1..+5] not zero-init at
per-bar entry. Slots written only on execution paths that fire (rc[2]
inside segment_complete trail-trigger, rc[3] inside positioned-non-
complete, rc[4] inside flat-with-features, rc[5] inside multiple
bonus paths, rc[1] inside CF block). Non-firing paths leave stale
values from previous batch's same out_off, contaminating SP4
reward_component_ema_kernel -> slots 64..68 -> SP11 mag-ratio canary
-> wrong controller weights.
Fix: zero-init rc[0..6) at the same location where r_<component>
locals are zero-initialized, so locals and buffer reset together.
Eliminates the entire stale-rc class of bugs.
Bug 2 (Important): cf_flip applied BEFORE inventory/churn/conviction
modifiers. Spec section 3.4.4 requires cf_flip LAST so subtractive
penalties operate on the right sign and multiplicative scales
attenuate before direction is flipped. Pre-fix on flipped samples,
inventory/churn ADDED to negated reward (penalty became bonus) and
conviction scaled the wrong-signed value. Asymmetric gradient signals
between flipped/non-flipped -> contaminated SP11 mag-ratio canary at
the cf axis.
Fix: move cf_flip to after conviction_scale, making it the last
modifier before out_rewards write. CF-block contract preserved
(reward at cf_off is still the post-flipped final value; CF block
already does `do_flip ? -reward : reward` to recover unflipped base).
Stale doc-strings: state_reset_registry.rs (2 sites) and
training_loop.rs (1 site) referenced pre-fix-up SP5_PRODUCER_COUNT=186
/ wiener_buffer=771. Updated to formula form
`(71 + SP5_PRODUCER_COUNT) × 3` so they don't drift on the next SP;
current value (post-B1b fix-up slot 360) is SP5_PRODUCER_COUNT=187,
buffer=774 floats.
dqn-wire-up-audit.md: appended SP11 B1b bug-hunt fix-up section
documenting both bugs, the structural fix, the doc-string sweep,
and verification (cargo check/build/tests).
cargo check + build clean; 6/6 SP11 GPU tests + 10/10 sp5_isv_slots
contract tests + 4/4 state_reset_registry tests still pass. The two
bugs were latent in the local smoke (sharpe 3 vs B1a 30 likely
traceable to either or both); structural fix sound. L40S smoke on
this commit will validate the empirical recovery.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per spec §4 amendment at 52c0b7521 on main: B1b smoke surfaced that
SP11 mag-ratio canary was reading slot 63 (REWARD_POPART_EMA_INDEX)
which is overloaded — pre-SP11 PopArt's normalization input (total
reward mag EMA) was the same value as popart-component magnitude
because composition was inline accumulation. B1b decomposition exposed
the overload; controller emitted w_pop ≈ 2.0 based on contaminated
ratio → 10× sharpe drop in smoke.
Resolution:
- Allocate ISV slot 360 = POPART_COMPONENT_MAG_EMA_INDEX
- Add popart_component_per_sample mapped-pinned buffer + write site
in experience_env_step at the r_popart assignment
- New popart_component_ema_kernel.cu writes slot 360 (single-block
block-tree-reduce per feedback_no_atomicadd)
- mag-ratio canary kernel signature changes from single
popart_ema_base_slot to (popart_specific_slot, cf_others_base_slot)
pair so it reads non-contiguous slot 360 + slots 64..68
- Reset registry: sp11_popart_component_mag_ema entry + dispatch arm
- Slot 63 (PopArt's input) UNCHANGED — pre-SP11 invariant preserved
ISV total: 360 → 361. SP5_SLOT_END = 361. SP5_PRODUCER_COUNT = 187.
cargo check + build clean; SP11 GPU oracle tests pass (6/6 including
updated mag_ratio test with 2 slot-index args); sp5_isv_slots layout
tests pass (10/10 with 185 unique slots / 187 linear span); state
reset registry tests pass (4/4 with new sp11_popart_component_mag_ema
entry + dispatch arm). Local multi_fold_convergence smoke gated on
data volume (175k bars on local fxcache vs 10-month walk-forward
requirement); validation deferred to L40S Argo run on PVC data per
the spec's pass criterion.
Per spec §3.5.3 amended at 7ddaf9c51 on main: experience_env_step
reward composition decomposed from 8+ inline accumulation sites into
explicit per-component locals (r_popart, r_cf, r_trail, r_micro,
r_opp_cost, r_bonus), then composed as Σ w_i × r_i with controller
weights from ISV[340..346).
Trail reward extraction (§3.5.4): trail-fire P&L now flows through
r_trail (forced-exit signal) instead of r_popart (voluntary-exit
signal). REWARD_TRAIL_WEIGHT_INDEX has real signal — controller can
weight forced-exit vs voluntary-exit P&L differently. rc[2] (the
prior structural-placeholder slot) now carries trail magnitude.
Universal post-composition modifiers (§3.4.4): drawdown / capital-
floor / inventory / churn / conviction-scale / cf-flip apply AFTER
the weighted Σ, unweighted. They are risk constraints and structural
operators, NOT learning components — agent cannot weigh them away.
Mean=1 normalization (B0, §3.4.3): weights normalize to mean=1 so per-
bar `w_active × r_active` ≈ pre-SP11 absolute scale on average.
Sentinel-defense: experience_env_step runs at start of epoch, SP11
controller runs at end (training_loop.rs ~3475). At fold 0 epoch 0
step 0 the controller has not yet emitted, so ISV[340..346) hold
sentinel 0. Defense: fmaxf(w_raw, 0.01) — same Invariant-1 hard floor
the controller enforces post-renorm. Cold-start scale = 1% of
pre-SP11; Pearl A bootstrap on first emit replaces sentinel.
cf_reward path: out_rewards[cf_off] now writes controller-weighted
cf reward (w_cf × r_cf with sentinel-defense). Loss-kernel
cf_weight=0.3f at mse:318/c51:789 (structural Q-blend, NOT reward
weight) UNTOUCHED per §3.5 amendment.
Mutual exclusivity preserved (popart / trail / micro / opp_cost):
exactly one path fires per bar; others stay 0. The cascade scalar
`reward` mirrors per-component locals so C.4/D.4b bonus blocks that
read in-progress trade reward (Q-cap pattern from
pearl_one_unbounded_signal_per_reward) keep bit-identical compounding.
After cascade, `reward` is overwritten with r_weighted; post-
composition modifiers operate on r_weighted as before.
This is the production-flip commit. Trainer is now on the SP11
controller end-to-end (modulo replay-time curiosity which lands in
B1c after Layer C audit per §3.5.5).
Verification:
- cargo check + build clean (1m32s release).
- 6/6 SP11 GPU oracle tests pass (none exercise env_step directly).
- 14/14 contract tests pass (sp5_isv_slots=10, state_reset_registry=4).
- Local smoke (RTX 3050 Ti, 20-epoch magnitude_distribution) verifies
HEALTH_DIAG sp11_reward weights drift epoch-over-epoch:
epoch 0: w_pop=1.000 w_cf=1.000 w_tr=1.000 ... (uniform sentinel-defense floor → mean=1)
epoch 3: w_pop=1.991 w_cf=2.036 w_tr=0.493 ... (controller redistributes)
epoch 9: w_pop=1.823 w_cf=1.887 w_tr=0.572 ... (mean ≈ 1.0 preserved, Σ ≈ 6)
EVAL_DIST bit-identical to B1a baseline (eq=0.803 eh=0.197 ef=0.000)
— pre-existing magnitude eval-collapse pathology
(project_magnitude_eval_collapse_kelly_capped) unchanged by B1b.
Audit doc updated (Invariant 7): docs/isv-slots.md SP11 section now
reflects Layer B status with B0/B1a/B1b/B1c rollout timeline.
Per spec §3.4.3 amended at 7ddaf9c51 on main: A2's Σweights=1
renormalization caused 6× reward-magnitude collapse on mutually-
exclusive components in experience_env_step (popart/micro/opp_cost
paths fire at most one per bar; weight 1/6 per fired path averages
1/6 of pre-SP11 reward magnitude).
Amended: weights normalize to mean(weights) = 1 (i.e., Σ = N = 6).
Each weight in [WEIGHT_HARD_FLOOR=0.01, MAX_WEIGHT=3.0]. Default
uniform = 1.0 each. Preserves pre-SP11 absolute scale on average.
Code change: reward_subsystem_controller_kernel.cu renormalization
step changes from `weights[c] = blends[c] / blend_sum` to
`weights[c] = min(MAX_WEIGHT, blends[c] × N / blend_sum)`. Anchors
N_COMPONENTS=6.0f and MAX_WEIGHT=3.0f added to the Invariant-1 const
float block at the top of the kernel.
3 A2 controller unit-test assertions updated:
- z_score_at_zero: weight_sum 1.0→6.0; per-component 1/6→1.0
- weights_renormalize_after_floor: assertion strengthened to
weight_sum ≤ N (cap binds in this pathological test where pre-cap
dominant weight ≈ 4.18 > MAX_WEIGHT=3.0); added per-component
≤ MAX_WEIGHT envelope check; added explicit cap-binding assertion
on dominant weight.
- saboteur_post_clamp_holds_min: weight assertions unaffected (this
test asserts only on s[6]/s[7], saboteur+curiosity are independent
of mean-vs-Σ choice).
Audit doc updated: docs/dqn-wire-up-audit.md gets a new
"SP11 B0 — controller renorm Σ=1 → mean=1 (2026-05-04)" section.
cargo check + release build clean. 6/6 SP11 GPU oracle tests pass on
RTX 3050 Ti. sp5_isv_slots (10/10) + state_reset_registry (4/4)
contract tests still pass.
Pre-requisite for B1b structural reward-composition refactor (§3.5.3)
which depends on the mean=1 semantic.
Code-quality review on 25eba79ad found two Important issues:
- launch_sp11_novelty_simhash_lookup + launch_sp11_novelty_simhash_update
were defined with #[allow(dead_code)] since B1 hasn't wired them yet —
feedback_no_hiding violation. Deleted both functions; kept the kernel
handle fields + cubin loads + MappedF32Buffer storage. B1 will inline
the launches at the actual call site (matching A0's deferral pattern
for the novelty-hash registry entry). Field doc-refs at the kernel
declarations updated so no stale name reference remains.
- Seed XOR-fold comment named what ("fold the high + low halves") but
not why ("to preserve entropy; truncation would silently discard
upper 32 bits"). Extended comment so a future reader switching to
truncation gets the warning.
Verified: cargo check clean (18 unrelated warnings, zero errors); 6/6
SP11 GPU oracle tests build under --features cuda (GPU-gated, ignored
on CPU runners); 10/10 sp5_isv_slots + 4/4 state_reset_registry
contract tests pass; allow(dead_code) count is exactly the 12
pre-existing attributes (was 14 before this fix; A2 added 2).
reward_subsystem_controller_kernel: 5 canaries → 10 outputs, true Z-score
(delta_ema/sqrt(var_ema)), sigmoid blending, weight renormalization to Σ=1,
saboteur post-clamp, curiosity permanent floor (0.2 × bound). Pearls A+D
chained on outputs per spec §3.4.1.
novelty_simhash_kernel: 42×16 random projection → 16-bit SimHash code,
1M-slot bucket count table for novelty signal `1/sqrt(1+count)`. Race-
tolerated update per feedback_no_atomicadd (under-counts bias novelty
UPWARD — safe direction).
novelty_simhash_proj_init_kernel: Philox-seeded GPU init for the
projection matrix (CPU is read-only per feedback_no_cpu_forwards).
HEALTH_DIAG `sp11_reward` line emits 10 outputs + improvement_z each
epoch. Reset registry: novelty hash table reset arm wired (closes the
A0 deferral); projection matrix is frozen at trainer init for run
lifetime, not reset.
All 20 SP11 slots populate every step. No consumer reads them yet —
training behavior unchanged from A1. 3 new GPU oracle tests pass on
RTX 3050 Ti (controller midpoint, weight renorm, saboteur clamp).
Spec: docs/superpowers/specs/2026-05-04-sp11-reward-as-controlled-subsystem.md §3.4 §3.5.2
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Code-quality review on 91b48bc7a found two issues in
sp11_producer_unit_tests.rs:
- :281 `let _ = REWARD_COMPONENT_MAG_RATIO_BASE` violated
feedback_no_hiding (silent dead-code suppression). Removed the
suppression and the unused import. The ISV landing slot for the
mag-ratio producer is exercised by Pearls A+D unit tests, not here.
- :190 saboteur_engagement test passed shared_mem_bytes=1024 for a
kernel that uses __shared__ (static, allocated at compile time),
not extern __shared__ (dynamic). Set to 0 matching the other two
test launches.
cargo check + all 3 GPU oracle tests still pass.
Adds val_sharpe_delta + saboteur_engagement + reward_component_mag_ratio
GPU producers for the SP11 reward-as-controlled-subsystem chain. Each is
a single-block producer chained with apply_pearls_ad_kernel for Pearls
A+D smoothing per pearl_first_observation_bootstrap.md +
pearl_wiener_optimal_adaptive_alpha.md. All three write to slots in
[350..360) which no consumer reads yet — Layer A is additive; consumer
migration lands atomically in Layer B.
A1.1 — val_sharpe_delta_compute_kernel.cu
Two-pass: writes raw delta + (delta - prev_delta_ema)^2 to scratch.
Chained Pearls A+D (n_slots=2) → ISV[VAL_SHARPE_DELTA_EMA_INDEX=350,
VAL_SHARPE_VAR_EMA_INDEX=351]. Host writes val_sharpe to mapped-pinned
history[1]; rotation handled in training_loop.rs at val emit boundary
(a literal already-computed value — no host-side compute, no htod_copy).
A1.2 — saboteur_engagement_compute_kernel.cu
Per-bar |Δreward| > 0.01 × ISV[PNL_REWARD_MAGNITUDE_EMA_INDEX] check
with block tree-reduce (no atomicAdd per feedback_no_atomicadd). The
per-bar Δreward signal is produced by experience_env_step's saboteur
perturbation site as `traded × |reward| × max(|eff_spread − 1|,
|eff_slip − 1|)` — a structural proxy for the cost-differential the
saboteur imposed on bars where the model traded. Single kernel-side
emit (no parallel reward computation), per spec §3.3.1.
Chained Pearls A+D → ISV[SABOTEUR_ENGAGEMENT_RATE_INDEX=358].
A1.3 — reward_component_mag_ratio_compute_kernel.cu
Reads ISV[REWARD_POPART_EMA_INDEX..+6) (the SP4 reward-component
magnitude EMAs), normalises to ratios, and mirrors popart magnitude
into scratch[6] as a side-output. ONE non-pointer parameter
(popart_ema_base_slot) — no _unused param per feedback_no_stubs.
Two chained Pearls A+D launches:
n_slots=6 → ISV[REWARD_COMPONENT_MAG_RATIO_BASE..+6)
n_slots=1 → ISV[PNL_REWARD_MAGNITUDE_EMA_INDEX=359]
(slots non-contiguous: 352..358 then 359.)
Wire-up (per feedback_wire_everything_up):
- 3 cubin entries appended to crates/ml/build.rs
- 3 kernel handles + val_sharpe_history_pinned (MappedF32Buffer[2]) +
saboteur_delta_reward dev-ptr cache fields on GpuDqnTrainer
- 3 launchers (launch_sp11_*) + 1 setter (set_sp11_saboteur_delta_reward_buf)
- saboteur_delta_reward_per_sample buffer field on GpuExperienceCollector
- experience_env_step kernel signature extended with the new buffer arg;
every call site in the same commit per feedback_no_partial_refactor
- training_loop.rs init wires collector→trainer setter; val emit boundary
invokes launch_sp11_val_sharpe_delta_compute; per-epoch metrics block
invokes launch_sp11_mag_ratio_compute then
launch_sp11_saboteur_engagement_compute (mag_ratio first so the
signal-relative threshold base is populated before the saboteur reader)
- SP5_SCRATCH_TOTAL grown 266 → 276 (10 new scratch slots: 2+1+7)
- docs/isv-slots.md SP11 section updated to reflect A1 producers
3 GPU oracle tests in crates/ml/tests/sp11_producer_unit_tests.rs
pass on RTX 3050 Ti via MappedF32Buffer fixtures (zero htod_copy /
dtoh_sync_copy / alloc_zeros — feedback_no_htod_htoh_only_mapped_pinned
compliant).
Note on Step 8a path: the plan offered two routes for the saboteur
Δreward producer — in-kernel diff emission OR a small dedicated
reader-of-existing-buffers. The existing reward path emits ONE reward
(not both with/without), so the dedicated-reader alternative was
infeasible. The in-kernel emission landed as a small write site at the
END of experience_env_step (after total_reward_per_sample is finalised),
threading saboteur_eff_spread/saboteur_eff_slip from the perturbation
site forward to the END via stack vars. Single new kernel parameter,
single new GPU-only buffer, single existing call site updated.
Spec: docs/superpowers/specs/2026-05-04-sp11-reward-as-controlled-subsystem.md §5
Plan: docs/superpowers/plans/2026-05-04-sp11-reward-as-controlled-subsystem.md (Task A1)
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Code-quality review on bf3a32d63 found two stale references that need
SP11 numbers:
- training_loop.rs:6672 + state_reset_registry.rs:891 — sp5_wiener_state
comments referenced the post-SP4/post-SP8 buffer sizes (543, 681);
post-SP11 is (71 + SP5_PRODUCER_COUNT) × 3 = 771 floats. Replaced the
literal sizes with formula form citing SP5_PRODUCER_COUNT directly so
this drifts less in the future.
- docs/isv-slots.md header — "Current ISV_TOTAL_DIM" said 171 (post-SP4
Task A1) while actual is 360. Updated header; SP11 section already
appended at the end of the file.
No logic changes. Cargo check + sp5_isv_slots / state_reset_registry
tests still pass.
Pure infrastructure. No producer kernels, no consumer reads. Existing
training paths trace identically because no consumer reads slots [340..360)
yet. Layout-fingerprint bumped to ISV_TOTAL_DIM=360.
Spec: docs/superpowers/specs/2026-05-04-sp11-reward-as-controlled-subsystem.md
feedback_no_htod_htoh_only_mapped_pinned (tests not exempt):
- All test fixtures converted from htod_copy/dtoh_sync_copy to
MappedF32Buffer with host_slice / host_slice_mut access.
- Novelty hash buffer + projection matrix changed from
cudarc::CudaSlice<f32> + alloc_zeros to MappedF32Buffer.
feedback_no_cpu_forwards (CPU is read-only):
- Projection matrix initialization changed from host-side StdRng +
host_slice_mut writes to a one-shot GPU init kernel
(novelty_simhash_proj_init_kernel) using Philox seeded from
config.seed. No host RNG, no host writes.
feedback_no_cpu_compute_strict (saboteur multiplication):
- B1 step 5 reverted from Rust-side `read_isv_slot * scale` to
GPU-side: pass base scale + ISV pointer + slot index to the
saboteur perturbation kernel; multiplication happens on-device.
feedback_trust_code_not_docs (grad-ratio terminology):
- Spec §3.3.1 "per-component grad EMA" was wrong — SP4 grad-balancer
is per-branch (4 slots), not per-reward-component. Renamed:
reward_component_grad_ratio_compute_kernel
→ reward_component_mag_ratio_compute_kernel
REWARD_COMPONENT_GRAD_RATIO_BASE
→ REWARD_COMPONENT_MAG_RATIO_BASE
Source: existing REWARD_POPART_EMA_INDEX..+6 (per-component reward
magnitude EMAs from SP4 reward_component_ema_kernel). Semantic
equivalent for the controller's exploit/diversify blend.
feedback_no_stubs (dead parameter):
- Removed `eps_div_idx_unused` from mag_ratio kernel signature.
Saboteur engagement: missing producer specified
- Spec §3.3.1's two-reward-arrays formulation replaced with single
`saboteur_delta_reward_buf` produced by the saboteur perturbation
kernel itself (single reward computation, diff emitted as side
output). Engagement kernel signature simplified to one input array.
PNL_REWARD_MAGNITUDE_EMA_INDEX (slot 359): producer wired
- mag-ratio kernel mirrors `isv[REWARD_POPART_EMA_INDEX]` to
scratch_out[6]; chained apply_pearls_ad targets slot 359.
Replay sample kernel location specified
- graph_utility_kernels.cu:71 (gather_f32_scalar). New sibling kernel
`gather_replay_reward_with_curiosity` defined; replaces the existing
scalar gather (no legacy alias per feedback_no_legacy_aliases).
novelty_simhash_lookup runs before, novelty_simhash_update after.
A2 controller test placeholders → full GPU oracle assertions
- Three controller tests (z=0 midpoint, weight renorm, saboteur clamp)
have full mapped-pinned fixtures with assertions on weight sum,
individual values, post-clamp bounds.
Plan now passes:
- feedback_no_htod_htoh_only_mapped_pinned (tests + production)
- feedback_no_cpu_forwards (CPU never writes/computes for GPU)
- feedback_no_cpu_compute_strict (all multiplications GPU-side)
- feedback_no_atomicadd (race-tolerated non-atomic, safety documented)
- feedback_no_partial_refactor (Layer B atomic; saboteur kernel sig
change touches all callers in the same commit)
- feedback_no_stubs (no dead parameters)
- feedback_trust_code_not_docs (corrected spec terminology)
- feedback_wire_everything_up (every new field has init + producer)
- feedback_no_legacy_aliases (old gather_f32_scalar replaced, deleted)
1447 lines, +268 from previous version.
User correction: curiosity is the *fix* for the ep1-peak overfitting
pathology, not a hazard to defend against. Reframed §7 from "Risks"
to "Design notes" — curiosity bound is a signal-relative scale, not
a defensive cap.
Fix the contradiction this exposed in the formula: previous
`curiosity_pressure = stagnant_or_worse * curiosity_bound` went to
zero when improving, which would cancel the always-on exploration
the §7 narrative now relies on. Replace with permanent-floor pattern
per pearl_blend_formulas_must_have_permanent_floor:
curiosity_floor = 0.2 * curiosity_bound (CURIOSITY_PERMANENT_FRACTION)
curiosity_dynamic = stagnant_or_worse * curiosity_bound
curiosity_pressure = max(curiosity_dynamic, curiosity_floor)
Now curiosity is always ≥ 20% of bound (anti-overfitting baseline)
and rises toward the bound when stagnant (stagnation breaker).
Updated unit-test guidance to assert pressure > 0 even at z=+10.
CURIOSITY_PERMANENT_FRACTION=0.2 added to Invariant-1 fraction list.
Saboteur-rising-with-improvement reframed as adversarial-load feature
rather than over-stress risk.
Brainstorm spec for SP11. Resolves the policy-stagnation pathology
surfaced in T10 train-multi-seed-xkjkb seed-0 ep0-14: model finds a
stable fixed point at ep1 (peak val sharpe 80.61), then OVERFITS to
it across remaining epochs (decline 80.61 → 70.58). Q-values grow but
val performance declines because reward function has no improvement
pressure.
Architecture (every input ISV-driven):
- Z-score-driven adaptation (no hardcoded "improving" threshold):
improvement_z = val_sharpe_delta_ema / max(val_sharpe_std_ema, EPS)
- 10 ISV outputs: 6 component weights + curiosity_pressure +
saboteur_intensity_mult + adaptive weight_floor + curiosity_bound
- 5 ISV canaries: val_sharpe_delta + val_sharpe_std (Z-score noise
estimate) + 6 per-component grad ratios + saboteur engagement +
PnL magnitude EMA (signal-relative curiosity bound)
- 4 new producer kernels (controller + 3 canary computers)
- Audit + migrate hardcoded cf_weight=0.3 in mse_loss_kernel.cu:318
and c51_loss_kernel.cu:789, plus other shaping multipliers
- NEW reward dimension: curiosity bonus, bounded by PnL magnitude
Per pearl_controller_anchors_isv_driven: every threshold replaced with
sigmoid(z) — no constants encode "what counts as improving". Per
pearl_blend_formulas_must_have_permanent_floor: every weight has
adaptive floor preventing zero-out. Per pearl_engagement_rate_self_
correction: saboteur intensity self-corrects via engagement rate canary.
Per pearl_cold_start_exit_signal_or: improvement signal OR'd from
multiple canaries so single-signal-failure doesn't stall controller.
New pearl authored alongside spec: pearl_reward_as_controlled_subsystem
— meta-principle that every reward path degree of freedom is a unified
controller output. Subsumes controller-anchor pearl at the reward layer.
Scope: 20 ISV slots, 4 producer kernels, audit + migration of
hardcoded reward shaping constants, 1 atomic commit.
~1300-1700 LOC; ~2.5-3 hours subagent work.
Success metric: val_sharpe[ep20] > val_sharpe[ep1] (Fix 33-38 baseline
peaked at ep1; SP11 should shift peak later as model continues
learning).
`total_return` from financials.rs:80-94 is log-space cumulative growth
across every per-bar step_return. With ~4M step_returns in a fold-
convergence run, even sub-bps positive bars compound to absurd
magnitudes (observed: 1.93e37%) when displayed as `{:+.2}%`. Math is
correct; display needs scientific notation.
Surfaced in T10 train-multi-seed-xkjkb seed-0 ep3 epoch summary while
SP10 chain validates structural fixes. Cosmetic-only change; no
training-path impact. Audit doc updated with Cosmetic 38.1 entry.
`total_return` from financials.rs:80-94 is the log-space cumulative growth
across every per-bar step_return: `exp(sum(ln(max(1+r, 1e-10)))) - 1`.
With ~4M per-bar step returns in a fold-convergence run, even sub-bps
positive bars compound to absurd magnitudes (observed: 1.93e37%) when
displayed as `{:+.2}%`. The math is correct; the display is broken.
Switching to `{:+.3e}%` shows the same magnitude compactly across the
full dynamic range without hiding the number. Comment added explaining
the underlying compute's HFT-inappropriate semantic so a future reader
knows the value is correct-but-meaningless rather than buggy.
Surfaced in T10 train-multi-seed-xkjkb seed-0 ep3 epoch summary while
the SP10 chain validates structural fixes. Cosmetic-only change; no
training-path impact.
T10 train-multi-seed-khr7c (commit 8a25b330f, post-Fix-37) showed val-Flat-
collapse persisting at the eval-side selector despite all SP9 controller
fixes — dir_entropy=0 / trade_count=1 in 214,654 bars because
experience_action_select branched on eval_mode and used argmax(E[Q]) at
eval. With Hold's E[Q] ≈ 0 and directional E[Q] = ε (small edge minus tx
costs), argmax wins Hold deterministically every bar.
Per pearl_thompson_for_distributional_action_selection (amended): the
rollout SELECTOR is unconditional Thompson at all times (training AND
eval); argmax is reserved for the Bellman TARGET Q computation only
(DDQN target). Per pearl_controller_anchors_isv_driven: the temperature
on the Thompson sample is ISV-driven from the SP9 intent_eval_divergence
canary. Per pearl_blend_formulas_must_have_permanent_floor: MIN_TEMP=0.5
is the permanent-stochasticity floor — the eval selector is NEVER fully
deterministic.
Atomic commit per feedback_no_partial_refactor:
* 1 new ISV slot @ [339..340) (EVAL_THOMPSON_TEMP_INDEX); ISV_TOTAL_DIM
339 → 340; SP5_PRODUCER_COUNT 165 → 166; layout fingerprint updated
* 1 new scratch slot @ [265..266) (SCRATCH_SP10_THOMPSON_TEMP);
SP5_SCRATCH_TOTAL 265 → 266
* intent_eval_divergence_compute_kernel.cu extended with 2 params
(divergence_target_isv_index, scratch_temp_idx) + new compute branch
temp = clamp(divergence/div_target, 0.5, 2.0); existing scratch_idx
renamed scratch_div_idx for semantic clarity
* experience_action_select: if (eval_mode) { argmax(E[Q]) } DELETED;
unconditional temperature-blended Thompson installed:
q_eff[d] = E[Q][d] + temp · (q_sample[d] − E[Q][d]); defensive clamp
to [0.5, 2.0] for cold-start before producer first observation. Other
branches (mag/ord/urg) keep their existing eps-greedy/Boltzmann logic
per pearl_thompson §3 exemption
* state_layout.cuh: new ISV_EVAL_THOMPSON_TEMP_IDX 339 define
* gpu_dqn_trainer.rs::launch_intent_eval_divergence_compute extended
with 2 kernel args + second apply_pearls_ad_kernel chain to smooth
the temperature into ISV[339]
* New FoldReset entry sp10_eval_thompson_temp + dispatch arm in
reset_named_state writing sentinel 0; Pearl A's first-observation
replacement fires on the new fold's first producer launch
* Test renamed test_eval_action_select_eval_argmax_picks_best →
test_eval_action_select_thompson_picks_proportionally; ISV buffer
setup with EVAL_THOMPSON_TEMP_INDEX=1.0 (pure Thompson); assertions
updated to ≥ 70% best-direction wins (was ≥ 99% under deterministic
argmax) and < 100% (selector is sampling)
* Pearl pearl_thompson_for_distributional_action_selection §4 amended;
MEMORY.md index entry updated
* Audit doc Fix 38 entry
Verification:
* SQLX_OFFLINE=true cargo check -p ml — clean (only pre-existing 18
warnings; no new errors or warnings introduced by Fix 38)
* SQLX_OFFLINE=true cargo test -p ml --lib state_reset — 4/4 pass
including contract test every_fold_and_soft_reset_entry_has_dispatch_arm
* SQLX_OFFLINE=true cargo test -p ml --lib sp5_isv_slots — 10/10 pass
including new sp10_thompson_temp_slot_above_sp9_block
* SQLX_OFFLINE=true cargo test -p ml --lib test_eval_action_select_thompson —
pass (RTX 3050 Ti local; τ=1.0 with clear Q gap → P(Long) ≥ 0.70 < 1.0)
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Brainstorm spec for SP10. Resolves the val-Flat-collapse pathology that
persisted through Fix 33-37: the eval-time argmax in experience_action_
select picks Hold deterministically every val bar (dir_entropy=0,
trade_count=1 in 214k bars) regardless of controller state.
Architecture:
- Delete `if (eval_mode)` argmax branch in direction-selector kernel
- Use temperature-blended Thompson: q_eff = E[Q] + τ × (Thompson - E[Q])
- τ = clamp(intent_eval_divergence / divergence_target, 0.5, 2.0)
- τ self-corrects: collapse → high τ; healthy → low τ; permanent 0.5 floor
- Reuses SP9's intent_eval_divergence_compute_kernel (extended, not new)
Per pearl_controller_anchors_isv_driven: τ is ISV-driven, no
hardcoded constants beyond Invariant 1 numerical anchors (clamp range).
Per pearl_blend_formulas_must_have_permanent_floor: MIN_TEMP=0.5 ensures
eval ALWAYS has stochasticity.
The pearl_thompson_for_distributional_action_selection was about Bellman
TARGET argmax (selector/target symmetry). It does NOT prohibit Thompson
at the rollout selector. SP10 amends the pearl to clarify.
Scope: 1 ISV slot, kernel modification (no new kernel — extend SP9's
producer), 1 consumer kernel rewrite, test update, pearl amendment,
audit doc Fix 38. ~300-500 LOC, single atomic commit per
feedback_no_partial_refactor.
Two targeted fixes to the SP9 Kelly cold-start warmup floor (Fix 37,
commit `48a8b9ee7`) for quirks observed in smoke-test-wrwkz on a
5-epoch L40S smoke. Atomic per `feedback_no_partial_refactor`.
**Quirk 1 — EMA target saturation:** ep1 HEALTH_DIAG showed
`stat_count_tgt=419 div_tgt=105002 temp_tgt=1.0 conf [stat=1.000
bhv=0.333 tmp=1.000] floor=0.0`. Pearl A sentinel-bootstrap on the
3 EMA target updaters (`kelly_*_target_ema_kernel.cu`) caused the
first observation to replace the sentinel directly, so
`target = first_obs`, `current/target = 1.0`, `confidence = 1.0`,
`floor = base × (1 − 1.0) = 0` — defeating the cold-start mechanism.
The principled fix per `feedback_isv_for_adaptive_bounds.md`:
"sufficient" is defined in absolute terms via Invariant-1 numerical
anchors. The 3 target ISV slots become constructor-written constants
(written ONCE in trainer constructor, identical pattern to existing
`config.cql_alpha → ISV[CQL_ALPHA_INDEX=48]` from Plan 1 Task 12).
- ISV[KELLY_SAMPLE_COUNT_TARGET_INDEX=333] = 100.0 (trades)
- ISV[KELLY_DIVERGENCE_TARGET_INDEX=334] = 2.0 (ratio)
- ISV[KELLY_TEMPORAL_TARGET_INDEX=335] = 5.0 (epochs)
**Quirk 2 — divergence ratio explosion:** ep1 showed
`divergence=70005`. Algebraically `intent_f / max(eval_f, 1e-6) =
0.07 / 1e-6 = 70 000` because `ISV[EVAL_DIST_F_INDEX=338]` was still
at Pearl A sentinel 0 before the first val window populated it.
Algebraically the warmup-floor kernel still derived
`behavioral_conf = 0` (correct semantic), but the synthetic 70 000
in HEALTH_DIAG masked the real signal flow.
Fix: explicit sentinel detection in
`intent_eval_divergence_compute_kernel.cu`:
if (eval_f < SENTINEL_THRESHOLD=1e-5)
divergence = SENTINEL_DIVERGENCE=1e6 // forces behavioral_conf=0
else
divergence = intent_f / eval_f
Same `behavioral_conf = 0` outcome but with explicit provenance —
HEALTH_DIAG now reads `divergence=1e6` until eval_f matures.
**Atomic structure:**
- DELETE 3 EMA target updater kernels (`.cu` files)
- DELETE 3 entries from `crates/ml/build.rs::kernels_with_common`
- DELETE 3 cubin static byte arrays + 3 struct fields + 3 cubin
loaders + 3 fields in trainer construction tuple in
`gpu_dqn_trainer.rs`
- DELETE 3 launches in `launch_sp9_kelly_warmup_floor` (chain
shrinks 5 → 2: q_var_mag_ema + main warmup-floor)
- DELETE 3 scratch slots; SP5_SCRATCH_TOTAL 268 → 265;
SCRATCH_SP9_EVAL_DIST_BASE slides 265 → 262
- ADD constructor-write of 3 Invariant-1 anchors (100.0, 2.0, 5.0)
immediately before layout-fingerprint write
- UPDATE 3 dispatch arms in `reset_named_state` to rewrite the
Invariant-1 anchors at fold boundary (NOT sentinel 0) — these
are constants, not stateful EMAs
- UPDATE 3 registry descriptions to reflect Invariant-1 anchor
semantic + smoke-test-wrwkz evidence
- UPDATE `intent_eval_divergence_compute_kernel.cu` with
sentinel-detect branch
- ISV slot layout UNCHANGED (3 target slots @ ISV[333..336)
remain); Wiener-buffer linear span (`SP5_PRODUCER_COUNT=165`)
UNCHANGED — 3 wiener triples for deleted producers become
reserved-unused like Pearl 6's [525..543) carve-out
- audit doc Fix 37.1 entry with full provenance + smoke evidence
**Verification:**
- `SQLX_OFFLINE=true cargo check -p ml` — clean (only pre-existing
18 warnings; no new errors or warnings).
- `SQLX_OFFLINE=true cargo test -p ml --lib state_reset` — 4/4 pass
including `every_fold_and_soft_reset_entry_has_dispatch_arm`.
- `SQLX_OFFLINE=true cargo test -p ml --lib sp5_isv_slots` — 9/9
pass including `sp9_slots_contiguous_above_sp8_block`.
**Expected smoke signature post-fix:**
- floor: starts ≈ base_floor × 1.0, decays as confidence accumulates
- divergence: 1e6 until first val window, small (≤ 5) afterward
- conf: stat=count/100, bhv=0 until eval_dist matures, tmp=ep/5
- combined_conf reaches 1.0 only when at least one axis genuinely
matures (typically temporal at epoch 5 in fold 0)
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per pearl_cold_start_exit_signal_or.md + pearl_controller_anchors_isv_driven.md:
the Kelly cap's `warmup_floor` and its release condition were both
regime-encoded constants. Single-axis statistical cold-start exit
(`total_trades >= 10` in trade_physics.cuh::kelly_position_cap) created
a chicken-and-egg deadlock: cap closed → no trades → no Kelly samples →
cap stays closed. T10 wsnc6 ep1-3 evidence (post-Fix 36):
intent_dist_f rising to 0.47 while eval_dist_f pinned to 0.00 with
kelly_f=0 across all epochs — train-vs-eval divergence as direct
symptom of the cold-start deadlock.
SP9 lifts the floor and its release condition fully onto ISV. The exit
condition is OR'd across statistical / behavioral / temporal axes per
the pearl — any single axis firing exits cold-start.
Atomic commit per feedback_no_partial_refactor:
- 9 new ISV slots @ [330..339):
- KELLY_WARMUP_FLOOR_INDEX (330) — adaptive floor consumed by
unified_env_step_core
- Q_VAR_MAG_EMA_INDEX (331) — historical EMA baseline for the
self-relative `base_floor` ratio (eliminates 0.5 hardcoded
half-Kelly)
- INTENT_EVAL_DIVERGENCE_INDEX (332) — behavioral-axis numerator
- 3 EMA targets (sample_count / divergence / temporal)
- 3 EVAL_DIST_{Q,H,F}_INDEX — replaces DtoH
read_eval_intent_magnitude_distribution()
- ISV_TOTAL_DIM 330 → 339; SP5_PRODUCER_COUNT linear span 156 → 165
- 9 new scratch slots @ [259..268); SP5_SCRATCH_TOTAL 259 → 268
- 7 new producer kernels (single-block cold-path, chained through
apply_pearls_ad_kernel for Pearl A bootstrap + Pearl D Wiener-α):
- eval_intent_dist_compute_kernel.cu — GPU reduction of
intent_mag_buf, replaces DtoH path per
feedback_no_cpu_compute_strict.md
- intent_eval_divergence_compute_kernel.cu
- q_var_mag_ema_compute_kernel.cu
- kelly_sample_count_target_ema_kernel.cu
- kelly_divergence_target_ema_kernel.cu
- kelly_temporal_target_ema_kernel.cu
- kelly_warmup_floor_compute_kernel.cu — main producer combining
all 8 ISV inputs
- consumer migration in trade_physics.cuh::unified_env_step_core:
health_safety_sp9 = fmaxf(health_safety_sp5,
isv[SP9_KELLY_WARMUP_FLOOR_INDEX])
threading through apply_kelly_cap's health_floor parameter; zero
sentinel → no-op via fmaxf so existing path's behavior is preserved
- DELETE read_eval_intent_magnitude_distribution() in
gpu_backtest_evaluator.rs (was DtoH read_all + host loop) per
feedback_no_htod_htoh_only_mapped_pinned.md; replaced by
intent_mag_dev_ptr_and_n() GPU accessor; metrics.rs:908 consumer
migrated from DtoH to launch + ISV mapped-pinned read
- 9 new state-reset registry FoldReset entries + 9 dispatch arms in
reset_named_state (Pearl A bootstrap on first observation)
- HEALTH_DIAG sp9_kelly_warmup line emits floor + divergence + 3
confidence axes + 3 EMA targets per feedback_no_hiding
- audit doc Fix 37 entry with full provenance, files touched, pearls
applied, verification
Constants: only KELLY_FLOOR_MIN_RATIO=0.25, KELLY_FLOOR_MAX_RATIO=1.0,
EPS_DIV=1e-6 — Invariant 1 numerical-stability anchors per
feedback_isv_for_adaptive_bounds.md. All thresholds, targets, and
the floor itself are signal-driven via ISV.
Pre-existing 18 cargo warnings unchanged. State-reset contract test
`every_fold_and_soft_reset_entry_has_dispatch_arm` passes for all 9
new entries; SP5 ISV slot layout test
`sp9_slots_contiguous_above_sp8_block` passes;
`SQLX_OFFLINE=true cargo check -p ml` clean.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Brainstorm spec for SP9 (Fix 37). Resolves the val-Flat-collapse pathology
surfaced in T10 train-multi-seed-wsnc6 ep1-3: training intent_dist_f=0.47
but eval_dist_f=0.00 because Kelly cap pinned eval mag to Quarter, with
trade_count=1 in 214k bars (chicken-and-egg: cap closed → no trades → no
Kelly samples → cap stays closed).
Architecture per pearl_controller_anchors_isv_driven and
pearl_cold_start_exit_signal_or:
final_kelly_f = max(measured_kelly_f, warmup_floor)
warmup_floor = base_floor × (1 − combined_confidence)
base_floor = clamp(q_var_mag / q_var_mag_ema, MIN, MAX)
combined_confidence = max(statistical, behavioral, temporal)
All thresholds ISV-driven via Pearl D Wiener-α; only Invariant 1 numerical
anchors (clamp ranges, EPS) remain as constants.
Scope: 9 new ISV slots, 6 new GPU producer kernels including a
mandatory eval_dist GPU migration (eliminating DtoH readback at
gpu_backtest_evaluator.rs:1353 per feedback_no_cpu_compute_strict).
~1000-1200 LOC, single atomic commit per feedback_no_partial_refactor.
Pearls authored as part of brainstorm:
- pearl_controller_anchors_isv_driven (regime-encoded constants)
- pearl_cold_start_exit_signal_or (OR'd signals across statistical /
behavioral / temporal axes)
Per pearl_controller_anchors_isv_driven.md: Fix 35 (CQL formula direction
flip) is necessary but insufficient — `const float MAX_BUDGET = 1.0f;` at
line 113 of loss_balance_controller_kernel.cu is regime-encoded the same
way the original CQL target_ratio formula was. Under healthy training
the cap is fine; under val-Flat-collapse it lets the controller saturate
budgets to 1.0 while the model is regressing into Flat.
The pearl prescribes the structural fix: pick the canary that fires
under the failure mode → train_active_frac (Long+Short / total during
training rollout). Lift it onto ISV with Pearls A+D smoothing; route it
through a producer kernel that derives per-(head, branch) cap via linear
interpolation FLOOR + active_frac × (CEIL - FLOOR); read the cap from
ISV in the controller kernel (no more hardcoded 1.0).
Atomic commit per feedback_no_partial_refactor:
- new kernel: train_active_frac_compute_kernel.cu — single-thread
reduction over monitoring_summary[5..17), writes scratch[250]; chained
apply_pearls_ad → ISV[TRAIN_ACTIVE_FRAC_INDEX=321]
- new kernel: loss_balance_max_budget_compute_kernel.cu — 8 threads
(2 heads × 4 branches), reads ISV[TRAIN_ACTIVE_FRAC_INDEX], writes
scratch[251..259); chained apply_pearls_ad ×8 →
ISV[LB_MAX_BUDGET_{CQL,C51}_BASE..+4)
- consumer: loss_balance_controller_kernel.cu — replace
`const float MAX_BUDGET = 1.0f;` with per-(head, branch) ISV reads,
defensive Pearl A bootstrap clamp
- 9 new ISV slots @ [321..330); ISV_TOTAL_DIM 321 → 330;
SP5_PRODUCER_COUNT linear span 147 → 156
- 9 new scratch slots @ [250..259); SP5_SCRATCH_TOTAL 250 → 259
- delete CPU train_active_frac compute (training_loop.rs:3392-3396) per
feedback_no_cpu_compute_strict; delete host field
`last_train_active_frac: f32`; replace with accessor reading from ISV
- 3 new state-reset registry entries (FoldReset; Pearl A bootstrap on
first observation) + dispatch arms in reset_named_state
- audit doc: Fix 36 entry with full provenance, files touched, pearls
applied, verification
Pre-existing 18 cargo warnings unchanged. State-reset contract test
`every_fold_and_soft_reset_entry_has_dispatch_arm` passes for all 3 new
entries; SP5 ISV slot layout test
`sp8_max_budget_slots_contiguous_and_above_activation_block` passes;
`SQLX_OFFLINE=true cargo check -p ml` clean.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
`target_ratio[CQL][b] = ANCHOR_CQL_RATIO * (1 - flatness[b])` had the
sign inverted: it pushed CQL HIGH when Q was flat (the collapse state)
and LOW when Q had variance. Per pearl_2_budget_kernel.cu, `flatness =
var_q / σ²`, so flatness HIGH means Q has variance — exactly when
overconfidence-risk-driven CQL pressure should kick in.
Symptom (T10-v3 train-multi-seed-x7sl2 logs): Once IQN Fix 34 woke
IQN-mag/ord/urg branches and produced real Q-targets, the controller's
inverted formula sent cql_budget to MAX_BUDGET=1.0 saturation on every
branch, the conservative pull kept Q flat, the model collapsed to
single-Flat-action eval (val_active_frac=0, dir_entropy=0,
trade_count=1, sharpe=0).
Fix is one operational character: `(1 - flatness)` → `flatness`. C51
formula was already correctly aligned (`flatness * ANCHOR_C51_RATIO`)
and unchanged.
Per pearl_controller_anchors_isv_driven.md: the inverted formula was
masked for months because the IQN-dead regime (Fix 34) suppressed
cql_raw on mag/ord/urg, never letting the controller engage with the
inverted target. Fixing the IQN regime exposed the dormant formula bug.
ANCHOR_CQL_RATIO=2.0 and MAX_BUDGET=1.0 remain hardcoded; Fix 36 will
make those ISV-driven via a GPU train_active_frac canary signal per
the pearl's "pick the canary that fires under the pathology" rule.
execute_training_pipeline consumes cached_target_h_s2_ptr via .take(), so
the SP6 Pearl 5 per-branch loop (4 sequential iqn calls per step) only
worked for branch 0. Branches 1, 2, 3 saw cached_ptr=None, returned a
"non-fatal" Err logged as a warning, and silently skipped apply_iqn_
trunk_gradient — only the direction branch's IQN quantile signal ever
reached the trunk, for months.
Surfaced in T10-retry train-multi-seed-ksjcm logs as repeated:
WARN: IQN parallel branch 1 step failed (non-fatal):
WARN: IQN parallel branch 2 step failed (non-fatal):
WARN: IQN parallel branch 3 step failed (non-fatal):
cached_target_h_s2_ptr is None.
Pre-existing since SP6 Pearl 5 (P4.T3-era multi-quantile IQN per-branch
τ schedules) — masked by the non-fatal warning classification.
Likely explains:
- cql_mag persistence at SP7 smoke (cql_mag=0.07 vs cql_dir=1.0):
CQL was the only learning signal mag had because IQN-mag was dead.
- SP7 c51 1000:1 controller suppression on mag/ord/urg may stabilize
at non-collapse values once IQN signal returns to those branches.
Fix is local: re-set cached ptr inside both 4-branch loops (parallel arm
near line 2123, sequential arm near line 2308). The ptr is the same
value (target_h_s2 is per-step, not per-branch), but the .take()
contract requires set-per-call. Defensive single-shot semantic preserved
— if a future bug skips the set, IQN errors loud rather than reusing
stale ptr.
Per feedback_no_partial_refactor: SP6 changed the call pattern 1→4 per
step but didn't migrate the cache contract. Per feedback_no_hiding: the
non-fatal classification hid the structural bug; escalation to hard
error deferred to next commit after post-fix T10 validates the warnings
disappear.
Audit doc updated with Fix 34 entry.
Without this trigger, cargo treated CUDA_COMPUTE_CAP-driven cubin
compilation as cached output of a non-tracked env var. The cargo-target
PVC is shared across H100 (sm_90) and L40S (sm_89) training jobs, so
swapping --gpu-pool between archs left stale cubins from the previous
build cached for any future sm_X variation.
Symptom (T10 train-multi-seed-rn559 today):
Failed to create DQN: rmsnorm cubin load: DriverError(
CUDA_ERROR_NO_BINARY_FOR_GPU,
"no kernel image is available for execution on the device")
The training pod scheduled on L40S (sm_89), the binary cache served
sm_90 cubins from a prior H100 build, and the rmsnorm cubin had no
sm_89 entrypoint.
Fix is one line: tell cargo to invalidate when CUDA_COMPUTE_CAP changes.
First post-fix build will rebuild all cubins (cargo conservatively
reruns when a new rerun-if-env-changed trigger is added without prior
state). Subsequent builds rebuild only on actual env changes.
Generalises the same lesson as the recent SP7 host-branch-in-captured-graph
fix: env-var-conditional code that doesn't declare its dependencies
freezes at first observation regardless of runtime input.
The SP7 controller wrote real budgets to ISV[BUDGET_CQL_BASE/C51_BASE] and
the activation flag correctly transitioned 0→1, but downstream SAXPY ops
still saw exactly 0.02/0.05 every step. Root cause: compute_adaptive_budgets
ran host-side Rust with `if cql_active >= 0.5 { real } else { bootstrap }`,
captured into the aux_child CUDA Graph at step 0 of each fold (when
cql_active is FoldReset sentinel 0.0). The bootstrap branch resolves to
literal 0.02/0.05 and freezes into kernel arg buffers; ~1000 graph replays
per fold use frozen scalars regardless of runtime ISV state.
Same disease as SP4 host-side EMA elimination 2026-05-01. Fix:
- New consume_lb_budget_kernel.cu (8 threads, 1 block): reads ISV slots
for activation flag + controller budget per (head, branch); applies
if/else dispatch on-device; writes trunk_mean + per-branch correction
to mapped-pinned scratch buf [10] = (cql_trunk + cql_corr×4 + c51_trunk
+ c51_corr×4). Same cubin houses dqn_{scale,saxpy}_f32_dev_ptr_kernel
variants reading alpha from a device pointer (no kernel-arg freeze).
- Buffer + launcher + cubin loading + 4 dev-ptr accessor methods
in gpu_dqn_trainer.rs.
- apply_c51_budget_scale / apply_cql_saxpy + per-branch variants take
alpha_dev_ptr: u64 instead of scalar f32. The "skip when ≈ 1.0"
optimization is dropped — the value lives on-device. Existing scalar
dqn_{saxpy,scale}_f32_kernel are NOT modified — they remain reused
by distill, VSN dW dilution, etc.
- compute_adaptive_budgets refactored to launch the dispatch kernel
inline (captured along with consumers; replay-time fresh values
every step) and stop returning host-resolved CQL/C51 scalars.
IQN/ENS keep their host-side resolution (out-of-scope per change
scope — no activation-flag dispatch).
- HEALTH_DIAG reads from lb_budget_effective_buf via new accessor
methods on FusedTrainingCtx; dead caches on GpuDqnTrainer
(last_{cql,c51}_budget_eff / _per_branch) deleted.
- Audit doc Fix 33. Memory pearl out-of-tree.
Bootstrap byte-equivalence preserved: kernel-arg constants
cql_bootstrap=0.02 / c51_bootstrap=0.05 remain byte-identical to the
prior host-side CQL_BOOTSTRAP_BUDGET / C51_BOOTSTRAP_BUDGET literals
and to loss_balance_controller_kernel.cu's COLD_START_FLOOR_* anchors.
Files: consume_lb_budget_kernel.cu (new), build.rs, gpu_dqn_trainer.rs,
fused_training.rs, training_loop.rs, dqn-wire-up-audit.md.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
After sp5 merge with -X theirs, two compile errors needed manual fix:
- gpu_her.rs lost the DevicePtrMut trait import
- gpu_tlob.rs Fix 20 regression test used the 4-arg adam_step signature
before sp5's Pearl 4 added β1/β2/ε params.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Conflicts auto-resolved with -X theirs strategy. sp5 has more
evolved MappedF32Buffer API + more recent work; main's via-pinned
cleanup landed in parallel as a lateral migration. Taking sp5's
side preserves the cleaner consolidated upload pattern.
The SP7 controller's CQL reference signal was reading cql_sx (post-SAXPY,
budget-scaled) which created a self-perpetuating deadlock at small
budget values: cql_sx_norm = budget × raw_grad → small budget → small
cql_sx → controller can't update → budget stays small.
The earlier offset 6 → 3 attempt failed because grad_decomp_launch_cql
was never effectively populating slot 3 — the snapshot pattern measures
‖grad_buf − snapshot‖, but apply_cql_gradient writes to cql_grad_scratch
(separate buffer), not grad_buf, so the snapshot delta is always 0.
This commit adds a real producer kernel cql_raw_norm_compute that reads
cql_grad_scratch directly and computes ‖raw_cql‖ over mag/dir/trunk
slices. Wired to fire AFTER apply_cql_gradient and BEFORE
apply_cql_saxpy, populating grad_decomp_result_pinned[3..6] with the
raw norm independent of cql_budget.
SP7 launcher updated to read from offset 3 (now: real raw CQL norm,
not the never-populated cql snapshot delta). HEALTH_DIAG label renamed
cql_sx → cql_raw to reflect the new contract; component index in the
cached grad_component_norms_* arrays switched 2 → 1.
The historical grad_decomp_launch_cql() call is removed — keeping it
would overwrite the slot with 0 after cql_raw_norm_compute fires. The
paired grad_decomp_snapshot_cql snapshot is left in place to scope the
diff to Path A; buffer cleanup (grad_snapshot_cql allocation +
grad_decomp_launch_cql definition) belongs in a follow-up commit per
feedback_no_partial_refactor.
Files: cql_raw_norm_kernel.cu (new, 97 LOC), build.rs,
gpu_dqn_trainer.rs (struct field + cubin static + load + launcher +
SP7 read offset 6→3), loss_balance_controller_kernel.cu (docstring +
arg comment), fused_training.rs (4 launch_cql_raw_norm call sites,
1 dead grad_decomp_launch_cql call removed), training_loop.rs
(HEALTH_DIAG label + index), audit doc Fix 31 SP7 Path A entry.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The SP7 smoke at 237b3dbfb showed all 8 (head, branch) combinations
stuck at bootstrap across 5 epochs, despite Q-variance and gradient
norms being non-zero. The activation flag mechanism is monotonic per
fold (once active, stays active), so 8/8 inactive across all observed
epochs implies the kernel never hits the active path — but we can't
confirm whether that's because (a) grad_decomp_result_pinned reads
zero at SP7's epoch-boundary call site, or (b) something else.
Added two HEALTH_DIAG lines:
- grad_decomp_pinned — reads bytes [0..12,24..36,36..48] of the
pinned buffer the SP7 kernel reads; surfaces iqn/cql_sx/c51 mag/
dir/trunk norms exactly as the kernel sees them.
- lb_active_per_branch — reads LB_{CQL,C51}_ACTIVE_BASE+0..4 from
ISV; surfaces the activation flag state per (head, branch).
Together these disambiguate "grad_decomp not populated at SP7 read
time" from "values populated but kernel cold-start branch fires for
other reasons".
Purely additive — no behavior change. Audit doc Fix 31 sub-bullet.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
