gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity
4,771 Commits 13 Branches 0 Tags
bfc3ffa9dc5fa5a14fa98dccab4347809dcd98fa
Commit Graph

2374 Commits

Author SHA1 Message Date
jgrusewski
bfc3ffa9dc diag(sp15-wave5): in-capture CAPTURE_PHASE_* checkpoints 
Smoke train-fp7xx printed all 16 Phase-6/7/8 checkpoints clean through
PER_PRIORITY_DONE. CAPTURE_DONE missing, exit changed 139→143 (SIGTERM)
— capture_training_graph hangs or takes ~40s past PER_PRIORITY_DONE
before Argo terminates the pod.

Adds 14 CAPTURE_PHASE_* checkpoints across the 12 child captures +
parent compose:
  BEGIN / PER_SAMPLE / COUNTERS / SPECTRAL / FORWARD / DDQN / AUX
  POST_AUX / ADAM_GRAD / ADAM_UPDATE / MAINTENANCE / IQL_MODULATE
  PER_PRIORITY / CHILDREN_STORED / PARENT_COMPOSED

Diagnostic-only.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-07 12:47:23 +02:00
jgrusewski
a8d6c33040 diag(sp15-wave5): extend STEP0_PHASE_* checkpoints into Phase 6/7/8 
Smoke train-xq9hg got past all 13 original checkpoints (BEGIN through
NAN_CHECKS_DONE) cleanly. SEGV is downstream in Phase 6/7/8.

Adds 16 more checkpoints: TLOB_BWD_ADAM, MAMBA2_BWD, MAMBA2_ADAM,
OFI_EMBED_BWD, OFI_EMBED_ADAM, PRUNING, BRANCH_GRAD_BALANCE, GRAD_NORM,
ADAM_OPS, Q_MAG_BIN, Q_DIR_BIN, ISV_UPDATE, MAINTENANCE, IQL_MODULATE,
PER_PRIORITY, CAPTURE.

Diagnostic-only — no behavior change.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-07 12:36:20 +02:00
jgrusewski
e9d9afbd61 diag(sp15-wave5): stderr checkpoints in run_full_step step-0 ungraphed path 
L40S smoke train-jfbzr (commit 23e9a1f78) segfaults at exit 139 in the
first run_full_step invocation, after "GPU training guard initialized
(epoch loop)" and before any HEALTH_DIAG output. Static analysis
debugger couldn't reproduce locally (test_data/futures-baseline/ lacks
.fxcache).

This commit adds eprintln! checkpoints at each phase boundary in the
ungraphed step-0 path of FusedTrainer::run_full_step (PER sample,
PopArt, counters, spectral norm, TLOB forward, forward_main, DDQN,
aux_ops, post_aux, NaN checks).

stderr is line-flushed (vs stdout buffered through tracing JSON
formatter), so the last printed checkpoint identifies the SEGV site.
Each fires once per fold's first step (graph capture absorbs subsequent
steps), so log overhead is minimal.

Diagnostic-only commit — no behavior change. Will be reverted after the
next L40S smoke localizes the SEGV and the root-cause fix lands.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-07 12:21:12 +02:00
jgrusewski
23e9a1f78c fix(cuda): compute_expected_q stride 13 vs denoise_target_q_buf size 12 — OOB writes threads 60-63 
Pre-existing latent bug surfaced by compute-sanitizer after the SP15
NULL-pointer fix at 2e37af29d unblocked the cascade. Threads 60-63
were writing past denoise_target_q_buf's end every batch.

Sanitizer evidence (pre-fix on 2e37af29d):
  Invalid __global__ write of size 4 bytes
    at compute_expected_q+0x2480
    by thread (60..63, 0, 0) in block (0, 0, 0)
    Access at <addr> is out of bounds (45/97/149/201 bytes after
    nearest allocation of size B*12*4 bytes)
  Host backtrace: GpuDqnTrainer::compute_denoise_target_q
    → submit_post_aux_ops → run_full_step

Root cause — semantic mismatch between writer stride and buffer size:
  - compute_expected_q writes q_values[i*total_actions + a] with
    total_actions = b0+b1+b2+b3 = 4+3+3+3 = 13 (4-direction factored)
  - denoise_target_q_buf was allocated b*12 (legacy 3+3+3+3 layout)
  - threads 60..63 wrote slot 12 of samples (60..63 % batch_size)
    past the buffer end every step

The downstream q_denoise_backward + denoise_loss_grad kernels read
Q_target[b * D + i] with hardcoded D=12 because the diffusion denoiser
MLP itself only has 12 output slots (W2[12,24] + b2[12]) — its 12 are
the q_coord_buf's post-cross-branch-attention narrowed output, NOT a
clean subset of the 13 raw Q-actions.

The exact same fix pattern was already applied to the sibling
q_var_buf_trainer allocation in the prior SP4 audit (see comment block
at gpu_dqn_trainer.rs:21620-21630 referencing the identical OOB at
threads 60..63); denoise_target_q_buf 8 lines below was missed because
it was guarded by the SP15 NULL-pointer ILLEGAL_ADDRESS that fault-
stopped the cascade before this OOB could fire — 2e37af29d removed the
upstream ILLEGAL_ADDRESS, surfacing the latent OOB.

Fix architecture (Option A — pad buffer to total_actions, pass stride
to consumer; same pattern as q_var_buf_trainer):
  1. Widen denoise_target_q_buf from b*12 to b*total_actions (= b*13)
     to match compute_expected_q's writer stride.
  2. Add refined_stride / target_stride / input_stride parameters to
     q_denoise_backward kernel; the kernel still computes D=12 per-
     sample (denoiser MLP fixed width) but addresses each input buffer
     at its own per-sample stride.
  3. Add refined_stride / target_stride parameters to denoise_loss_grad
     kernel (same pattern; used by launch_q_denoise_backward_cublas).
  4. Update both Rust launchers (launch_q_denoise_backward,
     launch_q_denoise_backward_cublas) to pass refined_stride=12 (q_coord
     and q_input are post-attention narrowed),
     target_stride=total_actions=13.
  5. denoise_q_input_buf STAYS at b*12 — it's a snapshot of q_coord_buf
     (also b*12) via snapshot_pre_denoise_q's DtoD copy; never written
     by compute_expected_q.
  6. Flat-scan consumers (SP4 target_q_p99_update producer + SP3 slot
     46 threshold-check + dqn_clamp_finite_f32) UNCHANGED — they
     consume the buffer flat via .len(); widening from b*12 to b*13 is
     monotone (one more valid Q-value per sample in the histogram).

Atomic per feedback_no_partial_refactor: 2 kernel signatures + buffer
allocation + 2 launch sites + struct doc comments + SP3 slot 46 doc +
audit doc — all in one commit. Every consumer of the writer's stride
migrates simultaneously.

Verification:
  - SQLX_OFFLINE=true cargo check -p ml --features cuda: clean
  - compute-sanitizer (RTX 3050 Ti): 0 Invalid __global__ errors at
    compute_expected_q post-fix (down from 4 per training step in
    baseline — re-verified on 2e37af29d to confirm the diff)
  - SQLX_OFFLINE=true cargo test -p ml --features cuda --lib: holds
    parent baseline 947 pass / 12 fail (same 12 pre-existing failures)

Files touched:
  crates/ml/src/cuda_pipeline/experience_kernels.cu
    — q_denoise_backward + denoise_loss_grad kernel signatures
  crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
    — alloc widen + 2 launcher updates + 4 doc comment updates
  docs/dqn-wire-up-audit.md
    — audit entry

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-07 11:39:34 +02:00
jgrusewski
2e37af29d4 fix(sp15-p1.3.b-followup-OOB): wire ISV bus + SP15 control pointers BEFORE first collect — fixes "load sp15_dd_state cubin" CUDA_ILLEGAL_ADDRESS 
Root cause: Wave 4.1b (s1_input_dim 102→103, eb9515e41) introduced
`bn_tanh_concat_dd_kernel` which reads `isv[DD_PCT_INDEX=406]` from the
collector's `isv_signals_dev_ptr`. Wave 4.1b's atomic refactor docstring
claimed "guarded by the captured-graph wiring in training_loop.rs:859
which sets the bus pointer BEFORE the first epoch's forward graph
capture" — but ordering audit shows that's wrong:

  Phase 1c (line 580): init_gpu_experience_collector
                        → collector.isv_signals_dev_ptr = 0 (NULL)
  Phase 2  (line 730): collect_gpu_experiences_slices
                        → bn_tanh_concat_dd_kernel reads isv[406]
                          = NULL+1624 = 0x658 → ILLEGAL_ADDRESS
  Phase 4  (line 859+): set_isv_signals_ptr (TOO LATE)

The error surfaces as "load sp15_dd_state cubin: ILLEGAL_ADDRESS" via
sticky-cascade (the OOB happened in the prior kernel; the next CUDA
call — dd_state cubin load — surfaces the cascade).

Reproduced locally via compute-sanitizer on
test_no_hang_single_epoch (RTX 3050 Ti). Sanitizer pinpointed
"Invalid __global__ read of size 4 bytes at bn_tanh_concat_dd_kernel
+0x4d0 ... Access at 0x658 is out of bounds" — exactly slot 406 * 4.

Fix: move the ISV / SP15 control pointer wiring (set_isv_signals_ptr,
set_sp15_alpha_warm_count_ptr, set_sp15_plasticity_target, PER
buffer's set_isv_signals_ptr + set_sp15_dd_trajectory_per_env_ptr +
set_sp15_per_env_dims) into init_gpu_experience_collector BEFORE the
collector is moved into self.gpu_experience_collector. This guarantees
the FIRST collect_experiences_gpu call (epoch 0, Phase 2) sees
non-NULL pointers — required because cudarc's graph capture bakes the
arg values into the captured graph at capture time, so subsequent
re-wiring has no effect on the captured replay path.

The per-epoch re-wiring at lines 855-957 of the outer epoch loop is
now redundant (idempotent setters re-applying the same stable
pointers) but kept in place per feedback_no_partial_refactor for
explicit per-epoch refresh semantics — these underlying pointers are
stable for the trainer's lifetime so re-setting is a no-op, and
removing the per-epoch call would scatter the wire-up logic across a
non-obvious pre-init / per-epoch split.

Verification:
  * 30/30 SP15 phase 1 oracle tests pass under compute-sanitizer
    memcheck with 0 errors (unchanged baseline)
  * test_no_hang_single_epoch under sanitizer: SP15 dd_state OOB is
    GONE; the first OOB now surfaces at compute_expected_q+0x2480
    (denoise_target_q_buf stride 12 vs total_actions=13 — a SEPARATE
    pre-existing bug previously masked by the SP15 OOB; reported but
    out of scope for this commit per single-atomic-fix discipline).
  * compute-sanitizer "Access at 0x658" matches the
    bn_tanh_concat_dd_kernel isv[406] read path; production training
    on L40S/H100 should now reach the next-layer issue.

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

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-07 11:15:30 +02:00
jgrusewski
81a9319d84 fix(sp15-wave3b-followup): migrate evaluate_baseline.rs 3 GpuBacktestEvaluator::new sites — Wave 3b missed the example binary 
Wave 3b (4320820ae) added a 6th window_lob_bars parameter to
GpuBacktestEvaluator::new and migrated 3 production lib call sites + 6
test call sites, but the lib-only `cargo check -p ml --features cuda`
validation didn't catch the 3 example-binary call sites in
evaluate_baseline.rs (lines 1344, 1641, 1802). Argo's ensure-binary
step compiles all 4 example binaries (hyperopt_baseline_rl,
train_baseline_rl, evaluate_baseline, precompute_features), and
evaluate_baseline failed with E0061 (wrong number of args).

This commit migrates all 3 sites to construct zero-OFI LobBar SoA
inline (eval-path lacks per-bar OFI features — same degradation
pattern as the PPO hyperopt adapter Wave 3b D4 resolution; OFI-impact
term degrades to 0 while commission + half-spread × position still
apply).

Validated: `cargo check -p ml --features cuda --all-targets` clean
(was --features cuda only before — now expanded to catch
example/test/bin compile-breaks).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-07 10:25:50 +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
eb9515e41c feat(sp15-wave4.1b): consumer migration — s1_input_dim 102→103, GRN w_s1 reshape, 4 forward + 3 backward sites 
Atomic consumer migration that flips production callers of bn_tanh_concat_kernel
over to Wave 4.1a's bn_tanh_concat_dd_kernel and bumps s1_input_dim from 102 to
103 across the entire trunk forward + backward path. Eliminates the documented
Wave 4.1a transient orphan.

Changes:
- s1_input_dim formula bump (bn_dim + portfolio_dim → bn_dim + portfolio_dim + 1)
  at compute_param_sizes, trainer ctor's CublasGemmSet, xavier_init_params_buf,
  the experience collector's CublasGemmSet, and CublasBackwardSet::new for the
  backward gemm cache. cuBLAS gemm caches re-key automatically (fresh HashMap).
- GRN w_a_h_s1[0] / w_residual_h_s1[4] reshape [shared_h1, 102] → [shared_h1, 103]
  via compute_param_sizes + xavier_init's fan_dims. Xavier-uniform init covers
  the new dd_pct column (bounded [0,1] — Xavier's small-magnitude assumption is
  appropriate; differs from SP14's aux_softmax_diff zero-init which was driven
  by the bidirectional ±1 range).
- bn_concat_dim() accessor +1 (TLOB backward row stride).
- 5 concat_dim local-var bumps (1 alloc + 3 forward + 2 backward + 1 in
  experience collector).
- 4 forward-call migrations to launch_sp15_bn_concat_dd: DDQN argmax pass
  (~27158), online forward (~27467), target forward (~27666), experience
  collector forward (~3853). Each takes self.isv_signals_dev_ptr; the kernel
  reads ISV[DD_PCT_INDEX=406] on-device and broadcasts.
- 3 GRN backward sites (main, ensemble, CQL) flow through encoder_backward_chain
  which uses s1_input_dim — bumped automatically. dd_pct column gradient is
  silently discarded by vsn_d_gated_state_portfolio_pad_kernel (reads
  [bn_dim..bn_dim+portfolio_dim) only) and bn_tanh_backward_kernel (reads
  [0..bn_dim) only). Correct: dd_pct sources from ISV bus, no learnable input.
- Legacy bn_tanh_concat_kernel field DELETED from trainer struct alongside its
  loader, tuple-element, destructuring, assignment (5 mechanical sites for the
  one dead field). Function tuple shrinks 44→43 elements. Kernel symbol stays
  in the cubin source for SP15 oracle parity tests.
- mag_concat / OFI concat audit verdict: DECOUPLED from s1_input_dim. They
  widen shared_h2, not the trunk INPUT dim.
- test_gpu_backtest_evaluator_state_dim_calculation migrated to assert
  STATE_DIM == 128 (was 96, stale per feedback_trust_code_not_docs).

Atomic per feedback_no_partial_refactor: every consumer of s1_input_dim and
bn_concat_buf row-stride migrated together. Eliminates Wave 4.1a transient-
orphan launcher per feedback_wire_everything_up. Legacy field deleted per
feedback_no_legacy_aliases.

Tests: cargo check clean (18 pre-existing unrelated warnings). Wave 4.1a
oracle parity test (bn_tanh_concat_dd_kernel_writes_dd_pct_column) still
passes. ML lib suite went from 945 pass / 14 fail (pre-Wave-4.1b baseline) to
947 pass / 12 fail post-Wave-4.1b — improved by +2 (state_dim_calculation
migration + ensemble checkpoint round-trip flake resolved).

Refs: SP15 Wave 4.1a (a8da1cb9c), pearl_no_host_branches_in_captured_graph,
feedback_no_partial_refactor, feedback_wire_everything_up,
feedback_no_legacy_aliases, feedback_isv_for_adaptive_bounds,
feedback_trust_code_not_docs.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-07 00:33:19 +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
d7f60d4dd7 feat(sp15-p1.6.b+1.7.b): wire dev-eval (Q8 final-fold) + test-eval (per-fold) into trainer 
Phase 1.6 (CLI flags + dev_features/holdout_features stash) and Phase
1.7 (set_test_data_from_slices observer + test_features stash) landed
the data-flow scaffolding; both deferred the actual eval consumer.

This task wires both atomically as parallel evaluator instances:
  - dev_evaluator: Option<GpuBacktestEvaluator> -- lazy-init after final
    fold, fires once against Q8 dev_features (when dev_quarters > 0)
  - test_evaluator: Option<GpuBacktestEvaluator> -- lazy-init per fold,
    fires inside the fold loop against the WF test slice (when
    fold.test_end > fold.test_start)

Architectural choice: parallel evaluator instances (NOT window-swap on
val_evaluator). Window-swap would require invalidating the CUDA graph
between val and dev/test runs -- fragile, and a direct violation of
pearl_no_host_branches_in_captured_graph. Parallel instances mirror
val_evaluator's lazy-init pattern (TLOB sync, ISV signal pointer,
training_mode = false toggle). Implementation lives behind a single
shared helper `launch_extra_eval` keyed on an `ExtraEvalKind` enum so
Dev / Test share TLOB / ISV / config setup verbatim.

HEALTH_DIAG additions:
  HEALTH_DIAG[N]: dev_eval dev_sharpe_net=... dev_calmar=... dev_max_dd=... dev_trades=...
  HEALTH_DIAG[N]: test_slice fold=K test_sharpe_net=... test_calmar=... test_max_dd=... test_trades=...

The *_sharpe_net key uses the fused-metrics-kernel cost-aware Sharpe
(post-Phase-1.1.b split -- already includes tx_cost_bps + spread_cost
via the env-step PnL feed); when Phase 1.2.b cost-net sharpe lands, the
key name is preserved so the aggregator-script contract holds.

Atomic per feedback_no_partial_refactor: both eval calls + both
evaluator fields + both HEALTH_DIAG lines + audit doc all in this
commit. No parallel paths, no feature flags. Dev_eval runs synchronously
via evaluate_dqn_graphed (one-shot, no async pipelining benefit since
it doesn't fire per-epoch); val path stays async.

Sealed Q9 holdout remains untouched -- Phase 4.3 will load Q9 via a
separate eval-only entry point (NOT train_walk_forward). The Phase 1.6
debug_assert sealed-slice guard catches accidental future refactors.

Verified: cargo check -p ml --features cuda clean; ml lib suite holds
946 pass / 13 fail baseline; the existing Phase 1.7 oracle test
set_test_data_from_slices_fires_observer_and_stashes still passes.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-06 20:55:23 +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
jgrusewski
b791bc8f7f refactor(sp15-p1.1.b): split sharpe out of fused backtest_metrics kernel; wire dedicated launch_sp15_sharpe_per_bar 
Phase 1.1 landed sharpe_per_bar_kernel.cu + launch_sp15_sharpe_per_bar
as orphan scaffolding because the val-side sharpe was inline in
backtest_metrics_kernel's 8-metric fusion (lines 208-211, 277), not
a host-side loop the spec sketch had assumed.

This refactor splits sharpe out:
  - backtest_metrics_kernel computes 7 metrics now (sortino, win_rate,
    max_dd, calmar, omega, VaR, CVaR; remaining counters unchanged).
    Output stride drops 14 -> 13; shmem 6 -> 5 reduction arrays.
  - gpu_backtest_evaluator calls launch_sp15_sharpe_per_bar against
    the same GPU-resident per-bar returns buffer, once per window
    (kernel is single-block by design; n_windows is small).
  - Annualization moves host-side: WindowMetrics.sharpe =
    raw_sharpe * annualization_factor.

Atomic per feedback_no_partial_refactor: kernel split + offset
rebase (every metric below sharpe shifted down by 1) + the lone
WindowMetrics.sharpe consumer (consume_metrics_after_event)
migrated in one commit. No parallel paths.

Output value of WindowMetrics.sharpe is preserved (verified to
1e-5 relative error against f64 closed-form via new oracle test
unified_sharpe_kernel_equivalence_under_annualization). All
existing Phase 1.1 oracle tests still pass; ml lib test suite
holds at the 945/13 baseline (no new regressions).

Eliminates the Phase 1.1 orphan launcher per feedback_wire_everything_up.
Sets up Phase 1.2.b cost-net sharpe to also use launch_sp15_cost_net_sharpe
on the cost-net returns buffer (separate task, separate commit).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-06 19:29:41 +02:00
jgrusewski
69b8fdb61a feat(sp15-p3.5.5): recovery curriculum — per-step DD_TRAJECTORY_DECREASING proxy 
Per spec §9.2 (3.5.5) post-amendment-2: replaces non-existent
episode-level metadata with per-bar signal that fires when
dd_pct(t) < dd_pct(t-1) AND dd_pct(t-1) > DD_TRAJECTORY_FLOOR — i.e.
transition is part of a recovery from non-trivial DD.

PER sampler (Phase 3.5.5.b follow-up) will read this and weight:
  sampling_weight = base × (1 + RECOVERY_OVERSAMPLE_WEIGHT × signal)
so recovery transitions get amplified gradient signal, completing the
downward-spiral break-out chain (3.5.2 reward asymmetry → 3.5.3
cooldown gate → 3.5.4 plasticity → 3.5.5 PER recovery curriculum).

3 ISV slots: 439 DD_TRAJECTORY_DECREASING, 440 RECOVERY_OVERSAMPLE_
WEIGHT (2.0 sentinel; ISV-driven from current dd_pct in follow-up),
441 DD_TRAJECTORY_FLOOR (0.02 sentinel; ISV-driven 25th percentile
of running dd_pct distribution in Phase 3.5.5.c follow-up per
feedback_isv_for_adaptive_bounds).

New sp15_dd_trajectory_prev_dd MappedF32Buffer (size 1) tracks
prev_dd across kernel calls — mirrors Task 3.5.3 sp15_cooldown_
consecutive_losses non-ISV mapped-pinned scratch pattern.

4 fold-reset registry entries + dispatch arms (3 ISV + 1 scratch).

Per established Phase precedent: kernel + launcher land first; PER
sampler integration and 25th-percentile floor producer are purely
additive follow-ups per feedback_no_partial_refactor.

Anchor test 2.10 recovery_after_streak (Phase 2C / Phase 3.5 paired) —
fully green via 3.5.2 + 3.5.4 + 3.5.5 combined.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-06 18:19:53 +02:00
jgrusewski
e0e0abfb28 feat(sp15-p3.5.4): plasticity injection trigger + warm-up tracker (weight-reset deferred) 
Per spec §9.2 (3.5.4) post-amendment-2 fix. TWO-STEP recovery:
  1. Fire when DD_PERSISTENCE > PLASTICITY_PERSISTENCE_THRESHOLD AND
     PLASTICITY_FIRED_THIS_FOLD == 0 → set fired flag, set warm-bars
     counter to M_warm (default 200). [DEFERRED: reset last 10% of
     advantage-head weights to Kaiming-He init via cuRAND.]
  2. Per-bar warm-bars decrement; action-selection layer (consumer wiring
     follow-up) reads max(COOLDOWN_BARS_REMAINING, PLASTICITY_WARM_BARS_
     REMAINING) and forces Hold while > 0.

Weight-reset DEFERRED to Phase 3.5.4.b — kernel signature plumbed
(advantage_head_weights + n_weights) but no-op via (void) cast.
Documented in audit doc.

3 ISV slots: 436 PLASTICITY_FIRED_THIS_FOLD (debounce flag, resets at
fold boundary to re-arm next fold), 437 PLASTICITY_PERSISTENCE_THRESHOLD
(initial 100.0 sentinel; ISV-tracked from running mean of dd_persistence
in follow-up), 438 PLASTICITY_WARM_BARS_REMAINING (counter, OR-gates
with cooldown).

3 fold-reset registry entries + dispatch arms.

Three GPU oracle tests pass: fires-when-persistence-exceeds-threshold
(warm_bars [198, 200] post-fire-and-decrement), debounced-within-fold
(no re-fire when fired=1; warm decrements 50→49), no-fire-below-threshold.

Anchor test 2.22 plasticity_cooldown_interlock (Phase 2C / Phase 3.5
paired) — green via 3.5.4.b follow-up + action-selection consumer wiring.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-06 17:37:56 +02:00
jgrusewski
649128e739 feat(sp15-p3.5.3): cooldown gate — K=5 force Hold for M=20 bars (initial sentinels) 
Per spec §9.2 (3.5.3). After K consecutive losing trades, force Hold for
M bars. Counter at COOLDOWN_BARS_REMAINING (slot 435) part of state —
model can reason about it.

Initial K=5, M=20 hardcoded sentinels. ISV-driven K via MEDIAN_STREAK_
LENGTH (slot 442) producer using two-heap median tracking is documented
Phase 3.5.3 follow-up. ISV-driven M from vol_normalizer time-to-mean-
reversion is also follow-up.

4 ISV slots: 433 K_THRESHOLD, 434 M_BARS, 435 BARS_REMAINING,
442 MEDIAN_STREAK_LENGTH. New sp15_cooldown_consecutive_losses
MappedF32Buffer tracks streak counter persistent across kernel calls.

5 fold-reset registry entries + dispatch arms (4 ISV slots + scratch
buffer reset).

Per spec post-amendment-2 fix: streak counter only updates on trade-close
events; per-bar non-close calls just decrement the cooldown counter. The
trigger gate fires only when a trade-close lands during an inactive
cooldown — re-arming mid-cooldown would extend the gate every closed
trade during the freeze, which is not the spec.

Per established Phase precedent: kernel + launcher land first; action-
selection wiring (force Hold while cooldown_remaining > 0) deferred to
follow-up commit per feedback_no_partial_refactor.

Anchor test 2.12 cooldown_engagement (Phase 2C / Phase 3.5 paired) —
green via this commit.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-06 17:14:53 +02:00
jgrusewski
aa91ce4d82 feat(sp15-p3.5.2): asymmetric reward under DD — gain × (1 + λ × dd_pct), pre-SP12-cap 
Per spec §9.2 (3.5.2). For gains: r_adjusted = r × (1 + λ × dd_pct).
For losses: unchanged. Multiplier applies BEFORE SP12 NEG/POS cap clamp;
saturating the cap for big recovery trades is behaviorally correct
per spec (encourages frequent small recoveries).

3 ISV slots: 430 DD_ASYMMETRY_LAMBDA (initial 0.5; ISV-tracked from
running DD variance via DD_DIST_VAR is follow-up), 431 R_GAIN_DD_BOOST
(diagnostic — most recent boost), 432 DD_DIST_VAR (running variance,
producer follow-up).

3 fold-reset registry entries + dispatch arms.

Per established Phase precedent: kernel + launcher land first; reward
composer site (apply multiplier BEFORE SP12 cap) deferred to follow-up
commit per feedback_no_partial_refactor.

Anchor test 2.10 recovery_after_streak (Phase 2C / Phase 3.5 paired) —
green via 3.5.2 + 3.5.3 + 3.5.4 + 3.5.5 combined.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-06 16:58:39 +02:00
jgrusewski
5d36f3238c feat(sp15-p3.5): confidence-aware Hold floor — bounded sigmoid 
Per spec §8.2 (3.5) post-amendment-2 fix. hold_floor = α × σ(k × (entropy − ε₀))
added to Q_hold pre-argmax/Thompson selection. Hold becomes uncertainty
expression, not distributional default.

α (HOLD_FLOOR_ALPHA slot 426) initial 0.5 sentinel — producer kernel
updating from rolling 95th percentile of |Q_dir| (NOT running max —
outlier-ratchet vulnerable per spec second-review #6) is documented
Phase 3.5 follow-up.
k (HOLD_FLOOR_K slot 427) initial 10.0 — producer from running variance
of entropy is follow-up.
ε₀ (HOLD_FLOOR_EPS0 slot 428) initial 1.0 — producer from 75th percentile
of entropy distribution (ENTROPY_DIST_REF slot 429) is follow-up.

4 fold-reset registry entries + dispatch arms.

Per established Phase precedent: kernel + launcher land first; action-
selection wiring (add hold_floor to Q_hold pre-argmax/Thompson) deferred
to follow-up commit per feedback_no_partial_refactor.

Anchor tests: 2.1 flat_market_holds + 2.6 regime_silences (Phase 2B
contracts) — green via this teaching's action-selection wiring.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-06 16:12:17 +02:00
jgrusewski
8bfc480d92 feat(sp15-p3.4): regret signal = r_discipline content + first-obs bootstrap EMA 
Per spec §8.2 (3.4). When policy holds AND a trade would have been
profitable past cost+trail, regret_bar = ideal_pnl_missed - cost_t.
EMA tracked at REGRET_EMA (slot 423) with first-observation bootstrap
(per pearl_first_observation_bootstrap) + simple exponential decay
(α=0.05; Wiener-α adaptive swap is a documented follow-up).

r_discipline_per_bar = -LAMBDA_REGRET × REGRET_EMA composed at the
reward-split site in the follow-up consumer commit per
feedback_no_partial_refactor.

3 ISV slots: 423 REGRET_EMA, 424 LAMBDA_REGRET (initial 1.0), 425
REGRET_GRAD_NORM. 3 fold-reset registry entries + dispatch arms.

Anchor test 2.6 regime_silences (Phase 2B contract) — green via
3.4 + 3.5 combined.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-06 16:02:59 +02:00
jgrusewski
7753cbef1b feat(sp15-p3.3): quadratic DD penalty + ISV-driven λ + threshold 
Per spec §8.2 (3.3). penalty = λ_dd × max(0, dd_current − dd_threshold)²

Asymmetric: zero below threshold, quadratic growth above. Encodes loss
aversion per pearl_audit_unboundedness_for_implicit_asymmetry.

3 ISV slots: 420 LAMBDA_DD (initial 1.0; ISV-tracked from grad-balance
in follow-up), 421 DD_THRESHOLD (initial 0.05 = 5% drawdown trigger),
422 DD_PENALTY_GRAD_NORM (initial 0.0).

3 fold-reset registry entries + dispatch arms.

Per established Phase precedent: kernel + launcher land first; reward
composition site (subtract penalty from r_total) deferred to follow-up
commit per feedback_no_partial_refactor.

Anchor test 2.5 drawdown_de_risks (Phase 2C / Phase 3.5 paired) — green
via Phase 3.5 mechanisms; this commit lands the penalty primitive.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-06 15:49:45 +02:00
jgrusewski
1eac41d644 feat(sp15-p3.2): explicit cost in r_quality on trade-close events 
Per spec §8.2 (3.2). Extends the Phase 3.1 composer kernel
`r_quality_discipline_split_kernel` with two new args (`float cost_t`,
`unsigned int trade_close_indicator`) and subtracts
`cost_t * (float)trade_close_indicator` from `r_quality` BEFORE the α
blend. Same `cost_t` scalar shape as the Phase 1.2 cost_net_sharpe
accumulator (commission + per-side spread + OFI-impact); the gate
`trade_close_indicator=1` on round-trip-close bars (rt_ind=1) and 0
otherwise so non-close bars receive a structural no-op identical to
the pre-Phase-3.2 behaviour. Model SEES the bill in the gradient
signal during training, not just in eval-time metrics.

Approach: kernel-signature-extension (NOT wrapper-kernel). The only
existing call site in the tree is the Phase 3.1 oracle test
(`training_loop.rs` has dispatch-arm reset wiring but does NOT yet
invoke the launcher per the Phase 3.1 commit's deferred-consumer
note), so the cascade is bounded to that single test — extending the
existing kernel is cleaner than a parallel wrapper that would have to
be retired the moment the Phase 3.1 deferred consumer migration
lands.

Anchor test 2.4 cost_sensitivity (Phase 2B contract) — green via this
commit + Phase 3.1 split structure (already landed in 2d226e6e7).

Per established Phase precedent: kernel/launcher signature change +
existing-test migration land atomically per
feedback_no_partial_refactor; production reward-composition wire-up
that feeds real `cost_t` from cost_net_sharpe is the same deferred
follow-up Phase 3.1 declared (no new debt added — both share one
follow-up commit).

cargo check -p ml --features cuda: clean (18 pre-existing warnings).
cargo test -p ml --test sp15_phase1_oracle_tests --features cuda
  -- --ignored r_quality_subtracts_explicit_cost: 1 passed.
cargo test -p ml --test sp15_phase1_oracle_tests --features cuda
  -- --ignored r_split_uses_sentinel_alpha_at_cold_start: 1 passed
  (Phase 3.1 sentinel test post-migration).
cargo test -p ml --features cuda: 946 passed / 13 failed
  (same 13 failures as parent 2d226e6e7; zero introduced).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-06 15:36:58 +02:00
jgrusewski
2d226e6e76 feat(sp15-p3.1): r_quality + r_discipline split with ISV-driven α + sentinel cold-start 
Per spec §8.2 (3.1) post-amendment-2 fix: ALPHA_SPLIT slot initialized
DIRECTLY to 0.5 in trainer constructor. Formula α = grad_norm_q /
(grad_norm_q + grad_norm_d + ε) takes over only after BOTH grad-norm
EMAs accumulate ≥ N_WARM=100 non-zero observations.

Two kernels in r_quality_discipline_split_kernel.cu (single cubin per
established 1:1-source-to-cubin pattern with multiple kernels):
  - r_quality_discipline_split_kernel: per-step composition + warm count
  - alpha_split_producer_kernel: per-step ALPHA_SPLIT update from grad ratio
    (gated on warm count to prevent premature formula activation)

3 ISV slots (417 ALPHA_SPLIT, 418 GRAD_NORM_QUALITY, 419 GRAD_NORM_DISCIPLINE)
+ sp15_alpha_warm_count [1] mapped-pinned scratch buffer on the trainer
struct. 4 fold-reset registry entries + dispatch arms (one for the
non-ISV warm-count buffer mirrors the sp11_novelty_hash host_slice_mut
pattern).

Per established Phase precedent: kernels + launchers land first; consumer
migration (per-step launches in training_loop.rs reward composition site)
deferred to a follow-up commit per feedback_no_partial_refactor.

Anchor tests: 2.4 cost_sensitivity + 2.6 regime_silences (Phase 2B
contracts) — green via Phase 3.4 regret + 3.2 cost; this commit lands
the split structure they depend on.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-06 15:26:56 +02:00
jgrusewski
7e63f83bf0 test(sp15-p2b): 17 behavioral tests (Group 1 + Group 2) as #[ignore] contracts 
Per spec §7.3 + plan v2 Phase 2B differential table. Each test documents
its behavioral contract and its enabling phase. Tests are #[ignore]'d
so the pre-commit-hook-gated suite does not fail; per-test #[ignore] is
removed by the commit that lands its enabling phase.

Group 1 trading behavior (7 tests): 2.1 flat_holds, 2.2 trend_directional,
2.3 mean_revert_reverses, 2.4 cost_sensitivity, 2.6 regime_silences,
2.7 stop_discipline, 2.18 action_latency.

Group 2 state/arch grounding (10 tests): 2.8 hold_vs_flat_semantics,
2.9 eval_fold_isolation, 2.13 eval_train_consistency, 2.14 magnitude_uses_
all_buckets, 2.15 fold_transition_stability, 2.16 q_value_bounded,
2.17 gradient_flow_to_all_heads, 2.19 kelly_calibration, 2.20 state_input_
grounding, 2.21 egf_gate_opens (re-export hook for sp14_oracle_tests.rs).

Group 3 (Phase 2C — 5 tests) lands paired with Phase 3.5 commits.

Harness invocations use trainer = () placeholder; Phase 2B.b extends
BehavioralResult/harness with per-bar actions, position-state inspection,
forced-action stepping, and per-head grad-norm snapshots. Each test's
ignore reason references which Phase X enabling work + harness extension
makes it green.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-06 15:05:34 +02:00
jgrusewski
ef373c34d7 feat(sp15-p1.7): consume the abandoned walk-forward test slice (stash + observer; eval invocation deferred) 
Per spec §6.7. The walk-forward generator emits `test_start..test_end`
per fold but the trainer at `mod.rs:1294` only consumed train+val — the
12.5% test slice was silently dropped, the model was never measured on
held-out data the train/val pipeline didn't see.

This commit lands the foundation: `set_test_data_from_slices` stashes
the per-fold range immediately after `set_val_data_from_slices`, gated
on `fold.test_end > fold.test_start` for defensively-empty slices. A
`set_test_data_observer` hook lets unit tests verify the wiring without
spinning up a full GPU eval pipeline.

The actual `evaluate_dqn_graphed` invocation against the stashed slice
plus the per-fold `HEALTH_DIAG test_slice fold=K test_sharpe_net=...`
emit is deferred to a follow-up commit per `feedback_no_partial_refactor`.
Wiring it through requires either standing up a second
`GpuBacktestEvaluator` instance (parallel to the val one at
`metrics.rs:550`) or refactoring the existing val evaluator to swap
window data between val and test eval — the val evaluator's lazy-init
path is fundamentally tied to the window passed at construction. Plus
TLOB weight sync, ISV signal wiring, and a `training_mode` toggle (no
such field exists yet on `DQNTrainer`).

This deferral matches the Phase 1.5 (kernel + launcher first, trunk
consumer follow-up) and Phase 1.6 (stash dev/holdout slices, eval
consumer follow-up) precedents on this branch. The stash + observer
surface is the analogous foundation; the L40S smoke once Task 1.7.b
lands will surface the per-fold `test_sharpe_net` HEALTH_DIAG line as
the canonical end-to-end verifier.

New oracle test `set_test_data_from_slices_fires_observer_and_stashes`
in `sp15_phase1_oracle_tests.rs` constructs a real trainer (sync init,
no GPU forward), registers an observer, exercises the API with a
synthetic [5000..6000) range, asserts the observer fires once with the
right bounds. Passes locally on RTX 3050 Ti.

`docs/dqn-wire-up-audit.md` extended with a Phase 1.7 entry documenting
what landed, what's deferred, the wire-up locations, and the rationale.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-06 14:53:05 +02:00
jgrusewski
ce019c72d2 feat(sp15-p1.6): --holdout-quarters + --dev-quarters CLI flags + sealed Q1-Q7/Q8/Q9 split 
Per spec §6.6 / Q6 train/dev/test split (defaults Q1-Q7 train, Q8 dev,
Q9 sealed final test):

- DQNHyperparameters: holdout_quarters + dev_quarters (default 1+1)
- crates/ml/examples/train_baseline_rl.rs: --holdout-quarters /
  --dev-quarters CLI flags forwarded to hyperparams (this is the actual
  training binary; bin/fxt/src/commands/train.rs is a gRPC client and
  services/ml_training_service/src/main.rs accepts training params via
  proto not CLI — see audit doc note).
- DQNTrainer::train_walk_forward slices training_data BEFORE fold
  generation; folds run on Q1..Q(9 - holdout - dev) only.
- DQNTrainer struct: dev_features/dev_targets/holdout_features/
  holdout_targets fields stash trailing slices for end-of-training dev
  eval and the Phase 4.3 separate eval-only workflow.
- debug_assert sealed-slice guard catches future refactors that
  re-introduce holdout into the training path.

Per established Phase 1 precedent (1.1-1.5: kernel/state lands first,
consumer wiring deferred to follow-up commit per
feedback_no_partial_refactor): CLI plumbing + slicing + dev/holdout
storage land in this commit. The post-final-fold dev evaluation call
(consumer of dev_features) is deferred to a follow-up commit and will
mirror Task 1.7's evaluate_dqn_graphed integration pattern. Phase 4.3
argo-eval-final.sh is the sole legitimate consumer of holdout_features
(separate eval-only workflow that does NOT call train_walk_forward).

cargo check -p ml --features cuda --example train_baseline_rl: clean
cargo check -p fxt: clean (no fxt changes needed; gRPC client only)

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-06 14:39:37 +02:00
jgrusewski
5309d4bee5 feat(sp15-p1.5): dd_pct foundational state input concat kernel — LAYOUT FINGERPRINT BREAK 
LAYOUT FINGERPRINT BREAK: pre-SP15 checkpoints WILL NOT LOAD after this
commit. Greenfield OK per spec Q1.

Per spec §6.5: dd_pct (slot 406, written by Task 1.3 dd_state_kernel)
gets concatenated to the trunk forward input as the last dim. Eval-time
policy SEES drawdown context on every forward pass; Phase 3 teachings
can condition on dd_pct directly via state, not just reward modulation.

New kernel `dd_pct_concat_kernel.cu` produces a [B, state_dim_padded + 1]
buffer whose leading state_dim_padded columns equal the input states_buf
and whose +1 last column equals isv[DD_PCT_INDEX=406] broadcast across
batch. Pure scatter-copy (no atomicAdd per feedback_no_atomicadd).

layout_fingerprint_seed extended with 'TRUNK_INPUT_DD_PCT=sp15_phase_1_5'
marker — FNV1a hash changes; old checkpoints fail to load with the
existing layout-mismatch error path (same fail-fast that fired on the
SP4 / SP14 layout breaks).

Phase 1.5 lands kernel + launcher + layout fingerprint marker + GPU
oracle test only. Trunk consumer migration (re-pointing forward_online
to consume the concat buffer + bumping s1_input_dim from 48 → 49 +
propagating through GRN encoder, VSN gate input, bottleneck path, and
backward dx scratch) is deferred to a follow-up atomic commit per
feedback_no_partial_refactor — matches the established Phase 1.1-1.4
precedent (kernels + launchers verify in isolation first; consumer
migration is a load-bearing change touching the GRN encoder, VSN
partition boundaries, fxcache schema, and backward gradient flow).

GPU oracle test `dd_pct_concat_kernel_writes_last_column` validates B=4,
raw state_dim=48, state_dim_padded=128: leading 128 columns of each
output row match the input states (including pad zeros), and the 129th
column equals isv[DD_PCT_INDEX]=0.42 broadcast across all 4 rows. Test
passes on local RTX 3050 Ti (sm_86) in 1.62s.

cargo test -p ml --lib --features cuda: 945 passed / 14 failed — same
14 failures pre-existing on the parent commit `c6fd4b4b2` (Task 1.4
partial baseline); zero introduced by this commit.

Per spec §6.5 step 7 (trunk-grounding behavioral test): Phase 4 L40S
smoke verifies dd_pct propagation at production scale via the existing
layout-fingerprint-mismatch fail-fast on cold-start of any pre-SP15
checkpoint. The follow-up Phase 1.5.b commit that lands the consumer
migration adds the explicit trunk-grounding KL test alongside the
forward_online wiring.

Touched:
- crates/ml/src/cuda_pipeline/dd_pct_concat_kernel.cu (new)
- crates/ml/build.rs (+1 cubin manifest entry)
- crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs (+SP15_DD_PCT_CONCAT_CUBIN
  + launch_sp15_dd_pct_concat + TRUNK_INPUT_DD_PCT layout fingerprint
  marker)
- crates/ml/tests/sp15_phase1_oracle_tests.rs (+1 GPU oracle test)
- docs/dqn-wire-up-audit.md (+1 Phase 1.5 entry)

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-06 14:31:16 +02:00
jgrusewski
c6fd4b4b2a feat(sp15-p1.4-partial): 4 constant-policy baselines (buyhold, hold_only, momentum, reversion) 
Per spec §6.4. Pure-CUDA kernels in baseline_kernels.cu — single cubin,
4 extern "C" __global__ functions sharing a templated compute_baseline_sharpe
helper. Trunk-shared baselines (random_dir_kelly slot 411, aux_only slot 413,
mag_quarter_fixed slot 414, trail_only slot 415) deferred to Task 1.4.b
follow-up — they need partial-policy-forward access from the main eval pass.

ISV slots written: 409 (buyhold), 410 (hold_only), 412 (naive_momentum),
416 (naive_reversion). Slots 411/413/414/415 stay at sentinel 0.0 until
follow-up commit.

Per established Phase 1 precedent: kernels + launchers land first;
per-eval-pass launches + HEALTH_DIAG baseline_deltas emit deferred to
follow-up commit per feedback_no_partial_refactor.

Anchor tests: buyhold positive on +drift, hold_only emits 0,
momentum + reversion sum near zero on mean-reverting series.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-06 14:10:20 +02:00
jgrusewski
9e84602486 feat(sp15-p1.3): drawdown state kernel — DD_CURRENT/MAX/RECOVERY/PERSISTENCE/CALMAR/PCT 
Per spec §6.3. Per-step kernel reads PS_PEAK_EQUITY (slot 7) and
PS_PREV_EQUITY (slot 9) from existing position state buffer (no new
equity slot needed). Writes 6 ISV slots (401-406). Calmar uses
max(dd_max, 1e-4) floor — eliminates the saturation-at-100 artifact
seen in train-dd4xl HEALTH_DIAG.

6 fold-reset registry entries + dispatch arms. 5 sentinel-0 stateful
outputs; calmar uses sentinel 1e-4 (same value as the kernel's floor)
so cold-start division uses the floor rather than ±inf.

Atomic split per feedback_no_partial_refactor.md (mirrors Task 1.1 +
1.2 precedent): kernel + launcher + registry land here; per-step
production wire-up + HEALTH_DIAG composer deferred to a follow-up
commit. Test 1.3 oracle on 6-step synthetic equity curve passes
locally on RTX 3050 Ti (sm_86) in 1.61s. cargo test -p ml --lib
--features cuda: 946 passed / 13 failed — same 13 pre-existing
failures as Task 1.2 baseline (a92ff28a9); zero introduced.
State-reset registry tests: 4/4 pass.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-06 13:56:57 +02:00
jgrusewski
a92ff28a98 feat(sp15-p1.2): cost-net sharpe kernel — commission + spread + OFI-impact 
Per spec §6.2 per-side semantics:
  cost_t = commission_per_rt × rt_ind[t]
         + half_spread[t] × |pos[t]| × side_ind[t]
         + ofi_lambda × |pos[t]| × |ofi[t]| × side_ind[t]

Commission charged at close; half-spread × |pos| at entry AND exit
(sums to one full spread per RT); OFI impact same per-side. Initial
λ=2.0e-4 in ISV[OFI_IMPACT_LAMBDA_INDEX=407] as Invariant-1 anchor
(constructor-write + FoldReset rewrite per feedback_isv_for_adaptive_bounds);
per-fold ISV refit may overwrite in later phases. Mean cost-per-bar
emitted to ISV[COST_PER_BAR_AVG_INDEX=408] (stateful kernel output;
FoldReset sentinel 0 + Pearl A first-observation bootstrap).

Same kernel reads LobBar fields from synthetic markets (Phase 2A) and
real fxcache LOB (prod) — dev/prod parity per Q3.

Phase 1.2 lands kernel + launcher + anchor seed + 2 registry entries
+ 2 dispatch arms only; consumer migration deferred to a follow-up
commit per feedback_no_partial_refactor (mirrors Phase 1.1 atomic
pattern). Oracle test cost_net_sharpe_round_trip_charges_full_spread
validates single round-trip cost = 2.00 / 10 bars = 0.20/bar with
mean_pnl_net = 0.80 on 1.69s RTX 3050 Ti (sm_86). Zero regressions
introduced (946 passed / 13 pre-existing failures, same as Task 1.1).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-06 13:46:03 +02:00
jgrusewski
3667cd1b04 feat(sp15-p1.1): unified sharpe kernel — single formula for train and val 
Replaces sharpe_ema (per-batch EMA, train) vs sharpe_annualised (val ×
sqrt(525600)) split. Single GPU kernel computes mean/std/sharpe via
2-pass block-tree-reduce; annualisation at the call site.

Per feedback_no_partial_refactor: consumer migration in metrics.rs /
training_loop.rs deferred to a follow-up commit; kernel + launcher land
in isolation first. Verified via 2 GPU oracle tests
(unified_sharpe_kernel_zero_mean, unified_sharpe_kernel_positive_drift)
passing on local RTX 3050 Ti (sm_86) in 1.78s.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-06 11:41:00 +02:00
jgrusewski
dff6e666ee feat(sp15-p2a.2): oracle policies + evaluator harness + pre-commit hook 
Phase 2A scaffolding complete (per spec §7.2):
- oracle.rs: OracleAction enum + 3 oracle policies (flat, drift, OU)
- harness.rs: BehavioralResult + evaluate_policy_on_market stub.
   Trainer-side methods (eval_actions_on_features, read_isv_for_test)
   are documented gap #7; Phase 2B tests add them per-test as needed.
- pre-commit hook: cargo test -p ml --test behavioral_suite gates
   argo-train.sh per spec §4.5 discipline rule.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-06 11:29:22 +02:00
jgrusewski
7d0a29dced feat(sp15-p2a.1): LobBar canonical ABI + 4 synthetic market generators 
Phase 2A scaffolding lands FIRST per spec §4.4 ABI contract. Phase 1.2
cost kernel reads LobBar; both dev synthetic and prod fxcache produce
LobBar — dev/prod parity per Q3.

Generators: flat_market, drift_market, ou_market, regime_switch_market
(seeded RNG for reproducible tests). Regime-switch test uses sticky
0.99/0.01 transitions (true regime persistence; spec's 50/50 was a
random walk, not a regime switch — corrected with code comment).

behavioral_suite test target wired into Cargo.toml; will run all
22 Phase 2 tests once they land in Phase 2B/2C.

Audit doc: SP15 Phase 2A.1 entry appended to docs/dqn-wire-up-audit.md
per Invariant 7 (component changes require audit-doc update).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-06 10:57:51 +02:00
jgrusewski
04ea5a0243 test(sp15-p1.0): scaffold sp15_phase1_oracle_tests.rs for Phase 1 
Empty mod gpu placeholder; Phase 1 tasks 1.1-1.7 will append per-task
oracle tests. No tests yet — file just must compile.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-06 10:53:34 +02:00
jgrusewski
c146c4fffd feat(sp15): scaffold sp15_isv_slots.rs with 46 slots [397..443) — ISV_TOTAL_DIM 396→443 
Per spec §4.3 allocation map. Pre-allocates disjoint slot ranges to
enable Approach B parallel sub-worktrees without index collisions:
  - Phase 0.B EGF retune: [397..401)
  - Phase 1.3 drawdown: [401..407)
  - Phase 1.2 cost: [407..409)
  - Phase 1.4 baselines: [409..417)
  - Phase 3.X-3.5.X teachings + recovery: [417..441)
  - Phase 3.5 deferred anchors: [441..443)

Layout fingerprint extended with all 46 slot names. Pre-SP15 checkpoints
will be incompatible (greenfield OK per Q1).

Two regression tests verify: (1) every slot < ISV_TOTAL_DIM, (2) layout
fingerprint locked at named indices. docs/isv-slots.md gets the SP15
section documenting the allocation map + greenfield sub-worktree plan.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-06 10:49:33 +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
60ad42676e fix(sp14): bump ISV_TOTAL_DIM 383 → 396 to cover SP14 EGF slots 
ROOT CAUSE of L1+L2 from Smoke A2-B: the ISV bus was sized for top
of SP13 (ISV_TOTAL_DIM=383) but B.1 allocated SP14 slots at 383-395.
Every SP14 read/write was OUT-OF-BOUNDS memory access. That's why:

- gate1 (slot 391) read as 0 always (OOB zero-init memory)
- post_open_min (slot 394) accumulated garbage values 9.5 → 28 → 46
- α_smoothed/α_raw values appeared to work but were undefined behavior

SP4/SP5 had a regression test (`all_sp4/5_slots_fit_within_isv_total_dim`)
that catches this exact failure mode at unit-test time. SP14 was missing
it — that gap let the bug ship across all 16 commits without being caught.

Changes:
- ISV_TOTAL_DIM: 383 → 396 (covers SP14 slots 383-395)
- layout_fingerprint_seed: extended with SP14 slot names + new
  ISV_TOTAL_DIM=396 marker (forces fingerprint hash bump per
  Invariant 8 — old checkpoints invalidated correctly)
- sp14_isv_slots.rs: 2 regression tests (mirror SP4/SP5 patterns)

Both tests pass. After this fix, SP14 EGF kernels will read/write
the correct slots; gate1 should actually flip open when aux_dir_acc
crosses target+0.03; gradient_hack_detect post_open_min stays bounded
in [0, 1] as designed.

NOT yet addressed (separate follow-up):
- warmup_gate hardcoded WARMUP_STEPS_FALLBACK=1000 violates
  feedback_isv_for_adaptive_bounds. Should be ISV-signal-driven OR
  removed entirely (k_aux/k_q already provide variance-driven warmup).
  Redesign post-re-smoke once bus-size fix is verified.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-05 21:59:47 +02:00
jgrusewski
e41dbb7d8a diag(sp14 B.12): per-epoch pearl_egf_diag HEALTH_DIAG emit 
Adds a new HEALTH_DIAG[{epoch}]: pearl_egf_diag line immediately after
the aux_moe block in the per-epoch metrics section of training_loop.rs.
Reads all 13 SP14 ISV slots [383..396) — α_smoothed, α_raw, β, k_aux,
k_q, var_aux, var_q, var_α, q_dis_short, q_dis_long, gate1 state,
post_open_min, lockout — via the established read_isv_signal_at pattern,
giving forensic visibility into EGF pearl state each epoch.

gate1/gate2 sigmoid outputs are intentionally omitted: recomputing them
host-side would violate feedback_no_cpu_compute_strict; the sigmoid
inputs are sufficient for a reader to infer the output values.

docs/isv-slots.md updated (Invariant 7): records B.12 HEALTH_DIAG wire-up.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
 2026-05-05 21:14:58 +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
dc3f948ee9 feat(sp14): B.10 — backward wire gradient gating by ALPHA_GRAD_SMOOTHED 
Critical safety mechanism that completes the EGF pearl: scales the wire
column of `dL/dx_concat [B, SH2 + 1]` (the gradient flowing FROM the
direction Q-head's first FC SGEMM TO `aux_softmax_diff`) by
`ISV[ALPHA_GRAD_SMOOTHED_INDEX = 393]`, computed by B.4's
`alpha_grad_compute_kernel` and orchestrated per-step in B.11.

`dL/dW[wire_col]` (Q-head's own weight gradient for the appended column)
is NOT scaled — the dW SGEMM `dY^T × x_concat` and the dX SGEMM
`dY × W^T` are independent, so scaling `dx[:, SH2]` AFTER both have
completed leaves dW unaffected. Q-head learns to USE the wire freely;
only the gradient PROPAGATING BACK to aux is gated.

Pre-B.11 (no producer wired) `ISV[393]` holds sentinel `0.0` →
wire force-closed (gradient zeroed) — the conservative safety state.
Post-B.11, B.4 writes the live gate output ∈ [0, 1] each step.

Closes the latent K-mismatch B.8/B.9 left in backward
============================================================

B.8 grew `w_b0fc` to `[adv_h, SH2 + 1]` end-to-end (Adam m/v +
spectral-norm vector + smoke fixtures); B.9 closed the forward dispatch
K-mismatch. The backward dW/dX SGEMMs for `d == 0` still used `K = SH2`
against the new `LDA = SH2 + 1` weight tensor — silently dropping the
last column of dW and zeroing the wire-col gradient. B.10 closes that
gap atomically with the wire-col scale per `feedback_no_partial_refactor`:

  * `backward_branch_dw` for `d == 0` now uses `(dir_qaux_concat_ptr, SH2 + 1)`
    instead of `(save_h_s2, SH2)` — matching the forward consumer
    pattern from B.9.
  * `backward_branch_dx` for `d == 0` now writes to
    `d_dir_qaux_concat [B, SH2 + 1]` with `K = SH2 + 1` instead of
    `scratch_d_h_s2 [B, SH2]` with `K = SH2`. Mirrors the magnitude
    branch's wider-buffer pattern.

New artifacts
=============

  * `sp14_scale_wire_col_kernel.cu`: one thread per batch row, scales
    `dx_concat[b, SH2]` by `isv[393]` IN-PLACE. NaN-safe per the
    `dqn_scale_f32_kernel` precedent (explicit `α==0 ⇒ 0` branch).
    Pure per-thread map, no atomicAdd, no shared memory.
  * `sp14_d_dir_qaux_concat: CudaSlice<f32>` `[B, SH2 + 1]` trainer-
    struct field. Dx SGEMM destination; the wire-col scale acts on
    this buffer; the strided accumulator copies the first SH2 columns
    into `bw_d_h_s2` after the scale.
  * `launch_sp14_scale_wire_col` launcher reads `self.isv_signals_dev_ptr`
    and the new buffer's raw_ptr.
  * `backward_full` signature grows two trailing `u64` args
    (`dir_qaux_concat_ptr`, `d_dir_qaux_concat_ptr`); both
    `backward_full` call sites (CQL aux + main online) wired
    atomically per `feedback_no_partial_refactor`.

Post-call orchestration at trainer level
========================================

  1. `launch_sp14_dir_concat_qaux(save_h_s2)` rebuilds the ONLINE
     concat in `sp14_dir_qaux_concat_scratch` (the forward pass had
     overwritten it with the TARGET concat at line ~25817). Same
     one-step-lag semantic preserved — `aux_nb_softmax_buf` is
     unchanged between forward and backward.
  2. `cuMemsetD32Async` zero of `d_h_s2` — pre-B.10 the direction
     branch (d==0) wrote it with beta=0; post-B.10 the dir-Q dX
     lives in `d_dir_qaux_concat` and is gated + accumulated AFTER
     `backward_full` returns, so the value-FC dx accumulator inside
     `backward_full` (beta=1) needs an explicit zero baseline.
  3. `backward_full` runs: dir branch → `d_dir_qaux_concat`,
     mag/ord/urg branches → their concat dX buffers, value-FC →
     `d_h_s2` (beta=1, on top of zeroed buffer).
  4. `launch_sp14_scale_wire_col` gates col SH2 of `d_dir_qaux_concat`.
  5. `accumulate_d_h_s2_from_concat` (beta=1) copies first SH2 cols
     of `d_dir_qaux_concat` into `d_h_s2`. Wire col stays in
     `d_dir_qaux_concat[:, SH2]`, untouched by this accumulator (its
     destination range is [0, SH2)). Pre-B.11 the wire is already
     zeroed by the sentinel-α gate; the orchestrator that propagates
     the gated wire-col gradient back to the aux head's softmax CE
     backward chain lives in B.11.
  6. mag/ord/urg accumulators continue with beta=1 (comments updated).

Wire status
===========

  * Forward dispatch: unchanged (B.9-complete).
  * Backward dispatch: GATED on both call sites (CQL aux + main online).
  * dW unchanged: the `dW = dY^T × x_concat` SGEMM writes
    `grad_buf[goff_w_b0fc..]` BEFORE the scale-wire-col launches;
    the scale operates ONLY on `d_dir_qaux_concat` (the dx buffer)
    AFTER both dW and dX SGEMMs complete.
  * Target net unaffected: Polyak EMA-only, no backward.
  * CudaSlice wrapper path: passes `0u64` for both new args, falls
    back to the legacy K=SH2 path. Consistent with the forward
    wrapper's diagnostic-only residual.

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)
  * Audit doc `docs/dqn-wire-up-audit.md` updated per Invariant 7.

After this commit, the EGF pearl is architecturally complete; the
orchestration of when/how the alpha_grad gates fire happens in B.11
(producer chain orchestrator).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-05 20:49:27 +02:00
jgrusewski
ecf4757c0d feat(sp14): B.9 — wire forward concat into direction Q-head SGEMM 
Closes the latent SGEMM K-mismatch left by B.8 (6715ab4ea):
`w_b0fc` had grown from `[adv_h, SH2]` to `[adv_h, SH2 + 1]` end-to-end,
but every direction-Q-head consumer's SGEMM still used `K = shared_h2`
against the new `LDA = SH2 + 1` weight tensor — safe ONLY because the
new column was zero-init in B.8 and Adam had not yet updated it. After
this commit the forward wire is FULLY ACTIVE; the SGEMM consumes
`sp14_dir_qaux_concat_scratch [B, SH2 + 1]` with `K = shared_h2 + 1`.

Direction Q-head input pointer: `h_s2_buf` → `sp14_dir_qaux_concat_scratch`.
K dim: `shared_h2` → `shared_h2 + 1`.

Concat kernel runs immediately before the direction Q-head SGEMM in
the same stream, enforcing `pearl_canary_input_freshness_launch_order`.
Mirrors the `launch_mag_concat_from` precedent: the aux head forward
that writes `aux_nb_softmax_buf` runs AFTER the per-step online
forward (line ~25599 in the new layout), so each forward consumes
the PREVIOUS step's aux predictions — same one-step-lag semantic as
mag_concat. Step 0 sees alloc_zeros (uniform 0.5/0.5 → diff = 0),
step 1+ sees the prior step's aux next-bar softmax.

Atomic-migration consumers (`feedback_no_partial_refactor`):

- `gpu_dqn_trainer.rs` — new `launch_sp14_dir_concat_qaux` method;
  online forward (line ~25583) and target forward (line ~25758)
  each precede their `forward_*_raw` call with a concat launch and
  pass `sp14_dir_qaux_concat_scratch.raw_ptr()`. Both replay paths
  (`replay_forward_ungraphed`, `replay_forward_for_q_values`
  ungraphed fallback) get the same wire — they use online weights
  and produce direction Q-values consumed by training/eval. Causal
  intervention sites (×2) and DDQN argmax pass `0u64` per spec
  (their direction Q outputs are either unread by the consumer or
  the spec accepts the K=SH2 fallback's residual one-step bias).

- `batched_forward.rs` — five `forward_*_raw` / `launch_vsn_glu_branch`
  signatures grow a trailing `dir_qaux_concat_ptr: u64`; new
  `d == 0 && dir_qaux_concat_ptr != 0` branch in every legacy
  ReLU-FC FC dispatch (multi-stream / sequential × online / target /
  F32-output) returning `(dir_qaux_concat_ptr, self.shared_h2 + 1)`.
  VSN-GLU branch path scatters `vsn_masked` into the first SH2 cols
  of the scratch, identical to the `d == 1/2/3` scatter pattern
  (the trailing aux_softmax_diff column was already written by the
  pre-VSN concat-kernel launch and survives the scatter). The
  `CublasGemmSet::new` heuristic-cache shape table grows by one
  unique tuple `(adv_h, batch, SH2 + 1, SH2 + 1)` so the first-call
  cublasLt heuristic search hits a fresh cache slot instead of the
  pre-B.8 `(adv_h, batch, SH2, SH2)` entry.

- `gpu_experience_collector.rs` / `value_decoder.rs` — pass `0u64`
  for the new arg (no aux-head dependency on those forwards;
  documented inline with rationale).

- `docs/dqn-wire-up-audit.md` — new SP14 Layer B B.9 entry per
  Invariant 7, documenting every new dispatch site, the
  diagnostic-path residual, and the launch-order constraint.

After this commit the forward wire is FULLY ACTIVE: aux-head
gradients flow back through the kernel's `s1 - s0` derivative into
`aux_nb_softmax_buf`'s logits, co-training the aux head with Q-loss.
Backward gradient flow is INTENTIONALLY UNGATED in this commit —
the EGF pearl gating (scale `dL/dx[wire_col]` by `α_grad_smoothed`
to prevent gradient-hacking) lands in B.10. Per
`feedback_no_partial_refactor`, this intermediate state is
functional (the model trains; aux gets co-trained by Q-loss) but
not yet behavior-protected by the gate.

Diagnostic-path residual (causal intervention, DDQN argmax, exp
collector, value decoder): the cuBLAS heuristic for `K=SH2, LDA=SH2`
against the underlying `[adv_h, SH2 + 1]` weight tensor reads the
first `adv_h * SH2` floats with stride SH2 — within bounds (no
OOB), produces stable-but-incorrect outputs for the residual paths.
Their direction Q outputs feed either (a) only-value-logit consumers
(causal sensitivity) or (b) downstream argmax-only consumers with
one-step-bias acknowledged by the spec (DDQN). The train-time wire
(online + target + replay) is fully closed.

Test: `SQLX_OFFLINE=true cargo check -p ml` clean (18 warnings,
pre-existing baseline). The smoke validation that the model
converges with the active forward wire happens in B.11 alongside
the captured-graph integration (B.10 gates backward first).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-05 20:26:23 +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