23 Commits

Author SHA1 Message Date
jgrusewski
e877e8aefe chore(sp18): archaeology cleanup — strip post-deletion markers + ISV_TOTAL_DIM giant docstring
Pure-comment follow-up to SP18 Phase 1 atomic deletion (3c318953a) and
Phase 2 device-fn introduction (8da8e2e58). Strips the "DELETED by SP18
Phase 1" markers + "RETIRED" docstrings + per-slot SP3→SP18 archaeology
that survived the focused deletion diff. Per feedback_no_legacy_aliases
("avoid backwards-compatibility hacks like removed-comment markers")
and CLAUDE.md ("don't add comments that explain what the code does"),
git history is the historical record — code shouldn't narrate its own
deletion.

Three categories cleaned:

1. ISV_TOTAL_DIM giant docstring in gpu_dqn_trainer.rs (Category 1):
   collapsed ~5000-char slot-by-slot SP3→SP18 history block into a
   13-line conceptual definition pointing readers at sp14_isv_slots.rs
   for the actual range definitions. Layout fingerprint seed
   (RESERVED_GAP_461_TO_468=SP18_P1_RETIRED;) is load-bearing and
   preserved.

2. RETIRED markers in sp14_isv_slots.rs + state_layout.cuh
   (Category 2): deleted "SP16 Phase 2 RETIRED 2026-05-09" header
   blocks, "HCS_* slots [462..468) RETIRED" markers, "SP13 [380..383)/
   SP16 [461..474) ALLOCATED but RETIRED per PP.4" prose, and the
   test-removal note before the surviving MHT block lock test.
   Tightened the SP16 T3 Wiener-α comment block to describe only the
   surviving MIN_HOLD_TEMPERATURE producer.

3. Stale doc-refs across the SP18 audit consumer set (Category 3):
   updated gpu_experience_collector.rs prose, sp5_isv_slots.rs SP13 v3
   description, gpu_dqn_trainer.rs slot-380 constructor block,
   state_reset_registry.rs HRC_*/MHT_*/TDB_* "mirrors slot N" refs +
   the registry-block introduction comment + the deleted-test marker
   at end of file, training_loop.rs launcher-call/diag-emit/
   dispatch-arms "DELETED" markers, build.rs cubin manifest "DELETED"
   marker, gpu_aux_trunk.rs trailing "HoldCostScaleUpdateOps DELETED"
   block, and tests/sp14_oracle_tests.rs leading "tests DELETED" marker.

Production-source grep verification (post-cleanup): zero matches for
HOLD_COST_SCALE_INDEX | HCS_TARGET | HCS_DIFF | HCS_PREV | HCS_SAMPLE
| HOLD_COST_CONTROLLER_GAIN | HOLD_COST_FLOOR_RATIO |
HOLD_COST_CEIL_RATIO | hold_cost_scale_update across crates/ml/src/
*.rs/*.cu/*.cuh.

Audit doc + script preserved per design:
- scripts/audit_sp18_consumers.sh grep patterns retain HOLD_COST_SCALE
  / HCS_* — those are the diagnostic tool.
- docs/sp18-wireup-audit.md historical sections retain references
  (audit doc IS the design-history record); fingerprint section
  regenerated to reflect post-cleanup hit counts (audit --check passes).
- docs/superpowers/specs/* + docs/superpowers/plans/* unchanged
  (spec/plan files retain HOLD_COST_SCALE references unchanged per
  feedback_trust_code_not_docs corollary — those are the past
  design-decision record).
- docs/dqn-wire-up-audit.md current-state header section gains a
  2026-05-09 archaeology-cleanup-pass entry citing the affected
  files + verification (per Pre-commit Invariant 7).

Verification:
- cargo check --workspace: clean (warnings only — pre-existing).
- cargo test -p ml --lib --no-run: compiles clean.
- cargo test --doc -p ml: 21 failures pre-existing (verified via stash;
  identical 21 failures on parent commit 8da8e2e58 — unrelated to
  this cleanup).
- bash scripts/audit_sp18_consumers.sh --check: passes.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-09 10:43:11 +02:00
jgrusewski
3c318953a7 feat(sp18 v2 P1.T3-T6): atomic delete SP13/SP16 hold_cost_scale chain
Deletes the entire reactive Hold-cost controller chain from SP13 P0a / SP16
P2 / SP16 T3 atomically — replaced by the SP18 D-leg structural opportunity-
cost reward landing in Phase 2.

Per DD7(c): observation chain preserved (slots 380, 382). Slot 380 becomes
constructor-init-only diagnostic at HOLD_COST_BASE=0.005, never updated.

18 consumer sites deleted (8 from plan checklist + 10 from A1-A10 audit
expansion):

Production code:
  - 3 reward subtraction sites (replaced with Phase 2 placeholder comments)
  - hold_cost_scale_update_kernel.cu (file deleted)
  - HoldCostScaleUpdateOps struct in gpu_aux_trunk.rs (~120 lines)
  - launch_hold_cost_scale_update in gpu_dqn_trainer.rs (forwarder)
  - Host controller block in training_loop.rs (~60 lines)
  - Slot constants [461..468) in sp14_isv_slots.rs
  - HOLD_COST_CONTROLLER_GAIN/FLOOR/CEIL constants in sp13_isv_slots.rs
  - state_layout.cuh mirror constants for [461..468)
  - build.rs cubin manifest entry
  - 7 fold-reset registry entries + 7 dispatch arms

Tests:
  - sp14_oracle_tests.rs lines 2173-2926 (11 tests + helpers, ~750 lines)
    — these include_bytes! the deleted cubin and cannot survive
  - lock_sp18_v2_pp4_retired_chain test (deleted, no inverse-contract
    replacement — locking a deletion is pointless ceremony per user call)

Layout fingerprint: range [461..468) is RESERVED gap (SP14-C.1 pattern;
ISV_TOTAL_DIM stays at 507 for checkpoint compatibility). Layout fingerprint
seed updated atomically: HOLD_COST_SCALE + HCS_* slots replaced with
RESERVED_GAP_461_TO_468=SP18_P1_RETIRED marker.

INTERIM STATE: 3 reward sites are now missing the per-bar Hold cost. This
is forbidden to L40S-dispatch until Phase 2 lands the
compute_sp18_hold_opportunity_cost device fn that replaces the deleted
subtraction. The audit doc records this hold.

Validation: cargo check --workspace clean; cargo test -p ml --lib passes
(state_reset_registry::every_fold_and_soft_reset_entry_has_dispatch_arm,
sp16_t3_wiener_welford_slot_layout_locked, sp18_combined_slot_layout_locked,
sp18_shrink_perturb_slot_layout_locked, sp18_fold_reset_entries_present —
all OK). Audit fingerprint: per-bar Hold-cost subtraction sites count
13 → 0 (smoking gun for complete deletion).

Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md
Audit: docs/sp18-wireup-audit.md (full A1-A10 expansion + post-deletion
       fingerprint)

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-09 09:51:03 +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
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
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
c0fc28e455 fix(sp14): delete warmup_gate — let variance-driven k_aux/k_q handle warmup (ISV-driven)
Per `feedback_isv_for_adaptive_bounds`, the hardcoded
`warmup_gate = (fold_step_counter / WARMUP_STEPS_FALLBACK).min(1.0)`
ramp violated the rule: adaptive bounds in ISV, never hardcoded
constants. The variance-driven k_aux/k_q sigmoid steepness already
provides warmup behavior intrinsically:

- High variance (cold-start, EMAs still moving) → k → K_MIN → flat
  sigmoid → gate ≈ 0.5 regardless of input. That IS the warmup.
- Low variance (settled) → k → K_BASE → sharp sigmoid → gates
  respond correctly to driver signals.

Adding a separate hardcoded step-counter multiplier on top was
double-counting + tuning-driven (the 1000-step threshold had no
principled basis). Removed entirely.

Removed (per `feedback_no_partial_refactor`, all atomically):
- `WARMUP_STEPS_FALLBACK` constant in `sp14_isv_slots.rs`
- `warmup_gate: f32` parameter in `alpha_grad_compute_kernel.cu`
- `gate1 * gate2 * warmup_gate` → `gate1 * gate2` in kernel
- `warmup_gate` arg from `launch_sp14_alpha_grad_compute`
- `fold_step_counter: usize` field on the trainer struct
- `fold_step_counter = 0` reset in `reset_for_fold`
- `fold_step_counter` init in trainer constructor
- `let warmup_gate: f32 = 1.0;` and `.arg(&warmup_gate)` from B.4
  oracle tests (4 launches: 2 in alpha_grad_schmitt_hysteresis,
  20 in alpha_grad_adaptive_beta loop)

Build: clean, 18 warnings (pre-existing baseline).
Tests: cargo test --no-run on sp14_oracle_tests succeeds.

Net result: EGF gate's warmup behavior now lives entirely in the
variance-driven k_aux/k_q sigmoid steepness controller (ISV slots
388/var_aux, 389/var_q). No hardcoded step counter. Honors
`feedback_isv_for_adaptive_bounds` and `pearl_controller_anchors_isv_driven`.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-05 22:07:55 +02:00
jgrusewski
857722e774 feat(sp14): B.11 — orchestrator wire-up for 3 EGF producer kernels + var_aux gap closure
Per-step launches (in graph capture order):
  1. Forward (existing)
  2. Action select (existing) → q_dir_logits available
  3. launch_sp14_q_disagreement_update → ISV[383, 384, 389]
  4. launch_sp14_alpha_grad_compute → ISV[385..395] (consumes q_disagreement)
  5. Backward (existing) — wire-col scale at B.10 reads ISV[393]

Per-epoch launch (end of epoch):
  6. launch_sp14_gradient_hack_detect → circuit breaker

α_short=0.3, α_long=0.05, α_var=0.05 per spec; warmup_gate derived from
steps_in_fold / WARMUP_STEPS_FALLBACK.

Var_aux producer gap closed (option C from B.4): alpha_grad_compute_kernel
now also writes ISV[VAR_AUX_INDEX=388] via Welford EMA against
(aux_dir_acc_short - aux_dir_acc_long). Adaptive k_aux is now functional
(was degenerate at K_BASE_AUX=20.0 constant pre-B.11). Closes the
"adaptive_k_aux currently degenerate" concern flagged in B.4 commit.

After this commit, the EGF pearl is FULLY ACTIVE end-to-end:
- Forward: aux signal feeds direction Q-head input (B.8/B.9)
- Backward: wire-col gradient gated by α_grad_smoothed (B.10)
- Producers: α_grad computed every step from real driver signals (B.11)
- Pre-B.11 force-closed gate (sentinel 0.0) → post-B.11 responsive gate

Build clean: 18 warnings pre-existing baseline, 0 new.
Tests: 4/4 P0b aux_w tests pass (no regression).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-05 21:11:26 +02:00
jgrusewski
6715ab4ea1 feat(sp14): B.8 — direction Q-head input dim SH2 → SH2+1 + fingerprint
Forward wire (weight-side only — dispatch consumer lands in B.9):
direction Q-head's first FC weight tensor `w_b0fc` (param-table index
17) input dim grows by 1 to accept aux_softmax_diff. The (rows, cols)
shape table at the trainer-init Xavier site is updated in lock-step
with `compute_param_sizes()`. Output dim (adv_h) and the bias `b_b0fc`
are unchanged — bias is per-output, not per-input.

`layout_fingerprint_seed()` entry renamed `PARAM_W_B0FC` →
`PARAM_W_B0FC_AUX1`, forcing the FNV-1a hash to bump per Invariant 8
(old checkpoints fail-fast at load via `check_layout_fingerprint`
instead of silently aliasing onto the new architecture). Per
`feedback_no_legacy_aliases`, no `_DEPRECATED` shim — straight
in-place rename.

New weight column zero-initialised in trainer construction (mirrors
the OFI column zeroing for w_b2fc/w_b3fc): the model starts ignoring
the new input and learns to use it through gradient descent. Xavier-
init on this column would inject day-0 noise the trunk would have to
denoise — let the EGF gate decide when the aux signal is trustworthy.

Atomic-migration consumers (`feedback_no_partial_refactor`) updated:

- `gpu_dqn_trainer.rs` — Xavier (rows, cols) table at index 17 grows
  to (adv_h, SH2+1); spectral-norm descriptor entry [4] for W_a1 grows
  in_dim from sh2 to sh2+1; spec_v_a1 power-iteration vector grows
  from sh2 to sh2+1 floats.
- `dqn/smoke_tests/gradient_budget.rs` — both `alloc_dueling`
  fixtures' slot 8 grow `cfg.adv_h * cfg.shared_h2` →
  `cfg.adv_h * (cfg.shared_h2 + 1)`.
- `docs/dqn-wire-up-audit.md` — new SP14 Layer B B.8 entry per
  Invariant 7.

Note: forward GEMM dispatch still uses `K = shared_h2` until B.9 lands
the concat-then-SGEMM consumer (per plan §2358). Until then the new
column reads as ignored padding; this is safe because (a) it's zero-
initialised, (b) GPU-only smoke tests are skipped on this CPU CI, (c)
the fingerprint bump invalidates any pre-SP14 checkpoint that would
attempt to load.

Test: `layout_fingerprint_bumps_after_sp14_wire` (CPU-only, in
`sp14_oracle_tests.rs`) hashes the pre-B.8 seed verbatim with the
single difference `PARAM_W_B0FC` (vs post-B.8 `PARAM_W_B0FC_AUX1`)
and asserts `LAYOUT_FINGERPRINT_CURRENT` differs — any silent revert
of the rename trips this test. Mirrors the `fingerprint_bumped_from_
pre_b1_1a` pattern from `sp13_layer_b_oracle_tests.rs`.

Verified:
- `SQLX_OFFLINE=true cargo check -p ml` clean, 18 warnings (baseline)
- `cargo test -p ml --test sp14_oracle_tests` 2 passed, 6 ignored (GPU)
- `cargo test -p ml --test sp13_layer_b_oracle_tests
   fingerprint_bumped_from_pre_b1_1a` still passes (sister test)

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-05 19:58:18 +02:00
jgrusewski
4527d8c852 feat(sp14): B.6 — dir_concat_qaux_kernel (pre-SGEMM forward-wire concat)
Mirrors mag_concat_qdir precedent (experience_kernels.cu:4560).
Concats [h_s2 ; aux_softmax_diff] into scratch buffer [B, SH2+1] for
the direction Q-head's first FC SGEMM.

aux_softmax_diff = softmax[b, 1] - softmax[b, 0] in [-1, +1] computed
inline; structurally bounded by softmax components per
pearl_bounded_modifier_outputs_require_structural_activation.

Pure per-thread map; no reduction; no atomicAdd. Does not read or write
ISV slots — purely data-movement. Launch order constraint: aux head
forward MUST complete before this concat reads aux_nb_softmax_buf
(enforced by orchestrator in B.10/B.11).

Test: dir_concat_qaux_correct verifies row-wise contiguous concat
+ correct softmax diff values for both 'down' (-0.8) and 'up' (+0.8)
synthetic aux predictions. B.3+B.4+B.5 regression: 5 GPU tests
unchanged (6 total pass).

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
2026-05-05 19:36:03 +02:00
jgrusewski
82fe6cea66 feat(sp14): B.5 — gradient_hack_detect_kernel (anti-mesa-opt circuit breaker)
Detects suspected gradient hacking: when gate1 is open AND aux_dir_acc
post-open-minimum drops > LOCKOUT_TRIGGER_DROP (0.05) below the Schmitt
open-threshold (target + SCHMITT_BAND = target + 0.03) AND q_disagreement
rises > LOCKOUT_TRIGGER_DIS_RISE (0.10) above the analytic random-alignment
baseline 0.5, simultaneously.

Action: force gate1_open_state = 0 (ISV[391]); set lockout_remaining = 2.0
epochs (LOCKOUT_EPOCHS). During lockout, gate1 stays force-closed
regardless of alpha_grad_compute_kernel output.

Tracks AUX_DIR_ACC_POST_OPEN_MIN (ISV[394]): running minimum of aux_dir_acc
since gate1 last opened; resets to 1.0 sentinel when gate closes naturally
or when circuit breaker fires.

Slot indices shifted +2 from original plan (SP13 closeout added
HOLD_RATE_TARGET=381 + HOLD_RATE_OBSERVED_EMA=382): Q_DIS_SHORT=383,
GATE1=391, POST_OPEN_MIN=394, LOCKOUT=395. Matches sp14_isv_slots.rs.

Single-thread state-machine kernel (threadIdx.x==0 guard); runs at end of
each epoch after alpha_grad_compute_kernel. No atomicAdd per
feedback_no_atomicadd.md.

1 oracle test: gradient_hack_circuit_breaker_fires verifies trigger
conditions (aux_drop=0.08 > 0.05, q_rise=0.15 > 0.10) cause lockout=2.0
and gate1 force-close=0.0.

B.3+B.4 regression: 4 GPU tests unchanged (5 total GPU pass, 1 host pass).

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
2026-05-05 19:29:22 +02:00
jgrusewski
49cdf90ecc feat(sp14): B.4 — alpha_grad_compute_kernel (EGF heart)
Single-thread state-machine kernel that is the heart of the Earned
Gradient Flow pearl. Reads driver signals from the global ISV bus,
runs Schmitt-trigger Gate 1, computes adaptive k_aux/k_q/β, evaluates
two sigmoids, multiplies with a host-supplied warmup gate, applies a
β-rate-limiter, and writes 7 outputs back to ISV.

Per-step pipeline:

1. Read aux_dir_acc (slot 373), q_disagreement (slot 383), Welford
   variance EMAs (388, 389, 390), persistent Schmitt state (391),
   alpha_smoothed_prev (393).
2. Compute adaptive k_aux = K_BASE_AUX/(1 + var_aux/VARIANCE_REF_AUX)
   and k_q analogously (B.2.5; floor at K_MIN = 1.0).
3. Run Schmitt-trigger Gate 1 state update (open at target+0.03,
   close at target-0.03; intentional discontinuity at transition is
   smoothed by the β rate-limiter downstream).
4. Evaluate Gate 1 sigmoid (aux competence, distance from threshold)
   and Gate 2 sigmoid (Q-aux disagreement vs analytic 0.5 baseline).
5. alpha_grad_raw = gate1 × gate2 × warmup_gate (structurally bounded
   to [0, 1] per pearl_bounded_modifier_outputs_require_structural_
   activation; no runtime clamp).
6. Update Welford variance of alpha_grad_raw → adaptive β (B.2.8;
   floor BETA_BASE = 0.5, ceiling BETA_MAX = 0.95).
7. alpha_grad_smoothed = β × prev + (1-β) × raw (rate-limited).
8. Write back 7 outputs: k_aux (385), k_q (386), β (387), var_alpha
   (390), gate1_state (391), alpha_raw (392), alpha_smoothed (393).

Sigmoid arguments clipped to [-30, 30] before __expf for fp32
overflow guard (precision-neutral; sigmoid saturates bit-equal at
those bounds).

Per pearl_bounded_modifier_outputs_require_structural_activation:
sigmoid composition produces structurally-bounded [0, 1] output.

KNOWN LIMITATION: as of B.4 landing, NO upstream kernel writes
ISV[388] (AUX_DIR_ACC_VARIANCE_EMA). The grep at status-report time
finds only the sp14_isv_slots.rs declaration. Effect: var_aux stays
at sentinel 0.0 forever, so k_aux is degenerate-but-non-fatal at
K_BASE_AUX (constant). Gate 1 still works, the sigmoid just doesn't
soften under noisy aux_dir_acc. To be resolved in B.11 producer-
chain orchestrator OR a separate fix-up task that adds a Welford-
variance update next to the existing AUX_DIR_ACC_SHORT_EMA producer.
var_q (389) IS written by q_disagreement_update_kernel (B.3), so
adaptive k_q is fully functional from B.4 onward.

Slot indices hardcoded inside the kernel via const int locals — must
match crates/ml/src/cuda_pipeline/sp14_isv_slots.rs (and 372/373
from sp13_isv_slots.rs). The plan originally documented 381/383/
384/385/386/387/388/389/390/391 for SP14 slots; the actual values
are +2 because SP13 closeout added HOLD_RATE_TARGET=381 +
HOLD_RATE_OBSERVED_EMA=382 after the plan was written.

Tests (RTX 3050 Ti pass; B.3's 2 tests still pass — no regression):

- alpha_grad_schmitt_hysteresis: 4-step trajectory verifies the
  closed→open→open→closed transition. Closed at aux=0.55 (below
  open=0.58); opens at aux=0.60; stays open at aux=0.54 (in
  hysteresis band [close=0.52, open=0.58]); finally closes at
  aux=0.50 (below close=0.52).
- alpha_grad_adaptive_beta: 20-oscillation regime verifies β grows
  above β_base=0.5 and remains bounded by β_max=0.95.

docs/dqn-wire-up-audit.md updated per Invariant 7 with full B.4
behaviour contract, per-step pipeline, single-thread launch
convention, sigmoid clip rationale, Schmitt discontinuity note,
and the var_aux Known Limitation.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-05 19:14:48 +02:00
jgrusewski
d3a35cc6e9 feat(sp14): B.3 — q_disagreement_update_kernel with K=4↔K=2 mapping
First of four producer kernels in the Earned Gradient Flow pearl chain.
Per-step computes the K=4↔K=2 mapped argmax mismatch between the Q-head's
4-way direction action (DIR_SHORT/DIR_HOLD/DIR_LONG/DIR_FLAT) and the
aux head's 2-way next-bar prediction (down/up), updates fast + slow
EMAs of the disagreement rate, and updates a Welford-style variance
EMA on the same signal in a single launch.

Mapping (per state_layout.cuh):
  DIR_SHORT (=0)  → aux down (=0)        contributes
  DIR_HOLD  (=1)  → masked                no contribution
  DIR_LONG  (=2)  → aux up   (=1)        contributes
  DIR_FLAT  (=3)  → masked                no contribution

Hold and Flat are masked because they represent "no NEW directional
commitment" (Hold = keep prior position; Flat = close all positions).
Penalising the aux head for not matching them would conflate
position-management actions with directional predictions.

Block tree-reduce on numerator + count separately, then divide and
update EMAs in a single thread (tid == 0). No atomicAdd per
feedback_no_atomicadd.md. Pure GPU compute per
feedback_no_cpu_compute_strict.md.

Pearl-A first-observation: when ISV[Q_DISAGREEMENT_SHORT_EMA=383] AND
ISV[Q_DISAGREEMENT_LONG_EMA=384] both equal sentinel 0.5f AND batch
has at least one valid contribution (total_cnt > 0), both EMAs are
replaced directly with the first observation per
pearl_first_observation_bootstrap.md. The 0.5f exact-match is safe
because 0.5 is exactly representable in IEEE 754 single precision
(mantissa = 1.0, exponent = -1). The Welford variance EMA at slot
389 drives the adaptive k_q sigmoid steepness consumed by the
alpha_grad_compute_kernel in B.4.

Slot indices are hardcoded inside the kernel via #define — must match
crates/ml/src/cuda_pipeline/sp14_isv_slots.rs (currently 383, 384, 389).
The kernel header documents the coupling explicitly.

Launch contract:
  grid_dim = (1, 1, 1)
  block_dim = (256, 1, 1)
  shared_mem_bytes = 2 * 256 * sizeof(float) = 2048

A shared_mem_bytes = 0 launch reads garbage and corrupts the EMA — the
kernel header documents the launcher requirement; oracle tests pass
2048 explicitly.

Files:
  - crates/ml/src/cuda_pipeline/q_disagreement_update_kernel.cu (NEW)
  - crates/ml/build.rs — register cubin in kernels_with_common
  - crates/ml/tests/sp14_oracle_tests.rs — append #[cfg(feature="cuda")]
    mod gpu with cubin handle + 2 GPU oracle tests
  - docs/dqn-wire-up-audit.md — SP14 B.3 section (Invariant 7)

Tests (both pass on RTX 3050 Ti):
  - q_disagreement_k4_k2_mapping: 8-row batch with 2 agreements,
    2 disagreements, 4 masked Hold/Flat → first-obs Pearl-A replaces
    both EMAs with batch_mean = 0.5 (asserted within 1e-4)
  - q_disagreement_all_hold_no_contribution: all-Hold edge case;
    total_cnt = 0 so batch_mean = 0/1 = 0; sentinel-bootstrap guard
    (total_cnt > 0) keeps EMA from collapsing to 0; current kernel
    blends to 0.35, test bound [0.0, 0.5] tolerant of either current
    blend or future skip-update refinement, asserts is_finite()

Wire-up status: producer kernel exists and is exercised only by the
oracle tests. The Rust launcher and graph-capture integration land in
B.7+ alongside the consumer (alpha_grad_compute in B.4 reads slots
383/384/389). This is one producer kernel of four (B.3 q_disagreement,
B.4 alpha_grad_compute, B.5 gradient_hack_detect, B.6 dir_concat_qaux);
known-orphan for the duration of the producer chain per
feedback_wire_everything_up.md (the same-commit wire-up rule is
relaxed for atomic chained-producer-consumer landings, with each
orphan documented in the audit doc; this orphan is acknowledged in
docs/dqn-wire-up-audit.md SP14 B.3 section).

Build + test verification:
  - SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --features cuda
    → clean (only the 18 pre-existing warnings)
  - SQLX_OFFLINE=true cargo test -p ml --test sp14_oracle_tests
    set_aux_weight → host-only A.2 still passes (no shared dependency)
  - SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml
    --test sp14_oracle_tests --features cuda q_disagreement
    -- --ignored --nocapture → 2 PASS

Slot indices reflect the SP13-closeout +2 shift documented in
sp14_isv_slots.rs (the original plan's 381/382/387 became 383/384/389
because HOLD_RATE_TARGET=381 and HOLD_RATE_OBSERVED_EMA=382 landed
between when the plan was written and B.1 was implemented).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-05 19:04:18 +02:00
jgrusewski
731cae4c80 fix(sp14): A.2 — lift set_aux_weight clamp to SP13 P0b range [0.15, 1.5]
The pre-fix clamp at gpu_dqn_trainer.rs:14722 was [0.05, 0.3] — the
SP11-era cap. P0b's controller computes aux_w in [0.15, 1.5] (base
0.5 × [0.3, 3.0]) but the setter silently chopped everything above
0.3, masking the deficit-amplification term entirely.

Smoke A trace confirmed:
  Fold 0/1: raw aux_w = 0.66-0.80 → clamped to 0.30 (deficit invisible)
  Fold 2:   raw aux_w = 0.164 (stagnation; below clamp) → 45% deficit

Post-fix: deficit-amp term `(1 + 5 × deficit)` actually expresses
through to the trainer. Fold 2 stagnation will get the designed
floor 0.15 instead of being capped at 0.3 — but the upper range
1.5 also opens, so deficit-amp can pull aux_w up when accuracy is
below target.

Constants imported from sp13_isv_slots.rs (AUX_W_BASE=0.5,
AUX_W_HARD_FLOOR_RATIO=0.3, AUX_W_HARD_CEIL_RATIO=3.0). No new
slots; existing constants exposed as the clamp bounds.

Test: set_aux_weight_clamp_range verifies constants resolve correctly.

A.1 (C51 atom-probability floor) deferred — Phase 0 verification
found the spec's stated `−log(p)/p · ∂p/∂z` divide does NOT exist
in c51_grad_kernel.cu at HEAD 037c24116; actual kernel uses the
numerically-stable `expf(lp) - proj`. The 1109 GRAD_CLIP_OUTLIER
events in Smoke A are real but their mechanism is different. Will
be re-spec'd as a separate task post-SP14.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-05 17:52:18 +02:00