Commit Graph

283 Commits

Author SHA1 Message Date
jgrusewski
3e9cefdbd9 fix(sp18-v2): restore-best precedes shrink-and-perturb at fold transition
`train_walk_forward` called `reset_for_fold().await?` directly at the
fold boundary, which applies `shrink_and_perturb(α=0.8, σ=0.01)` to
whatever `params_flat` holds. At end-of-fold, `params_flat` carries the
LATEST-EPOCH decayed weights, NOT the best-Sharpe snapshot saved
mid-fold. Fold N+1 then started from `0.8 × decayed + noise` and
carried the within-fold edge-decay forward.

Insert `restore_best_gpu_params()` BEFORE `reset_for_fold` so S&P
operates on the best-Sharpe snapshot. Cold-start guard swallows the
"no snapshot saved" error on the first fold (or any fold without a
Sharpe improvement) — matches pre-fix behavior on cold start, no
regression.

State interaction with `reset_for_fold`: restore-best writes
`params_flat` ONLY (online weights); reset_for_fold then runs S&P on
the restored online, hard-syncs target ← restored online, and resets
Adam state on every optimizer. Non-conflicting.

`best_params_snapshot` is constructor-init `None` on FusedTrainingCtx
and is NEVER cleared at fold boundary. The trainer's `self.best_sharpe`
IS reset to NEG_INFINITY in `reset_for_fold` so the first improvement
in fold N+1 will overwrite the snapshot. Until then, the snapshot holds
whichever fold most recently saved a peak — intended training-scoped
behavior. Strict within-fold semantics flagged as a follow-up
consideration.

Behavioral test (CPU oracle) pins the math contract:
  - shrink_and_perturb(α, σ) on X = α × X + (1−α) × N(0, σ)
  - with-fix vs without-fix differ by α × (W_best − W_curr)
  - cold-start (best == current) is bit-identical between orderings

GPU-level kernel coverage stays in compile_training_kernels smoke and
the L40S 30-epoch validation gating SP18 Phase 1.

Pre-commit Invariant 7: docs/dqn-wire-up-audit.md updated with full
rationale, state-interaction analysis, lifecycle notes, and the
preserved cross-pearl invariants.

Per:
- feedback_no_partial_refactor (call-ordering + test + audit-doc atomic)
- feedback_no_legacy_aliases (reuses existing API unchanged)
- feedback_wire_everything_up (restore_best_gpu_params gains second
  cold-path consumer)
- pearl_no_host_branches_in_captured_graph (runs outside graph capture)
2026-05-09 01:36:52 +02:00
jgrusewski
27f8e332da plan(sp18 v2): Phase 0 Task 0.2 — B-leg V_SHARE trajectory + TD-error magnitude diagnostic
Adds the SP18 v2 Phase 0 B-leg observability scaffold per the plan's
Task 0.2 spec:

- New kernel `td_error_mag_ema_kernel.cu` (single-block × 256 threads):
  reads `td_errors_buf [B]` (already populated by C51/MSE loss for PER
  priority recomputation, post-train-step), block tree-reduces
  `mean(|td_errors[b]|)` (no atomicAdd), and EMA-blends into
  `ISV[TD_ERROR_MAG_EMA_INDEX=493]` via Pearl-A first-observation
  bootstrap (sentinel 0.0 → REPLACE) + fixed α=`WELFORD_ALPHA_MIN=0.4`
  per `pearl_wiener_alpha_floor_for_nonstationary`. The TDB_* Welford
  accumulators in slots [498..504) are RESERVED for the Phase 4
  q_next_target Wiener-α chain — not used in Phase 0.

- Cubin manifest entry in `crates/ml/build.rs` + `TD_ERROR_MAG_EMA_
  CUBIN` re-export in `gpu_dqn_trainer.rs`.

- `sp18_td_error_mag_ema_kernel` field on `GpuDqnTrainer` + cubin load
  on the trainer's stream + `launch_sp18_td_error_mag_ema_update()`
  cold-path launcher + `read_sp18_td_error_mag_ema()` convenience
  wrapper.

- New `sp18_v_share_history: [[f32; 4]; 5]` field on `DQNTrainer` —
  fixed-size ring buffer of the last 5 epochs of per-branch V_SHARE
  EMA readings (slots [478..482) per SP17 Phase 3.2). Initialised to
  `[[NaN; 4]; 5]`; epochs 0–3 emit `nan` as the slope and skip the
  ISV write; epoch 4 onward computes `(EMA[now] - EMA[now-4]) / 4`
  per branch and writes the dir-branch slope to
  `ISV[V_SHARE_TREND_DIAG_INDEX=496]`.

- Two new HEALTH_DIAG lines in `training_loop.rs` at the per-epoch
  boundary (right after the SP18 reward_decomp line):

    HEALTH_DIAG[N]: v_share_traj [dir_slope=X mag_slope=Y ord_slope=Z urg_slope=W]
    HEALTH_DIAG[N]: td_error_pre [magnitude_ema=X]

  V_SHARE slope is host-side computation against the ring buffer
  (`(now - now_m4) / 4` per branch). TD-error magnitude is post-blend
  ISV slot 493 read (producer fires inside `read_sp18_td_error_mag_
  ema()`). Pre-fix baseline for the B-DD9 ratio gate
  (`avg(|TD-error|) ratio post-fix / pre-fix ∈ [0.5, 5.0]`).

- New GPU oracle tests in `crates/ml/tests/sp18_hold_reward_oracle_
  tests.rs`:
  * `td_error_mag_ema_pearl_a_bootstrap` — synthetic td_errors with
    closed-form mean(|td|)=1.125; pre-populate slot at sentinel;
    assert post-launch slot equals the mean (Pearl-A direct-replace).
  * `td_error_mag_ema_blend_post_bootstrap` — synthetic td_errors
    with mean(|td|)=0.5; pre-populate slot at non-sentinel 1.0;
    assert blend equals `(1 - 0.4) × 1.0 + 0.4 × 0.5 = 0.8`.

Pure observability — no production-path consumer in this commit. No
reward changes, no Bellman target changes, no kernel modifications to
the action-selection or training paths. Per `feedback_no_partial_
refactor` the kernel + cubin manifest + buffer + launcher + ring
buffer + HEALTH_DIAG emit + GPU oracle tests all land atomically.

Verification:
  SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check --workspace clean.
  All 4 GPU oracle tests pass on RTX 3050 Ti (2.09s):
    reward_decomp_per_action_gpu_oracle, reward_decomp_empty_bin,
    td_error_mag_ema_pearl_a_bootstrap, td_error_mag_ema_blend_post_bootstrap.
  Existing slot lock + state_reset_registry tests still pass.

Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md
      § Phase 0 Task 0.2.
Audit: docs/dqn-wire-up-audit.md § "SP18 v2 Phase 0 Task 0.2".

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-09 01:18:15 +02:00
jgrusewski
15b50ac38f plan(sp18 v2): Phase 0 Task 0.1 — D-leg per-action reward decomposition diagnostic
Adds the SP18 v2 Phase 0 D-leg observability scaffold per the plan's
"Phase 0 — diagnostic emit (NO functional change)" task:

- New kernel `reward_decomp_diag_kernel.cu`: block tree-reduce
  (4 blocks × 256 threads, one block per direction-axis bin) reading
  `reward_components_per_sample [N×6]` + `actions_out [N]` and emitting
  5 per-bin stats (mean r_micro / mean r_opp_cost / mean r_popart /
  mean |reward| / fire_rate) into a 20-float row-major output. Bin
  order: Hold(0)→Long(1)→Short(2)→Flat(3); col order: micro→opp→
  popart→abs→fire. Empty-bin guard emits 0.0 (NOT NaN) per the
  consumer-side KILL CRITERION arithmetic contract.

- Cubin manifest entry in `crates/ml/build.rs` + `REWARD_DECOMP_DIAG_
  CUBIN` re-export in `gpu_dqn_trainer.rs`.

- 20-float `MappedF32Buffer sp18_reward_decomp_diag_buf` field on
  `GpuDqnTrainer` + accessor pair (`sp18_reward_decomp_diag_dev_ptr`
  for the writer-side launcher; `read_sp18_reward_decomp_diag` for the
  HEALTH_DIAG reader). Buffer is constructor-zeroed so cold-start
  HEALTH_DIAG emits a deterministic zero block.

- `sp18_reward_decomp_diag_kernel` field on `GpuExperienceCollector` +
  cubin load on the collector's stream + `launch_sp18_reward_decomp_
  diag(n, b1, b2, b3, out_dev_ptr)` launcher. Wired in
  `training_loop.rs` at the per-step boundary, BEFORE
  `launch_reward_component_ema_inplace` (which `memset_zeros` the
  source buffer after consuming it) per `pearl_canary_input_freshness_
  launch_order`.

- New per-epoch HEALTH_DIAG line emit at the existing per-epoch
  boundary (after the SP17 dueling line):

    HEALTH_DIAG[N]: reward_decomp [hold(micro=X opp=Y popart=Z abs=W
                    fire=F) long(...) short(...) flat(...)]

  Reads the mapped-pinned 20-float diag buffer directly via the
  collector→trainer host_ptr — no DtoH copy.

- New `crates/ml/tests/sp18_hold_reward_oracle_tests.rs`:
  * `reward_decomp_per_action_cpu_oracle` (CPU oracle pinning the
    per-bin reduction math against a 4-sample synthetic batch).
  * `reward_decomp_per_action_gpu_oracle` (GPU oracle, ignored unless
    `--ignored`; asserts kernel matches CPU oracle bit-for-bit within
    1e-6 f32 budget).
  * `reward_decomp_empty_bin_emits_zero_not_nan` (empty-bin contract
    guard).

Pure observability — no production-path consumer in this commit. No
reward changes, no Bellman target changes, no kernel modifications to
the action-selection or training paths. Per
`feedback_no_partial_refactor` the kernel + cubin manifest + buffer +
launcher + production wire-up + HEALTH_DIAG emit + GPU oracle test all
land atomically.

Verification:
  SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check --workspace clean.
  CPU oracle test passes; GPU oracle + empty-bin guard both pass on
  RTX 3050 Ti (2.13s).

Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md
      § Phase 0 Task 0.1.
Audit: docs/dqn-wire-up-audit.md § "SP18 v2 Phase 0 Task 0.1".

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-09 01:06:48 +02:00
jgrusewski
b6b17d46bb feat(sp17-3.2): V_share + advantage_clip_bound producers + extended emit
Phase 3.2 lands the remaining two SP17 dueling-Q diagnostic producers
atomically with kernel + launcher + Rust wrapper + extended HEALTH_DIAG
emit + GPU oracle tests per `feedback_wire_everything_up`.

V_share[d] = |E[V]| / (|E[V]| + |E[A_centered, picked]|)
  where picked = argmax_a Σ_z A_raw[i, a, z] (max-Q semantic — tractable
  per-batch without depending on actions_history_buf which is collector-
  time state stale relative to the cuBLAS forward at HEALTH_DIAG cadence).
  Pearl-A bootstrap (sentinel 0.5) + α=WELFORD_ALPHA_MIN=0.4 + bilateral
  [0, 1] clamp per `pearl_symmetric_clamp_audit`. 4 blocks × 256 threads.

advantage_clip_bound = p99(|A_centered|) × ADVANTAGE_CLIP_SAFETY_FACTOR=1.5
  via sp4_histogram_p99 (block tree-reduce + per-warp tile binning, NO
  atomicAdd per `pearl_fused_per_group_statistics_oracle`). EMA α=0.01
  slow per-fold + bilateral clamp [0.1, 100.0] per
  `pearl_symmetric_clamp_audit`. Pearl-A bootstrap (sentinel 1.0).
  Single block × 256 threads + flat |A_centered| scratch buffer
  (mapped-pinned, sized to B × Σ_d b_d × NA).

Observability-only — the actual clipping wire-up is Phase 5 follow-up.
The Phase 1 mean-zero contract (commits eabcf8d52..6f53d676f) makes
A_centered a meaningful signal; this commit observes it.

Extended HEALTH_DIAG line:

  HEALTH_DIAG[N]: dueling [v_share=(d=X m=Y o=Z u=W)]
                          [a_var=(d=A m=B o=C u=D)] [clip=K]

GPU oracle tests on RTX 3050 Ti (all pass, 13/13 SP17 tests):
- v_share_per_branch_matches_closed_form: synthetic V=2.0 + linear A
  per branch; closed-form V_share = 2/(2 + |K_d × (n_d-1)/2|);
  ε=1e-4. Pearl-A bootstrap REPLACES on first launch.
- advantage_clip_bound_tracks_p99_safety: synthetic A with action-
  dominant + per-(i,z) jitter (the jitter is REQUIRED — pathologically
  lockstep values undercount in sp4_histogram_p99's non-atomic warp
  tile binning per the kernel's documented "1/(256×32) loss for
  uniformly distributed signals" qualifier; concentrated values violate
  the assumption. Real |A_centered| in production is continuous, so
  this is a test-data-only effect.) ε=0.20 (jitter + linear histogram
  quantization).

Plan: docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md
      Phase 3.2.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 23:02:35 +02:00
jgrusewski
1e70cd5e59 feat(sp17-3.1): A_var_ema per-branch producer + HEALTH_DIAG emit
Phase 3 of SP17 dueling-Q identifiability — first of three diagnostic
producers landed atomically with kernel + launcher + Rust wrapper +
HEALTH_DIAG emit + GPU oracle test per `feedback_wire_everything_up`.

Per branch d ∈ {dir, mag, ord, urg}:
  Var_d = (1/(B × n_d × NA)) Σ_{i, a, z} (A[i, a, z] − mean_a A[*, z])²

Block tree-reduce (no atomicAdd, `feedback_no_atomicadd`); 4 blocks ×
256 threads. Pearl-A first-observation bootstrap (sentinel 0.0 →
REPLACE on first launch); steady-state α = WELFORD_ALPHA_MIN=0.4 per
`pearl_wiener_alpha_floor_for_nonstationary` — the structural-control
floor preserves catch-up bandwidth without storing 24 Welford
accumulator slots for a cold-path-cadence diagnostic.

Cold-path emit: single launch per HEALTH_DIAG cadence (epoch boundary)
right after `v_a_means`. New line:

  HEALTH_DIAG[N]: dueling [a_var=(d=X m=Y o=Z u=W)]

The line will be extended with V_share + advantage_clip_bound in
Phase 3.2, then finalised in Phase 3.3.

GPU oracle test on RTX 3050 Ti: synthetic A constructed so each branch
d has a closed-form Var(A_centered); kernel readback matches expected
value within ε=1e-4 (f32 rounding budget for ~8×n×51 accumulator
length).

Plan: docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md
      Phase 3.1.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 22:45:40 +02:00
jgrusewski
079a06e485 test(sp17): asymmetric A → deterministic argmax under MIN_TEMP=0.5
Verifies that when one direction's centered advantage dominates
strongly enough, action selection picks it deterministically across
all N=1000 Philox-keyed runs.

Plan specifies "Thompson temp = 0.0 → pure argmax E[Q]", but the
kernel floors thompson_temp at MIN_TEMP=0.5 per
pearl_blend_formulas_must_have_permanent_floor — pure τ=0 isn't
ISV-accessible. With τ=0.5 the blend is q_eff = 0.5·E[Q] +
0.5·q_sample; deterministic argmax across all τ=0.5 draws requires
the E[Q] gap to exceed atom_span/2.

Construction: NA=3 atoms [-0.1, 0, +0.1] (atom_span=0.2). A_dir
puts +300 at z=2 for Long, -100 at z=2 for others (per-atom mean
zero ⇒ centering preserves shape). Long centered E[Q] ≈ +0.1, others
≈ -0.05 (gap 0.15 > 0.1). q_sample[Long] = +0.1 with prob ≈ 1
(softmax(300) ≈ delta at z=2); q_sample[other] ∈ {-0.1, 0} (zero
prob for +0.1). q_eff[Long] = 0.1; q_eff[other] ≤ -0.025. Long
strictly wins all draws.

If the centering breaks (Long no longer dominant under centered E[Q])
or τ blends a non-Long sample over Long, this test fires.

Plan: docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 22:25:51 +02:00
jgrusewski
f31a6b7ff0 test(sp17): symmetric A ⇒ identical per-direction E[Q]
Verifies that under uniform A (== 0 across all action × atom slots),
every direction has identical centered E[Q] regardless of V.

The plan's original wording probes this through Thompson selector
("uniform action distribution"), but the kernel's first-wins-strict-
`>` argmax over four i.i.d. Thompson samples produces a structurally
non-uniform distribution under symmetric A even with correct centering
(closed-form earlier-bias predicts ≈[44%, 26%, 18%, 11%] across
Short/Hold/Long/Flat from the tie statistics). The Thompson distribution
is V-dependent through tie statistics — NOT a centering regression.

Restated as the structural pre-Thompson property: with A=0 and
V arbitrary, centered logit = V + 0 is identical across all directions
⇒ per-direction E[Q] identical to ε=1e-5. The Thompson selector reads
these centered logits; if A=0 produced non-zero per-direction E[Q]
spread, *that* would be the centering regression — exactly what this
test catches.

Probed via compute_expected_q (reads back per-action E[Q] directly,
no Thompson noise as red herring). V-non-uniform sanity check confirms
the kernel reads V (non-zero E[Q] when V ≠ 0).

Plan: docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 22:23:19 +02:00
jgrusewski
c2dd6917e0 test(sp17): A-centered invariance under V shift
Verifies the architectural contract: a uniform additive shift to V
across atoms cannot leak into the post-centering distribution. The
plan specifies reading A_centered directly, but A_centered is a
register-local quantity inside compute_expected_q; this test
restates the property as the equivalent behavioral assertion that
adding a uniform constant to V leaves every per-action E[Q]
identical (softmax translation invariance).

A regression that accidentally reduced over (V + A) instead of A
alone would shift the per-atom mean by V_SHIFT and corrupt the
centered logits; the per-action E[Q] would diverge by O(1), failing
the ε=1e-4 assertion.

Plan: docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 22:16:16 +02:00
jgrusewski
10fafe8e60 test(sp17): V-invariance behavioral test
Verifies the architectural contract: V depends only on state, not on
the specific advantage tensor. Two A tensors with same per-atom mean
produce different post-centering shapes ⇒ different per-action E[Q]
(centering is shape-sensitive), but the V-only diagnostic readout is
identical across the two runs (V cannot leak in via the per-atom mean
reduction).

This is the structural property that makes dueling work — without it,
the model conflates "state value" with "action value" and gradient
updates corrupt V via raw-A noise.

Plan: docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 22:15:08 +02:00
jgrusewski
ffa6fda868 feat(sp17): centered A in barrier/ib_gradient_direction (Commit D)
User design call DD11: migrate the two aux-CQL gradient kernels in
c51_loss_kernel.cu to the SP17 mean-zero advantage contract. The
pre-SP17 comment "skip advantage-mean centering — small epsilon vs
correctness simplicity" at barrier_gradient_direction:1212 is REMOVED;
its empirical-without-verification rationale doesn't hold under the
SP17 contract that compute_expected_q + c51_loss + c51_grad +
mag_concat_qdir + Thompson + quantile_q_select already enforce.

barrier_gradient_direction:
- Per-thread `a_mean_per_atom[NUM_ATOMS_MAX]` reduction over b0_size
  direction actions.
- Forward Q-value computation reads `v_row[z] + (adv_a[z] - mean[z])`.
- Backward gradient recompute uses the same centered logits in the
  per-action softmax probability accumulation. The Jacobian's symmetric
  -1/b0_size per-atom offset cancels in the barrier's relative-push
  gradient direction (max up, 2nd down — both targets see same offset).
- Stale "skip centering" comment deleted; replaced with SP17 explanation.

ib_gradient_direction:
- Same per-atom mean reduction at function entry.
- Forward Q computation + backward dq_dlogit recompute both flow through
  centered probabilities. Variance var_q is invariant under common
  per-atom shifts (math: shift cancels in (Q(a) - mean_q)^2), so var_q
  numerics are bit-equivalent — but the gradient flows through the
  centered probability `p(z|a)` for consistency with c51_grad backward.

NUM_ATOMS_MAX=128 ceiling guard added to both kernels (mirror of
experience_kernels.cu); early-exit `if (num_atoms > NUM_ATOMS_MAX) return`
matches the pattern already used by these kernels for unrelated
zero-op guards.

GPU oracle test (RTX 3050 Ti, 6/6 PASS):
  barrier_gradient_direction_uses_centered_advantage — A=0, V=0 ⇒
  centered logits all zero ⇒ uniform softmax ⇒ E[Q]=0 across actions ⇒
  q_gap=0 ⇒ barrier fires at min_req=0.05 ⇒ asserts total |grad| > 1e-6.
  Regression detector: any centering breakage produces non-finite or
  zero gradients ⇒ test fails loudly. ib_gradient_direction shares the
  identical per-atom mean reduction pattern so the same test covers
  both kernels structurally.

Verification:
  cargo check --workspace                                          → clean
  cargo test sp17_dueling_oracle_tests --features cuda
    -- --ignored                                                    → 6/6 PASS

⚠ INTERIM STATE: c51_loss_batched + c51_grad_kernel already-centered
sites still need Commit E annotation pass to mark the existing Jacobian
+ per-d=1 magnitude-std as SP17-compliant.

Plan: docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 22:00:07 +02:00
jgrusewski
3107edb8f7 feat(sp17): Thompson V-wire-in (Commit C)
User design call DD10 option (b): Thompson direction selector now reads
softmax(V[z] + (A[a, z] - mean_a A[*, z])) instead of pre-SP17
softmax(A[a, z]). Without V wired in, action selection responded to a
DIFFERENT distribution than compute_expected_q's E[Q] (the Bellman-target
ranking) — the very pathology SP17 is fixing at the action layer.

Architectural change closes the contract gap atomically across:

Kernel signature (`experience_action_select`):
- Added `const float* __restrict__ v_logits_dir` after `b_logits_dir`.
  NULL is invalid (no fallback) — kernel hard-requires V.
- New per-thread `a_mean_per_atom_dir[THOMPSON_MAX_ATOMS]` reduction
  computes mean A across the 4 direction actions per atom z (sample-local
  register, no atomicAdd).
- Pass 1 (E[Q] for temperature blend + conviction): builds per-action
  `combined_logits_d[z] = V[z] + (A[a,z] - mean_a A[*,z])` and feeds
  `softmax_c51_inline` instead of raw `b_logits_dir + d * n_atoms`.
- Pass 2 (Thompson sample): same combined-logits rebuild per direction.
- `softmax_c51_inline` device helper UNCHANGED — keeping centering at
  caller maintains a narrower contract; thompson_test_kernel and other
  potential callers stay unaffected.
- THOMPSON_MAX_ATOMS=128 ceiling preserved; __trap() on overflow.

QValueProvider trait extension:
- `compute_q_and_b_logits_to` now takes `v_logits_out_ptr: u64` and
  DtoD-copies `on_v_logits_buf` per sub-iteration (atomic per
  feedback_no_partial_refactor — every consumer migrates in lockstep).

Trainer + evaluator wire-up:
- `gpu_dqn_trainer.rs::on_v_logits_buf_ptr() -> u64` (new pub fn, mirror
  of existing `on_b_logits_buf_ptr`).
- `gpu_backtest_evaluator.rs::chunked_v_logits_buf` field allocated
  [chunk_n * NA + 32*3] (cuBLAS tail-safety pad).
- Both call sites (collector + evaluator) pass v_logits arg in launch.

GPU oracle test (RTX 3050 Ti, 5/5 PASS):
  thompson_direction_select_reads_v_logits — runs production cubin
  twice on identical A logits with V=[0,0,0] vs V=[10,0,0]; asserts
  q_gap_v_dominant < 50% of q_gap_v_zero. If V is being IGNORED
  (regression), both runs produce IDENTICAL q_gaps and the test fails
  with a clear message. Uses MappedF32Buffer / MappedI32Buffer per
  feedback_no_htod_htoh_only_mapped_pinned.

Note on the plan's "raw argmax = Hold but centered argmax = Long" test:
  Mathematical analysis shows softmax-with-constant-shift preserves
  action ordering (mean subtraction adds the same per-atom constant to
  every action's logits), so the plan's specific assertion isn't
  algebraically constructable with simple A/V. The replacement test
  (V-dependence of q_gap) is more sensitive — it fails on the actual
  regression case (V ignored ⇒ identical q_gaps) the plan was trying
  to detect.

Verification:
  cargo check --workspace                                          → clean
  cargo test sp17_dueling_oracle_tests --features cuda
    -- --ignored                                                    → 5/5 PASS

⚠ INTERIM STATE: aux-CQL barrier_gradient_direction +
ib_gradient_direction still read raw advantage. Commit D closes them;
Commit E annotates the pre-SP17 c51_loss/c51_grad already-centered sites.

Plan: docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 21:52:36 +02:00
jgrusewski
4802879494 feat(sp17): centered A in mag_concat_qdir (Commit B)
Migrates the magnitude-branch's direction-Q conditioning input to the
SP17 mean-zero contract. mag_concat_qdir builds a per-(sample, action)
softmax over V[z] + A_dir[a, z] across THREE inner passes (max,
sum_exp, E[Q]); all three now read `dir_logits_b[..] - a_mean_per_atom[z]`.

Without this migration the magnitude trunk-input would see raw E[Q_dir]
while c51_loss/c51_grad backward consumes centered A — mixed contract
ruled out by `feedback_no_partial_refactor`.

The Plan-4 q_rms adaptive scale (ISV[96]) is preserved unchanged;
centering changes the per-action E[Q_dir] values that feed it but the
scale formula itself is structurally orthogonal. Backward into the
direction logits flows through c51_grad (already mean-zero) so no bw
change is needed in this kernel.

GPU oracle test: B=1, SH2=2, NA=3, B0=4 with the same synthetic where
mean_a = [0.75, 0, 0.75] is non-zero. Asserts:
  (a) h_s2[0..SH2] copied verbatim into concat prefix
  (b) tail slots = centered_E[Q_dir] × (1.0 / q_rms) to ε=1e-5
  (c) tail[0] < 0 / tail[3] > 0 sign asymmetry (regression detector)

Verification (RTX 3050 Ti):
  cargo check --workspace                                          → clean
  cargo test sp17_dueling_oracle_tests --features cuda
    -- --ignored                                                    → 4/4 PASS

⚠ INTERIM STATE: Thompson direction-select + aux-CQL barrier/ib still
read raw advantage. Commits C-D close them; Commit E annotates the two
already-centered c51_loss/c51_grad pre-SP17 sites.

Plan: docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 21:38:03 +02:00
jgrusewski
9b6e94854f feat(sp17): centered A in quantile_q_select (Commit A)
Migrates `quantile_q_select` (uncertainty-driven action selection from
C51 atom CDFs) to the SP17 mean-zero advantage contract. The kernel
builds a per-(sample, branch, action) softmax over `V[z] + A[a, z]`
across THREE inner passes (max, sum_exp, CDF); all three now read
`adv_a[z] - a_mean_per_atom[z]`.

Plan flagged this as a hidden 6th consumer. Task 1.4/1.5 as authored
covered only c51_loss + c51_grad + mag_concat_qdir + Thompson; the
post-Task-1.2 audit found `quantile_q_select` reads raw advantage at
lines 5704/5709/5720 across all 4 branches and was missed entirely.

Sample-local register reduction (`float a_mean_per_atom[NUM_ATOMS_MAX]`,
NA_MAX=128 mirroring compute_expected_q); no atomicAdd, no shared mem,
no cross-thread sync per `feedback_no_atomicadd`. Device-side __trap()
on num_atoms overflow per `feedback_no_quickfixes`.

GPU oracle test: N=1, NA=3, B0=4 with V=[0,0,0] and A_raw chosen so the
per-atom mean is [0.75, 0, 0.75]. Test runs the production cubin with
iqn_readiness=0 and util_ema=1.0 and asserts each per-action q90 matches
the CPU oracle to ε=1e-5. Two structural-asymmetry assertions fail
loudly if centering regresses (action 0 q90 ≤ 0, action 3 q90 ≥ 0).
Mapped-pinned per `feedback_no_htod_htoh_only_mapped_pinned`.

Verification (RTX 3050 Ti):
  cargo check --workspace                                          → clean
  cargo test -p ml --test sp17_dueling_oracle_tests --features cuda
    -- --ignored                                                    → 3/3 PASS

⚠ INTERIM STATE: mag_concat_qdir + Thompson + aux-CQL barrier/ib still
read raw advantage. Commits B-D close them; Commit E annotates the two
already-centered c51_loss/c51_grad pre-SP17 sites. No L40S dispatch
escapes feat/sp17-dueling until every consumer is migrated per
`feedback_no_partial_refactor`.

Plan: docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 21:33:06 +02:00
jgrusewski
eabcf8d529 feat(sp17): mean-zero identifiability in compute_expected_q
Inserts per-atom mean-A reduction inside the per-branch outer loop:
Q[a, z] = V[z] + (A[a, z] - mean_a A[*, z])
The post-centering atom-level identifiability holds: Σ_a A_centered[a, z] = 0
for every (sample, branch, atom).

Per Wang et al. 2016 Dueling Networks: mean-zero is smoother and more
stable than max-zero. Atom-level subtraction (not expectation-level)
preserves distributional shape per action.

Sample-local register reduction; no atomicAdd, no shared memory, no
cross-thread sync. NUM_ATOMS_MAX=128 mirrors existing THOMPSON_MAX_ATOMS;
device __trap() on overflow.

⚠ INTERIM STATE: c51_loss_kernel + c51_grad_kernel + mag_concat_qdir
still read RAW (un-centered) advantage logits. Phase 1 Tasks 1.3-1.5
migrate them in lockstep within this branch BEFORE any L40S dispatch.
A partially-migrated state is forbidden per feedback_no_partial_refactor.

Plan: docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 21:13:16 +02:00
jgrusewski
7b13324bcb test(sp17): CPU oracle for mean-zero centered Q computation
Locks the mean-zero identifiability math contract independent of GPU
implementation. Subsequent Phase 1 GPU oracle tests compare kernel
output against this contract bit-for-bit.

Plan: docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 21:07:24 +02:00
jgrusewski
cc4746b48d plan(sp17): HEALTH_DIAG v_a_means baseline (PRE-CENTERING)
Adds per-epoch readout of the V/A breakdown so we can observe the
pre-SP17 baseline distribution before the identifiability projection
lands. Also instruments the kill criterion for Phase 0:

  HEALTH_DIAG[N]: v_a_means [v=X a_dir=Y a_mag=... a_ord=... a_urg=...]

If train-multi-seed-b5gmp's Q(Flat) over-attribution is V-driven, V
should be elevated (~0.4+) while a_dir is small. If V is balanced and
A is doing the work, dueling cannot help and we abort SP17.

Block-tree-reduce kernel \`v_a_means_diag_kernel\` (no atomicAdd per
\`feedback_no_atomicadd\`); MappedF32Buffer per
\`feedback_no_htod_htoh_only_mapped_pinned\`.

Plan: docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 20:52:35 +02:00
jgrusewski
25ff1d4195 fix(sp16): floor Wiener-α at 0.4 for non-stationary control loops
Wiener-optimal α minimizes MSE for stationary stochastic signals.
The min_hold_temp / hold_cost_scale loops track a target driven by
the adapting policy itself — non-stationary by construction.

Empirical evidence: train-multi-seed-b5gmp sha 641aa0dfd, Fold 1
collapse Sharpe 88.65 → 55.55 because α decayed to 0.248 by HD4
and could no longer track the climbing overrun signal.

Pearl: pearl_wiener_alpha_floor_for_nonstationary (memory).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 20:14:18 +02:00
jgrusewski
641aa0dfde fix(sp16-t3): Wiener-optimal adaptive α per pearl — hold_cost_scale + min_hold_temperature
Per train-multi-seed-hjzss validation: SP16 T1+T2 chain was structurally
landed but BEHAVIORALLY INERT in 5-epoch smoke (bit-identical to pfh9n
baseline through epoch 3). Root cause: hardcoded `alpha = 0.05f` in both
producer kernels violates feedback_isv_for_adaptive_bounds AND prevents
convergence in short runs (~60 epochs needed from cold start).

Fix per pearl_wiener_optimal_adaptive_alpha:
  α = diff_var / (diff_var + sample_var + ε)

Where sample_var = running variance of target signal (Welford accumulator)
and diff_var = running variance of consecutive one-step differences.

Cold-start: target jumps 1.0 → 6.4 → 7.0 → high diff_var → α ≈ 0.6+
            → near-bootstrap responsiveness in epochs 1-3
Steady-state: signal stabilizes → diff_var drops → α decays naturally
              → smoothing emerges without hardcoded constant

Adds 12 new ISV slots (6 per producer):
- HCS_TARGET_MEAN/M2, HCS_DIFF_MEAN/M2, HCS_PREV_TARGET, HCS_SAMPLE_COUNT
- MHT_TARGET_MEAN/M2, MHT_DIFF_MEAN/M2, MHT_PREV_TARGET, MHT_SAMPLE_COUNT

ISV_TOTAL_DIM 462 → 474.

Both kernels migrated atomically. Pearl-A bootstrap preserved (sentinel
on prev_blended triggers REPLACE; cold-start α=1.0 when N<3 samples).
Defensive bounds [WELFORD_ALPHA_MIN=0.01, WELFORD_ALPHA_MAX=0.95] on the
Wiener-derived α to guard against denormal/underflow corner cases.

HEALTH_DIAG[N] emit extended with `alpha=...` and `sample_count=...` for
direct trajectory observation in validation smoke.

Behavioral tests verify:
- α high during signal jumps (>0.3 at epoch 3 post-cold-start)
- α low in steady state (mean tail α<0.4 under converging signal)
- Pearl-A bootstrap fires on first observation (Welford state advances
  regardless of REPLACE branch)
- α stays within [WELFORD_ALPHA_MIN, WELFORD_ALPHA_MAX] over 50 epochs
  (post-cold-start; cold-start α=1.0 by design)
- No 0.05f hardcoded literal remains in blend math (regression-locked
  via host-only string scan)

5 GPU + host tests pass: sp16_phase3_alpha_high_during_signal_jump,
alpha_low_in_steady_state, pearl_a_bootstrap_first_obs,
alpha_naturally_bounded, no_hardcoded_alpha. sp14 + sp15 oracle suites
unchanged (34 GPU tests + 4 host tests).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 18:56:58 +02:00
jgrusewski
1a3bcf97b8 feat(sp16-p2): adaptive Hold cost scale via ISV[461]
Per train-multi-seed-pfh9n post-mortem: observed_hold_rate climbed 0.25 → 0.52
across training while cost penalty (~0.006) was 100× smaller than per-bar
reward magnitudes (popart=0.97, cf=0.65). Hold action was effectively free,
allowing Q(Hold) to dominate via structural low-variance bias.

Fix: scale Hold cost adaptively. ISV[HOLD_COST_SCALE_INDEX=461] tracks:
  scale = clamp(1.0 + 24.0 × max(0, observed - target) / max(target, 0.01), 1.0, 25.0)

Effective cost at 100% overrun (observed=2× target): 0.006 × 25 = 0.15,
competitive with per-bar reward magnitudes (~0.01-0.1).
At/below target: scale = 1.0 (no extra penalty).

Pearl-A bootstrap + Welford slow EMA (α=0.05). Mirrors T1's
MIN_HOLD_TEMPERATURE pattern (same input signals: ISV[382] observed,
ISV[381] target).

Producer: hold_cost_scale_update_kernel.cu — single-thread cold-path,
per-epoch boundary, AFTER MIN_HOLD_TEMPERATURE in training_loop.rs.

Consumer migration (atomic per feedback_no_partial_refactor):
3 sites in experience_kernels.cu — segment_complete branch (line ~3089),
per-bar positioned-Hold branch (line ~3553), per-bar flat-Hold branch
(line ~3617). Cold-start fallback: scale=1.0 when slot ≤ 0 or
out-of-bounds (bit-identical pre-Phase-2 cost magnitude).

ISV_TOTAL_DIM: 461 → 462.

Behavioral tests (5/5 PASS on RTX 3050):
- sp16_phase2_hold_cost_scale_climbs_with_overrun
- sp16_phase2_hold_cost_scale_at_target_is_one
- sp16_phase2_hold_cost_scale_under_target_is_one
- sp16_phase2_hold_cost_scale_bounds_clamp
- sp16_phase2_hold_cost_scale_pearl_a_bootstrap

Regression: SP14 oracle suite 30/30 PASS, SP15 phase 1 oracle suite
36/36 PASS.

Instrumentation: HEALTH_DIAG[N]: hold_cost_scale_diag obs/tgt/norm/scale.

Per feedback_isv_for_adaptive_bounds + feedback_no_partial_refactor.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 15:56:46 +02:00
jgrusewski
2c469af2f7 fixup(sp16-p1): correct stale build.rs comment + remove tautological test assertion
Code-quality review of 0426ce888 flagged two Important doc/test items:

1. crates/ml/build.rs lines 859-882 — comment described OLD slot-373
   dir-acc-skill mapping; replaced with short description of new
   slot-382/381 hold-rate-overrun mapping.

2. crates/ml/tests/sp14_oracle_tests.rs lines 1898-1903 — tautological
   `(A || ¬A)` assertion (dead test code); deleted. Preceding assert at
   1892-1897 already covers the real check.

No runtime impact — doc + test cleanup only.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 15:34:01 +02:00
jgrusewski
0426ce8887 fix(sp16-p1): MIN_HOLD_TEMPERATURE signal swap to hold-rate overrun + slot 330 non-bug closure
Per train-multi-seed-pfh9n post-mortem follow-up: slot 460 stuck at 50
in Fold 1 was NOT a launch-lifecycle bug. Producer fires per-epoch but
kernel had early-return guard on AUX_DIR_ACC_SHORT_EMA (slot 373) at
sentinel 0.5. Slot 373 reset on fold boundary; aux dir-acc EMA either
didn't fire or settled within ε of 0.5 → kernel kept early-returning.

Fix: drop slot 373 dependency entirely. Drive temperature from observed
hold-rate vs target overrun:

  overrun = max(0, observed_hold_rate - target_hold_rate)
  overrun_norm = clamp(overrun / max(target, 0.01), 0, 1)
  new_temp = TEMP_MIN + (TEMP_MAX - TEMP_MIN) × overrun_norm
  blended_temp = Welford EMA α=0.05 with Pearl-A bootstrap

When over-holding: temp HIGH → exit ramp permissive (matches design intent).
When at/under target: temp LOW → exit penalty strict.

Survives fold reset: hold-rate measurement starts fresh with real data
immediately, no chained-input-sentinel masking.

Slot 330 (KELLY_WARMUP_FLOOR) investigated and confirmed NON-BUG: producer
behaves correctly per pearl_kelly_cap_signal_driven_floors cross-fold-
persistence. floor=0 post-warmup is correct steady state.

Behavioral tests:
- sp16_phase1_min_hold_temp_climbs_with_hold_overrun
- sp16_phase1_min_hold_temp_strict_when_at_target
- sp16_phase1_min_hold_temp_strict_when_under_target
- sp16_phase1_min_hold_temp_no_longer_reads_slot_373

Instrumentation: HEALTH_DIAG[N]: min_hold_temp_diag obs/target/overrun/temp.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 15:24:12 +02:00
jgrusewski
fb24614f07 feat(sp16-p0): Q-by-action HEALTH_DIAG diagnostic for Hold-bias observation
Per train-multi-seed-pfh9n post-mortem: structural-Q-bias hypothesis
(Adam's m/sqrt(v) prefers low-variance Hold over noisy direction
Q-targets) needs direct per-action Q observations to verify or refute.
WR plateaued at ~0.46 across both folds while Hold% climbed 15% → 52%
and Q-mean climbed 0.19 → 0.41 — but per-action Q was unobservable.

Adds HEALTH_DIAG emit each epoch:

  HEALTH_DIAG[N]: q_by_action [hold=X long=Y short=Z flat=W]

Reads via host-side averaging of `q_out_buf [B, total_actions]`
direction-branch columns [0..b0=4]. No new kernel needed — q_out_buf
is already populated row-major by compute_expected_q. Modeled directly
on the existing Task 0.3 magnitude-bucket diagnostic (same dtoh-and-
average cold path).

Action-index ordering canonical, see state_layout.cuh:123-126:
  DIR_SHORT=0, DIR_HOLD=1, DIR_LONG=2, DIR_FLAT=3.
Emit slot order is [hold, long, short, flat] (Hold first because the
hypothesis is about Hold's ascent dominating direction Q-magnitudes).

New surface:
- gpu_dqn_trainer.rs: q_dir_means_cached field + update_q_dir_means_cached
  method + q_direction_action_means accessor + free static helper
  compute_q_dir_means_from_host_buf (factored for testability).
- fused_training.rs: FusedTrainingCtx wrappers parallel to q_mag pair.
- training_loop.rs: emit block adjacent to q_var_per_branch.

Behavioral test: sp16_phase0_q_by_action_diagnostic_reads_four_action_means
seeds known per-column constants, asserts canonical-dir-idx → emit-slot
mapping at 1e-3 tolerance, and includes sentinel-leak guard (non-direction
columns at 99.0 fail any wrong index→slot mapping with margin >12).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 14:34:59 +02:00
jgrusewski
9fb980da2b fixup(class-a-audit-batch-4b): invert MIN_HOLD_TEMPERATURE dir_acc → temp mapping
Post-commit kernel-formula audit of `compute_min_hold_penalty` at
trade_physics.cuh:567-577 showed the original `temp = TEMP_MIN +
(TEMP_MAX - TEMP_MIN) × skill` mapping was BACKWARDS for the
WR-plateau scenario this 8-commit chain targets.

Formula recap:
  soft_factor = deficit / (deficit + T)
  HIGH T → soft_factor → 0 → forgiving (no exit penalty)
  LOW T  → soft_factor → 1 → sharp (max exit penalty)

The deleted schedule's `START=50 → END=5` therefore meant
"permissive early, strict late" (classic explore-then-exploit), not
"permissive when committed, strict when uncertain" as the
substituted mapping assumed.

WR-plateau scenario: dir_acc ≈ 0.46-0.48 (model committed but
wrong) — the model needs a permissive (HIGH T) exit ramp to bail
out of bad committed directions, not maximum exit penalty (LOW T)
locking it into them.

Inverted the mapping using `confusion = 1 - skill`:
  - dir_acc ≈ 0.5 (random / plateau) → confusion=1 → temp=50
    (permissive, plateau exit ramp)
  - dir_acc → 1.0 (saturated skill)  → confusion=0 → temp=5
    (strict, force informed commitment)

This preserves the deleted schedule's epoch-0 anchor (T=50
cold-start) while reacting to actual realized skill instead of an
epoch-time anchor.

Files touched (atomic per `feedback_no_partial_refactor`):
- min_hold_temperature_update_kernel.cu — formula + header doctring.
- sp14_isv_slots.rs — slot-doc comment expanded with new mapping
  + fixup-rationale paragraph.
- gpu_aux_trunk.rs — `MinHoldTemperatureUpdateOps` docstring.
- gpu_dqn_trainer.rs — cubin docstring + ISV_TOTAL_DIM doc-comment.
- tests/sp14_oracle_tests.rs — Tests 3 + 4 flipped (high dir_acc
  now expects temp=TEMP_MIN; low dir_acc now expects TEMP_MAX);
  Tests 1, 2, 5 numerically unchanged (sentinel guard fires before
  the mapping; midpoint dir_acc=0.75 has skill=confusion=0.5 so
  pre/post-fixup numerics match).
- docs/dqn-wire-up-audit.md — Item 4 entry rewritten with the
  fixup rationale + mapping table updated.

Verification:
  cargo check -p ml --tests --all-targets   PASS
  cargo test sp14_audit_4b (lib)            4/4 PASS (slot layout
                                                   locks unchanged)

GPU oracle re-run deferred to next L40S smoke (kernel was rebuilt
in-place so cubin contents change; the 5 oracles encode the new
mapping and will exercise the new cubin).

Per `pearl_first_observation_bootstrap.md` (sentinel→target
replace; sentinel anchors unchanged) +
`pearl_symmetric_clamp_audit.md` (bilateral clamp on target_temp
preserved) + `feedback_no_partial_refactor.md`.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 11:28:00 +02:00
jgrusewski
0b9ea77dc4 fix(class-a-audit-batch-4b): plan_threshold floor adaptive + MIN_HOLD_TEMPERATURE ISV-driven
Per Class A audit-fix Batch 4-B (final 2 of 4 deferred items from
P1-wiring/P1-producer). Completes the 8-commit WR-plateau intervention
chain. Validation deferred to next L40S smoke.

Item 3: plan_threshold adaptive floor (Design Y - inline producer)
  - NEW slot PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX=459
  - Writes slow-EMA shadow of `0.5 * readiness_ema` from inside the
    existing update kernel (no new file, no new launch).
  - Pearl-A first-observation bootstrap (sentinel 0.1 matches pre-fix
    hardcoded value for bit-identical cold-start) + Welford alpha=0.005
    slow EMA.
  - Bilateral clamp [0.05, 0.50] (probability units) per
    pearl_symmetric_clamp_audit.
  - Consumer reads isv[459] as the floor in the same launch's final
    fmaxf; cold-start sentinel REPLACES with threshold_target so the
    pre-fix `fmaxf(0.1, 0.5*ema)` semantic is preserved bit-identical
    for any readiness EMA above 0.20.

Item 4: MIN_HOLD_TEMPERATURE -> ISV-driven (driving signal: dir_acc skill)
  - NEW slot MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460
  - NEW kernel min_hold_temperature_update_kernel.cu (single-thread
    cold-path, per-epoch boundary launch).
  - Driving signal: dir_acc skill = clamp((short_ema - 0.5)/0.5, 0, 1)
    from ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]. When committing skillfully
    (high dir_acc) -> temp HIGH (permissive). When at random baseline
    (~0.5) -> temp LOW (sharp commitment pressure). Substituted for
    the audit-spec's `dir_entropy_deficit` because no dir_entropy ISV
    slot exists - dir_acc skill is the closest semantically-equivalent
    signal that preserves the spec intent.
  - Pearl-A bootstrap (sentinel 50.0 matches the deleted
    MIN_HOLD_TEMPERATURE_START=50 anchor) + alpha=0.05 mid-cadence EMA.
  - Bounds [5, 50] (matches the deleted schedule range).
  - Decouples temperature from epoch number - the old schedule pinned
    LOW temp (sharp) at end of training, exactly when a WR-plateaued
    model needed forgiveness to escape.
  - DELETED: state_layout.cuh::MIN_HOLD_TEMPERATURE_{START, END, DECAY}
    #defines + training_loop.rs::min_hold_temperature_for_epoch helper
    function (kept docstring tombstone explaining the deletion). Both
    call sites migrated to the new ISV reader. Per
    feedback_no_legacy_aliases + feedback_no_partial_refactor.

ISV_TOTAL_DIM: 459 -> 461.

Cumulative WR-plateau fix series (final commit, #8):
- 8f218cab2 (Class C bug 1 + P0-B)
- 316db416b (P0-C MIN_HOLD_TARGET)
- 394de7d43 (P0-A REWARD_POS/NEG_CAP)
- c4b6d6ef2 (P1 wiring var_floor)
- 657972a4b (P0-A downstream DD penalty + MIN_HOLD_PENALTY_MAX)
- 87d597d5d (P1 producer Bayesian priors)
- 7e9a8f6ef (Batch 4-A DD saturation floor + legacy DELETE)
- this commit (Batch 4-B plan_threshold floor + MIN_HOLD_TEMPERATURE)

Verification:
  cargo check -p ml --tests --all-targets   PASS
  cargo test sp14 (lib)                     16/16 PASS (slot layout locks)
  cargo test sp15 (lib)                     2/2   PASS (no regression)
  cargo test state_reset_registry (lib)     4/4   PASS (dispatch coverage)
  cargo test sp14_oracle_tests (GPU)        24/24 PASS (8 new + 16 pre)
  cargo test sp15_phase1_oracle_tests (GPU) 36/36 PASS (no regression)

Per feedback_isv_for_adaptive_bounds + feedback_no_partial_refactor +
feedback_no_legacy_aliases.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 11:13:54 +02:00
jgrusewski
7e9a8f6ef1 fix(class-a-audit-batch-4a): DD saturation floor adaptive + legacy DD path Case A
Per Class A audit-fix Batch 4-A (deferred from P1-wiring/P1-producer due to
audit-doc errors). Fixes 2 of 4 deferred items; Batch 4-B handles plan_threshold
floor + MIN_HOLD_TEMPERATURE in a separate commit.

Item 1: DD saturation floor (the upper end of the DD ramp at trade_physics.cuh:154
in apply_margin_cap, NOT line 548 as the audit doc claimed — that line is a
magnitude action constant; the actual saturation floor lives in apply_margin_cap)
  - NEW slot DD_SATURATION_FLOOR_ADAPTIVE_INDEX=458
  - Producer dd_saturation_floor_update_kernel.cu — p75(per-env DD_MAX) × 1.5
    via Welford `mean + Z_75 × sigma` estimator with `max(p75, mean)` robustness
    guard, mirrors P0-A REWARD_POS_CAP producer pattern (Pearl-A bootstrap +
    Welford α=0.01)
  - Cold-start fallback: 0.25f (DD_SATURATION_FLOOR_DEFAULT in state_layout.cuh)
  - Bounds: [0.10, 0.50] (Category-1 dimensional safety)
  - Distinct from SP15_DD_THRESHOLD_INDEX=421 (the SP15 quadratic DD-penalty
    *trigger* threshold, a *lower* bound; this slot is the *upper* end of the
    linear position-size scaling ramp dd_scale = max(0.05, 1.0 − dd_frac/floor))
  - Threaded `isv_signals_ptr` into `apply_margin_cap` with NULL-tolerant
    cold-start fallback to DD_SATURATION_FLOOR_DEFAULT
  - 4 oracle tests (Pearl-A bootstrap, no-DD guard, bounds clamp, Welford EMA)

Item 2: Legacy compute_drawdown_penalty path → Case A (DELETED)
  - Decision rationale: SP15's quadratic asymmetric DD penalty
    (compute_sp15_final_reward_kernel.cu:154 via sp15_dd_penalty helper) runs
    unconditionally as a post-modifier on the SP11-composed reward with
    ISV-driven λ_dd (slot 420) and DD threshold (slot 421). Layering the legacy
    linear-ramp penalty inside the SP11 composer on top of the SP15 quadratic
    creates double-counting of DD shaping — exactly the code-smell the Class A
    audit was designed to eliminate. Per `feedback_no_legacy_aliases.md` and
    `feedback_no_partial_refactor.md`.
  - Atomic deletion across:
      - `compute_drawdown_penalty` device function (trade_physics.cuh)
      - Single call site at experience_kernels.cu:3822
      - `dd_threshold` and `w_dd` kernel arguments
      - `w_dd` Rust config field (gpu_experience_collector.rs +
        trainers/dqn/config.rs DQNHyperparameters)
      - `w_dd` profile section + dispatch (training_profile.rs RewardSection,
        OptimizableParameterRanges, FixedRewardParameters, ParamLookup
        dispatch, profile→hyperparam mapping, test assertion)
      - `w_dd *= rki` risk-intensity multiplier (config.rs)
      - `w_dd` TOML keys (dqn-hyperopt.toml × 2, dqn-localdev.toml,
        dqn-production.toml, dqn-smoketest.toml)
      - Stale doc comments on hyperopt/adapters/dqn.rs + config.rs
        risk_intensity field
  - `config.dd_threshold` SURVIVES (still consumed by `launch_sp15_dd_state`
    as the dd_budget for DD_PCT scaling). Documented in field comment.

ISV_TOTAL_DIM: 458 → 459 (Item 1 adds 1 slot; Item 2 is pure deletion)

Cumulative WR-plateau fix series (this is commit 7):
- Class C bug 1 + P0-B (8f218cab2)
- P0-C (316db416b)
- P0-A (394de7d43)
- P1 wiring (c4b6d6ef2) — 1 of 4 wireable
- P0-A downstream (657972a4b)
- P1 producer (87d597d5d)
- audit-fix 4-A (this commit)

Verification: 16 sp14_oracle_tests pass (incl. 4 new), 36 sp15_phase1_oracle_tests
pass, 12 sp14_isv_slots layout tests pass, 4 state_reset_registry tests pass
(every-FoldReset-arm-has-dispatch contract holds), workspace cargo check clean.

Per feedback_isv_for_adaptive_bounds + feedback_no_partial_refactor +
feedback_no_legacy_aliases.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 10:41:07 +02:00
jgrusewski
87d597d5d7 fix(class-a-p1-producer): adaptive Bayesian Kelly priors per-fold-end
Replace 4 hardcoded Bayesian priors in Kelly cap calculations
(prior_wins=2.0, prior_losses=2.0, prior_sum_wins=0.01,
prior_sum_losses=0.01) with ISV-driven slow-EMA values fed by a new
producer kernel that aggregates realized PS_KELLY_* fields across envs
from the same portfolio_state buffer kelly_cap_update_kernel reads
from at the same per-epoch boundary.

Slots [454..458):
- KELLY_PRIOR_WINS, KELLY_PRIOR_LOSSES,
  KELLY_PRIOR_SUM_WINS, KELLY_PRIOR_SUM_LOSSES

ISV_TOTAL_DIM bumps 454 -> 458.

Producer:
- kelly_bayesian_priors_update_kernel.cu (per-fold-end / per-epoch
  boundary, block-tree-reduce in shmem, no atomicAdd). Pearl-A
  first-observation bootstrap (sentinels 2.0/2.0/0.01/0.01 match
  pre-P1-Producer hardcoded values for bit-identical cold-start) +
  alpha=0.005 slow EMA. Bounds counts in [0.5, 100], sums in [0.001,
  1.0] per feedback_isv_for_adaptive_bounds. Launched RIGHT BEFORE
  launch_kelly_cap_update so that kernel sees the freshly-blended
  priors.

Consumers migrated atomically (same commit):
- kelly_cap_update_kernel.cu:39-42 -> ISV[454..458) with cold-start
  fallback via kelly_prior_or_default helper (range guard).
- trade_physics.cuh::kelly_position_cap:304-307 -> NULL-tolerant
  isv_signals_ptr threaded through apply_kelly_cap (single caller in
  unified_env_step_core line 898 already had the bus pointer; mirrors
  the existing kelly_f_smooth / kelly_warmup_floor_sp9 patterns).

State reset:
- 4 FoldReset registry entries (sp14_p1_kelly_prior_*) +
  4 dispatch arms in training_loop.rs::reset_named_state.

Oracle tests (sp14_oracle_tests.rs, GPU-gated #[ignore]):
- Pearl-A bootstrap (sentinel -> REPLACE with aggregated targets)
- No realized trades -> ISV preserved bit-exactly
- Bounds clamp on extreme aggregates
- Slow EMA blend after bootstrap

CPU tests passing:
- sp14_p1_kelly_prior_slot_layout_locked
- all_sp14_p1_slots_fit_within_isv_total_dim
- every_fold_and_soft_reset_entry_has_dispatch_arm (C.10 lesson)
- layout_fingerprint_bumps_after_sp14_wire

DEFERRED — Item 2 (MIN_HOLD_TEMPERATURE EMA): the audit-spec said
"MIN_HOLD_TEMPERATURE = 0.5f hardcoded somewhere" but the actual code
has MIN_HOLD_TEMPERATURE_{START=50.0f, END=5.0f, DECAY=20.0f} as a
PER-EPOCH ANNEALING SCHEDULE driven by min_hold_temperature_for_epoch
in training_loop.rs:68-73. The kernel takes T as a runtime scalar
specifically to enable Phase 2 ISV-driven lift "without recompiling
cubin" per the SP12 v3 design comment. The audit's claim of "0.5f
hardcoded" does not match reality; the right Phase 2 signal is a
separate spec decision and should not be guessed at per
feedback_no_quickfixes. Reporting back per the prompt's hard rule #8.

Per feedback_isv_for_adaptive_bounds + feedback_no_partial_refactor +
pearl_controller_anchors_isv_driven.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 09:38:51 +02:00
jgrusewski
394de7d434 feat(class-a-p0a): REWARD_POS/NEG_CAP → ISV-driven adaptive caps from realized return distribution
Per Class A audit ranking, the highest-suspected-impact fix for the
months-long WR-stuck-at-46-48% plateau across 11 superprojects.

Hardcoded REWARD_POS_CAP=+5.0f / REWARD_NEG_CAP=-10.0f
(state_layout.cuh:266-267) was structurally clipping the upper tail of
realized alpha. Controller signals (sharpe EMA, var_q, q_gap) all
derive from this CAPPED buffer, so the controller cannot select for
trades it cannot see. Selectivity gradient evaporates — small wins
clip to +5 alongside large wins also clipping to +5.

The state_layout comment lines 261-264 explicitly deferred Phase 2
(ISV-driven) "IF Phase 1 validation reveals adaptive need". Phase 1
has been running 11 SPs without budging WR — adaptive need revealed.

Architecture:
- 2 new ISV slots [452..454): REWARD_POS_CAP_ADAPTIVE,
  REWARD_NEG_CAP_ADAPTIVE.
- Producer kernel reward_cap_update_kernel.cu: block-tree-reduce
  Welford `mean + Z_99 × sigma` p99 estimator over winning realized
  returns + conservative `max(p99, max_win)` takeover, × 1.5 safety
  factor → POS cap. NEG cap = -2 × POS cap (preserves Kahneman 2:1
  asymmetry per pearl_audit_unboundedness_for_implicit_asymmetry —
  asymmetry stays, but moved from hardcoded scalar to producer-time
  multiplier; single source of truth, no consumer applies the 2× ratio
  itself).
- Pearl-A first-observation bootstrap from sentinel (5.0 / -10.0,
  matching pre-P0-A hardcoded values for bit-identical cold-start).
  Welford EMA α=0.01 thereafter (slow blend — reward distribution is
  the foundation of training and shouldn't move fast).
- Bounds: POS in [1, 50], NEG in [-100, -2] (Category-1 dimensional
  safety per feedback_isv_for_adaptive_bounds, NOT tuning).
- 3 consumer sites migrated atomically per
  feedback_no_partial_refactor: experience_kernels.cu:3112-3114
  (segment_complete cap), compute_sp15_final_reward_kernel.cu:163
  (Stage 4 helper invocation), sp15_reward_axis_helpers.cuh:211
  (sp15_apply_sp12_cap device fn signature change to take isv ptr).
- Cold-start fallback: when ISV slot at sentinel OR outside
  [REWARD_POS_CAP_MIN_BOUND=1, REWARD_POS_CAP_MAX_BOUND=50],
  consumers fall back to original macros (still defined in
  state_layout.cuh).
- 2 new device-ptr accessors on the experience collector
  (step_ret_per_sample_dev_ptr, trade_close_per_sample_dev_ptr) —
  reuses existing per-sample buffers; no new buffer allocated.
- Per-epoch boundary launch (cold path) in training_loop.rs alongside
  launch_aux_horizon_chain.
- Reset registry entries + dispatch arms in reset_named_state per the
  C.10 lesson (missing dispatch causes runtime crash).
- Layout fingerprint seed updated: ISV_TOTAL_DIM 452→454 +
  AUX_PRED_HORIZON_BARS=450 + AVG_WIN_HOLD_TIME_BARS=451 (previously
  missing from seed) + REWARD_POS_CAP_ADAPTIVE=452 +
  REWARD_NEG_CAP_ADAPTIVE=453.

Per feedback_isv_for_adaptive_bounds: every adaptive bound in ISV.

Verification:
- cargo check -p ml --tests --all-targets: clean (19 pre-existing
  warnings, 0 new).
- sp14_isv_slots tests: 8/8 pass (4 layout + 4 fits-within).
- sp14_oracle_tests with --features cuda --ignored: 8/8 pass (4
  existing q_disagreement/dir_concat + 4 new P0-A tests covering
  Pearl-A bootstrap, no-winning-trades preservation, bounds clamping
  to [1,50], Welford α=0.01 EMA blend).
- sp15_phase1_oracle_tests with --features cuda --ignored: 36/36 pass
  (no regression from sp15_apply_sp12_cap signature change).

Cumulative WR-plateau fix series:
- Class C bug 1 (8f218cab2): replay buffer intent→realized.
- Class A P0-B (8f218cab2): Kelly warmup floor wiring.
- Class A P0-C (316db416b): MIN_HOLD_TARGET adaptive.
- Class A P0-A (this commit): REWARD_POS/NEG_CAP adaptive — restores
  upper-tail alpha to the controller's view.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 08:45:43 +02:00
jgrusewski
10e647c141 test(sp14-c9): synthetic smoke for aux trunk gradient chain + C.8/C.9 audit close-out
C.8 (ISV-driven aux trunk Adam β1/β2/ε/LR/grad-clip) was already complete in C.5a
commit c90de9859 — all 5 ISV reads and fold-boundary StateResetRegistry defaults were
wired atomically with the Adam launcher. No new code required; noted in audit doc.

C.9 adds `aux_trunk_learns_synthetic_uptrend` to aux_trunk_oracle_tests.rs:
- B=16, ENC=32, H1=32, H2=16, SH2=32, H_HEAD=32, K=2, 100 steps
- Backward kernel invocations corrected to match actual signatures:
  - aux_trunk_bwd_dh_pre(d_logits, w3, w2, h_aux1, h_aux2, dh_pre2, dh_pre1, B, H1, H2, SH2)
    shmem = H2 floats (sh_dh2_pre cache), NOT (H1+H2+SH2)
  - aux_trunk_bwd_dW_reduce called 3×: dW3/dW2/dW1 each with (A, B_grad, dW_out, B, Krows, Jcols)
  - aux_trunk_bwd_db_reduce called 3×: db3/db2/db1 each with (B_grad, db_out, B, Jcols)
- Head params trained via host-side SGD (test orchestration only; reads mapped-pinned partials)
- Trunk params trained via dqn_adam_update_kernel (GPU Adam)
- Pass gate: CE loss < 0.1 AND dir_acc > 0.95 after 100 steps
  Near-random baseline (ln(2)≈0.693) = broken gradient chain, L40S dispatch blocked

Memory pearl pearl_separate_aux_trunk_when_shared_starves.md added and indexed.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
2026-05-08 03:30:44 +02:00
jgrusewski
0e61de408f feat(sp14-c6): h_s2_aux_rms_ema producer — ISV[449] per-collector-step
Single-block 256-thread CUDA kernel computing RMS(h_s2_aux [B, SH2])
and EMA-blending the step observation into ISV[H_S2_AUX_RMS_EMA_INDEX=449]
directly. Pearl-A first-observation bootstrap embedded in kernel body
(sentinel 0.0 → replace); fixed α=0.05 EMA blend thereafter.

ISV slot 449 is outside the SP4/SP5 wiener buffer linear span so the
scratch+apply_pearls_ad_kernel path is not available — self-contained
Pearl-A logic mirrors the avg_win_hold_time_update_kernel precedent
(slot 451). No atomicAdd; shmem block-tree-reduce only. Launched after
aux_trunk_forward in the collector per-step hot path.

- h_s2_aux_rms_ema_kernel.cu — new CUDA kernel (81 lines)
- build.rs — cubin manifest entry
- gpu_dqn_trainer.rs — H_S2_AUX_RMS_EMA_CUBIN static
- gpu_aux_trunk.rs — HS2AuxRmsEmaOps struct + launch()
- gpu_experience_collector.rs — field + constructor + hot-path launch
- aux_trunk_oracle_tests.rs — h_s2_aux_rms_ema_pearl_a_bootstrap test
- dqn-wire-up-audit.md — Phase C.6 audit entry

cargo check -p ml --tests: clean (only pre-existing warnings)
Oracle test: 1 new test added (requires GPU to run)

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
2026-05-08 03:17:30 +02:00
jgrusewski
3b71d21834 feat(sp14-c): aux prediction horizon ISV-driven (multi-bar pivot)
Aux's original label was (p_{t+1} > p_t) — pure HFT-scale microstructure
noise that's unlearnable at our HFT-MFT trading frequency. Migrated to
(p_{t+H} > p_t) where H is read from ISV[AUX_PRED_HORIZON_BARS_INDEX=450].

Adaptive producer drives H from observed avg winning hold time:
- Pearl-A first-observation bootstrap: replace sentinel H=60 directly
  on first valid observation
- Steady-state Wiener-α EMA blend, slow (α=0.01) for stable horizon
  (no target-variance EMA available, fallback per
  pearl_wiener_optimal_adaptive_alpha)
- "No winning trades yet" guard keeps sentinel until first valid observation

Lookahead truncation: labels at t where t+H >= total_bars are masked
(sentinel -1, loss-reduce skips). The existing aux_next_bar_loss_reduce
in aux_heads_kernel.cu already supports the -1 mask convention via the
B_valid count — no new valid_mask parameter needed.

Step 5b finding: Case B — existing per-sample buffers
(hold_at_exit_per_sample, trade_profitable_per_sample) populated by
unified_env_step_core, but no aggregate ISV slot. Added new aggregator
slot AVG_WIN_HOLD_TIME_BARS_INDEX=451 + new producer kernel
avg_win_hold_time_update_kernel.cu (block-tree-reduce, no atomicAdd).
ISV_TOTAL_DIM bumped 450 → 452.

ATOMIC migration per feedback_no_partial_refactor: both label kernels
(aux_sign_label_kernel.cu trajectory + aux_sign_label_per_step_kernel.cu
per-rollout-step) migrated together to the new
(targets, bar_indices, isv, isv_h_idx, out_labels, total, total_bars)
signature. The lookahead host-passed scalar argument is removed; H is
read from ISV inside the kernel (broadcast value, single read per
thread, on-device clamp [1, 240]).

Producer chain (per-epoch boundary): new
GpuDqnTrainer::launch_aux_horizon_chain orchestrates
avg_win_hold_time_update → aux_horizon_update sequentially alongside
launch_kelly_cap_update at the existing epoch-boundary slot in
training_loop.rs.

Trunk math (C.2/C.3/C.4) unchanged — separate aux trunk is label-
agnostic. Validation in C.10 will use H=60 cold-start; the adaptive
producer drives H from real winning-trade observations.

Tests (8 oracle, 5 new + 3 preserved):
- aux_trunk_forward_matches_numpy_reference (C.3) ✓
- aux_trunk_backward_gradient_check (C.4) ✓
- aux_trunk_backward_does_not_write_dx (C.4) ✓
- aux_sign_label_h_bar_horizon (NEW) ✓
- aux_sign_label_lookahead_mask (NEW) ✓
- aux_horizon_pearl_a_bootstrap (NEW) ✓
- aux_horizon_converges_to_steady_target (NEW) ✓
- aux_horizon_holds_sentinel_with_no_winning_trades (NEW) ✓

8/8 pass on RTX 3050 Ti.

Phase C.4b of SP14 Layer C separate-aux-trunk refactor.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 01:37:39 +02:00
jgrusewski
5d584dc751 feat(sp14-c): aux trunk backward kernel + gradient check + stop-grad invariant test
Backward propagates dh_s2_aux through w3/w2/w1 with block-tree-reduce
(no atomicAdd per feedback_no_atomicadd). Critical: kernel set does NOT
write dx_in — encoder gradient remains Q-shaped only. Stop-grad
invariant verified via parameter-list structural enforcement (kernels
literally cannot reference an `dx_in_out` pointer they don't accept) +
kernel source inspection that strips comments and asserts no `dx_in`
write pattern.

Three kernels in aux_trunk_backward_kernel.cu:
  - aux_trunk_bwd_dh_pre: per-sample, computes dh_aux2_pre [B, H2] +
    dh_aux1_pre [B, H1] using ELU' from POST-activation form
    (`(y > 0) ? 1 : (1 + y)` mirrors aux_elu_bwd_from_post in
    aux_heads_kernel.cu).
  - aux_trunk_bwd_dW_reduce: generic outer-product reduce
    `dW[k, j] = sum_b A[b, k] * B[b, j]`. One block per output
    element, shmem-tree reduce over batch. Used 3× (dW3, dW2, dW1).
  - aux_trunk_bwd_db_reduce: generic batch-reduce `db[j] = sum_b
    B[b, j]`. One block per output element. Used 3× (db3, db2, db1).

Memory-efficient: no per-sample partials (avoids B×163,072 floats for
production topology). Per-element reduction means O(P) blocks each
doing O(B) work in shmem.

Rust wrapper AuxTrunkBackwardOps in gpu_aux_trunk.rs orchestrates seven
launches in fixed sequence (capture-friendly, no host branches per
pearl_no_host_branches_in_captured_graph). All three CudaFunction
handles pre-loaded once at construction. Field added to GpuDqnTrainer
alongside aux_trunk_forward_ops; constructor mirrors C.3 pattern.

Tests (all pass on RTX 3050 Ti, sub-ULP forward, 1.33e-2 max rel-err
backward gradient at smallest sampled gradient):
  - aux_trunk_forward_matches_numpy_reference (C.3 — preserved).
  - aux_trunk_backward_gradient_check (NEW): central-difference
    numerical gradient at 16 sampled dW3 indices vs analytic from
    backward kernel. Loss = 0.5 * ||h_s2_aux||^2 so dh_s2_aux =
    h_s2_aux. EPS=1e-3, B=4, ENC=H1=H2=AUX=32 (33 forwards in ~2s).
    REL_TOL = 2e-2 (f32 finite-difference noise floor for
    small-gradient tail; production topology is dimension-independent
    given runtime args).
  - aux_trunk_backward_does_not_write_dx (NEW): reads kernel source,
    strips C-style comments (so design-discussion text mentioning
    `dx_in` doesn't false-positive), asserts no `dx_in` / `dx_in_out`
    symbol survives in code. Complements the structural enforcement
    (kernel signatures don't accept `dx_in_out` pointer).

Phase C.4 of SP14 Layer C separate-aux-trunk refactor. Module is
additive — wire-up into collector backward chain + Adam updates lands
in Phase C.5 (atomic).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 01:16:37 +02:00
jgrusewski
cb6bca4629 feat(sp14-c): aux trunk forward kernel + Rust wrapper + oracle test
3-layer MLP forward (Linear→ELU→Linear→ELU→Linear). Pre-loaded
CudaFunction for graph-capture safety per pearl_no_host_branches_in_captured_graph.
Oracle test verifies bit-for-bit match against numpy reference within
1e-4 tol. Saves h_aux1 and h_aux2 to global memory for backward.

Phase C.3 of SP14 Layer C separate-aux-trunk refactor (plan:
docs/superpowers/plans/2026-05-07-sp14-layer-c-separate-aux-trunk.md).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 00:56:23 +02:00
jgrusewski
4f372a49a8 refactor(sp14-c): atomic α machinery deletion + aux trunk ISV slot allocation
Phase C.1 of SP14 Layer C separate-aux-trunk refactor. Single atomic
commit per feedback_no_partial_refactor and feedback_no_legacy_aliases —
no DEPRECATED stage.

Deleted (per C.0 audit a7a162d1f):
- alpha_grad_compute_kernel.cu
- sp14_scale_wire_col_kernel.cu
- gradient_hack_detect_kernel.cu (EGF circuit breaker — α-coupled)
- 10 ISV slot constants (K_AUX_ADAPTIVE, K_Q_ADAPTIVE,
  BETA_RATE_LIMITER_ADAPTIVE, AUX_DIR_ACC_VARIANCE_EMA,
  ALPHA_GRAD_RAW_VARIANCE_EMA, GATE1_OPEN_STATE, ALPHA_GRAD_RAW,
  ALPHA_GRAD_SMOOTHED, AUX_DIR_ACC_POST_OPEN_MIN,
  GRADIENT_HACK_LOCKOUT_REMAINING) plus 16 supporting sentinels and
  structural-anchor constants
- 31 reference sites across 7 files (collector, trainer, training_loop,
  fused_training, batched_backward, build.rs cubin manifest,
  state_reset_registry)
- 3 oracle tests (alpha_grad_adaptive_beta, alpha_grad_schmitt_hysteresis,
  gradient_hack_circuit_breaker_fires)

Added:
- 6 aux trunk control plane ISV slots [444..450):
  AUX_TRUNK_LR (444), AUX_TRUNK_BETA1 (445), AUX_TRUNK_BETA2 (446),
  AUX_TRUNK_EPS (447), AUX_TRUNK_GRAD_CLIP (448), H_S2_AUX_RMS_EMA (449)
- 6 reset registry entries (5 Invariant-1 anchors + Pearl-A first-
  observation EMA)
- 6 reset_named_state dispatch arms (mirrors SP5 Layer A pattern)
- ISV_TOTAL_DIM bumped 444 → 450

Preserved:
- dir_concat_qaux_kernel.cu (Coupling A: forward feature wire)
- q_disagreement_update_kernel.cu (diagnostic-only)
- q_disagreement_* ISV slots (383, 384, 389) — HEALTH_DIAG consumer

Build clean (cargo check -p ml --tests --all-targets, RTX 3050 Ti);
4 surviving oracle tests pass (dir_concat_qaux_correct + 3 q_disagreement_*).
ISV layout: deleted α slots left as RESERVED gap (NOT compacted) for
checkpoint fingerprint compatibility.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-08 00:32:40 +02:00
jgrusewski
0ec791734d fix(sp15-cubin-preload): pre-load 6 SP15 launchers' CudaFunction handles in experience collector
Bug audit Pattern 2 — 6 SP15 launchers in gpu_dqn_trainer.rs were doing
load_cubin + load_function PER CALL inside the launcher body, called
from gpu_experience_collector.rs's per-rollout-step body (~thousands
of times per epoch x 4096 envs x 1000 timesteps).

Same architectural bug class as commits 5d63762ab (bn_tanh_concat_dd)
and 1396b62ec (sp15_baseline + cost_net). load_cubin/load_function
are host-side driver API calls — not capturable inside graph capture,
and prone to CUDA_ERROR_ILLEGAL_ADDRESS in subprocess child contexts
(documented failure mode in 1396b62ec).

Migration:
- launch_sp15_dd_state, launch_sp15_dd_state_reduce,
  launch_sp15_dd_trajectory_decreasing, launch_sp15_alpha_split_producer,
  launch_sp15_final_reward, launch_sp15_plasticity_injection —
  signatures changed to accept &CudaFunction parameter.
- Collector struct gains 6 new pre-loaded CudaFunction fields,
  populated once in collector::new() (mirrors SP14 EGF kernel loading
  added in commit 09202aa99).
- All 6 hot-path call sites updated to pass &self.exp_<kernel>_kernel
  reference (5 in collect_experiences_gpu, 1 in plasticity-injection
  trigger path).
- Oracle tests in sp15_phase1_oracle_tests.rs pass through a small
  load_sp15_kernel test helper that inlines the per-call load (graph
  capture isn't a concern in oracle scaffolds; production callers use
  struct-cached handles via the collector).
- docs/dqn-gpu-hot-path-audit.md gains Fix 29 documenting the
  pattern, the 6 migrated sites, deferred evaluator launchers, and
  out-of-scope orphan launchers.

Per feedback_no_partial_refactor: all 6 launchers + their callers
migrate atomically. Per pearl_no_host_branches_in_captured_graph:
zero load_cubin in collector per-rollout-step body after this commit.

Deferred (separate scope): launch_sp15_sharpe_per_bar and
launch_sp15_position_history_derivation are evaluator-path callers
(GpuBacktestEvaluator) — fixing them requires adding fields to a
different struct, deferred to a separate atomic commit. Orphan
launchers (regret_signal, cooldown) have no production callers; left
untouched.

Cubin static visibility: all 6 affected SP15 cubins were already
pub static (collector imports work without visibility promotion).

Verification: cargo check clean (only pre-existing warnings); 7/7
sp14_oracle_tests pass; 36/36 sp15_phase1_oracle_tests pass.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-07 23:36:07 +02:00
jgrusewski
9d0c124cee fix(sp14-egf): gate q_disagreement EMA update on total_cnt > 0 — fixes training-time decay-to-zero of rollout signal
Root cause from train-6fcml 5-epoch trajectory (commit 5608b866b after
producer cadence migration): HEALTH_DIAG[0] (post-experience-collection)
showed q_dis_s=0.0595 q_dis_l=0.1329 var_q=0.00091 — meaningful rollout
signal. HEALTH_DIAG[1+] (post-training, per-step launches) all showed
q_dis_s=0.0000 q_dis_l=0.0000 var_q=0.00000 — signal decayed to zero
inside ONE epoch.

The kernel's ISV write block ran unconditionally even when total_cnt
(non-masked-row count after Hold/Flat masking) was 0. Empty-batch
launches blended `batch_mean = 0/1 = 0` into the EMA, decaying the
rollout signal to 0 over ~178 training steps × 0.7^n. Per-step training
launches read replay batches whose Q-direction picks are dominated by
Hold/Flat (the natural distribution); so total_cnt = 0 was the common
case, not a corner case.

Fix (atomic, single kernel):
- Wrap the ISV write block in `if (total_cnt > 0.0f) { ... }`. When the
  training batch has no non-masked rows, the kernel is a no-op for that
  step — EMAs stay at the prior step's values. Stream-ordered launches
  still run; only the ISV write is skipped.
- Remove redundant `&& (total_cnt > 0.0f)` clause from the `is_first`
  bootstrap check (now guaranteed by the outer gate).

Per pearl_first_observation_bootstrap semantics: "no observation"
preserves prior; only "first observation" replaces sentinel. Decay-on-
empty was inconsistent with both rules.

Other EGF-chain kernels audited:
- alpha_grad_compute_kernel.cu — operates on persistent ISV state,
  no batch concept; var_aux/var_alpha Welford updates use `diff` of
  persistent EMAs, not batch means. No empty-batch path. SAFE.
- aux_dir_acc_reduce_kernel.cu — emits out_6[0..3] with sentinel
  fallback (0.5) when denom==0; downstream apply_fixed_alpha_ema then
  blends 0.5 toward EMA. The sentinel is the random-baseline (target
  threshold lies above it), so empty-batch pulls EMA toward harmless
  baseline rather than zero. Different semantics from q_disagreement
  (which has 0 — far below baseline 0.5). SAFE.
- gradient_hack_detect_kernel.cu — single-thread state machine on
  persistent ISV, no batch. SAFE.

Verification:
- 6 existing sp14_oracle_tests pass.
- New q_disagreement_empty_batch_preserves_ema test asserts bit-exact
  preservation of pre-seeded EMAs (0.0595, 0.1329, 0.0009 — the
  train-6fcml HEALTH_DIAG[0] values) across an all-Hold batch. Catches
  the regression the existing all-hold test missed (its bound
  `[0.0, 0.5]` accepted both decay-to-blend and preserve-prior; the new
  test is strict bit-equality).
- L40S smoke validation pending — train-6fcml symptoms (alpha_smoothed
  stuck at 0.0002, gate1 closed forever) expected to resolve.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-07 20:54:07 +02:00
jgrusewski
1396b62ec6 fix(sp15-wave5-followup): pre-load sp15_baseline + cost_net cubins — fixes hyperopt-trial CUDA_ERROR_ILLEGAL_ADDRESS
Root cause: 5 SP15 evaluation launchers (cost_net_sharpe + 4 baseline_*)
were doing `load_cubin` + `load_function` PER-CALL inside
`GpuBacktestEvaluator`'s eval hot loop. Pattern is fragile across CUDA
context lifetimes — works in single-pass train-best context (smoke
train-9bcwm verified), fails in hyperopt-trial child stream context
(workflow train-xggfc trial 1 failed at "load sp15_baseline_kernels
cubin: ILLEGAL_ADDRESS"; after the host-side load corrupted the trial's
context, trials 2-20 all cascade-failed at "Fork CUDA stream for trial").

Fix (atomic, matches 5d63762ab precedent for bn_tanh_concat_dd):
- Add 5 `CudaFunction` fields on `GpuBacktestEvaluator`.
- Pre-load both `SP15_BASELINE_KERNELS_CUBIN` (4 functions) and
  `SP15_COST_NET_SHARPE_CUBIN` (1 function) once in
  `GpuBacktestEvaluator::new()`, alongside the existing `env_module` /
  `metrics_module` loads.
- Change all 5 launcher signatures in `gpu_dqn_trainer.rs` to take
  `&CudaFunction` instead of doing per-call cubin load.
- Update all 5 call sites in `gpu_backtest_evaluator.rs:2877..2922` to
  pass `&self.sp15_*_kernel`.
- Update 4 oracle test call sites in
  `crates/ml/tests/sp15_phase1_oracle_tests.rs` (cost_net + 4 baselines)
  to pre-load and pass the kernel handle directly, matching 5d63762ab's
  pattern for `DQN_UTILITY_CUBIN`.
- Update Invariant 7 audit doc (`docs/dqn-wire-up-audit.md`).

Verification: `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml
--tests --all-targets` clean (warnings only, no errors).

Refs: train-xggfc failure 2026-05-07T13:11:43, 5d63762ab precedent,
pearl_no_host_branches_in_captured_graph (this is the eval analogue),
feedback_no_partial_refactor, feedback_wire_everything_up.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-07 16:47:41 +02:00
jgrusewski
5d63762ab3 fix(sp15-wave4.1b-OOB-followup): pre-load bn_tanh_concat_dd_kernel — fixes forward-capture SEGV
Root cause: SP15 Wave 4.1b (eb9515e41) migrated the trainer's 3 production
bottleneck-concat call sites (online forward, target forward, DDQN argmax)
from the pre-loaded `bn_tanh_concat_kernel: CudaFunction` field to a
`launch_sp15_bn_concat_dd` free function that resolved the symbol via
`load_cubin` + `load_function` ON EVERY CALL. That pattern is incompatible
with CUDA Graph capture: `cuModuleLoadData` and `cuModuleGetFunction` are
HOST-side driver API calls that allocate memory and mutate driver state,
and they are NOT capturable inside a `cuStreamBeginCapture` region. Calling
the launcher from `submit_forward_ops_main` (graph-captured forward child)
caused a SEGV at exit 139 — the loader raced with the capture-mode driver
state, the host-side corruption surfaced as a segfault before
`CAPTURE_PHASE_FORWARD_DONE` could print.

Diagnostic evidence (L40S smoke `train-vg5f7` on commit `bfc3ffa9d`):
ALL 16 step-0 ungraphed checkpoints printed clean. Capture begins:
  CAPTURE_PHASE_BEGIN
  CAPTURE_PHASE_PER_SAMPLE_DONE
  CAPTURE_PHASE_COUNTERS_DONE
  CAPTURE_PHASE_SPECTRAL_DONE
Then: SEGV. The next checkpoint that didn't print was
CAPTURE_PHASE_FORWARD_DONE -> SEGV is inside `submit_forward_ops_main`'s
`forward` child capture. The same function ran cleanly ungraphed in
step 0 because the loader-host-branch was harmless without active
capture.

The experience collector's equivalent caller (gpu_experience_collector.rs:4127)
was already doing this correctly — pre-loaded `bn_tanh_concat_fn` field
populated at construction. Wave 4.1b's mistake was asymmetric: it kept the
collector's pre-load pattern but introduced an on-demand loader for the
trainer's call sites.

Fix (atomic, restores graph-safe contract):
- Add `bn_tanh_concat_dd_kernel: CudaFunction` field on `GpuDqnTrainer`
  (back-fills what Wave 4.1b removed, with updated docstring naming the
  new dd_pct-aware kernel).
- Pre-load the symbol in `compile_training_kernels` from the same `module`
  as `bn_tanh_bw` / `bn_bias_grad` (forward-child captured replay group).
  Tuple grows 43 -> 44 CudaFunctions; struct ctor wires the field.
- Change `launch_sp15_bn_concat_dd` signature to take `&CudaFunction` as
  parameter; remove the per-call `load_cubin` + `load_function`. All 3
  trainer call sites pass `&self.bn_tanh_concat_dd_kernel`.
- Promote `DQN_UTILITY_CUBIN` from `pub(crate)` to `pub` so the oracle
  tests in `crates/ml/tests/sp15_phase1_oracle_tests.rs` can pre-load
  the kernel handle (the launcher no longer hides this for them).
- Update the 2 test call sites (kernel-level oracle + Wave 4.1c behavioral
  KL test) to load the kernel handle once up-front and pass it through.

Verification:
- `cargo check -p ml --tests` clean (release + dev profile).
- `cargo test -p ml --test sp15_phase1_oracle_tests -- --ignored`:
  36/36 pass, including the Wave 4.1c behavioral KL test that exercises
  two end-to-end forward passes through the modified launcher.

Refs: SP15 Wave 4.1b (eb9515e41), pearl_no_host_branches_in_captured_graph,
feedback_no_partial_refactor, feedback_wire_everything_up.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-07 13:18:14 +02:00
jgrusewski
c16b3b5a80 feat(sp15-p1.3.b-followup-B): per-(env,t) dd_trajectory + PER sampler — fixes Wave 4.3 uniform-batch limitation
Phase 1.3.b-followup (5b394f103) landed the main per-env redesign but
deferred dd_trajectory + per_insert_pa migration as Followup-B because
the contract shape is per-(env, t) buffer [N*L], not per-env tile.

This commit:
  - Reshapes dd_trajectory_decreasing_kernel to per-env grid
    [n_envs, 1, 1] x [1, 1, 1]; writes dd_trajectory_per_env[env]
  - Migrates per_insert_pa to per-transition lookup via
    env_id = (j % (n_envs * lookback)) / lookback; reads
    dd_trajectory_per_env[env_id]
  - Allocates two collector-owned per-env tiles
    (sp15_dd_trajectory_per_env + sp15_dd_trajectory_prev_dd_per_env);
    deletes the trainer-owned [1] sp15_dd_trajectory_prev_dd scratch +
    its setter wiring (replaced by unconditional collector ownership,
    mirrors sp15_dd_state_per_env pattern)
  - Extends dd_state_reduce_kernel to mean-aggregate per-env
    trajectory tile -> ISV[DD_TRAJECTORY_DECREASING_INDEX=439] for
    HEALTH_DIAG diagnostic preservation
  - Wires per-env tile dev_ptr + (n_envs, lookback) dims into
    GpuReplayBuffer via new setters (set_sp15_dd_trajectory_per_env_ptr,
    set_sp15_per_env_dims) called from training_loop
  - Migrates 4 dd_trajectory + 1 PER oracle tests to per-env contract
  - Adds NEW behavioral test
    per_sampler_weights_per_transition_not_uniform_across_batch:
    n_envs=2 batch with env-0 trajectory=1, env-1 trajectory=0; asserts
    priorities[env-0 slots]=3.0 and priorities[env-1 slots]=1.0 in the
    SAME insert batch (Wave 4.3 would produce uniform 3.0 OR uniform
    1.0 across all 8 priorities — the test directly fails the old
    implementation)
  - Layout fingerprint marker DD_TRAJECTORY_PER_ENV=sp15_phase_1_3_b_followup_B
    (greenfield checkpoints OK per spec Q1)

Fixes the Wave 4.3 PER limitation: ISV[DD_TRAJECTORY_DECREASING_INDEX=
439] was read ONCE per insert call and applied uniformly to ALL
n_envs * lookback * 2 transitions in the batch, so the recovery
oversample was statistically biased (whichever env wrote ISV[439]
most recently determined the boost for ALL inserted transitions).
Per-transition lookup fixes this — each transition's weight reflects
its own env's recovery context.

Atomic per feedback_no_partial_refactor: kernel reshape + 2 collector-
owned per-env tiles + trainer struct cleanup + reduction kernel
extension + per_insert_pa kernel signature change + 3 new GPU PER
replay buffer fields/setters/accessors + per_insert_pa launch site
update + collector launch sequence update + state-reset registry
rename (1->2 entries) + 2 dispatch arms + training_loop wiring
delete/replace + 4 dd_trajectory test migrations + 1 PER test
migration + 1 NEW behavioral test + 2 dd_state_reduce callsite
null updates + audit doc all in this commit.

Closes the per-env DD redesign chain end-to-end. SP15 reward shaping
+ recovery-curriculum PER oversample now correctly apply per-env
context everywhere they're consumed; no more "whichever env wrote
last wins" paths.

Verified: all dd_state, dd_trajectory, per_sampler, final_reward,
plasticity oracle tests green; ml lib HOLDS the Phase 1.3.b-followup
baseline (947 pass / 12 fail) on RTX 3050 Ti — no new regressions.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-07 10:08:49 +02:00
jgrusewski
5b394f1035 feat(sp15-p1.3.b-followup): per-env DD redesign — Path A env-0-canonical → Path B per-env tile + reduction
Closes the Phase 1.3.b deferred per-env redesign per
feedback_no_partial_refactor. Path A (env-0-canonical, commit 132609724)
made every downstream consumer read env-0's DD context; production envs
each have their own DD trajectory, so when env-3 was at 30% DD and
env-0 was at ATH, the model learning from env-3's transitions saw
dd_pct=0 and silently skipped the recovery shaping. Path B threads
each env's actual DD context through the reward shaping atomically:

(1) dd_state_kernel reshape — grid [n_envs, 1, 1], one thread per env,
    writes 6 scalars per env to a new per-env tile dd_state_per_env
    [n_envs * 6]. No more ISV scalar writes.
(2) NEW dd_state_reduce_kernel — single-block tree-reduce (no
    atomicAdd; BLOCK=256 with strided initial pass for n_envs up to
    32768 on H100). Mean-aggregates per-env tile → 6 scalar ISV slots
    [401..407) for HEALTH_DIAG diagnostic. Max-aggregates DD_PERSISTENCE
    → new slot DD_PERSISTENCE_MAX_INDEX=443 for plasticity injection
    trigger (one shared advantage-head ⇒ global firing ⇒ max-aggregate
    is the only correct rule). ISV_TOTAL_DIM 443→444; SP15_SLOT_END
    443→444; SP15_SLOT_COUNT 46→47.
(3) compute_sp15_final_reward_kernel migration — per-(i,t) per-env DD
    lookup via env_id = (idx % (N*L)) / L. Helpers sp15_dd_asymmetric_reward
    and sp15_dd_penalty migrated to take dd_pct + dd_current as scalar
    parameters (the kernel reads from the per-env tile, threads scalars
    in). On-policy + CF threads at the same (i,t) read the SAME tile
    entry (one DD trajectory per env, shared across slot kinds).
(4) plasticity_injection_kernel migration — persistence read switched
    from ISV[404] (mean) to ISV[443] (max). One set of advantage
    weights ⇒ ANY env exceeding the threshold should arm the gate.
(5) Per-env tile owned by GpuExperienceCollector (not the trainer) —
    the collector knows alloc_episodes (= n_envs); the trainer's
    batch_size is a different quantity. Reset to zero via the
    sp15_dd_state_per_env registry-arm dispatch.
(6) Per-step launch order: dd_state → dd_state_reduce →
    alpha_split_producer → final_reward, all on the same stream
    (CUDA serialises producer→consumer without explicit event sync).
(7) HEALTH_DIAG semantic shift (documented breaking change): slots
    401-406 now report cross-env mean, not env-0 value. For n_envs=1
    smoke configs the mean equals env-0's value (bit-stable migration).
(8) Layout fingerprint break: added markers DD_PERSISTENCE_MAX=443;
    ISV_TOTAL_DIM=444; DD_STATE_PER_ENV=sp15_phase_1_3_b_followup.
    Pre-followup checkpoints will not load (greenfield OK per spec Q1).
(9) 6 oracle tests migrated + 1 NEW behavioral test
    `dd_state_per_env_diverge_independently` — two-env config (env-0
    in recovery, env-1 deepening) verifies independent trajectories
    + mean-aggregate + max-aggregate semantics.

Phase 1.3.b-followup-B (separate split): dd_trajectory_decreasing_kernel
+ per_insert_pa migration is NOT in scope. The current Wave 4.3 reads
ISV[439] at insert-batch time (epoch end) — applied uniformly to ALL
inserted transitions (a known pre-existing limitation). Proper fix
requires per-(env, t) trajectory buffer [N*L] + env_id-aware lookup
in per_insert_pa via env_id = (j % (N*L)) / L. That's a different
contract change; splitting preserves no-partial-refactor within each
migration.

Atomic per feedback_no_partial_refactor: all 5 consumers of single-env-
canonical DD slots (final_reward kernel + plasticity kernel +
HEALTH_DIAG diagnostic + 6 oracle tests + new behavioral test) migrate
to per-env tile lookup in this commit; the new DD_PERSISTENCE_MAX
ISV slot lands with its sole consumer (plasticity).

Verified: SQLX_OFFLINE=true cargo check -p ml --features cuda clean;
cargo check -p ml --features cuda --tests clean;
CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests
--features cuda -- --ignored: 17/17 oracle tests green (2 dd_state
incl. new per-env behavioral + 5 plasticity + 3 dd_trajectory + 6
final_reward + 1 per_sampler); cargo test -p ml --features cuda --lib:
947 pass / 12 fail HOLDS the 483cef454 baseline (no new regressions).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-07 09:35:58 +02:00
jgrusewski
483cef454c feat(sp15-p3.5.4.c): production caller + OR-gate consumer for plasticity injection — closes 3.5.4 end-to-end
Wave 4.2 (ef08611d3) landed cuRAND Kaiming-He weight reset as orphan-
with-tests-only. Phase 3.5.4.c creates the production caller and the
action-selection consumer atomically per spec lines 4346-4448.

Production caller in gpu_experience_collector.rs step 4-pre (BEFORE
experience_action_select):
  - fan_in = cfg.adv_h = 128 (CORRECTED from Wave 4.2's audit-doc
    error claiming 6528 = adv_h × num_atoms; that's total weight
    count of one branch's projection, not per-unit input dim. Per
    feedback_trust_code_not_docs the audit-doc error is also fixed
    inline in this commit.)
  - n_weights = branch_0_size × num_atoms × adv_h = 4 × 51 × 128 =
    26112 (default config). Resets last 10% = 2611 directional
    advantage-head weights via cuRAND curand_normal × sqrt(2/128) ≈
    0.125 stddev (was wrongly documented as 0.0175).
  - Branch: w_b0out (directional) only — plasticity is about
    escaping stuck-in-flat regimes; targeted at directional choice.
  - Seed: mix_seed(Phase-3.5.4.c-unique base) ^ t per-step
    deterministic source; same (FOXHUNT_SEED, t) always yields the
    same Kaiming-He samples per kernel thread.
  - Trainer exposes target via new pub fn sp15_w_b0out_target()
    returning (dev_ptr, n_weights, fan_in); collector consumes via
    new set_sp15_plasticity_target(); training_loop wires the two.

OR-gate consumer in experience_action_select:
  - New kernel arg float plasticity_m_warm threaded into the
    cooldown-mask code path.
  - Wave 1.B's cooldown_active = (cooldown_remaining > 0) extended
    to cooldown_active = (cooldown_remaining > 0) ||
                         (plasticity_warm_remaining > 0).
  - Flat-on-fire-bar detection — option (a), warm ≥ m_warm − 1.5f
    per the trigger-then-decrement convention. The kernel sees
    warm = m_warm − 1 on the fire bar; subsequent warm-up bars
    observe warm ≤ m_warm − 2. New if (plasticity_fire_bar) branch
    BEFORE the existing else if (cooldown_active) branch — the
    fire-bar gets DIR_FLAT, subsequent warm bars get DIR_HOLD via
    the OR-gate cooldown.

Two-step recovery semantics:
  - Bar T (fire): plasticity launches → ISV[436]=1, ISV[438]=
    m_warm then decrements to m_warm−1. action_select detects
    fire-bar → dir_idx = DIR_FLAT.
  - Bars T+1..T+M_warm−1 (warm): action_select OR-gate forces
    dir_idx = DIR_HOLD.
  - Bar T+M_warm: warm transitions to 0, OR-gate inactive, normal
    Thompson/argmax resumes.

When the production caller is unwired (test scaffold path):
plasticity_m_warm passes 0.0f, the kernel's fire-bar predicate is
dead (warm == 0 too), and the OR-gate degenerates to cooldown-only
— bit-identical to the pre-3.5.4.c behaviour. Existing 3 SP15
3.5.b/3.5.3.b action_select oracle tests pass unchanged at
m_warm = 0.0f.

New behavioral oracle test plasticity_fires_force_flat_then_cooldown_
holds verifies the full sequence on RTX 3050 Ti with M_warm = 5
(shrunk from production's 200 for test runtime): bar 0 → DIR_FLAT,
bars 1-3 → DIR_HOLD, bar 4 (warm boundary) → DIR_LONG. ISV side-
effects (fired flips 0→1 then debounces, warm decrements with
underflow guard) verified bar-by-bar.

Eval/backtest path passes m_warm = 0.0f because plasticity is a
training-only mechanism (during deterministic eval the weights must
remain frozen at their checkpoint values).

Atomic per feedback_no_partial_refactor: kernel + caller + consumer +
trainer plumbing + 4 launcher-call-site updates (1 production
collector + 1 eval-path + 2 test scaffolds) + new behavioral test +
audit-doc fan_in inline correction + new audit-doc entry all in this
commit. No fallback. SP15 Phase 3.5 recovery-dynamics chain is now
end-to-end production-wired (3.5.2 + 3.5.3 + 3.5.4 + 3.5.4.c +
3.5.5 + 3.5.5.b all firing).

Verified: cargo check -p ml --features cuda clean; cargo check -p ml
--features cuda --tests clean; CUDA_COMPUTE_CAP=86 cargo test -p ml
--test sp15_phase1_oracle_tests --features cuda -- --ignored
--nocapture 34 of 34 SP15 oracle tests green (5 plasticity + 3
action_select + 3 cooldown + 23 others); cargo test -p ml --features
cuda --lib HOLDS the Wave 4.3 baseline (946 pass / 13 fail) on RTX
3050 Ti — no new regressions.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-07 08:58:17 +02:00
jgrusewski
19f6cce510 feat(sp15-wave4.3 / 3.5.5.b): PER sampler integration — recovery transitions oversampled
Phase 3.5.5 (69b8fdb61) landed dd_trajectory_kernel writing
ISV[DD_TRAJECTORY_DECREASING_INDEX=439] per-step but DEFERRED the PER
sampler integration AND the production launch. Wave 4.3 wires both
atomically per feedback_wire_everything_up.

Investigation finding — Case B (existing TD-error-driven priority
sampler): the GPU PER architecture in
crates/ml-dqn/src/gpu_replay_buffer.rs already weights replay by
per-transition priority (raw priority + priorities_pa = priority^alpha
feeding the prefix-sum sampler). Recovery boost slots in cleanly as a
multiplier on priorities[idx] at insert time — no architectural
refactor, no new sampling-path machinery.

Recovery-oversample formula:
  effective = base_priority × (1.0 + ISV[440] × ISV[439])
  priorities[idx] = effective
  priorities_pa[idx] = effective^alpha

When a transition is in a recovery (DD shrinking from non-trivial DD,
ISV[439]=1.0), it lands in the buffer at 1 + ω(2.0) × δ(1.0) = 3.0×
the baseline max_priority, then sampled 3.0^0.6 ≈ 1.93× more often
than baseline (alpha=0.6 PER convention) until per_update_pa
overwrites priority based on TD-error and normal PER takes over.
Recovery transitions get amplified gradient signal — the agent learns
recovery dynamics over typical "average-DD" Bellman noise.

Atomic landings (per feedback_no_partial_refactor):
  * crates/ml-dqn/src/per_kernels.cu — per_insert_pa kernel signature
    extended (+priorities, +isv pointers); kernel writes both columns
    of the (priority, priority^alpha) row pair; nullable ISV ptr falls
    back to boost_factor=1.0
  * crates/ml-dqn/src/gpu_replay_buffer.rs — new isv_signals_dev_ptr
    field + setter + accessor + priorities_pa_slice accessor; redundant
    pre-3.5.5.b scatter_insert_f32 broadcast of max_priority REMOVED
    (now subsumed by per_insert_pa's priorities[idx]=effective store)
  * crates/ml/src/cuda_pipeline/gpu_experience_collector.rs — new
    sp15_dd_trajectory_prev_dd_dev_ptr field + setter; per-step
    launch_sp15_dd_trajectory_decreasing invocation gated on both
    ISV ptr + prev_dd ptr being non-zero, placed immediately after
    launch_sp15_dd_state in the env-step loop
  * crates/ml/src/trainers/dqn/trainer/training_loop.rs — two new
    wiring blocks for the collector's prev_dd ptr and the replay
    buffer's ISV ptr, mirroring the existing
    set_isv_signals_ptr / set_sp15_alpha_warm_count_ptr plumbing
  * crates/ml/tests/sp15_phase1_oracle_tests.rs — new oracle test
    per_sampler_weights_recovery_transitions_higher verifying both
    columns of the priority-buffer write equal exactly the
    boosted/baseline values (3.0 vs 1.0; 3.0^0.6 vs 1.0^0.6) and the
    boosted/baseline ratio = 3.0 within 1e-5 — the recovery oversample
    factor by construction

Phase 3.5.5 deferred consumer eliminated per
feedback_wire_everything_up. This closes the SP15 Phase 3.5
recovery-dynamics chain (3.5.2 + 3.5.3 + 3.5.4 + 3.5.4.b + 3.5.5 +
3.5.5.b all wired). DD_TRAJECTORY_FLOOR (slot 441) and ISV-driven
RECOVERY_OVERSAMPLE_WEIGHT (slot 440) producers remain documented
follow-ups per feedback_isv_for_adaptive_bounds — kernel reads from
slots rather than literals so consumer migration is a no-op when the
producers land.

Verified on RTX 3050 Ti (CUDA 12.9, sm_86):
  * cargo check -p ml-dqn / -p ml --features cuda clean
  * 3 of 3 dd_trajectory oracle tests still green
  * 1 of 1 new per_sampler oracle test green (3.0 vs 1.0 priority
    bias, ratio = 3.000... within 1e-5)
  * 8 of 8 (1 ignored) gpu_residency replay-buffer + adamw smoke
    tests green
  * 4 of 4 PER smoke tests green
  * Full ml lib suite IMPROVES Wave 4.2 baseline: 947 pass / 12 fail
    (was 946 pass / 13 fail; the previously-failing
    test_dqn_checkpoint_round_trip now passes — likely the redundant
    pre-3.5.5.b scatter_insert_f32 was racing with the immediate
    per_insert_pa overwrite under particular timing conditions; the
    merged single-kernel write closes that race)

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-07 01:38:13 +02:00
jgrusewski
ef08611d3f feat(sp15-wave4.2 / 3.5.4.b): cuRAND Kaiming-He weight reset for plasticity injection
Phase 3.5.4 (e0e0abfb2) landed the trigger + warm-up tracker but
deferred the actual weight reset (kernel accepted advantage_head_weights
+ n_weights but no-op'd via (void) cast). Wave 4.2 lands the real
reset.

When DD_PERSISTENCE exceeds threshold AND not yet fired this fold,
the kernel now resets the last 10% of advantage-head weights to
Kaiming-He init: Normal(0, sqrt(2/fan_in)) sampled via cuRAND
curand_normal() with per-thread state initialized from a host-passed
seed (Option A: per-call curand_init(seed, tid, 0, &state) for full
determinism — bit-identical samples for identical (seed, tid) pairs).

Single kernel, two phases:
  1. Every thread independently re-evaluates fire_now from the same
     ISV reads (DD_PERSISTENCE / threshold / fired_flag); the trigger
     condition is a pure function of these reads so all threads
     converge without cross-block synchronisation. Block 0 / thread 0
     also runs the trigger + warm-bars decrement (single-thread ISV
     write path).
  2. If fire_now: each tid < reset_count writes Kaiming-He sample to
     advantage_head_weights[reset_start + tid]. Per-thread independent
     write, no atomic, no reduction (feedback_no_atomicadd clean).

New launcher params: fan_in (i32), seed (u64). Grid:
((n_weights/10 + 255) / 256).max(1) x [256, 1, 1]. The .max(1) floor
ensures block 0 always exists even when n_weights/10 == 0.

cuRAND device functions (curand_init, curand_normal) are inlined into
the cubin by nvcc from <curand_kernel.h> in the standard CUDA toolkit
include path — no host-side cuRAND linker dependency required, no
build.rs link change needed.

Oracle test plasticity_injection_kernel_resets_last_10pct_kaiming_he
verifies: (a) ISV[fired] flipped 0->1, (b) ISV[warm] = m_warm - 1,
(c) first 90% bit-identical to 1.0, (d) last 10% all moved off 1.0,
(e) sample mean |mean| < 0.05, (f) sample std within +-20% of
sqrt(2/fan_in), (g) determinism re-check produces bit-identical
samples for identical seed.

3 existing trigger tests migrated to the new launcher signature; the
debounced + no-fire variants additionally assert weight-stability
(early-out path skips the reset region).

Atomic per feedback_no_partial_refactor: kernel + launcher + 4 tests
+ audit doc + build.rs comment + 2 docstrings land together.

Eliminates Phase 3.5.4 deferred consumer per feedback_wire_everything_up
(the (void) casts are gone).

Action-selection consumer wiring (Phase 3.5.4.c) remains the separate
follow-up per the established Phase 3.5.X pattern.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-07 01:15:39 +02:00
jgrusewski
a54f53e4ed feat(sp15-wave4.1c): behavioral KL test — dd_pct trunk integration shifts policy distribution
Closes out Wave 4.1 (Phase 1.5.b consumer migration). Wave 4.1a (a8da1cb9c)
landed bn_tanh_concat_dd_kernel that fuses dd_pct into the trunk input;
Wave 4.1b (eb9515e41) wired s1_input_dim 102→103 through GRN reshape + 4
forward + 3 backward call sites. Wave 4.1c proves the wiring actually
changes the policy: a synthetic GPU forward composing launch_sp15_bn_concat_dd
with cublasSgemm_v2 against random Xavier-init weights W[proj_h=4, 103]
yields measurably different action distributions when ISV[DD_PCT]=0.0 vs
0.10 — observed mean KL=1.158e-4 (max=3.005e-4) vs threshold 1e-6
(~100× headroom).

Why a synthetic projection vs the real GRN trunk: the seeded forward_trunk_for_test
helper from Wave 4.1a noted (lines 2304-2319) that exposing the trainer's
trunk forward for tests would require either (a) a public surface change on
DQNTrainer exposing internal cuBLAS handles + GRN scratch + weights (the
trainer's fused_ctx is pub(crate) and only initialised inside the training
loop at training_loop.rs:547 — DQNTrainer::new returns with fused_ctx: None),
or (b) duplicating the trunk's cuBLAS setup in a test (≥200 lines of buffer
plumbing). Both options are architecturally heavier than the test's purpose
justifies. Per the spec dispatch ("the test's purpose is 'non-zero KL proves
the wire is connected' not 'verifies trained behavior'"), the synthetic
single-layer projection is the right scope: it exercises the new column-102
weights on the dd_pct value — exactly the path the real GRN's Linear_a first
GEMM takes for w_a_h_s1[:, 102] (the dd_pct column added by Wave 4.1b's
reshape).

Test contract:
- Two passes through launch_sp15_bn_concat_dd + cublasSgemm_v2 differ ONLY in
  ISV[DD_PCT_INDEX=406] (0.0 at-ATH vs 0.10 in-DD).
- Inputs (bn_hidden, states) deterministic; weights deterministic via LCG
  seed=42 with Xavier-uniform bound = sqrt(6 / (103+4)) ≈ 0.237.
- KL > 1e-6 (set 100× below the observed magnitude so a real wiring break
  fires this test, not silently passing).

What this test does NOT verify: the full GRN composition (ELU/GLU/LN/residual)
propagating dd_pct through h_s2 + the branch advantage heads. That end-to-end
behavior is exercised by the L40S smoke + production training runs.

Phase 1.5.b orphan launcher chain fully eliminated per
feedback_wire_everything_up: kernel landed (4.1a) → consumer migration
(4.1b) → behavioral verification (4.1c) — three atomic commits, the
3a/3b/3c split-pattern matching Wave 3's a/b decomposition. The Wave 4.1a
transient orphan window opened in a8da1cb9c → closed in eb9515e41 →
behavioral coverage added here.

Wave 4.1a's seeded helpers consumed: kl_divergence (used) and
minimal_trainer_for_tests (retained but unused — the seeded comment
correctly identified that exposing the trunk forward via the trainer
surface is non-trivial, so the helper waits for a future cargo-cult test
that needs trainer construction without GPU forward, e.g. weight-shape
introspection).

Touched: crates/ml/tests/sp15_phase1_oracle_tests.rs (+1 module
sp15_wave_4_1c_behavioral with 1 ignored test, 2 helper fns, 1 assertion
block — purely additive, no kernel or production-code changes), docs/dqn-wire-up-audit.md
(Wave 4.1c entry at top of audit doc).

Verified:
- SQLX_OFFLINE=true cargo check -p ml --features cuda --tests clean (18
  pre-existing unrelated warnings, no new warnings).
- CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests
  --features cuda -- --ignored bn_concat dd_pct --nocapture: 2 of 2 oracle
  tests green (Wave 4.1a bn_tanh_concat_dd_kernel_writes_dd_pct_column +
  Wave 4.1c dd_pct_trunk_input_shifts_policy_distribution).
- cargo test -p ml --features cuda --lib: 947 passed / 12 failed —
  exactly matches Wave 4.1b baseline (test addition is in the --test
  integration target, not lib target).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-07 00:52:26 +02:00
jgrusewski
a8da1cb9cf feat(sp15-wave4.1a): bn_tanh_concat appends dd_pct column from ISV — bottleneck-aware Phase 1.5 consumer migration
The standalone dd_pct_concat_kernel from Phase 1.5 was bottleneck-
incompatible — it operated on raw [B, 128] state, but production trunk
consumes [B, s1_input_dim] = [B, 102] post-bottleneck. Wave 4.1a fixes
this at the kernel level; Wave 4.1b lands the consumer migration
(s1_input_dim 102→103, GRN w_s1 reshape, 3 forward + 3 backward sites).

Spec correction (per feedback_trust_code_not_docs): the spec's
'state_dim 48→49' is stale terminology pre-STATE_DIM 48→112→128
evolution. Production s1_input_dim is bottleneck_dim + (STATE_DIM −
market_dim) = 16 + (128 − 42) = 102. Wave 4.1b will bump this to 103.

NEW bn_tanh_concat_dd_kernel in dqn_utility_kernels.cu:
  - Fuses dd_pct append into the same launch as bn_tanh + portfolio
    concat (output shape [B, bn_dim + portfolio_dim + 1])
  - Reads isv[DD_PCT_INDEX=406] (set by Wave 1.3.b dd_state_kernel
    per-step), broadcasts the scalar across batch as the appended
    last column

DELETED standalone dd_pct_concat_kernel.cu + launch_sp15_dd_pct_concat
+ cubin manifest entry per feedback_no_legacy_aliases (zero production
callers — only test consumer; bottleneck-on path is canonical).

Test helpers added (used by Wave 4.1c behavioral KL test).
Phase 1.5 oracle test migrated to bn_tanh_concat_dd_kernel contract:
test name bn_tanh_concat_dd_kernel_writes_dd_pct_column passes on
RTX 3050 Ti.

Layout fingerprint already covers Phase 1.5 via the existing
TRUNK_INPUT_DD_PCT=sp15_phase_1_5; marker — pre-SP15 checkpoints
already break.

fxcache schema_hash auto-bumps from file content hashes (per task
P5T5 Phase F mechanism); no manual schema bump needed.

Atomic per feedback_no_partial_refactor for the kernel-signature
contract change. Consumer wiring (s1_input_dim propagation, GRN
reshape, forward/backward call sites) deferred to Wave 4.1b's atomic
commit per the established 3a/3b split precedent — kernel + launcher
land first.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-07 00:06:35 +02:00
jgrusewski
4320820ae2 feat(sp15-wave3b): host-side wire-up — eliminates 5 orphan launchers via GpuBacktestEvaluator constructor signature change
Second half of the Wave 3 val-cost-streams refactor (3a kernel-side
foundation landed at e968f4ded). Atomically migrates the
GpuBacktestEvaluator::new contract; 3 call sites + 6 test call sites
+ 5 WindowMetrics fields + 11 new buffers + launch sequence wiring
all in this commit.

Constructor signature change: GpuBacktestEvaluator::new gains
window_lob_bars: &[Vec<LobBar>] parameter alongside existing
window_prices + window_features. Three production call sites migrated
atomically:
  - trainers/dqn/trainer/metrics.rs:651 (val_evaluator construction)
  - trainers/dqn/trainer/metrics.rs:1166 (extra_eval Dev/Test)
  - hyperopt/adapters/dqn.rs:1493 (full LobBar with real OFI)
  - hyperopt/adapters/ppo.rs:1364 (zero-OFI LobBar — PPO lacks per-bar
    OFI features; cost-net OFI-impact term degrades to 0; commission +
    half-spread × position still apply)

11 new mapped-pinned buffers on GpuBacktestEvaluator:
  Input (3): close_prices_buf, half_spread_buf, ofi_scalar_buf
  Derivation (3): position_history_buf, side_ind_buf, rt_ind_buf
  Output (5): cost_net_sharpe_buf, baseline_{buyhold,hold_only,
    momentum,reversion}_sharpe_buf

5 new WindowMetrics fields (host-annualised via × annualization_factor):
  - sharpe_cost_net (1.2.b cost-net sharpe)
  - baseline_{buyhold,hold_only,momentum,reversion}_sharpe (1.4.b)

Per-window eval flow now: existing fused metrics kernel → 1.1.b sharpe →
position_history_derivation (one launch over all windows) →
cost_net_sharpe (per-window) → 4 × baseline_* (per-window).

commission_per_rt: host-computed constant per D3 resolution, formula
config.tx_cost_bps × 0.0001 × config.initial_capital, passed by-value
at each baseline + cost_net launcher invocation.

Wave 3a kernel signature follow-up: cost_net_sharpe_kernel side_ind /
rt_ind switched from unsigned int* → float* so the cost-net kernel
chains directly with the f32 streams emitted by
position_history_derivation_kernel (no u32→f32 adapter buffer; bit-pun
mismatch fixed). cost_net oracle test migrated MappedU32Buffer →
MappedF32Buffer accordingly.

Atomic per feedback_no_partial_refactor: constructor sig change + 3
production + 6 test call site migrations + 5 orphan launchers
eliminated + WindowMetrics field additions + cost_net kernel sig fix
+ audit doc all in this commit.

Closes 1.2.b + 1.4.b + position_history_derivation orphan launchers
per feedback_wire_everything_up.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-06 23:39:45 +02:00
jgrusewski
e968f4ded9 feat(sp15-wave3a): kernel-side foundation — baseline output buffers + position_history derivation
Wave 3a half of the val-cost-streams refactor (3b host-side wire-up
follows). Atomically migrates the kernel-side contracts; 5 launchers
remain orphan transiently awaiting 3b production callers.

Baseline kernels (1.4):
  - 4 baseline_*_kernel signatures gain 'out: float*' parameter writing
    per-window [mean, std, raw_sharpe] (matches 1.1.b sharpe_per_bar shape)
  - ISV writes to slots 409, 410, 412, 416 removed entirely
    (per-window output is correct for WindowMetrics consumption;
    ISV-scalar writes were spec scaffolding for a single-fold-aggregate
    version that 1.4.b's per-window contract supersedes)
  - 4 ISV slot constants removed from sp15_isv_slots.rs
  - state_reset_registry: NO entries to remove (verified via grep —
    the 4 slots never had registry entries / dispatch arms in the first
    place; they were single-fold-aggregate scalars defaulted at every
    fold start by the constructor-write that initialises the ISV bus).
    Task 4 from the dispatch is a no-op; the
    every_fold_and_soft_reset_entry_has_dispatch_arm regression test
    continues to pass unchanged.
  - 4 oracle tests migrated to output-buffer assertion
  - layout_fingerprint_seed string updated (4 retired entries removed,
    4 trunk-shared entries retained; layout-break-class change)

New action_decoding_helpers.cuh:
  - Extracts factored_action_to_dir_idx + factored_action_to_position
    __device__ helpers (the latter is a higher-level position state-
    machine helper not previously available)
  - Mirrors trade_physics.cuh::decode_direction_4b semantics exactly so
    on-policy and counterfactual paths agree on factored-action meaning
  - Single source of truth for action→direction→position mapping;
    consumers #include the header

New position_history_derivation_kernel.cu (post-loop derivation for
cost_net_sharpe consumer in Wave 3b):
  - Reads actions_history_buf, reconstructs per-bar position_history
    (-1/0/+1), side_ind (1.0 on position change), rt_ind (1.0 on
    transition-to-flat from non-flat) via sequential walk (single
    block per window, no atomicAdd per feedback_no_atomicadd)
  - New launcher launch_sp15_position_history_derivation in
    gpu_dqn_trainer.rs
  - New cubin manifest entry in build.rs
  - 1 oracle test covering 8-bar Short→Hold→Long→Hold→Flat→Long→Flat→
    Short sequence; expected position/side_ind/rt_ind triples match
    hand-computed values

Atomic per feedback_no_partial_refactor for the ISV-contract change
(every consumer of slots 409/410/412/416 migrated in this commit; their
consumers were the 4 oracle tests, all migrated). The orphan launcher
transient state for the 5 baselines + derivation kernel is explicitly
the 3a/3b split point — production callers land in 3b's
GpuBacktestEvaluator::new constructor signature change.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-06 23:03:30 +02:00
jgrusewski
334b496647 feat(sp15-wave2): fused post-SP11 reward-axis composer (layered architecture)
Closes the deferred-consumer gap left by Phase 3.1
(r_quality_discipline_split_kernel, commit 5d36f3238), Phase 3.3
(dd_penalty_kernel), and Phase 3.5.2 (dd_asymmetric_reward_kernel) —
three SP15 reward-axis kernels that landed only as standalone scalar
producers awaiting deferred consumer wiring per
feedback_no_partial_refactor.md.

Wave 2 chooses Option β (layered, composable post-modifier) over
Option α (replace SP11 entirely): SP11 B1b stays canonical
"trader-quality" composer that writes out_rewards; the SP15 reward-axis
composition becomes a fused PER-(i,t) post-modifier read-modify-writing
the same buffer in place. This preserves SP11's z-score mag-ratio
contract (canary tests untouched) while landing all three deferred SP15
consumers atomically with zero parallel paths.

Architecture (Q1-Q4 user resolutions):
  Q1: SP12 caps stay as state_layout.cuh macros (REWARD_NEG_CAP=-10,
      REWARD_POS_CAP=+5), NOT lifted to ISV — spec'd constants per the
      SP12 v3 design, not adaptive bounds.
  Q2: Both on-policy + CF slots get the same DD-aware shaping. CF reward
      = w_cf × r_cf (SP11 controller weight already applied) is composed
      via the same helpers as on-policy. DD context is per-step, not
      per-action.
  Q3: New slot_completed_normally[N*L] flag preserves SP11's early-return
      semantics (data-end at experience_kernels.cu:2142 → reward=0.0;
      blown-account at :2203 → reward=-10.0). Fused kernel skips slots
      where flag==0.
  Q4: alpha_split_producer_kernel OWNS the warm-count increment (per-step
      scalar producer; the fused kernel has 2*N*L threads and would
      over-tick by that factor if it owned the increment).

Per-step launch sequence in gpu_experience_collector.rs (after Phase
1.3.b's dd_state launch): experience_env_step → alpha_split_producer
(reads grad-norm slots 418/419, writes ALPHA_SPLIT slot 417, increments
warm-count) → compute_sp15_final_reward (fused per-(i,t) over [N*2*L]
slots, applies α-blend → DD asymmetric → DD penalty → SP12 cap, writes
back to out_rewards in place).

experience_env_step signature change — 2 new output params:
  r_discipline_out: float* [N*L] — per-step REGRET_EMA mirror, written
    at end of normal reward composition.
  slot_completed_normally_out: int* [N*L] — 0 default at entry, set to
    1 only on the path that reaches out_rewards[out_off] = reward.

3 deleted kernel files:
  - r_quality_discipline_split_kernel.cu (composer + producer; replaced
    by renamed alpha_split_producer_kernel.cu keeping ONLY the producer
    with the moved warm-count increment).
  - dd_penalty_kernel.cu (replaced by sp15_dd_penalty __device__ helper).
  - dd_asymmetric_reward_kernel.cu (replaced by sp15_dd_asymmetric_reward
    helper).

3 new files:
  - alpha_split_producer_kernel.cu (per-step scalar producer of α from
    grad-norm ratio, with warm-count increment moved here per Q4).
  - sp15_reward_axis_helpers.cuh (4 __device__ inline helpers:
    sp15_alpha_blend, sp15_dd_asymmetric_reward, sp15_dd_penalty,
    sp15_apply_sp12_cap).
  - compute_sp15_final_reward_kernel.cu (fused per-(i,t) parallel over
    [N*2*L] slots — α-blend + DD-asymmetric + DD-penalty + SP12 cap,
    skips early-return sentinel slots).

3 deleted launchers + 3 deleted CUBIN statics in gpu_dqn_trainer.rs:
  - launch_sp15_r_quality_discipline_split (composer scalar variant).
  - launch_sp15_dd_penalty.
  - launch_sp15_dd_asymmetric_reward.
  - SP15_R_QUALITY_DISCIPLINE_SPLIT_CUBIN.
  - SP15_DD_PENALTY_CUBIN.
  - SP15_DD_ASYMMETRIC_REWARD_CUBIN.

2 new launchers + 2 new CUBIN statics:
  - launch_sp15_final_reward + SP15_FINAL_REWARD_CUBIN.
  - SP15_ALPHA_SPLIT_PRODUCER_CUBIN (the retained launch_sp15_alpha_split_producer
    now loads this).

GpuExperienceCollector: 2 new CudaSlice fields
(r_discipline_per_sample, slot_completed_normally_per_sample) +
sp15_alpha_warm_count_dev_ptr field + set_sp15_alpha_warm_count_ptr
setter + Step 5b launch block.

State reset registry — no new entries: r_discipline +
slot_completed_normally are per-step ephemeral (defaulted at every
kernel entry); cross-fold leakage impossible. Existing Phase 3.1/3.3/
3.5.2 ISV slot entries cover the rest.

6 new oracle tests in sp15_phase1_oracle_tests.rs::mod gpu drive
compute_sp15_final_reward_kernel directly with hand-crafted buffers:
  - final_reward_alpha_blend_at_cold_start (Stage 1)
  - final_reward_dd_penalty_above_threshold (Stage 3)
  - final_reward_dd_asymmetric_gain (Stage 2 + R_GAIN_DD_BOOST diag)
  - final_reward_dd_asymmetric_loss (Stage 2 asymmetric guard)
  - final_reward_sp12_cap_clamps_both_directions (Stage 4)
  - final_reward_skips_early_return_slots (Q3 sentinel preservation)

9 deleted scalar-kernel oracle tests:
  - r_split_uses_sentinel_alpha_at_cold_start
  - r_quality_subtracts_explicit_cost
  - dd_penalty_quadratic_above_threshold + dd_penalty_zero_below_threshold
  - dd_asymmetric_reward_gain_amplified_by_dd_pct +
    dd_asymmetric_reward_loss_unchanged + dd_asymmetric_reward_no_op_at_ath
  (the new fused-kernel tests cover the same behavioral surface
  end-to-end through the in-place RMW path).

Verified: SQLX_OFFLINE=true cargo check -p ml --features cuda clean (18
unrelated warnings); all 29 SP15 phase1 oracle tests pass on RTX 3050 Ti
(includes the 6 new fused-kernel tests + 17 retained tests + 6 old);
SP11 mag-ratio canary tests still untouched (no canary-test renames or
deletions); ml lib suite holds 946 pass / 13 fail = baseline.

Hard rules: feedback_no_partial_refactor (3 phases' deferred consumers +
3 deletions + 3 new files + 6 new tests + audit doc all in this commit;
no parallel paths, no feature flags), feedback_wire_everything_up
(closes 3 SP15 phase orphan launchers atomically), feedback_no_legacy_aliases
(deletions land in same commit as replacement; no compatibility shim),
feedback_no_atomicadd (fused kernel is per-(i,t) parallel, pure scalar
arithmetic), pearl_audit_unboundedness_for_implicit_asymmetry (gain-only
DD multiplier + asymmetric NEG/POS caps preserve loss aversion),
pearl_symmetric_clamp_audit (SP12 cap is bilateral via fmaxf/fminf even
though bounds are intentionally asymmetric per spec),
pearl_no_host_branches_in_captured_graph (new launches happen inside
collect_experiences_gpu::launch_timestep_loop per-step, OUTSIDE the
experience-fwd CUDA Graph capture region — same precedent as Phase
1.3.b's dd_state launch).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-06 22:22:33 +02:00
jgrusewski
f01a292f6f feat(sp15-p3.5.b+3.5.3.b): wire hold_floor (inline) + cooldown mask into experience_action_select
Phase 3.5 (hold_floor_kernel) + Phase 3.5.3 (cooldown_kernel) landed
the producer + state machinery; both deferred the action-selection
consumer wiring. This task wires both atomically.

Architectural decision: hold_floor is now an INLINE __device__
computation inside experience_action_select reading ISV slots
426/427/428/429 directly. The standalone hold_floor_kernel.cu +
launch_sp15_hold_floor + HOLD_FLOOR_CUBIN are deleted — launching a
kernel to write one f32 just to read it back was unnecessary. ISV
slots + state_reset_registry entries remain; only the launch path is
removed per feedback_wire_everything_up + feedback_no_legacy_aliases.

Entropy source: per-step Shannon entropy of softmax(e_dir) computed
inline from the 4 e_dir floats already in registers (Pass 1 of the
Thompson direction selector). High entropy = uncertain policy → Hold
gets the floor lift; low entropy = confident policy → floor ≈ 0.

q_eff_dir scratch preserves e_dir for downstream consumers
(out_conviction, out_q_gaps, out_magnitude_conviction) — adding
hold_floor there would corrupt the Kelly-cap warmup floor with a
meta-confidence mask.

cooldown mask: when ISV[COOLDOWN_BARS_REMAINING=435] > 0,
action_select hard short-circuits to dir_idx = DIR_HOLD before
Pass 2 — sidesteps the temperature-blend numerics where a
finite-sentinel-on-non-Hold approach would let pure-Thompson (τ=1)
samples dominate the masked direction. Cooldown supersedes
hold_floor — when forcing Hold the floor is moot.

3 new oracle tests:
  - action_select_applies_hold_floor_inline (no cooldown)
  - action_select_forces_hold_during_cooldown
  - action_select_no_force_hold_when_cooldown_zero

Atomic per feedback_no_partial_refactor: action_select changes +
hold_floor_kernel deletion + cubin manifest update + 3 oracle tests +
audit doc all in this commit. No parallel paths, no feature flags.

Eliminates Phase 3.5 + Phase 3.5.3 deferred consumers. The
cooldown_kernel itself remains (it maintains the consecutive_losses
streak + decrements COOLDOWN_BARS_REMAINING per bar); only its
consumer is now wired.

Verified: cargo check -p ml --features cuda clean; ml lib suite
holds 946 pass / 13 fail = baseline; all 6 oracle tests pass on
RTX 3050 Ti (3 pre-existing cooldown + 3 new action_select).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-06 21:18:31 +02:00
jgrusewski
132609724e feat(sp15-p1.3.b): wire dd_state per-step launch + drop equity-recompute bug + env-0 canonical observable
Path A of the blocked 1.3.b investigation: fixes two architectural
issues atomically and wires the launcher.

(1) Bug fix: dd_state_kernel.cu was recomputing new_equity =
PS_PREV_EQUITY + pnl_step and writing it back, but experience_env_step
already maintains PS_PREV_EQUITY (experience_kernels.cu:3473-3475) —
wiring as-is would silently double-accumulate equity every step.
Kernel now READS PS_PREV_EQUITY / PS_PEAK_EQUITY only; does not
modify them. pnl_step parameter dropped from both kernel and
launcher signatures.

(2) Per-env shape decision: kernel is single-thread/single-block;
production has N envs but DD ISV slots [401..407) are scalars.
Picks 'env 0 as canonical observable' — kernel reads
pos_state[0 * PS_STRIDE + ...]. Per-env redesign (per-env tiles +
reduction kernel) deferred to Phase 1.3.b-followup if L40S smoke
shows single-env DD aggregation is insufficient.

(3) Wire-up: launch added at gpu_experience_collector.rs step 5b in
launch_timestep_loop, immediately after env_step writes PS_PREV_EQUITY,
outside the exp-fwd graph capture region (which ends at line ~3829,
well before env_step). Atomic per feedback_no_partial_refactor:
kernel signature change + oracle test update + launcher call site
update all in this commit.

Eliminates the Phase 1.3 orphan launcher per feedback_wire_everything_up.
Downstream Phase 3.3 / 3.5.2 / 3.5.4 / 3.5.5 readers will receive live
DD values when their consumer wiring lands.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-06 20:06:57 +02:00
jgrusewski
eda1eccb1a merge(sp15): bring phase2a (LobBar + behavioral test scaffold + 17 tests) into phase1
Brings in Phase 2A.1 (LobBar canonical ABI + 4 synthetic market generators),
Phase 2A.2 (oracle + harness + pre-commit hook), and Phase 2B (17 #[ignore]
behavioral test contracts) so the phase1 honest-numbers branch has access
to the LobBar (price, half_spread, ofi) ABI needed for Phase 1.2.b cost-net
sharpe consumer wiring.

Path 2 of the BLOCKED 1.2.b investigation: the cost-net kernel needs
GPU-resident streams (half_spread, ofi, rt_ind, side_ind, position) that
do not exist on phase1; Phase 2A.1's LobBar provides the canonical ABI.

Conflict resolution: docs/dqn-wire-up-audit.md — both branches prepended
entries; merged by keeping all three (Phase 2A.1 from phase2a, Phase 1.6,
Phase 1.7 from phase1). Phase 2A.1 entry placed above Phase 1.6 / 1.7 to
keep this region's audit ordering consistent (newer-first locally).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-06 19:43:46 +02:00