da5e564ccfad43071f579a824ef122dc7017d837
557 Commits
| Author | SHA1 | Message | Date | |
|---|---|---|---|---|
jgrusewski
|
3b71d21834 |
feat(sp14-c): aux prediction horizon ISV-driven (multi-bar pivot)
Aux's original label was (p_{t+1} > p_t) — pure HFT-scale microstructure
noise that's unlearnable at our HFT-MFT trading frequency. Migrated to
(p_{t+H} > p_t) where H is read from ISV[AUX_PRED_HORIZON_BARS_INDEX=450].
Adaptive producer drives H from observed avg winning hold time:
- Pearl-A first-observation bootstrap: replace sentinel H=60 directly
on first valid observation
- Steady-state Wiener-α EMA blend, slow (α=0.01) for stable horizon
(no target-variance EMA available, fallback per
pearl_wiener_optimal_adaptive_alpha)
- "No winning trades yet" guard keeps sentinel until first valid observation
Lookahead truncation: labels at t where t+H >= total_bars are masked
(sentinel -1, loss-reduce skips). The existing aux_next_bar_loss_reduce
in aux_heads_kernel.cu already supports the -1 mask convention via the
B_valid count — no new valid_mask parameter needed.
Step 5b finding: Case B — existing per-sample buffers
(hold_at_exit_per_sample, trade_profitable_per_sample) populated by
unified_env_step_core, but no aggregate ISV slot. Added new aggregator
slot AVG_WIN_HOLD_TIME_BARS_INDEX=451 + new producer kernel
avg_win_hold_time_update_kernel.cu (block-tree-reduce, no atomicAdd).
ISV_TOTAL_DIM bumped 450 → 452.
ATOMIC migration per feedback_no_partial_refactor: both label kernels
(aux_sign_label_kernel.cu trajectory + aux_sign_label_per_step_kernel.cu
per-rollout-step) migrated together to the new
(targets, bar_indices, isv, isv_h_idx, out_labels, total, total_bars)
signature. The lookahead host-passed scalar argument is removed; H is
read from ISV inside the kernel (broadcast value, single read per
thread, on-device clamp [1, 240]).
Producer chain (per-epoch boundary): new
GpuDqnTrainer::launch_aux_horizon_chain orchestrates
avg_win_hold_time_update → aux_horizon_update sequentially alongside
launch_kelly_cap_update at the existing epoch-boundary slot in
training_loop.rs.
Trunk math (C.2/C.3/C.4) unchanged — separate aux trunk is label-
agnostic. Validation in C.10 will use H=60 cold-start; the adaptive
producer drives H from real winning-trade observations.
Tests (8 oracle, 5 new + 3 preserved):
- aux_trunk_forward_matches_numpy_reference (C.3) ✓
- aux_trunk_backward_gradient_check (C.4) ✓
- aux_trunk_backward_does_not_write_dx (C.4) ✓
- aux_sign_label_h_bar_horizon (NEW) ✓
- aux_sign_label_lookahead_mask (NEW) ✓
- aux_horizon_pearl_a_bootstrap (NEW) ✓
- aux_horizon_converges_to_steady_target (NEW) ✓
- aux_horizon_holds_sentinel_with_no_winning_trades (NEW) ✓
8/8 pass on RTX 3050 Ti.
Phase C.4b of SP14 Layer C separate-aux-trunk refactor.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
5d584dc751 |
feat(sp14-c): aux trunk backward kernel + gradient check + stop-grad invariant test
Backward propagates dh_s2_aux through w3/w2/w1 with block-tree-reduce
(no atomicAdd per feedback_no_atomicadd). Critical: kernel set does NOT
write dx_in — encoder gradient remains Q-shaped only. Stop-grad
invariant verified via parameter-list structural enforcement (kernels
literally cannot reference an `dx_in_out` pointer they don't accept) +
kernel source inspection that strips comments and asserts no `dx_in`
write pattern.
Three kernels in aux_trunk_backward_kernel.cu:
- aux_trunk_bwd_dh_pre: per-sample, computes dh_aux2_pre [B, H2] +
dh_aux1_pre [B, H1] using ELU' from POST-activation form
(`(y > 0) ? 1 : (1 + y)` mirrors aux_elu_bwd_from_post in
aux_heads_kernel.cu).
- aux_trunk_bwd_dW_reduce: generic outer-product reduce
`dW[k, j] = sum_b A[b, k] * B[b, j]`. One block per output
element, shmem-tree reduce over batch. Used 3× (dW3, dW2, dW1).
- aux_trunk_bwd_db_reduce: generic batch-reduce `db[j] = sum_b
B[b, j]`. One block per output element. Used 3× (db3, db2, db1).
Memory-efficient: no per-sample partials (avoids B×163,072 floats for
production topology). Per-element reduction means O(P) blocks each
doing O(B) work in shmem.
Rust wrapper AuxTrunkBackwardOps in gpu_aux_trunk.rs orchestrates seven
launches in fixed sequence (capture-friendly, no host branches per
pearl_no_host_branches_in_captured_graph). All three CudaFunction
handles pre-loaded once at construction. Field added to GpuDqnTrainer
alongside aux_trunk_forward_ops; constructor mirrors C.3 pattern.
Tests (all pass on RTX 3050 Ti, sub-ULP forward, 1.33e-2 max rel-err
backward gradient at smallest sampled gradient):
- aux_trunk_forward_matches_numpy_reference (C.3 — preserved).
- aux_trunk_backward_gradient_check (NEW): central-difference
numerical gradient at 16 sampled dW3 indices vs analytic from
backward kernel. Loss = 0.5 * ||h_s2_aux||^2 so dh_s2_aux =
h_s2_aux. EPS=1e-3, B=4, ENC=H1=H2=AUX=32 (33 forwards in ~2s).
REL_TOL = 2e-2 (f32 finite-difference noise floor for
small-gradient tail; production topology is dimension-independent
given runtime args).
- aux_trunk_backward_does_not_write_dx (NEW): reads kernel source,
strips C-style comments (so design-discussion text mentioning
`dx_in` doesn't false-positive), asserts no `dx_in` / `dx_in_out`
symbol survives in code. Complements the structural enforcement
(kernel signatures don't accept `dx_in_out` pointer).
Phase C.4 of SP14 Layer C separate-aux-trunk refactor. Module is
additive — wire-up into collector backward chain + Adam updates lands
in Phase C.5 (atomic).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
cb6bca4629 |
feat(sp14-c): aux trunk forward kernel + Rust wrapper + oracle test
3-layer MLP forward (Linear→ELU→Linear→ELU→Linear). Pre-loaded CudaFunction for graph-capture safety per pearl_no_host_branches_in_captured_graph. Oracle test verifies bit-for-bit match against numpy reference within 1e-4 tol. Saves h_aux1 and h_aux2 to global memory for backward. Phase C.3 of SP14 Layer C separate-aux-trunk refactor (plan: docs/superpowers/plans/2026-05-07-sp14-layer-c-separate-aux-trunk.md). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
4926fb7c65 |
feat(sp14-c): allocate aux trunk params (~132K) + Adam m/v state
3-layer MLP: encoder_out_dim → 256 → 128 → AUX_HIDDEN_DIM. Kaiming-He weights (Box-Muller from LCG-uniform), zero biases, separate Adam m/v buffers (12 state tensors). Allocated in trainer constructor — collector borrows via raw_ptr at wire-up time per existing OFI-embed / q-attn ownership pattern (no parallel param mirror needed). Not yet wired to forward/backward — pure allocation per Phase C.2 design. Topology dimensions resolved against actual codebase: - encoder_out_dim = config.shared_h1 (= SH1 = 256 in production) - AUX_HIDDEN_DIM = config.shared_h2 (= SH2 = 256, matches existing aux head's input dim per aux_heads_kernel.cu:118 `h_s2 [B, SH2]`) Total params: 65,536 + 256 + 32,768 + 128 + 32,768 + 256 = 131,712. Audit doc updated per Invariant 7. Phase C.2 of SP14 Layer C separate-aux-trunk refactor (plan: docs/superpowers/plans/2026-05-07-sp14-layer-c-separate-aux-trunk.md). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
4f372a49a8 |
refactor(sp14-c): atomic α machinery deletion + aux trunk ISV slot allocation
Phase C.1 of SP14 Layer C separate-aux-trunk refactor. Single atomic
commit per feedback_no_partial_refactor and feedback_no_legacy_aliases —
no DEPRECATED stage.
Deleted (per C.0 audit
|
||
|
jgrusewski
|
a7a162d1ff |
docs(sp14-c-preflight): catalogue α-machinery launch sites pending atomic deletion
Phase C.0 of SP14 Layer C separate-aux-trunk refactor. Pure audit doc entry — no source code change. Establishes "before" state of α machinery (Coupling B) launch sites + ISV slots, ahead of atomic deletion in Phase C.7. Files slated for deletion: alpha_grad_compute_kernel.cu, sp14_scale_wire_col_kernel.cu. ISV slots: VAR_AUX/VAR_Q/VAR_ALPHA/ ALPHA_RAW/ALPHA_SMOOTHED/GATE1_OPEN_STATE/GATE2_OPEN_STATE (7 total). Preserved: q_disagreement_* slots [383..390) + producer (diagnostic-only); dir_concat_qaux_kernel (Coupling A: forward feature wire, survives unchanged). Plan: docs/superpowers/plans/2026-05-07-sp14-layer-c-separate-aux-trunk.md Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
fa1f299147 |
fix(sp7-cadence): migrate SP5 Pearl 2 budget + SP7 loss-balance controller from process_epoch_boundary to per-step submit_aux_ops
Bug audit finding #2 (post train-d2b2s diagnostic — same Pattern 1 class as SP14 B.11 commit |
||
|
jgrusewski
|
411a304731 |
fix(aux-head-regime): stop-gradient on aux_regime_backward dh_s2 — completes today's stop-gradient pair
Mirrors commit |
||
|
jgrusewski
|
872bd73927 |
fix(aux-head): stop-gradient on aux's h_s2 input — fixes aux-loss-rises-during-training pathology
Root cause from train-v8ztm 9-epoch HEALTH_DIAG aux next_bar_mse trajectory: - Ep 0: 0.352 (learnable signal — below random baseline ln(2)≈0.693) - Ep 9: 0.717 (above random baseline — aux is now WORSE than random) - aux_dir_acc_long stuck at 0.19 (anti-correlated with truth) Aux head's backward gradient was flowing back to shared trunk activation h_s2 via dh_s2_out write at aux_heads_kernel.cu:599-613. Q-loss gradient on h_s2 dominates (larger magnitude, structurally different objective: cumulative discounted reward vs next-bar direction). h_s2 evolves to support Q's task; aux's CE loss climbs as h_s2 features become anti-aligned with direction prediction. Fix: stop-gradient. Aux reads h_s2 via forward, trains its own w1/b1/w2/b2 from CE loss, but does NOT propagate to h_s2. Q-loss is the sole shaping force on h_s2. Aux must adapt to whatever h_s2 happens to be — if the representation has direction signal, aux's params will extract it; if not, aux can't learn (separate-trunk Option 2 deferred for that case). This was the SEVENTH fix in today's chain (after 6 SP14 EGF cadence/ gate/saturation fixes). The EGF was a scaffold over a broken aux head; fixing aux first is the architectural prerequisite for EGF to route useful signal. Verification: cargo check clean; sp14_oracle_tests 7/7 pass. Validation: aux next_bar_mse should now DECREASE during training in the next L40S smoke (vs the rising-from-0.35-to-0.72 pattern in v8ztm). Deferred follow-up: aux_regime_backward has the same architecture (propagates dh_s2 to trunk). Same fix is a candidate once next_bar result validates the approach. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
c260dca8bd |
fix(sp14-β): remove training-time EGF launches from submit_aux_ops — collector path is canonical
Atomic step 5 of β migration. SP14 EGF producer chain now fires
ONLY from the experience collector (commit
|
||
|
jgrusewski
|
c691bd381a |
feat(sp14-β): wire collector-native SP14 producer chain after aux forward
Step 4 of β migration: 3 SP14/SP13-EGF producers fire per-rollout-step
in collector using collector-owned kernel handles + collector stream.
Reads rollout-time q_values (post-expected-Q, pre-IQR/ensemble/noise)
+ exp_aux_nb_softmax. Writes to shared ISV.
Producer order preserved (matches trainer submit_aux_ops chain):
1. SP13 dir-acc reduce → 2 fixed-α EMAs → aux_pred to ISV[375]
2. SP14 q_disagreement_update (reads aux softmax + q_values)
3. SP14 alpha_grad_compute (pure ISV state machine)
Same kernel gate (commit
|
||
|
jgrusewski
|
296ba92282 |
feat(sp14-β): wire collector-side aux-head forward + label producer per-rollout-step
Step 3 of β migration: collector now runs aux_next_bar_forward on rollout state every step. Label producer (new thin variant aux_sign_label_per_step_kernel) derives sign(price[t+1] - price[t]) per env using bar = episode_starts[ep] + t. Aux predictions feed the EGF kernel chain (step 4), NOT the Q-head's input (rollout Q-head still sees raw h_s2, dir_qaux_concat_ptr remains 0u64). Placement: AFTER captured forward graph, BEFORE expected_q kernel. Same-stream serial ordering reads exp_h_s2_f32 populated by forward_online_f32 inside the captured graph. Cold-start gated on trainer_params_ptr != 0 to skip the test-scaffold path where the trainer hasn't wired its params yet. Files added: aux_sign_label_per_step_kernel.cu (66 lines). Files modified: build.rs (+8 lines, register cubin), gpu_experience_collector.rs (+106 lines: struct field, cubin static, load in new(), per-step launch block). Compile clean; sp14_oracle_tests 2/2 non-GPU pass. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
88eb7aa241 |
feat(sp14-β): allocate rollout-sized aux buffers + AuxHeadsForwardOps in collector
Step 2 of β migration: 5 new buffers sized to alloc_episodes (vs trainer's batch_size). AuxHeadsForwardOps instance is collector- owned and stream-bound to collector stream via AuxHeadsForwardOps::new(&stream). Param tensors shared with trainer via existing f32_weight_ptrs_from_base path. Buffer sizing: exp_aux_nb_hidden_buf [alloc_episodes × 32], exp_aux_nb_logits/softmax_buf [alloc_episodes × 2], exp_aux_nb_label_buf [alloc_episodes] i32, exp_aux_dir_acc_buf [6] mapped-pinned (post-B1.1a 6-float layout matches trainer). Compile clean; sp14_oracle_tests 2/2 non-GPU pass (7 GPU tests ignored on RTX 3050 Ti host). No aux forward yet — buffers allocated, ready for wire. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
09202aa991 |
feat(sp14-β): pre-load SP14 EGF kernel handles in experience collector (Option B)
Step 1 of β migration — Option B (collector-native): collector loads its own CudaFunction handles for the 3 SP14 producer kernels plus their sub-kernels (4 total: aux_dir_acc_reduce, aux_pred_to_isv_tanh, q_disagreement_update, alpha_grad_compute). Mirrors SP13 hold_rate pattern at gpu_experience_collector.rs:1820. Cleaner than cross-component launcher calls (avoids trainer-stream / collector-stream race; no signature surgery on the existing trainer launchers). Cubin static decls flipped to pub(crate) so the collector can re-load on its own stream. Compile clean; sp14_oracle_tests pass (2/2 non-GPU; GPU-gated 7 ignored on RTX 3050 Ti host). No launches yet — additive infrastructure commit. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
9d0c124cee |
fix(sp14-egf): gate q_disagreement EMA update on total_cnt > 0 — fixes training-time decay-to-zero of rollout signal
Root cause from train-6fcml 5-epoch trajectory (commit
|
||
|
jgrusewski
|
5608b866b6 |
fix(producer-cadence): migrate 4 more per-step ISV producers from process_epoch_boundary to per-step hot path
Continuation of SP14 B.11 cadence fix (commit
|
||
|
jgrusewski
|
200f05fcef |
fix(sp14-B.11): move EGF producer chain into per-step training loop — fixes per-epoch staleness
Root cause from train-v8ztm 10-ep validation (commit |
||
|
jgrusewski
|
1396b62ec6 |
fix(sp15-wave5-followup): pre-load sp15_baseline + cost_net cubins — fixes hyperopt-trial CUDA_ERROR_ILLEGAL_ADDRESS
Root cause: 5 SP15 evaluation launchers (cost_net_sharpe + 4 baseline_*) were doing `load_cubin` + `load_function` PER-CALL inside `GpuBacktestEvaluator`'s eval hot loop. Pattern is fragile across CUDA context lifetimes — works in single-pass train-best context (smoke train-9bcwm verified), fails in hyperopt-trial child stream context (workflow train-xggfc trial 1 failed at "load sp15_baseline_kernels cubin: ILLEGAL_ADDRESS"; after the host-side load corrupted the trial's context, trials 2-20 all cascade-failed at "Fork CUDA stream for trial"). Fix (atomic, matches |
||
|
jgrusewski
|
5d63762ab3 |
fix(sp15-wave4.1b-OOB-followup): pre-load bn_tanh_concat_dd_kernel — fixes forward-capture SEGV
Root cause: SP15 Wave 4.1b ( |
||
|
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> |
||
|
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> |
||
|
jgrusewski
|
e9d9afbd61 |
diag(sp15-wave5): stderr checkpoints in run_full_step step-0 ungraphed path
L40S smoke train-jfbzr (commit
|
||
|
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 |
||
|
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, |
||
|
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 (
|
||
|
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 |
||
|
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 (
|
||
|
jgrusewski
|
19f6cce510 |
feat(sp15-wave4.3 / 3.5.5.b): PER sampler integration — recovery transitions oversampled
Phase 3.5.5 (
|
||
|
jgrusewski
|
ef08611d3f |
feat(sp15-wave4.2 / 3.5.4.b): cuRAND Kaiming-He weight reset for plasticity injection
Phase 3.5.4 (
|
||
|
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 ( |
||
|
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 (
|
||
|
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>
|
||
|
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
|
||
|
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>
|
||
|
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
|
||
|
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> |
||
|
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>
|
||
|
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> |
||
|
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> |
||
|
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>
|
||
|
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> |
||
|
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>
|
||
|
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> |
||
|
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> |
||
|
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> |
||
|
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> |
||
|
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> |
||
|
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
|
||
|
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>
|
||
|
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> |