20b1163ffcb67eb2186f679bfa8dec1c4abe0f12
1161 Commits
| Author | SHA1 | Message | Date | |
|---|---|---|---|---|
jgrusewski
|
bc3ebb5fb0 |
docs(claude): foxhunt agents & skills implementation plan — 14 tasks
14-task plan covering: foundation hook router (Task 1), 5 auditor agents (Tasks 2-5, 13), 5 workflow skills (Tasks 6-10), 2 maintenance skills (Tasks 11-12), end-to-end acceptance check (Task 14). Tasks 2-5, 6-8, 9-10, 11-12 fan out in parallel. Task 13 (sp-critical-reviewer) composes the four domain auditors and is built last. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
9e6637c065 |
docs(claude): foxhunt-aware agents & skills design — phase 1 (12 items)
Pattern-cluster auditors + workflow-skill hybrid that internalize accumulated project wisdom (15 feedback_*, 30+ pearl_*, 12+ project_* memory files; 30 SP specs) into foxhunt-namespaced agents and skills under .claude/agents/foxhunt/ and .claude/skills/foxhunt/. Phase 1: 5 auditor agents (isv-discipline, gpu-contract, reward-controller, code-hygiene, sp-critical-reviewer) + 5 workflow skills (sp-spec-writer, isv-slot-scaffolder, pearl-distiller, argo-deploy-helper, smoke-pilot) + 2 maintenance skills (memory-curator, stale-worktree-cleaner). Hook integration is warn-only PostToolUse via a single thin router; agents read memory but only pearl-distiller and memory-curator may write into memory/. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
26deaa5004 |
feat(sp21): T2.2 Phase 1.5 + Phase 2 — entry_q tracking + enrichment real-tape wireup (atomic)
Closes the multi-step T2.2 work. Phase 1.5 captures entry_q (predicted
Q at trade open) on the per-trade tape; Phase 2 wires
enrichment::run_enrichments to the real GpuBacktestEvaluator
per-trade tape and deletes the fake-trade synthesizer
(extract_eval_trades_from_metrics) per feedback_no_stubs.
Plan amendment from on-paper design:
- Audit caught plan's "portfolio_state[ps+6] is unused" claim was
wrong — slot 6 is in active use as cum_return. Replaced with a
dedicated entry_q_state_buf [n_windows] separate from
portfolio_state. Doesn't touch shmem layout, gather kernel, or
model state-dim.
- E5 design question resolved as option (2): quartile-spread Sharpe
with ISV-driven significance anchor read from
ISV[VAL_SHARPE_VAR_EMA_INDEX=351] (per
pearl_controller_anchors_isv_driven). Hardcoded 1.5σ/0.5σ
thresholds eliminated; the noise floor is the val_sharpe variance
EMA already produced per epoch by the early-stopping pipeline.
- Single-step evaluate() launcher passes NULL q_values_per_window
(forward_fn closure exposes only action indices, not Q-values);
same NULL-tolerant pattern as exploration_scale_ptr et al. The
production val pipeline (chunked path) DOES wire real Q-values.
Files (atomic per feedback_no_partial_refactor):
- backtest_env_kernel.cu: 4 new args at end of both kernels
(entry_q_state, q_values_per_window, num_actions,
per_trade_predicted_q_out); pre_entry_q snapshot + close emit +
open/reverse capture.
- gpu_backtest_evaluator.rs: entry_q_state_buf + per_trade_predicted_q_buf
allocated; both reset in reset_evaluation_state; both launchers
migrated; EvalTrade.predicted_q field added; read_per_trade_tape
populates it.
- trainer/enrichment.rs: EvalTrade is now pub(crate) use re-export
from gpu_backtest_evaluator (drops ensemble_var); E5 refactored
to quartile-spread Sharpe with ISV-driven anchor;
extract_eval_trades_from_metrics deleted; HindsightExperience
retained (Phase 6 wiring).
- trainer/training_loop.rs: enrichment block replaced with
evaluator.read_per_trade_tape() + ISV slot 351 read; val_bars /
real_trade_count / real_total_pnl / real_win_rate plumbing
dropped.
- trainer/{mod,constructor,metrics}.rs: last_val_metrics field
removed (last consumer gone — feedback_no_hiding).
- docs/dqn-wire-up-audit.md: 2026-05-11 audit entry.
Verification (passing):
- SQLX_OFFLINE=true cargo check -p ml --tests --features cuda
- sp20_aggregate_inputs_test (12/12)
- sp20_phase1_4_wireup_test (2/2)
- sp20_emas_compute_test (4/4)
- sp20_controllers_compute_test (7/7)
After-this scope (Phase 3-7 + Phase 8 in T2.2 multi-phase):
- E1 q_correction → ISV slot consumer
- E4 per-branch LR scaling via per-group Adam
- E6 winner indices → PER priority bumps
- E7 hindsight → replay buffer injection
- E8 curriculum weights → segment sampling
- Signal-drive remaining controller GAINS (0.9/1.1 in E5; 2.0/0.5/-0.5 in E2)
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
c274b99ea9 |
docs(sp21): T2.2 Phase 1.5 + Phase 2 continuation plan + amendment
Adds the continuation plan for SP21 T2.2 Phase 1.5 (entry_q tracking
in portfolio_state slot 6) + Phase 2 (wire enrichment to real
per-trade tape from gpu_backtest_evaluator). Self-contained plan
that a fresh session can pick up cold:
- State-at-session-start summary with all 8 prior commits
- Phase 1.5 design (storage slot, capture site, plumbing,
EvalTrade extension)
- Phase 2 design (training_loop wire-up, enrichment refactor,
E5 alternate-signal options with recommendation)
- Hard rules carried from prior session (no NULLs, no stubs,
atomic per feedback_no_partial_refactor)
- Verification plan (5 GPU oracle test suites)
- Open design question (E5 alternate signal) with recommendation
- Multi-phase continuation map (Phases 3-7 + Phase 8)
- Final cascade verification (smoke run criteria)
Also amends the parent SP21 plan
(2026-05-10-sp21-train-eval-coherence-isv-defrost.md) with the
T2.2 multi-phase scope section that documents the full per-trade-tape
expansion the user committed to (option 3 — full wiring across
multiple sessions instead of the original recommended option (b)
aggregate-stats refactor).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
7d538d9304 |
feat(sp21): T2.2 Phase 1 Step B — per-trade tape buffers + readback (atomic)
Replaces Step A's NULL launcher passes with real device buffers.
Both kernel launch sites in gpu_backtest_evaluator.rs (backtest_env_step
and backtest_env_step_batch) now pass real per-trade tape pointers.
The kernel's per-trade emission block fires unconditionally on close
events — single-threaded per-window writes preserve event ordering and
enable race-free counter increment without atomicAdd.
New constants + types:
- MAX_TRADES_PER_WINDOW = 200_000 (typical eval window bar count;
per-window memory: 4 SoA buffers × 4 bytes × 200k = 3.2MB)
- pub struct EvalTrade with 5 fields: bar_index, pnl, holding_bars,
direction, magnitude. Does NOT include predicted_q / ensemble_var
— those need entry-time captures (entry_q, entry_var in portfolio
state) deferred to Phase 1.5.
New struct fields on GpuBacktestEvaluator:
- per_trade_pnl_buf: CudaSlice<f32> [n_windows × MAX_TRADES]
- per_trade_holding_bars_buf: CudaSlice<u32>
- per_trade_bar_index_buf: CudaSlice<u32>
- per_trade_dir_mag_buf: CudaSlice<u32> (packed dir/mag)
- per_trade_count_buf: CudaSlice<u32> [n_windows]
New methods:
- reset_per_trade_tape (folded into reset_evaluation_state): zeros
the count buffer at the start of each eval window. SoA value
buffers don't need zeroing — read up to count[w] only.
- pub fn read_per_trade_tape(&self) -> Result<Vec<EvalTrade>, MLError>:
reads count buffer first (cheap), early-returns empty if no trades,
else reads 4 SoA buffers (~16MB DtoH at PCIe ≈ 1ms) and flattens
window-major into chronological Vec<EvalTrade>.
Phase 2 follow-up (next commit) — wire read_per_trade_tape to the
enrichment caller in training_loop.rs:1510-1568, replacing
extract_eval_trades_from_metrics (the fake-trade synthesizer).
Phase 1.5 follow-up (if Phase 2 keeps E1+E5) — add entry_q + entry_var
to portfolio_state at trade open, extend per-trade tape with 6th/7th
SoA buffers.
Affected files:
- crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs
(constants, struct fields, alloc, construction, 2 launcher sites,
reset_evaluation_state addition, read_per_trade_tape method)
Verification:
- cargo check -p ml --tests: passes (warnings only)
- GPU oracle tests: behavior preserved by construction (existing
WindowMetrics aggregator unaffected — separate kernel)
Plan reference: docs/plans/2026-05-10-sp21-train-eval-coherence-isv-defrost.md
T2.2 multi-phase scope; Phase 1 Step B closure.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
5d01190e19 |
feat(sp21): T2.2 Phase 1 Step A — per-trade tape kernel ABI (atomic)
Extends backtest_env_step + backtest_env_step_batch kernel signatures
with 6 new NULL-tolerant args for per-trade tape emission:
- float* per_trade_pnl_out
- unsigned int* per_trade_holding_bars_out
- unsigned int* per_trade_bar_index_out
- unsigned int* per_trade_dir_mag_out (packed hi-16 dir, lo-16 mag)
- unsigned int* per_trade_count
- int max_trades_per_window
Each kernel snapshots entry_price + hold_time BEFORE the
unified_env_step_core call, then post-call mirrors
record_kelly_trade_outcome's close-detection predicate
(is_exit || is_reversal with positive entry_price + valid pre-trade
equity guards) to recompute the realized return for per-trade tape
emission. Single thread per window writes its own slot — race-free
counter increment without atomicAdd per feedback_no_atomicadd.
Step A only — Step B (host alloc + readback + enrichment integration)
follows in a separate commit. This commit is the kernel ABI extension
with NULL launchers, bit-identical to pre-Phase-1 behavior.
NULL-tolerant ABI extension is the codebase's canonical phased-rollout
pattern. Same shape as alpha_per_env / is_win_per_env additive contracts
in sp20_aggregate_inputs_kernel. Per feedback_no_partial_refactor's
"consumer migrates with contract change" rule, the consumer (launcher)
MIGRATES here by passing NULL — output behavior preserved bit-identically.
Step B (next commit):
- Allocate per-trade buffers in GpuBacktestEvaluator constructor
- Pass real device pointers from launcher
- Reset per-window counter at fold start
- Host-side readback into Vec<EvalTrade>
- Wire to enrichment caller, replace extract_eval_trades_from_metrics
Affected files:
- crates/ml/src/cuda_pipeline/backtest_env_kernel.cu (both kernels)
- crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs (both launchers)
Verification:
- cargo check -p ml --tests: passes (warnings only)
- GPU oracle tests: NULL-tolerance contract preserved by construction
(per_trade_count == NULL gates the entire emission block)
Plan reference: docs/plans/2026-05-10-sp21-train-eval-coherence-isv-defrost.md
T2.2 multi-phase scope section.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
4ab1c132e8 |
feat(sp21): T2.1+T2.4 — Q-value early-stop + MIN_HOLD zombies deleted
Closes two architectural-debt items from SP21 Tier 2.
T2.1 — check_early_stopping(avg_q_value) deleted entirely:
- Combined two failed mechanisms: (a) Q-value floor — not a learning
signal (high Q can mean edge OR value explosion, indistinguishable);
(b) Sharpe plateau with hardcoded `improvement < 0.01` threshold,
structurally meaningless against typical val-sharpe deltas O(1-10).
- Both subsumed by the SP21 T1.1a+T1.1b val-loss patience early-stop
with signal-driven min_delta from VAL_SHARPE_VAR_EMA.
- Legacy `old_should_stop` branch + function body deleted.
- Per feedback_no_legacy_aliases.
T2.4 — MIN_HOLD_TARGET / MIN_HOLD_PENALTY_MAX #defines deleted:
- Investigation: macros referenced ONLY in comments and the defining
line itself — no actual code use. The SP12 v3 production callers
were removed in SP20 Phase 2 Task 2.2.
- HEALTH_DIAG line at training_loop.rs:5159 updated to drop the dead
30.0/3.0 literals.
- Scope boundary: MIN_HOLD_TEMPERATURE_* chain is NOT a zombie —
actively wired (kernel producer + SP16 controller consumer).
- Per feedback_no_legacy_aliases.
T2.5 — PER hyperparams disposition (no code change):
- per_alpha=0.6, per_beta_start=0.6 are paper-canonical (Schaul et al.).
Per the SP21 plan recommendation, kept fixed for SP21. Filed for a
separate SP if later identified as a leverage point.
Affected files:
- crates/ml/src/trainers/dqn/trainer/metrics.rs:435-488
(check_early_stopping body deleted)
- crates/ml/src/trainers/dqn/trainer/training_loop.rs (7253 caller +
7291-7322 old_should_stop branch + 5159-5167 HEALTH_DIAG line)
- crates/ml/src/cuda_pipeline/state_layout.cuh:317-318
(#defines deleted)
Verification:
- cargo check -p ml --tests: passes (warnings only)
Cumulative SP21 Tier 2 status: T2.1 ✓, T2.4 ✓, T2.5 ✓ (deferred-doc).
T2.2+T1.3 (enrichment.rs constants soup, ~400 LOC) remaining.
Plan reference: docs/plans/2026-05-10-sp21-train-eval-coherence-isv-defrost.md
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
c90553953c |
feat(sp21): T1.2+T1.4 — enrichment real metrics + backtracking signal-driven (atomic)
Closes the remaining SP21 Tier 1 hardcoded-constant items.
T1.2 — enrichment fed real metrics, not placeholders:
- was: extract_eval_trades_from_metrics(_, 60000.0, 0.0, 0.5, ...)
with hardcoded trade_count=60000, total_pnl=0.0, win_rate=0.5
- now: reads from self.last_val_metrics: Option<[f32; 14]> populated
by val backtest pass at metrics.rs:868. Layout [2]=win_rate,
[4]=total_trades, [7]=total_pnl. Cold-start fallback (None)
is (0.0, 0.0, 0.0) — preferable to fabricated 60000-trade
signal that biased E2/gamma/ensemble from epoch 0.
- Per feedback_no_todo_fixme + feedback_no_stubs.
T1.4 — backtracking thresholds signal-driven:
- Three hardcoded thresholds in run_backtracking_epoch_end replaced
with sigma = sqrt(ISV[VAL_SHARPE_VAR_EMA_INDEX=351]) derivatives:
a) Save trigger (improvement_rate > 0.01) → > 0.5σ.
The 0.01 fired on every epoch (any tiny change > 0.01);
0.5σ requires a meaningful move (typical sigma O(1-10)).
b) Plateau-detection frozen check (abs(delta) < 0.01) → < 0.5σ.
The 0.01 ~never fired; 0.5σ correctly identifies stagnation.
c) Route acceptance (>= min_improvement_rate=0.1) → >= 1.5σ.
Stricter than save-trigger as designed.
- BacktrackingState::min_improvement_rate field deleted — replaced
by per-call signal-driven computation. Floor 0.5 covers cold-start
before var_ema bootstraps from sentinel per
pearl_blend_formulas_must_have_permanent_floor.
- Per feedback_isv_for_adaptive_bounds + feedback_adaptive_not_tuned.
Affected files:
- crates/ml/src/trainers/dqn/trainer/training_loop.rs:1510-1530
(T1.2 enrichment) + :7458-7530 (T1.4 sigma + 3 threshold sites)
- crates/ml/src/trainers/dqn/trainer/mod.rs:108,142
(T1.4 min_improvement_rate field deletion)
Verification:
- cargo check -p ml --tests: passes (warnings only)
- cargo test -p ml --lib early_stopping: 8/8 pass
Cumulative SP21 Tier 1 status: T1.1a ✓, T1.1b ✓, T1.2 ✓, T1.4 ✓,
T2.3 ✓ — Tier 1 closed. Tier 2 (check_early_stopping(avg_q_value)
deletion + enrichment.rs constants soup) is next.
Plan reference: docs/plans/2026-05-10-sp21-train-eval-coherence-isv-defrost.md
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
74c7a80114 |
feat(sp21): T1.1a+T1.1b+T2.3 — signal-driven early-stopping (atomic)
Closes the patience-based early-stopping bug that ran xmd6b 30 epochs
past peak val performance (epoch 2: val_Sharpe=90, total_pnl=0.44 →
epoch 30: val_Sharpe=26, total_pnl=0.16) — a 70% loss of alpha to
training-induced overfitting.
Two intertwined bugs, fixed atomically:
T1.1a (wrong-source) at training_loop.rs:7234:
- was: self.early_stopping.should_stop(-log_output.epoch_sharpe, epoch)
reading the TRAINING ROLLOUT Sharpe (Thompson-noisy, in-sample,
oscillates even when the model is frozen)
- now: self.early_stopping.should_stop(log_output.val_loss, min_delta, epoch)
reading the deterministic-backtest val_loss
- The comment 5733 lines earlier (line 1499) explicitly says "Use
val_Sharpe (deterministic backtest), NOT epoch_sharpe" — patience
path was the inconsistency, backtracking already honored it.
T1.1b (hardcoded threshold) in early_stopping.rs:
- was: EarlyStopping::new(patience, min_delta) with min_delta=0.001
constructor-set, struct field, structurally meaningless
against the val_loss noise floor (typical val_sharpe deltas
are O(1-10), so 0.001 essentially never gates)
- now: EarlyStopping::new(patience), should_stop(val_loss, min_delta,
epoch) with min_delta computed per-call from
ISV[VAL_SHARPE_VAR_EMA_INDEX=351] as
sqrt(var_ema).max(0.5)
- Floor 0.5 covers cold-start before var_ema bootstraps from sentinel
per pearl_blend_formulas_must_have_permanent_floor.
- Per feedback_isv_for_adaptive_bounds + feedback_adaptive_not_tuned.
T2.3 (test signature update) absorbed:
- 6 existing unit tests migrated to new should_stop signature.
- 1 NEW test (test_min_delta_can_change_per_call) verifying per-call
threshold change works correctly.
- EarlyStopping::min_delta struct field deleted.
- Atomic per feedback_no_partial_refactor.
Affected files:
- crates/ml/src/trainers/dqn/early_stopping.rs (struct + tests)
- crates/ml/src/trainers/dqn/trainer/constructor.rs (new() arg)
- crates/ml/src/trainers/dqn/trainer/training_loop.rs (call site)
Verification:
- cargo check -p ml --tests: passes
- cargo test -p ml --lib early_stopping: 8/8 pass
Behavioral expectation post-fix: xmd6b-shape runs (val_Sharpe rising
31→90 epochs 0-2, declining 90→26 epochs 3-30) will trigger early-stop
near the peak. With patience=5 and var_ema bootstrapping by epoch 2-3,
the controller detects "no improvement of ≥ 1σ for 5 consecutive
epochs" by ~epoch 7-8 and stops, saving ~22 epochs of overfitting.
Plan reference: docs/plans/2026-05-10-sp21-train-eval-coherence-isv-defrost.md
Tier 1 status: T1.1a ✓, T1.1b ✓, T2.3 ✓ (this commit). T1.2 + T1.4 next.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
4d4cd996db |
feat(sp21): T3.3 ISV defrost — hold_reward_ema (atomic Phase 3.2 wireup)
Closes the Phase 3.2 forward-reference loop in the SP20 aggregator.
Previously `out_inputs->per_bar_hold_reward = 0.0f` was hardcoded; the
per-bar Hold opportunity-cost producer existed
(experience_kernels.cu:3823 — `per_bar_opp_cost = -aux_conf × cost_scale`)
and wrote to `hold_baseline_buffer` and `r_micro` directly, but never
reached HOLD_REWARD_EMA. Result: HOLD_REWARD_EMA frozen at sentinel 0.0
across all observed epochs; the SP20 reward centering loop
(`r_micro += per_bar_opp_cost - HOLD_REWARD_EMA`) stayed uncentered,
biasing the policy's reward signal away from zero-mean.
T3.3 wireup mirrors the T2.2 alpha refactor: a per-env scratch buffer
that the producer writes at every step (alongside the existing
`hold_baseline_buffer` write), and the aggregator reads at the same
step. Sums opp_cost over Hold-direction envs only; emits the mean as
`per_bar_hold_reward`. The HOLD_REWARD_EMA's gate (`hold_fraction > 0.5f`
from T3.2) preserves the strict-majority semantic from before.
Atomic across producer site, kernel signature, aggregator, launcher,
collector, and tests (per feedback_no_partial_refactor):
- experience_kernels.cu — new `float* per_bar_opp_cost_per_env`
kernel arg, NULL-tolerant write at line ~3823.
- sp20_aggregate_inputs_kernel.cu — new arg, 6th shmem stripe
(`sh_opp_cost_sum`), per-thread Hold-gated accumulation, tree
reduction extended to 6 stripes, output `per_bar_hold_reward =
opp_cost_sum / hold_count` when hold_count > 0 else 0.
- sp20_aggregate_inputs.rs — launcher signature, dynamic_shmem_bytes
5→6 stripes, doc table, internal test.
- gpu_experience_collector.rs — new `per_bar_opp_cost_per_env:
CudaSlice<f32>` field, alloc, struct construction, kernel-arg
pass at both experience_env_step and sp20_aggregate_inputs
launches (raw_ptr).
- sp20_aggregate_inputs_test.rs — helper renamed
`run_kernel_with_is_win_and_opp_cost`, 4 existing sites pass
NULL, 2 NEW oracle tests verifying real producer + NULL fallback.
- sp20_phase1_4_wireup_test.rs — NULL fallback at the wireup site;
HOLD_REWARD_EMA-stays-at-zero assertion remains valid.
Verification:
- cargo check -p ml --tests: passes (warnings only)
- cargo test -p ml --test sp20_aggregate_inputs_test --features cuda
-- --ignored: 12/12 GPU oracle tests pass on RTX 3050 Ti, including
both new T3.3 tests (per_bar_hold_reward_means_over_hold_envs_only,
null_per_bar_opp_cost_emits_zero).
- cargo test -p ml --test sp20_phase1_4_wireup_test --features cuda
-- --ignored: 2/2 pass under the new NULL-tolerant contract.
Plan reference: docs/plans/2026-05-10-sp21-train-eval-coherence-isv-defrost.md
Tier 3 status: T3.1 ✓, T3.2 ✓, T3.3 ✓ (this commit), T3.4 withdrawn,
T3.5 cascade-pending, T3.6 withdrawn.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
1790a31b66 |
feat(sp21): T3.1+T3.2 ISV defrost — wr_ema + hold_pct_ema (atomic)
SP21 Tier-3 foundation: defrost two production-frozen ISV slots that
pinned at 0 across every observed training epoch (d7bj7, xmd6b). Both
bugs in the same aggregator kernel, same structural shape: per-step
binary majority-vote indicators feeding fractional EMAs.
T3.1 — WR_EMA defrost (sp20_aggregate_inputs_kernel.cu):
- was: is_win_out = (2 * wins_count >= closed_count) ? 1 : 0
binary "majority won this step" indicator
- now: win_fraction_out = (float)wins_count / (float)closed_count
fractional in [0, 1]
- With actual val WR≈0.46, the binary signal was 0 most of the time,
pinning WR_EMA at 0 across 30+ epochs in xmd6b. EMA now converges
to population win rate.
T3.2 — HOLD_PCT_EMA defrost (same kernel, same pattern):
- was: action_is_hold_out = (hold_count * 2 > n_envs) ? 1 : 0
strict-majority indicator
- now: hold_fraction_out = (float)hold_count / (float)n_envs
- Cascade: HOLD_COST_SCALE controller compared hold_pct_ema=0 to
tgt±0.05, always saw < lower, ramped × 0.95 → clamped at floor
0.01 (the hold_cost_scale=0.0100 observation in d7bj7 logs).
T3.5 expected to cascade-fix in next training run.
- HOLD_REWARD_EMA gate preserves strict-majority semantic via
hold_fraction > 0.5f test in the consumer kernel.
Atomic across struct fields, kernel logic, Rust mirror, byte
serialization, doc tables, and 4 test files (per
feedback_no_partial_refactor):
- sp20_aggregate_inputs_kernel.cu (struct + 2 computations)
- sp20_emas_compute_kernel.cu (struct + reader + gate)
- sp20_aggregate_inputs.rs (doc table)
- sp20_emas_compute.rs (Rust struct + serialize + 3 tests)
- sp20_aggregate_inputs_test.rs (reader sig + 4 test assertions)
- sp20_emas_compute_test.rs (5 struct literal updates)
- sp20_phase1_4_wireup_test.rs (HOLD_PCT_EMA expected 0.625)
Verification:
- cargo check -p ml --tests: passes (warnings only)
- cargo test -p ml --lib sp20_emas: 6/6 unit tests pass
Pearl candidate: binary-majority aggregator over a fractional
underlying signal cannot serve as input to a fractional-target EMA.
The previous fix (commit
|
||
|
jgrusewski
|
6df44e4c6d |
docs(sp21): plan + revert sp20-aux-h-fixed experiment
Reverts the forced H=30 diagnostic in aux_horizon_update_kernel.cu
(commit
|
||
|
jgrusewski
|
c78c4766ca |
experiment(sp20-aux-h-fixed): force H=30 to test bootstrap failure hypothesis
Throwaway diagnostic edit on aux_horizon_update_kernel.cu. Tests whether the adaptive horizon's self-defeating loop (WR=50% → H~2 → noise horizon → 53.5% acc → WR=50%) is the actual bottleneck. If aux_dir_acc rises to 58%+ at fixed H=30, bootstrap failure confirmed and the proper fix is constraint-based (min hold time during exploration). Audit-doc updated. DO NOT merge to mainline — branch is throwaway after the experiment. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
d6bfad7033 |
docs(sp20): consolidate Phase 5 audit-doc + close-out
Replaces the per-commit audit-doc entries from commits 1-3 with a single
consolidated Phase 5 close-out entry. The consolidated entry covers:
- The full design rationale: gate the REWARD (not the target); avoids
needing q_mean_a entirely. At low aux confidence, r_used → 0 ⇒
Bellman target collapses to gamma * Q(s', a').
- All 5 components: trainer buffer, FusedTrainerCtx accessor, training
loop wire-up, kernel signature + gate, launcher arg.
- NULL-tolerance contract: aux_conf_at_state == NULL OR isv_signals
== NULL ⇒ gate = 1.0 (identity).
- Default-state semantics: alloc_zeros 0.0 sentinel → gate ≈ 0.12 →
reward mostly suppressed pre-population (graceful degradation).
- reward_bias interaction (composes cleanly — gate damps reward
pre-projection, reward_bias lifts target Q-mean per-branch).
- Plan accuracy errata: the user spec's "add aux_conf_at_state_buf
field to GpuBatch struct" was unnecessary — GpuBatch doesn't
carry the SP13 B1.1b aux_sign_labels_ptr either; both follow the
"trainer-only buffer + direct-gather" pattern.
- Test coverage: 3 CPU math tests + 1 GPU behavioral integration test.
- Confidence: medium-high that the gate fires correctly on real data;
end-to-end smoke validation deferred (no smokes dispatched per
controller instruction).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
ab4a7db33c |
feat(sp20): c51_loss launcher aux_conf arg + Phase 5 gate tests
Threads `self.aux_conf_at_state_buf` into the `c51_loss_batched` launch
in `GpuDqnTrainer::launch_c51_loss`. Position matches the kernel's
appended trailing arg from the previous commit.
Tests added in `crates/ml-dqn/src/gpu_replay_buffer.rs::tests`:
- `aux_gate_high_confidence_passes_full_target` (CPU pure-math):
gate(aux_conf=0.5, threshold=0.10, temp=0.05) > 0.99 proves
high-confidence reward pass-through.
- `aux_gate_low_confidence_attenuates_reward` (CPU pure-math):
gate(aux_conf=0.02, threshold=0.10, temp=0.05) < 0.20 proves
the uncertain-state neutralizer semantic.
- `aux_gate_temp_floor_keeps_gate_finite` (CPU pure-math):
sweeps {temp, aux_conf, threshold} and asserts finite gate ∈ [0,1]
across the ISV-controllable parameter range — proves the
fmaxf(temp, 1e-3) floor keeps the kernel numerically safe.
- `aux_conf_direct_to_trainer_gather_populates_destination` (GPU
behavioral): wires a fresh CudaSlice<f32> as the trainer
destination, inserts 8 transitions with strictly-positive distinct
aux_conf values, samples 1, asserts the trainer destination
buffer post-sample holds a value from the inserted set (NOT the
alloc_zeros sentinel) — proves the direct-gather wiring actually
populates the trainer buffer with non-trivial data.
All 3 CPU math tests + 1 GPU integration test pass on RTX 3050.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
96b76d9298 |
feat(sp20): c51_loss_batched aux_conf_at_state reward gate
Adds the Phase 5 consumer kernel-side gate. New kernel arg `const float* __restrict__ aux_conf_at_state` appended to `c51_loss_batched`'s signature. Gate computation runs once per sample at the kernel-entry reward-setup site (after the #27 ensemble- disagreement adjustment), then the gated `reward` propagates through every branch's `block_bellman_project_f` call without per-branch changes. Formula: gate = sigmoid((aux_conf - threshold) / temp) reward = gate * reward where: threshold = ISV[AUX_CONF_THRESHOLD_INDEX=518] temp = max(ISV[AUX_GATE_TEMP_INDEX=519], 1e-3) Mathematical interpretation: at low aux confidence (gate→0), `r_used → 0`, so the Bellman target becomes `gamma * Q(s', a')`. The Q value at the current state collapses toward `gamma * Q(s', a')` — model gets no reward feedback on uncertain transitions. Effectively "don't update Q on uncertain transitions" — the "uncertain-state neutralizer" semantic from the Phase 3 Task 3.4 audit doc spec §4.4. NULL-tolerant: `aux_conf_at_state == NULL` OR `isv_signals == NULL` ⇒ gate skipped (identity, no-op = pre-Phase-5 behaviour). Test scaffolds without a wired aux head still work. Out of scope: `iqn_dual_head_kernel.cu` — IQN is the auxiliary loss, C51 is production. Gating IQN is more complexity for marginal gain. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
d3a057af4f |
feat(sp20): allocate trainer aux_conf_at_state_buf + wire PER direct-gather
Phase 5 plumbing — consumer-side wire-up. Allocates `aux_conf_at_state_buf: CudaSlice<f32>` ([batch_size]) on `GpuDqnTrainer`, exposes the raw_ptr via `aux_conf_at_state_buf_ptr()` and the `FusedTrainerCtx` delegating accessor `trainer_aux_conf_at_state_buf_ptr()`, and invokes `GpuReplayBuffer::set_trainer_aux_conf_ptr` at both fused_ctx init sites in training_loop.rs (init + re-init, atomic per `feedback_no_partial_refactor`). The PER `gather_f32_scalar` now writes the SAMPLED bar's per-batch aux_conf directly into the trainer's f32 buffer on every step — same direct-to-trainer pattern as the SP13 B1.1b `aux_nb_label_buf` (i32) wire-up immediately above the new call. The c51_loss_batched reward gate (lands in the next commit) reads this buffer to compute `gate = sigmoid((aux_conf - ISV[AUX_CONF_THRESHOLD]) / ISV[AUX_GATE_TEMP])` and applies `r_used = gate * reward` at the Bellman projection. `alloc_zeros` cold-start: 0.0 sentinel → at threshold ≈ 0.10 and temp ≈ 0.05 the gate is `sigmoid(-2) ≈ 0.12` → reward is mostly suppressed pre-population. This is the "graceful degradation" semantic from the Phase 3 Task 3.4 audit doc spec §4.4. Once PER's direct-gather populates from the producer ring on the first sample step, the per-bar aux_conf values drive the gate as designed. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
d1a8ec206f |
fix(sp20): bump MAX_UPLOAD_BYTES 2GB → 8GB + audit doc
L40S (48GB) and H100 (80GB) have plenty of headroom; the 2GB cap was a conservative leftover that tripped on workflow zgjgc with 17.8M imbalance bars (3.2GB). 8GB budget leaves ~28GB free on L40S after model + activations + workspace. Updates both call sites: - DqnGpuData::upload_slices (training data, the failing one on zgjgc) - PpoGpuData::upload (market data, same constant) Error message format strings updated 2.0 GB → 8.0 GB. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
6289710463 |
merge(sp20): restore is_win_per_env across struct/launcher/registry after Phase 3 cherry-pick
Cherry-picking Phase 3 ( |
||
|
jgrusewski
|
dd9c70df98 |
feat(sp20): Phase 3 Task 3.4 — replay buffer schema for per-bar aux_conf
Lands the producer→ring→consumer schema plumbing for per-bar aux_conf
(the K=2 peak-softmax-above-uniform confidence value the Phase 5
Aux→Q gate consumes at the Bellman target). Atomic across all struct
boundaries per `feedback_no_partial_refactor`.
Data flow (TWO struct boundaries, not one — plan errata Gap 11):
ExperienceCollector → GpuExperienceBatch → insert_batch()
(.aux_conf field) (8th arg)
→ ring → sample() → GpuBatchPtrs → Trainer
(.aux_conf_ptr) (Phase 5 consumer)
Phase 5 lands the consumer (Aux→Q gate kernel at the Bellman target);
this commit is the producer-side schema plumbing only.
Spec §4.4 Phase 5 forward-reference contract:
At each replay batch step:
aux_conf = ptrs.aux_conf_ptr[k] # SAMPLED bar
threshold = ISV[AUX_CONF_THRESHOLD_INDEX] # [0.01, 0.20]
temp = ISV[AUX_GATE_TEMP_INDEX]
gate = sigmoid((aux_conf - threshold) / temp) # ∈ [0, 1]
Q_target = gate × Q_full + (1 - gate) × mean_a Q(s, a)
Producer (experience_env_step):
- 1 new kernel arg `aux_conf_per_sample: float* [N×L×cf_mult]`.
- Per-bar write at kernel entry: `aux_conf_per_sample[out_off] =
sp20_compute_per_bar_aux_conf_k2(aux_logits + i*K)`.
- CF slot write: `aux_conf_per_sample[cf_off] =
aux_conf_per_sample[out_off]` (CF state shares the same env state
so shares the same aux signal).
- NULL-tolerant: defaults to 0.0f (K=2 uniform-prior sentinel).
GpuExperienceCollector:
- +`aux_conf_per_sample: CudaSlice<f32>` field.
- alloc with `total_output * cf_mult` for both on-policy + CF.
- Threads as kernel arg of experience_env_step.
- Clones into `GpuExperienceBatch.aux_conf` at end of
collect_experiences_gpu (size `total = base_total × 2`).
GpuExperienceBatch: +`aux_conf: CudaSlice<f32>` field.
GpuReplayBuffer (crates/ml-dqn):
- +`GpuBatchPtrs.aux_conf_ptr: u64`.
- +ring storage `aux_conf: CudaSlice<f32>` [capacity].
- +sample buffer `sample_aux_conf: CudaSlice<f32>` [max_batch_size].
- +`trainer_aux_conf_ptr: u64` for direct path.
- +`set_trainer_aux_conf_ptr` setter + `trainer_aux_conf_ptr`
accessor (mirrors `set_isv_signals_ptr` pattern).
- `insert_batch` adds 8th arg `aux_conf_in: &CudaSlice<f32>` —
scatters via `scatter_insert_f32` (reuses the rewards/dones
scatter kernel).
- `sample_proportional` adds gather: direct-to-trainer if
`trainer_aux_conf_ptr != 0`, else fallback into
`sample_aux_conf`. Both paths populate `GpuBatchPtrs.aux_conf_ptr`.
Atomic caller updates (insert_batch arg from 7 to 8):
- training_loop.rs:2522 (production)
- gpu_residency.rs:77 (smoke)
- performance.rs:144 (smoke)
- training_stability.rs:154+201 (smoke ×2)
- gpu_per_integration_test.rs:128 (integration)
- sp15_phase1_oracle_tests.rs:2170+2189+2356 (oracle ×3)
- 3 inline tests in gpu_replay_buffer.rs
Tests:
- test_aux_conf_schema_round_trip (GPU): inserts 8 transitions with
distinct aux_conf [0.0, 0.05, ..., 0.35]; samples 1; asserts the
gathered slot value matches one of the inserted values
(round-trip integrity invariant).
- test_aux_conf_trainer_ptr_wiring_contract (CPU): asserts
set_trainer_aux_conf_ptr / trainer_aux_conf_ptr setter+accessor
behavior matches the SP15 Phase 3.5.5.b mirror pattern.
Verification:
SQLX_OFFLINE=true cargo check -p ml --tests --features cuda # green
SQLX_OFFLINE=true cargo check -p ml-dqn --tests --features cuda # green
SQLX_OFFLINE=true cargo test -p ml --lib sp20 # 21/21 pass
SQLX_OFFLINE=true cargo test -p ml-dqn test_aux_conf_trainer_ptr_wiring_contract # CPU pass
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
06dbb78ffc |
feat(sp20): Phase 3 Task 3.2 — Hold opportunity-cost dual emission (Component 2)
Lands the per-bar Hold opportunity-cost dual emission at
experience_env_step's per-bar branches (positioned-non-event ~3650 +
flat-non-event ~3737), atomically with the trade-close consumer
wiring at the segment_complete branch (~3289 — replaces the Phase 2
Task 2.2 forward-reference placeholder `hold_baseline = 0.0f`) per
`feedback_no_partial_refactor`.
Spec §4.2 dual-emission contract:
per_bar_opp_cost = -aux_conf × cost_scale
Path 1 (Hold-only): r_micro/r_opp_cost += per_bar_opp_cost
- ISV[HOLD_REWARD_EMA]
(centered for Q-target stability)
Path 2 (always): hold_baseline_buffer[env, t % 30] = per_bar_opp_cost
(uncentered, consumed at trade close)
At trade close:
hold_baseline = sp20_sum_hold_baseline_over_trade(buf, env, 30,
current_t,
segment_hold_time)
alpha = R_event - hold_baseline # replaces Phase 2 placeholder
The summation walks backwards `min(30, segment_hold_time)` slots from
`(current_t - 1) % 30` in env i's row of the per-env circular buffer.
The trade-close bar's segment_complete branch does NOT write to the
buffer, so the most recent per-bar write is at current_t - 1.
Layout decision (plan errata Gap 9): per-env stride
`[N_envs × HOLD_BASELINE_BUFFER_SIZE = 30]` row-major. The kernel is
per-env-parallel; a single global ring would interleave bars across
envs and break trade-attribution semantic.
Trade-open tracking decision (plan errata Gap 10): NO new state slot
needed. The existing `segment_hold_time` (= saved_hold_time captured
before PS_HOLD_TIME reset at experience_kernels.cu:2618) plus
current_t suffice — walk backwards in the buffer.
Replaces SP18 D-leg `compute_sp18_hold_opportunity_cost` calls at
experience_kernels.cu:3672 and :3783. The device function in
trade_physics.cuh:655 is RETAINED — still called by
sp18_hold_opp_test_kernel.cu (oracle test surface). Per
`feedback_no_stubs`: only production callers migrate; the helper
stays for the test surface.
New device helpers (sp20_hold_baseline.cuh):
- sp20_compute_per_bar_aux_conf_k2(logits) — bit-identical to the
formula in sp20_stats_compute_kernel.cu Pass A.
- sp20_sum_hold_baseline_over_trade(buf, env, size, current_t,
segment_hold_time) — walks backwards with modulo wrap-around.
New buffer + reset infrastructure:
- GpuExperienceCollector.hold_baseline_buffer field
- HOLD_BASELINE_BUFFER_SIZE = 30 constant (re-exported)
- StateResetRegistry FoldReset entry + invariant test
- reset_named_state dispatch arm
Six new kernel args appended to experience_env_step signature:
aux_logits_per_env, hold_baseline_buffer, hold_buffer_size, aux_k,
hold_cost_scale_idx, hold_reward_ema_idx. All NULL-tolerant.
HOLD_REWARD_EMA forward-reference: ISV[516] is updated by
sp20_emas_compute_kernel reading per_bar_hold_reward from
sp20_aggregate_inputs_kernel which currently emits 0.0f as a Phase
3.2 forward-ref placeholder (line 292). The parallel `sp20-phase-2-fix`
agent owns wiring the real producer; Task 3.2's centering math
references the ISV slot (not a hardcoded 0), so when the parallel
branch lands, EMA starts updating and centering becomes load-bearing
automatically. No additional Task 3.2 work needed post-merge.
Tests (sp20_hold_baseline_test.rs, 4 GPU oracle tests):
1. aux_conf_k2_bounds_and_fixed_points — bounds, uniform/saturated/
spec-warmup/spec-confident/symmetry fixed points.
2. sum_hold_baseline_over_trade_indexing — basic indexing,
hold_time>size clamp, defensive guards, modulo wrap-around.
3. sum_hold_baseline_per_env_stride — env-stride correctness across
row-major buffer.
4. dual_emission_hold_vs_non_hold — full Path 1 + Path 2 contract,
mirrors plan §3.2 reference fixture.
Verification:
SQLX_OFFLINE=true cargo check -p ml --tests --features cuda # green
SQLX_OFFLINE=true cargo test -p ml --lib sp20 # 21/21 pass
SQLX_OFFLINE=true cargo build -p ml # cubin compile
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
13ce783853 |
test(sp20): post-fix runtime diagnostic — read-back accessors + producer test scaffold
Adds infrastructure for verifying the SP20 Phase 2 Task 2.2-fix
(`is_win_per_env`) producer side at runtime. The consumer-side oracle
test passes; the producer-side wire-up has not been independently
verified end-to-end. Production HEALTH_DIAG observed `wr_ema=0.0000`
across all training epochs even after `64bbbe418` landed, with
`alpha_ema` evolving normally — strongly suggesting either (a) all
profitable trade closes had `hold_time == 0` so the producer write
didn't fire, or (b) the producer write fired but `(segment_return >
0.0f)` evaluated to false in the production data.
Code-inspection audit (this commit, summarized in audit-doc):
- Verified 81 kernel args (launcher) match 81 kernel signature args
for `experience_env_step`. `is_win_per_env` at position 81;
`alpha_per_env` at position 78. cudarc `PushKernelArg<&mut
CudaSlice<T>>` impl pushes `&arg.cu_device_ptr` (stable address
for the lifetime of the chained statement).
- Verified buffer allocation at construction (alloc_zeros), pub(crate)
field, FoldReset entry + dispatch arm, kernel-arg pass at both
experience_env_step and sp20_aggregate_inputs launch sites.
- Verified producer site — both alpha + is_win writes inside the SAME
`if (segment_complete && segment_hold_time > 0.0f)` block at
experience_kernels.cu:3196, both NULL-guarded, race-free.
- Verified consumer site — aggregate kernel reads `is_win_per_env[env]`
+ `alpha_per_env[env]` side-by-side inside `if (tc != 0)`.
- Verified stream ordering — both kernels run on `self.stream` in the
same `for t in 0..timesteps` loop; experience kernel at line 6069,
aggregate at 6577. Stream-implicit producer→consumer ordering.
**No code-level bug surfaced by inspection.** The `is_win_per_env`
write is structurally identical to the `alpha_per_env` write 14
lines earlier in the same gating block.
Files modified:
- `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` —
`read_is_win_per_env_for_test()` and `read_alpha_per_env_for_test()`
pub fns. Mirror `read_epoch_dsr_state()` pattern: dtoh memcpy via
the collector's stream, returns Vec<i32> / Vec<f32>. Cost-free in
production (only called from tests) and benign for invariants.
- `crates/ml/tests/sp20_is_win_producer_test.rs` — production-path
producer test scaffold. Builds 6-float-per-bar synthetic OHLCV
matching `experience_env_step`'s tgt[0..6] contract, runs
`collect_experiences_gpu`, reads back both buffers, asserts the
three structural invariants:
1. alpha_per_env has at least one non-zero slot (precondition —
the synthetic upward-trend data must produce ≥ 1 trade close).
2. is_win_per_env has at least one slot == 1 (regression-pin for
the Task 2.2-fix producer wire-up).
3. Co-occurrence: every env where alpha > 0 must have is_win == 1
(R_event > 0 ⇒ segment_return > 0 ⇒ is_win = 1; structural
guarantee from the kernel's two writes at the same gate).
Currently `#[ignore]` because the collector's cuBLAS forward chain
requires a real `trainer_params_ptr` (NUM_WEIGHT_TENSORS=163
weights). Activation requires either trainer-params plumbing in the
test scaffold OR moving the test inline as `#[cfg(test)]` in
gpu_experience_collector.rs.
- `docs/dqn-wire-up-audit.md` — audit-doc entry per Invariant 7.
Documents the inspection scope, the unverified producer-side
pre-conditions, and the open-issue framing for the wr_ema=0.0000
discrepancy.
Verification:
- SQLX_OFFLINE=true cargo check -p ml --tests passes
- 9 state_reset_registry tests pass (incl.
sp20_is_win_per_env_registered_fold_reset)
- The new test file compiles cleanly under `--features cuda`
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
64bbbe4181 |
fix(sp20): WR_EMA pinning bug — wire segment-level is_win_per_env producer + kernel-body read
Closes the fix gap left by the failing-test commit. The aggregation kernel body now reads `is_win_per_env[env]` (when non-NULL) in the `tc != 0` branch in place of the broken per-bar `step_ret > 0` predicate; the producer site at `experience_kernels.cu::segment_ complete` writes the segment-level `(segment_return > 0.0f) ? 1 : 0` flag at the same location as `alpha_per_env[i] = alpha`. Atomic changes: - sp20_aggregate_inputs_kernel.cu: when is_win_per_env != NULL, read the per-env i32 flag for wins_count; legacy `sr > 0` predicate preserved as NULL-tolerant fallback for oracle-test scaffolds. - experience_kernels.cu: new `int* is_win_per_env` kernel arg; segment_complete branch writes `(segment_return > 0.0f) ? 1 : 0` at the same site as `alpha_per_env[i] = alpha`. Race-free per-i write (mirrors the alpha producer pattern). NULL-tolerant. - gpu_experience_collector.rs: `is_win_per_env: CudaSlice<i32>` `[alloc_episodes]` field, alloc_zeros at construction, kernel-arg pass at experience_env_step launch site, raw_ptr() pass at sp20_aggregate_inputs launch site. - state_reset_registry.rs: FoldReset entry + assertion test (sp20_is_win_per_env_registered_fold_reset). The every_fold_and_soft_reset_entry_has_dispatch_arm test passing confirms the dispatch arm is present. - training_loop.rs: `"is_win_per_env" => memset_zeros` arm. - docs/dqn-wire-up-audit.md: 2026-05-10 fix-commit entry. Bug pattern for the memory pearl: per-bar mark-to-market step_return is NOT the trade's segment-level win/loss outcome at close bars because tx_cost (deducted via `*cash -= cost` in trade_physics.cuh::execute_trade) plus the per-bar Δprice tick swamp out segment-accumulated unrealized gains. The segment-level realized P&L (`(realized_pnl + unrealized_at_exit - trade_start_pnl) / prev_equity`) is the correct outcome scalar; the producer in segment_complete already had it (used by sp20_compute_event_reward's sign check for alpha), the aggregation kernel just wasn't reading the right thing. Bug latent since SP20 Phase 1.4 / 2.2 landed on 2026-05-09; caught 2026-05-10 by HEALTH_DIAG observation in a multi-seed L40S smoke (alpha_ema populated correctly while wr_ema pinned at 0.0000 across 17 epochs). Verification: - SQLX_OFFLINE=true cargo check -p ml --tests passes - 21 SP20 lib tests pass; 9 state_reset_registry tests pass - The new GPU oracle test wr_ema_uses_segment_pnl_via_is_win_per_env_predicate passes after this commit (would FAIL on prior commit's kernel-body) - The NULL-tolerance pin null_is_win_per_env_falls_back_to_legacy_step_ret_predicate preserves the pre-fix oracle-test contract Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
5a29c37cd8 |
test(sp20): failing test pinning WR_EMA bug — wins predicate must use segment P&L not per-bar step_ret
Plumbs new is_win_per_env i32 device-buffer arg through the
sp20_aggregate_inputs kernel signature and Rust launcher, NULL-tolerant
to preserve the legacy `step_ret_per_env > 0` predicate for existing
oracle-test scaffolds. Production caller passes 0 (NULL) in this
commit; commit 2 wires the per-env producer + collector buffer + reset
entry atomically with the kernel-body change that honors the arg.
Adds two new GPU oracle tests in sp20_aggregate_inputs_test:
- wr_ema_uses_segment_pnl_via_is_win_per_env_predicate (FAILS without
fix): step_ret all negative (tx-cost dominated close bars — the
realistic production case), is_win_per_env = [1,1,1,0]. With the
legacy predicate wins_count = 0 ⇒ is_win = 0 (the production bug).
With the fix wins_count = 3 ⇒ is_win = 1.
- null_is_win_per_env_falls_back_to_legacy_step_ret_predicate: pins
the NULL-tolerance contract so existing oracle tests + Phase 1.4
wireup test scaffolds keep working unchanged.
Bug root cause: WR_EMA pinned at 0 across all training epochs in
production (HEALTH_DIAG observed `wr_ema=0.0000` while `alpha_ema`
evolved with real values). The aggregation kernel's win predicate was
`step_ret_per_env[env] > 0.0f` at the close bar, which is the per-bar
mark-to-market `(new_value - prev_equity) / prev_equity`, NOT the
trade's segment-level outcome. At close bars `step_ret = position *
(close_t - close_{t-1}) - tx_cost`; the per-bar tick is small while
tx_cost is fixed → `step_ret < 0` is dominant across closing bars
even for trades that closed profitably overall. Result: wins_count =
0 always, is_win = 0 always, WR_EMA pinned at 0, LOSS_CAP pinned at
-1.0 (cold-start cap, never ramping to -2.0 per spec §4.1).
Fix arrives in the next commit: experience_kernels.cu segment_complete
branch writes `is_win_per_env[i] = (segment_return > 0.0f) ? 1 : 0`
(the segment-level realized-return signed-flag), passed through to
the aggregation kernel which uses the new per-env i32 buffer in place
of the per-bar `sr > 0` predicate when wired.
Per `feedback_no_partial_refactor` the kernel signature change + all
callers (production + 2 test files) migrate atomically in this commit;
the kernel-body behavioral change + producer wire-up + reset registry
+ audit-doc entry land atomically in the immediately following commit.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
e1aef53373 |
docs(sp20): plan accuracy errata for Phase 2 (Tasks 2.0/2.1/2.2/2.3/2.4)
Append a "Plan Accuracy Errata" section to the SP19+20 plan
documenting the 6 deviations from the original Phase 2 plan that
emerged during implementation. Each entry captures the gap, the
decision made, and the rationale, so future implementers see what
was actually built vs what was specified.
Gaps documented:
1. Task 2.0 (label-at-open infrastructure) was NEW — split out from
Task 2.2 to land buffer + alloc + reset + write site atomically
before Task 2.2's consumer ships. Sign convention mapping detail
captured (kernel emits {0, 1, -1}, SP20 spec uses {-1, 0, +1}).
2. Per-env alpha plumbing was not in Task 2.2 plan scope — built in
Task 2.2 (NOT deferred to Phase 4) to preserve the
`feedback_no_partial_refactor` contract atomicity. Touches 5
files in one commit.
3. Per-bar SP18 D-leg sites — KEEP for Phase 2 per `feedback_no_stubs`.
Plan's wording "DELETE [these helpers]" was overbroad; only the
trade-close-site call is deleted. The phantom
`compute_sp12_reward_with_cost` doesn't exist as a function (the
SP12 v3 reward is the inlined block).
4. `sp20_compute_event_reward` placement — new dedicated
`sp20_reward.cuh` header (NOT inside `experience_kernels.cu`,
NOT inside `trade_physics.cuh`). Mirrors the
compute_asymmetric_capped_pnl / compute_min_hold_penalty
header-only pattern; needed for GPU oracle test wrapper to share
the function bit-for-bit per `feedback_no_cpu_test_fallbacks`.
5. Task 2.3 was subsumed by Task 2.2 — the existing Path C chain
consumes the new `alpha` field automatically once `alpha_per_env`
is wired; no separate `sp20_emas_compute` producer call needed.
6. `min_hold_*` kernel-arg trio cleanup deferred to Task 2.4 — the 3
kernel args were deleted in Task 2.2 (per `feedback_no_hiding`)
but the upstream producer chain (`min_hold_temperature_update_kernel`,
ISV[460], `read_min_hold_temperature_from_isv`, `config.min_hold_*`)
deferred to Task 2.4 because it touches SP14 ISV slot registry +
StateResetRegistry + ISV layout fingerprint bump.
Implementation-level details remain in `docs/dqn-wire-up-audit.md`
Task 2.0 / 2.1 / 2.2 entries; this errata is the plan-level
"what was actually built vs what was specified".
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
3f9b34030c |
feat(sp20): Phase 2 Task 2.2 — SP12 v3 reward block → 4-quadrant + alpha plumbing (atomic)
Lands the SP20 4-quadrant reward at the experience_env_step::
segment_complete site, atomically with per-env alpha plumbing through
the SP20 aggregation kernel per `feedback_no_partial_refactor`:
* `experience_env_step.cu` — replace SP12 v3 inlined block (vol-norm
calc + base_reward + SP18 D-leg trade-close call + asymmetric cap +
min-hold penalty + r_popart/r_trail cascade, ~260 LoC) with the SP20
4-quadrant reward (~80 LoC) computing:
R_event = sp20_compute_event_reward(segment_return,
label_at_open_per_env[i],
ISV[LOSS_CAP_INDEX])
alpha = R_event - 0.0 (Phase 3.2 hold_baseline placeholder)
r_used = alpha - ISV[ALPHA_EMA_INDEX] (advantage centering)
alpha_per_env[i] = alpha (consumed by sp20_aggregate)
r_popart / r_trail = r_used (SP11 component split preserved)
Adds 3 new kernel args at the end (preserves Task 2.0 args):
float* alpha_per_env, int loss_cap_idx, int alpha_ema_idx
Deletes 3 SP12 v3 args (min_hold_target, min_hold_penalty_max,
min_hold_temperature) per `feedback_no_hiding` "wire up or delete".
* `sp20_aggregate_inputs_kernel.cu` — add `alpha_per_env` input arg
(NULL-tolerant), 5th shmem stripe (sh_alpha_sum, was 4 stripes),
per-thread close-gated accumulation, 5-way tree-reduce loop body,
output write `alpha = mean / closed_count`. Replaces the Phase 1.4
hardcoded `out_inputs->alpha = 0.0f` placeholder atomically.
* `sp20_aggregate_inputs.rs` launcher — +1 arg `alpha_per_env_dev`
(passes 0/NULL for tests). Shmem byte-count: 4 stripes → 5 stripes.
* `gpu_experience_collector.rs` — new `alpha_per_env: CudaSlice<f32>`
field [alloc_episodes], allocated next to `label_at_open_per_env`,
threaded into env_step launch + sp20_aggregate launch. Removes the
3 obsolete config.min_hold_* `.arg(...)` calls from env_step launch.
* `state_reset_registry.rs` — `alpha_per_env` FoldReset entry +
invariant test `sp20_alpha_per_env_registered_fold_reset`.
* `training_loop.rs::reset_named_state` — `alpha_per_env` dispatch arm
via `stream.memset_zeros` (mirrors `label_at_open_per_env` pattern).
* Pin-test updates per spec (atomic with the contract change):
- `sp20_aggregate_inputs_test::alpha_and_per_bar_hold_reward_are_
phase_2_and_3_2_placeholders` →
`null_alpha_producer_keeps_phase_1_4_placeholder_contract` (asserts
NULL alpha producer ⇒ alpha=0.0 backward-compat).
- `sp20_phase1_4_wireup_test::alpha_and_hold_reward_emas_stay_at_
zero_phase_2_3_2_placeholders` →
`..._with_null_producers` (same NULL-tolerance pin).
- +2 NEW tests in `sp20_aggregate_inputs_test`:
* `alpha_mean_over_closed_envs_phase_2_contract` — 4 envs (3
close, 1 doesn't), asserts close-gated mean (env 3's alpha=99
ignored).
* `alpha_zero_when_no_close_phase_2_contract` — no-close steps
emit alpha=0.0 regardless of stale alpha_per_env content.
* `dqn-wire-up-audit.md` — Task 2.2 entry documents the deleted SP12 v3
block, the per-env alpha plumbing, the retained device functions
per `feedback_no_stubs`, and the deferred Task 2.4 cleanup of the
orphaned min_hold_temperature_update_kernel producer chain.
Per `pearl_event_driven_reward_density_alignment`: per-bar SP18 D-leg
sites at experience_kernels.cu:3722 / :3833 are RETAINED pending
Phase 3.2 Component 2 replacement. Per `feedback_no_stubs`:
compute_sp18_hold_opportunity_cost / compute_asymmetric_capped_pnl /
compute_min_hold_penalty device functions are RETAINED (still called
by per-bar SP18 sites + sp12_reward_math_test_kernel test wrapper).
Per spec §4.1: alpha_ema centering (`R_used = alpha - alpha_ema`) is
load-bearing for cold-start learning at WR=46% with raw EV ≈ -0.06,
NOT just Q-target stability. Documented in the kernel comment block.
Verification:
SQLX_OFFLINE=true cargo build -p ml # cubin compile
SQLX_OFFLINE=true cargo check -p ml --tests --features cuda # tests compile
SQLX_OFFLINE=true cargo test -p ml --lib sp20 # 20/20 lib tests
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
eaaab152fc |
feat(sp20): Phase 2 Task 2.1 — sp20_compute_event_reward + 8 GPU oracle tests
Lands the SP20 Component 1 reward math as a header-only `__device__
__forceinline__` function in a new `sp20_reward.cuh` header, plus the
GPU oracle test scaffold for the 4-quadrant truth table per spec §4.1.
* `sp20_reward.cuh` — `sp20_compute_event_reward(close_pnl,
label_at_open_sign, loss_cap)` function. Returns bounded scalar
`R_event ∈ [loss_cap, +1.0]`. 4-quadrant table:
+1, +1 → +1.0 (right reason, right outcome)
+1, -1 → +0.5 (wrong reason, right outcome)
-1, -1 → -0.5 (right reason, wrong outcome)
-1, +1 → loss_cap (wrong reason, wrong outcome — ISV-driven)
0, * → 0.0 (no information — close_pnl == 0)
Sentinel `label_at_open_sign == 0` (Task 2.0 contract) maps to
dir_match=false ⇒ wrong-reason quadrant. Documented as the safer
default in the header comment.
* `sp20_event_reward_test_kernel.cu` — standalone single-thread test
wrapper, mirrors the `sp12_reward_math_test_kernel.cu` /
`thompson_test_kernel.cu` pattern. Outputs to a mapped-pinned [1]
f32 with `__threadfence_system()` for PCIe-visible coherence.
* `tests/sp20_event_reward_test.rs` — 8 GPU oracle tests
(`#[ignore = "requires GPU"]`-gated):
1. quadrant_right_reason_right_outcome
2. quadrant_wrong_reason_right_outcome
3. quadrant_right_reason_wrong_outcome
4. quadrant_wrong_reason_wrong_outcome_loss_cap_at_minus_1
5. quadrant_wrong_reason_wrong_outcome_loss_cap_at_minus_2
6. zero_pnl_returns_zero
7. sentinel_label_winning_trade_lands_shoulder (Task 2.0 contract)
8. sentinel_label_losing_trade_lands_loss_cap (Task 2.0 contract)
* `experience_kernels.cu` — `#include "sp20_reward.cuh"` so Task 2.2's
trade-close site replacement has the function in scope.
* `build.rs` — `sp20_event_reward_test_kernel.cu` cubin entry.
No production callers in this commit. Task 2.2 atomically:
- replaces the SP12 v3 reward block at experience_kernels.cu:3216-3380
with the SP20 4-quadrant reward;
- adds the `is_close` consumer of `label_at_open_per_env[i]` (Task
2.0's per-env scratch);
- threads per-env alpha through the SP20 aggregation kernel,
making `alpha_ema` non-zero post-trade-close (replacing the
Phase 1.4 0.0 forward-reference placeholder).
Per `feedback_no_partial_refactor.md` the device function ships
BEFORE the production caller so Task 2.1's GPU oracle tests can
exercise the math in isolation; the partial refactor that breaks
the no-partial-refactor rule is "feature behind the function with
no caller", which this commit's tests resolve in scope.
Per `feedback_no_cpu_test_fallbacks.md` GPU oracle only — no CPU
reference impl. Per `feedback_no_htod_htoh_only_mapped_pinned.md`
all CPU↔GPU buffers are mapped-pinned (`MappedF32Buffer`).
Verification:
SQLX_OFFLINE=true cargo check -p ml --tests --features cuda # clean
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
29a2615f6a |
feat(sp20): Phase 2 Task 2.0 — per-env trade-open label-sign infra
Lands the producer-side infrastructure for SP20 Phase 2's 4-quadrant
reward kernel (Task 2.2):
* New `[alloc_episodes]` `i32` device buffer `label_at_open_per_env`
on `GpuExperienceCollector`, allocated alongside `episode_starts_buf`.
* Two new `experience_env_step` kernel args (`aux_label_per_env` input,
`label_at_open_per_env` output), threaded into the launch site.
* Write site at the `entering_trade` branch maps the upstream aux
next-bar label (`{0,1,-1}` from `aux_sign_label_per_step_kernel`) to
the SP20 spec sign convention (`{-1,+1,0}`) per spec §4.1.
* StateResetRegistry entry `label_at_open_per_env` (FoldReset) +
`reset_named_state` dispatch arm via `stream.memset_zeros`.
* Registry invariant test `sp20_label_at_open_per_env_registered_fold_reset`
pins the FoldReset category + key description anchors.
* Audit-doc entry documents the design rationale (why per-env buffer
vs. derived-at-trade-close read), sign mapping, FoldReset semantics,
kernel-arg threading, and Task 2.1/2.2 forward references.
Producer-only this commit. The consumer (`is_close` branch reading
`label_at_open_per_env[i]` and passing to `sp20_compute_event_reward`)
lands atomically with the SP12 v3 reward block replacement in Task 2.2,
per `feedback_no_partial_refactor`.
NULL-tolerant kernel arg pair lets test scaffolds without aux-head
wiring continue to work; the FoldReset-zeroed buffer reads as sentinel
0 at the consumer (which Task 2.2 maps to "no info" / wrong-reason
quadrant — safer default than over-rewarding a no-signal trade).
Verification:
SQLX_OFFLINE=true cargo check -p ml # clean
SQLX_OFFLINE=true cargo test -p ml --lib sp20 # 18/18 (+1 new)
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
abd7e533bc |
fix(architectural): volume_bar_size in cache key + OFI front-month filter
## Two architectural cleanups, both surfaced by the wgdc8 experiment ### Part 1: volume_bar_size in cache key Mirrors the imbalance_bar_threshold/ewma_alpha fix from `f7718b376`. The volume bar size constant (100 contracts/bar) was previously hardcoded and not in the fxcache key. Tuning it would have hit the same fossilization bug as imbalance_bar_threshold did pre-fix. Changes: - `Hyperparams.volume_bar_size: u64` field added (default 100, matches `DEFAULT_VOLUME_BAR_SIZE` for backwards compat). - TrainingProfile loader reads `volume_bar_size` TOML key. - `calculate_dbn_cache_key_full` signature 7 → 8 args. Hashed via `to_le_bytes()`. Test `test_cache_key_includes_volume_bar_size` added; passes alongside the 5 existing tests. - 4 callers updated atomically (per `feedback_no_partial_refactor`): `discover_and_load`, `data_loading.rs:146`, `train_baseline_rl.rs:599`, `precompute_features.rs:259,720`. - `data_loading.rs:279` now passes `self.hyperparams.volume_bar_size` to `build_volume_bars` instead of the hardcoded `DEFAULT_VOLUME_BAR_SIZE`. - New `--volume-bar-size` CLI arg on both binaries (default 100). - New Argo workflow params `volume-bar-size` (default "100") and `data-source` (default "mbp10") on both `train-template.yaml` and `train-multi-seed-template.yaml`. Threaded into precompute + trainer invocations. - `scripts/argo-train.sh` exposes `--volume-bar-size <n>` and `--data-source <s>` for ad-hoc overrides. ### Part 2: OFI front-month filter (latent bug fix) `crates/ml/examples/precompute_features.rs:539-557` (the OFI/VPIN/Kyle's Lambda computation branch when MBP-10 + trades data is available) was loading trades unfiltered for per-bar microstructure feature computation. The volume bar formation path filters front-month per-file (line 354), but the OFI path did not. Effect pre-fix: during contract rollover windows (e.g., ESZ24 → ESH25), OFI per-bar microstructure features included trades from BOTH contracts simultaneously, distorting VPIN, Kyle's Lambda, and trade imbalance signals. Severity in production: small (front-month dominates ES.FUT volume by 10-100×) but real and present in every prior MBP-10+trades production run. Fix: mirror the per-file `filter_front_month` call from the volume bar path. Volume bar formation and OFI computation now both see the same in-month trade tape. Added log line shows raw vs filtered count per file for transparency. ## Why bundled Both fixes touch trade-data plumbing in `precompute_features.rs` and the fxcache key contract. Per `feedback_no_partial_refactor`, related architectural cleanups land atomically. Both surfaced from the same wgdc8 audit; bundling avoids two cache-key-invalidating commits in sequence (each would force full fxcache regen). ## Compatibility - `volume_bar_size` defaults to 100 → existing wgdc7-equivalent runs reproduce, but with a *new* fxcache key (the f7718b376-era cache file is unreachable; harmless, can GC manually). - OFI fix is strictly more correct; no opt-out needed. Existing models trained on contaminated OFI features may show slight feature distribution drift on first cache regen — expected, not a regression. - `data_source = "ohlcv"` Argo param now possible; routes precompute through volume bar branch directly. wgdc8 experiment uses this to test bar resolution sensitivity at volume_bar_size=500 (5× DEFAULT). Tests: 6/6 feature_cache tests pass. Workspace + examples compile clean. Audit-doc: `docs/dqn-wire-up-audit.md` updated. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
f7718b3761 |
fix(architectural): include bar formation params in fxcache key + actually USE imbalance bars
## The bug (audit 2026-05-09) `crates/ml/src/feature_cache.rs:calculate_dbn_cache_key_full` hashed only `(symbol, data_source, dbn_filenames+sizes)` — NOT `imbalance_bar_threshold` or `imbalance_bar_ewma_alpha`. Combined with `precompute_features.rs:346` unconditionally calling `build_volume_bars` regardless of `data_source`, this meant: 1. 14 audited production runs (Apr 11–May 8) all collided on the same fxcache key (`a3f933aa...` / `c07c960a...`) regardless of TOML `imbalance_bar_threshold` value 2. The imbalance-bar code path was reachable only via fxcache MISS, which never happens in production because `ensure-fxcache` always populates first 3. Every "tuning" of `imbalance_bar_threshold` across 16+ SP runs was a silent no-op — the system was actually running volume bars at DEFAULT_VOLUME_BAR_SIZE (100 contracts/bar) ## The fix (this commit) **Part A — cache key includes bar formation params:** - `calculate_dbn_cache_key_full` signature: 5 args → 7 args. Two new f64 params hashed via `to_le_bytes()`. - 4 callers updated atomically (per `feedback_no_partial_refactor`). - 2 new unit tests (`test_cache_key_includes_bar_threshold`, `test_cache_key_includes_bar_alpha`) pin the contract. **Part B — precompute_features actually USES data_source:** - New CLI args `--imbalance-bar-threshold` (default 0.5) and `--imbalance-bar-ewma-alpha` (default 0.1) on both train_baseline_rl and precompute_features. - `precompute_features.rs:346` now branches: when `data_source == "mbp10"` AND `mbp10_data_dir.is_some()`, calls `mbp10_to_imbalance_bars` instead of `build_volume_bars`. **Argo plumbing:** - `train-template.yaml` + `train-multi-seed-template.yaml`: new workflow parameters threaded into BOTH precompute and trainer invocations so both compute the same fxcache key. - `scripts/argo-train.sh`: new CLI flags for ad-hoc overrides. - ensure-fxcache regen path: removed `rm -f /feature-cache/*.fxcache` (with bar-params now in key, parallel experiments coexist). ## Effects going forward - Tuning `imbalance_bar_threshold` actually changes bar density - Configuring `data_source = "mbp10"` actually produces imbalance bars - Multiple parallel experiments at different thresholds coexist on PVC - `dqn-production.toml: imbalance_bar_threshold = 0.5` no longer ignored Default values match prior production behavior → existing wgdc7-equivalent runs reproduce, just with a *new* fxcache key (the old volume-bar cache file is still on disk but won't be hit; harmless, can GC manually). Audit-doc: `docs/dqn-wire-up-audit.md` updated with full context. Tests: 5/5 feature_cache tests pass, full workspace + examples compile clean. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
235e838422 |
feat(sp20): Phase 1.4 (Path C) — kernel-arg refactor + aggregation kernel + wire-up
Atomic Phase 1.4 commit per `feedback_no_partial_refactor`. Wires the
SP20 fused-producer chain (Stats → Aggregate → EMAs → Controllers) into
GpuExperienceCollector's per-rollout-step path with one new aggregation
kernel + Phase 1.2 EMA kernel-signature refactor + production wire-up
+ tests + audit/spec/plan amendments, all in one commit.
## Path C decision rationale
The Phase 1.2 EMA kernel originally took 9 scalar value-args. The Phase
1.4 wire-up site (per-rollout-step in GpuExperienceCollector) needs to
feed per-env GPU-resident trade signals (trade_close_per_sample,
step_ret_per_sample, hold_at_exit_per_sample, packed factored
actions_out) into the kernel — forbidden via host sync
(`feedback_cpu_is_read_only`) and via memcpy_dtoh/htod
(`feedback_no_htod_htoh_only_mapped_pinned`). Path C refactors the EMA
kernel to take a device struct pointer (`SP20EmaInputs* ema_inputs`) +
adds an aggregation kernel that writes the struct on the GPU. Phase 1.2
had zero production consumers yet, so the sig change + Phase 1.4 wire-up
land atomically.
## What this commit contains
1. **EMA kernel signature refactor** (Phase 1.2 → Path C):
`sp20_emas_compute_kernel.cu` swaps 9 scalar args for a single
`const SP20EmaInputs* __restrict__ ema_inputs` device pointer.
Math semantics bit-identical. Rust `EmaInputs` mirrored as
`#[repr(C)]` byte-for-byte; new `pack_inputs_into_f32_view` helper
for tests + collectors that fill the struct via a mapped-pinned
f32-aliased buffer.
2. **New aggregation kernel** (`sp20_aggregate_inputs_kernel.cu`):
Per-env arrays (sliced to current rollout step) + sp20_stats outputs
(p50/std) + aux_dir_acc_reduce_kernel output [0] → SP20EmaInputs
struct. Block tree-reduce (4 stripes × bdim × 4 bytes shmem); no
atomicAdd. Aggregation rules per spec §4.5:
- is_close = OR over envs
- is_win = (≥ 0.5 of closed envs were wins)
- trade_duration = round(mean(hold_at_exit) over closed envs)
- action_is_hold = strict majority (count*2 > n_envs) HOLD
- alpha = 0.0 (Phase 2 forward ref — reward kernel)
- per_bar_hold_reward = 0.0 (Phase 3.2 forward ref — Hold-cost dual)
- aux_logits_p50/std/dir_acc forwarded from upstream
3. **Production wire-up in GpuExperienceCollector**:
4 kernel handles + 4 mapped-pinned buffers (struct fields, alloc,
init); per-rollout-step launch sequence after env_step (step 5c):
`Stats → Aggregate → EMAs → Controllers`. All on the same stream;
stream-implicit producer→consumer ordering. Gated on
`isv_signals_dev_ptr != 0 && trainer_params_ptr != 0` (matches the
existing SP14-β EGF chain pattern).
4. **Fold-boundary reset**:
3 new RegistryEntry records (sp20_ema_inputs_buf,
sp20_emas_internal_buf, sp20_emas_obs_count_buf) + 13 new dispatch
arms in training_loop.rs (10 SP20 ISV slot resets + 3 buffer resets).
Closes the pre-existing `every_fold_and_soft_reset_entry_has_dispatch_arm`
regression (was failing on fresh check after the SP20 ISV slot
registrations landed in commit
|
||
|
jgrusewski
|
e96f7fcdd1 |
fix(sp20): replay buffer records magnitude INTENT (Bellman consistency)
Parallel to the 2026-05-08 Class C direction fix but on the **magnitude axis**.
Class C correctly chose REALIZED for direction because env-gated trades (capital
floor / trail / broker cap → Flat) mean no trade happens — recording as Flat is
honest. For magnitude, Kelly cap clamping does NOT gate the trade — the trade
still happens, just at a smaller size. `actual_mag_core` is the bucketed
magnitude derived in trade_physics.cuh:1099-1117 from |position|/max_position;
when the policy intends Full and Kelly clamps to 0.35× max (bucket → Quarter),
recording Quarter loses the policy's intent. Q(s, mag=Full) never receives the
reward gradient from Full-intent trades that actually executed.
One-line swap inside the Class C encoding block at experience_kernels.cu:2689-
2697: `actual_mag_core * b2_size * b3_size` → `mag_idx * b2_size * b3_size`.
`mag_idx` is the local intent magnitude already decoded at line ~2460 from the
original action_idx BEFORE any Kelly/CVaR/Q-gap clamping. Direction axis keeps
Class C semantics (`actual_dir_core` → gated trades record as Flat). Order/
urgency remain intent passthrough.
`actual_mag_core` remains consumed downstream by Task 2.X per-magnitude trade-
lifecycle instrumentation at the seg_mag_bin site below — direction-branch Q
learns from realized; magnitude-branch Q learns from intent. The two axes have
different env-enforcement semantics so they take different actions at the
replay-write site.
Verification: new CPU oracle test crates/ml/tests/sp20_magnitude_intent_record_
test.rs pins the encoding contract with 4 invariants (Full→Quarter scenario,
all-pairs sweep, order/urgency passthrough, direction Class C unchanged). Per
pearl_tests_must_prove_not_lock_observations the test asserts invariants
("recorded mag == intent mag, regardless of realized") not observed values, so
it cannot become a bug-lock if the kernel's compute path changes — only if the
encoding contract itself is broken. All 4 tests pass; cargo check clean.
Audit-doc entry added to docs/dqn-wire-up-audit.md as the SP20 magnitude intent
fix, parallel to the Class C direction fix above with the asymmetry between
the two axes spelled out.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
4e21c38b85 |
fixup(sp20): Phase 1.1 K=3 → K=2 retarget (production aux head is K=2)
The Phase 1.1 sp20_stats_compute kernel (
|
||
|
jgrusewski
|
5ded1cb4b9 |
feat(sp20): Phase 1.3 sp20_controllers_compute kernel
Component 5 / Kernel 2 of the SP20 fused-producer chain — the
deterministic derivation step. Reads the EMAs Phase 1.2's
sp20_emas_compute wrote (4 ISV + 4 internal scratch slots) and
derives 6 controller outputs into ISV slots:
- LOSS_CAP (510) = -1 - clamp((wr-0.50)/0.05, 0, 1)
- HOLD_COST_SCALE (513) two-sided ramp around TARGET_HOLD_PCT ±0.05
- TARGET_HOLD_PCT (514) = clamp(0.8 - aux_p50_ema*1.5, 0.1, 0.8)
- N_STEP (517) = clamp(round(trade_dur_ema), 1, 30) (f32)
- AUX_CONF_THRESHOLD (518) = clamp(aux_dir_acc_ema-0.50, 0.01, 0.20)
- AUX_GATE_TEMP (519) = max(aux_conf_std_ema, 0.01)
TARGET_HOLD_PCT computed before HOLD_COST_SCALE because the
HOLD_COST_SCALE controller reads tgt; implemented as a local f32
written to ISV then re-used for the controller (no redundant ISV
read-back). HOLD_COST_SCALE is the only output that READS its own
previous ISV value (in-place update); deadband path writes the
unchanged previous value verbatim per feedback_no_stubs.
Single-block, single-thread kernel — captureable in the per-step
CUDA Graph per pearl_no_host_branches_in_captured_graph. ISV +
internal pointers are mapped-pinned device pointers (existing
buffers from Phase 1.2); kernel emits __threadfence_system() for
PCIe-visible coherence.
Pearls + invariants honoured:
- pearl_controller_anchors_isv_driven: every output's anchor /
target / cap derives from EMAs; only spec-frozen ramp / clamp
parameters are compile-time constants.
- feedback_isv_for_adaptive_bounds: these 6 outputs ARE the
adaptive bounds.
- feedback_no_atomicadd, feedback_no_cpu_compute_strict,
feedback_no_htod_htoh_only_mapped_pinned, feedback_no_stubs.
- feedback_no_partial_refactor: kernel + launcher + tests + build
entry land atomically; production wire-up lands in Phase 1.4.
- pearl_tests_must_prove_not_lock_observations: 7 GPU oracle
tests assert clamp boundaries, formula correctness, and
bidirectional ramp behavior — NOT specific observed values.
Tests (all #[ignore = "requires GPU"], pass on RTX 3050 Ti sm_86):
1. loss_cap_ramp_boundaries — 4 wr_ema sweeps incl. clamps.
2. n_step_bounds — round + clamp [1, 30].
3. aux_conf_threshold_bounds — clamp [0.01, 0.20].
4. aux_gate_temp_floor — max(std, 0.01).
5. target_hold_pct_inverse_relation — 0.8 − p50*1.5 with clamps.
6. hold_cost_scale_two_sided_ramp — up/down/deadband + clamps.
7. all_six_outputs_written_in_one_launch — guards missed writes.
Plus 3 launcher unit tests (slot range, slot uniqueness, ISV input
slot range).
Verification:
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --features cuda
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml \
--test sp20_controllers_compute_test --features cuda \
-- --ignored --nocapture
→ 7/7 GPU oracle tests pass; 3/3 launcher unit tests pass.
Phase 1.4 will land the production caller atomically alongside
Kernel 1 + Kernel 3 launches on the same stream in order
Stats → EMAs → Controllers per the SP20 design.
|
||
|
jgrusewski
|
71aade18cf |
feat(sp20): Phase 1.2 sp20_emas_compute kernel
Component 5 / Kernel 1 of the SP20 design — central state-tracker
that updates 8 Wiener-α EMAs per training step:
- 4 ISV slots: ALPHA_EMA (511), WR_EMA (512),
HOLD_PCT_EMA (515), HOLD_REWARD_EMA (516)
- 4 internal: trade_duration_ema, aux_conf_p50_ema,
aux_conf_std_ema, aux_dir_acc_ema (private
scratch consumed by Phase 1.3 controllers)
Per-EMA i32 observation counters (mapped-pinned obs_count[8])
handle the pearl_first_observation_bootstrap sentinel transition
correctly even when 0.0 is a legitimate observation (e.g.,
first-loss WR=0). count==0 ⇒ replace, count>0 ⇒ Wiener-blend at
α = 0.4 (WIENER_ALPHA_FLOOR per
pearl_wiener_alpha_floor_for_nonstationary).
Phase 1.2 lands kernel + launcher + 5 tests (4 GPU oracle, 1
floor lock) + build entry + audit doc atomically per
feedback_no_partial_refactor. Production wire-up (Phase 1.4)
deferred — kernel is dead code until then.
Verified on RTX 3050 Ti (sm_86):
- 4 GPU oracle tests pass: first_observation_replaces_sentinel,
wiener_alpha_converges_to_long_run_mean,
hold_reward_ema_gated_on_hold_bars,
per_step_emas_fire_unconditionally
- 4 launcher unit tests pass (constants + struct sanity)
- 1 floor-lock test pass (WIENER_ALPHA_FLOOR == 0.4)
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
de922c6a4a |
feat(sp20): Phase 1.1 sp20_stats_compute kernel
Component 5 / Kernel 3 of the SP20 fused-producer chain. Single-block BLOCK=256 kernel reads `aux_logits [B, 3]` (the SP14-C aux head's 3-class direction logits) and emits `[aux_conf_p50, aux_conf_std]` into a `MappedF32Buffer<2>`, where the per-row signal is `aux_conf[i] = max_c softmax(logits[i, *])[c] - 1/3`. p50 uses the inlined `sp4_histogram_p99` pattern (per-warp tile binning + cumulative-from-bottom, no atomicAdd per `feedback_no_atomicadd`); std uses two block tree-reductions sharing one shmem tile sequentially. One fused kernel streams `aux_logits` once for both stats per `pearl_fused_per_group_statistics_oracle`. Phase 1.4 wires the production launch site atomically with the rest of the SP20 reward chain per `feedback_no_partial_refactor`. This commit lands kernel + Rust launcher + GPU oracle tests + build entry + audit-doc entry together so the kernel is independently verifiable on RTX 3050 Ti (sm_86) and L40S (sm_89) before the EMA + controller producers (Phase 1.2 + 1.3) reference its outputs. Tests verify: - uniform logits → aux_conf = 0 → [p50, std] = [0, 0] - varied confidence (logit ramp 0 → 3) → matches CPU oracle - heterogeneous half-hot half-uniform → matches CPU oracle - empty batch → degenerate-guard writes [0, 0] All 4 GPU oracle tests + 4 launcher unit tests pass on RTX 3050 Ti. Test data uses per-row variance to avoid the `pearl_sp4_histogram_warp_tile_undercount` lockstep-uniform trap (concentrated values within one bin_width race the per-warp non-atomic increments). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
4249ebc961 |
feat(sp20): register 10 ISV slots in StateResetRegistry
Sentinel = 0.0 per pearl_first_observation_bootstrap. First observation of each EMA replaces the sentinel directly, no blending. Slots [510..520): - loss_cap (510): adaptive loss cap for reward clamp - alpha_ema (511): Wiener-α EMA for loss_cap producer - wr_ema (512): win-rate EMA driving loss_cap adaptive ramp - hold_cost_scale (513): hold penalty cost multiplier - target_hold_pct (514): hold-engagement target - hold_pct_ema (515): hold-engagement EMA - hold_reward_ema (516): hold-action reward EMA - n_step (517): multi-step TD horizon adapter - aux_conf_threshold (518): auxiliary task confidence threshold - aux_gate_temp (519): auxiliary task gating temperature |
||
|
jgrusewski
|
f5eed1fa79 | spec(sp20): reserve 10 ISV slots [510..520) for WR-first reward optimization | ||
|
jgrusewski
|
99332003a4 |
plan(sp19+20): WR-first reward implementation plan
Implements the spec at docs/superpowers/specs/2026-05-09-sp19-20-wr-first-design.md
(commit
|
||
|
jgrusewski
|
9d9ca3e6e7 |
spec(sp19+20): apply Q1(b) — Hold-reward EMA for Q-scale comparability
Q1 design decision (b): add explicit hold_reward_ema to center the per-bar Hold reward, so Q(Hold) and Q(trade) targets are scale-comparable. The marginal Q(trade) > Q(Hold) preference now comes from data variance in each state (high-aux states pull Q(Hold) more negative), not from structural scale asymmetry that depends on cost_scale magnitude. Q2 decision: keep 4-quadrant fixed (no ramping partials). Already in spec. Changes: - §4.2 Hold opp-cost: dual emission documented — R_per_bar_centered (= R_per_bar - hold_reward_ema) for Q-target/replay tuple, R_per_bar uncentered for hold_baseline_buffer (Component 1 baseline) - §4.5 Kernel 1: hold_reward_ema added (per-step on Hold-state bars only) - §5 data flow: per-bar reward path shows the centered/uncentered split - ISV slots: 9 → 10 (HOLD_REWARD_EMA_INDEX added) - §8 footprint: Component 2 LoC 70 → 90 (+20 for dual emission) - Total LoC estimate: 1620 |
||
|
jgrusewski
|
defbd0abe1 |
spec(sp19+20): patch 7 review issues
P1 (must-fix bugs/gaps): 1. ASYM_RATIO_INDEX → LOSS_CAP_INDEX with explicit formula in §4.1 and §4.5 (was double source-of-truth — formula in §4.1, ISV slot orphaned) 2. Replay buffer schema change (per-bar aux_conf) added to §8 implementation footprint (~50 LoC additional, was invisible in original spec) 3. hold_baseline_buffer size specified = LOOKAHEAD_HORIZON_MAX = 30 bars (§4.2) 4. "4-tier gate" typo → "5-tier gate" in §7 P2 (clarifications): 5. alpha_ema centering documented as load-bearing for cold-start learning (§4.1) — not just for Q-target stability. EV is slightly negative at WR=46% with uninformed SP19 label; advantage-style centering rescues cold-start. 6. Q-scale asymmetry between Hold (uncentered) and trade (alpha_ema centered) documented as INTENTIONAL design choice (§4.2) — produces marginal Q(trade) > Q(Hold) preference that counteracts the Q(Hold) attractor. Behavioral test sp20_pure_noise validates the no-signal Hold default still works. P3 (tightening): 7. Behavioral test thresholds tightened (§4.6): - sp20_pure_trend: WR > 70% → > 90%, PF > 2.0 → > 3.0 - sp20_pure_noise: Hold% > 80% → > 95%, trades < 50 → < 20 |
||
|
jgrusewski
|
730337375f |
spec(sp19+20): WR-first reward + multi-horizon label utilization
Combines SP19 (multi-horizon labels, already landed) + SP20 (WR-first reward) into one spec per pearl_no_deferrals_for_complementary_fixes. Goal: WR ≥ 55%, textbook PF ≥ 2.0, walk-forward stable, per-regime stable. Six components, atomic ship per feedback_no_partial_refactor: 1. Reward kernel (event-driven, 4-quadrant, asymmetric clamp ramped from wr_ema, multi-horizon directional ground-truth check) 2. Hold opportunity-cost (per-bar, dual emission for real reward + Hold baseline buffer) 3. n-step credit distributor (uniform over trade duration, fixes SP18 B-leg self-bootstrap bug at gpu_experience_collector.rs:4154) 4. Aux→Q confidence gate (sigmoid threshold, mean_a Q baseline avoids Hold-everywhere punishment) 5. 3 fused producer kernels (sp20_emas, sp20_controllers, sp20_stats) driving 9 ISV slots in [510..520) 6. 7 behavioral tests on RTX 3050 Ti, gating L40S deployment ~1,550 LoC total. Spec includes data flow, error handling philosophy, 4-tier test gate, success criteria with explicit failure modes. Cross-references: - pearl_event_driven_reward_density_alignment (design principle) - pearl_audit_unboundedness_for_implicit_asymmetry (asymmetric clamp) - pearl_separate_aux_trunk_when_shared_starves (aux gate leverages SP14-C) - project_metric_pipeline_inflation_audit (WR honest, Sharpe honest, goal grounded) - project_goal_wr_55_pf_2 (the goal this spec implements) |
||
|
jgrusewski
|
d39005c6f4 |
feat(sp19 commit b): producer-side multi-horizon reward blend at fxcache write time
Path (B) Commit B — load-bearing semantic change for the SP19 multi-
horizon reward augmentation. At fxcache-write time, blend 1-bar / 5-bar
/ 30-bar log-returns into `tgt[1]` (`preproc_next`) using equal-thirds
weights with `1/sqrt(N)` vol-scale correction:
blended = (1/3) * r_1
+ (1/3) * r_5 / sqrt(5)
+ (1/3) * r_30 / sqrt(30)
The vol-scale correction applies INSIDE the blend (before weighting) so
each horizon's log-return is at unit-volatility-equivalent under
random-walk assumption — drops naturally into the existing `tgt[1]`
preprocessing pipeline (consumed in `experience_kernels.cu:1556` and
`cuda_pipeline/mod.rs:508`) without double-scaling.
Producer trim: valid range shrinks by `LOOKAHEAD_HORIZON_MAX = 30`
bars since `r_30` needs `close[i + WARMUP + 30]`. Walk-forward fold
construction is dataset-relative, so trimming at producer time shifts
every fold's tail by 30 bars — one-time cost per fxcache regen.
Both producer call sites (`precompute_features.rs` for the
feature-precompute binary, `data_loading.rs` for the DBN-fallback
in-process path) change atomically per `feedback_no_partial_refactor`.
Kernel consumers are unchanged — only `tgt[1]`'s value composition
differs from the consumer's perspective.
ISV slots [507..510) (allocated in Commit A) are NOT consumed at
producer time — `precompute_features` runs BEFORE training so the
ISV bus isn't initialised when the blend happens. Producer hardcodes
1/3 each via `SP19_HORIZON_BLEND_WEIGHTS`. Future Path (A) refactor
would bump TARGET_DIM to carry per-horizon log-returns and read
ISV at training-step time; for the empirical hypothesis test
("does multi-horizon blend lift WR?") equal-thirds is sufficient.
fxcache schema invalidation: `FXCACHE_VERSION` 8 → 9. Existing
`.fxcache` files (1-bar-only `preproc_next`) fail validation at load
and trigger Argo's ensure-fxcache regen step. First L40S run after
this commit takes ~10-15 min longer for the regen — one-time cost,
expected.
Touches:
- crates/ml/examples/precompute_features.rs: SP19_HORIZON_BLEND_WEIGHTS
constant + LOOKAHEAD_HORIZON_MAX trim + blend in tgt[] writer +
early-bail-out check on minimum dataset size.
- crates/ml/src/trainers/dqn/data_loading.rs: same blend +
trim in DBN-fallback path; `last sample targets itself` block
removed (the trimmed range guarantees feature-vector and target
lengths match without a sentinel last row).
- crates/ml/src/fxcache.rs: FXCACHE_VERSION 8 → 9 + v9 docstring entry.
- crates/ml/tests/multi_horizon_reward_blend_test.rs (NEW): CPU-only
oracle behavioural test — 4 cases including known returns, trim
contract, zero-close short-circuit, constant-price blend.
- docs/dqn-wire-up-audit.md: Concerns subsection appended to the
SP19 Commit B entry already drafted in Commit A (pre-existing
test_fxcache_empty + test_dqn_checkpoint_round_trip flakiness
documented for Invariant 7).
Verification:
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check --workspace clean
cargo test -p ml --test multi_horizon_reward_blend_test 4/4 pass
cargo test -p ml --test fxcache_roundtrip_test test_fxcache_f32_roundtrip passes (TARGET_DIM unchanged)
cargo test -p ml --lib ... 13 baseline failures (pre-existing); zero new regressions
bash scripts/audit_sp18_consumers.sh --check exit 0
Pre-existing test_fxcache_empty failure: hand-crafts a 64-byte header
but the actual header is 72 bytes; reader bails early on
"failed to fill whole buffer" instead of "bar_count=0". Test setup
bug, NOT a regression. Pre-Commit B failure count: 13. Post-Commit B
failure count: 13 (the test_dqn_checkpoint_round_trip test is flaky
and toggles independently of this change — verified by stashing and
re-running 3× pre-Commit B).
Atomic-refactor invariant satisfied: Commit A (slot reservations) +
Commit B (producer-side blend) land on the same branch with no L40S
dispatch between. Per task instruction the branch stays unpushed
pending user review.
DBN-fallback path: applied identically in `data_loading.rs:497-528`.
Both producers use the same `SP19_HORIZON_BLEND_WEIGHTS` constant and
the same `LOOKAHEAD_HORIZON_MAX = 30` trim.
Vol-scale correction site: applied inside the blend (before weighting)
in BOTH producers. The existing `tgt[1]` preprocessing pipeline does
NOT double-scale — `experience_kernels.cu:1556` and
`cuda_pipeline/mod.rs:508` consume `tgt[1]` as a unit-scale log-return
exactly as before the blend was added; the blended value drops in
without further scaling.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
37964ae2a4 |
feat(sp19 commit a): ISV slot reservations [507..510) for multi-horizon reward blend
Path (B) Commit A — additive infrastructure for the SP19 producer-side
multi-horizon reward augmentation. Three ISV slots reserve indices for
a future Path (A) refactor that would let a controller drive per-batch
horizon weights; producer (Commit B) hardcodes equal-thirds 1/3 each at
fxcache-write time so the slot reservations are dormant in this branch.
Pure additive — changes NO runtime behaviour. The dispatch arms write
the equal-thirds sentinel at every fold boundary, but no kernel
consumes the slots. This is forward-compatible reservation per
feedback_isv_for_adaptive_bounds, NOT a half-fix; the producer wiring
is complete (next commit), the slot reservations are documented
forward-compatibility for the TARGET_DIM-bumping refactor.
Slot allocation:
- 507 REWARD_HORIZON_WEIGHT_1BAR_INDEX sentinel = 1/3
- 508 REWARD_HORIZON_WEIGHT_5BAR_INDEX sentinel = 1/3
- 509 REWARD_HORIZON_WEIGHT_30BAR_INDEX sentinel = 1/3
Touches:
- crates/ml/src/cuda_pipeline/sp14_isv_slots.rs: 3 slot constants +
sentinel + range markers + sp19_reward_horizon_slot_layout_locked +
all_sp19_slots_fit_within_isv_total_dim lock tests.
- crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs: ISV_TOTAL_DIM 507 →
510 + layout_fingerprint_seed extension (3 SLOT_* lines + a
SP19_PRODUCER_HARDCODED_HORIZON_BLEND token registering producer-time
consumption).
- crates/ml/src/cuda_pipeline/state_layout.cuh: 3 #define mirrors of
the Rust slot indices + SENTINEL_REWARD_HORIZON_WEIGHT_DEFAULT macro.
- crates/ml/src/trainers/dqn/state_reset_registry.rs: 3 RegistryEntry {
FoldReset } with the "SP19 Path (B) reservation" marker + lock-test
expansion 24 → 27 entries.
- crates/ml/src/trainers/dqn/trainer/training_loop.rs: 3 dispatch arms
in reset_named_state writing SENTINEL_REWARD_HORIZON_WEIGHT_DEFAULT.
- docs/dqn-wire-up-audit.md: SP19 Commit A entry documenting the
reservation rationale + verification steps.
Verification:
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check --workspace clean
cargo test -p ml --lib sp19_reward_horizon_slot_layout_locked passes
cargo test -p ml --lib all_sp19_slots_fit_within_isv_total_dim passes
cargo test -p ml --lib sp18_fold_reset_entries_present passes (27)
cargo test -p ml --lib every_fold_and_soft_reset_entry_has_dispatch_arm passes
bash scripts/audit_sp18_consumers.sh --check exit 0
Atomic-refactor invariant (HARD — feedback_no_partial_refactor): NO
L40S DISPATCH between Commit A and Commit B. The producer-side blend
+ fxcache version bump + behavioural test land in the next commit on
this branch.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
ce841eb56e |
fix(audit): lookahead housekeeping — purge gap + per-fold NormStats
Closes recs 2 + 3 in lookahead-bias-audit-2026-04-28.md.
Fix 2A (rec 2): add walk_forward::PURGE_BARS = 5 constant and insert
val_start = train_end + PURGE_BARS in all 4 fold-construction sites:
generate_walk_forward_indices, _from_timestamps, _windows (date-based,
drops first PURGE_BARS bars of val), and gpu_walk_forward::generate_folds
(both stratified variants preserve the gap on boundary shift). 5-bar
width matches max per-bar feature lookback (autocorr lags 1/5/10);
compile-time per feedback_isv_for_adaptive_bounds since the
feature-pipeline lookback is itself compile-time.
Fix 2B (rec 3): per-fold NormStats fit from train-only data.
- walk_forward.rs: add from_features_slice + denormalize/denormalize_batch
helpers (inverse of normalize_batch for clamp-bounded round-trip).
- train_baseline_rl.rs fold loop: load fxcache sidecar norm_stats.json,
denormalise back to RAW, refit per-fold via from_features_slice,
renormalise full dataset, re-upload via init_from_fxcache, save
per-fold norm_stats_fold{N}.json. DBN-fallback path keeps legacy
behaviour (no sidecar to denormalise from) with a warn! log.
Ensemble trainer block is documented follow-up (separate code path).
Behavioral tests:
- test_norm_stats_per_fold_fit_train_only: synthetic train-mean=1.0 /
val-mean=9.0 dataset; from_features_slice recovers train-only mean to
ε=1e-5 while from_features returns global 5.0
- test_norm_stats_denormalize_roundtrip: non-clamped features round-trip
bit-identically
- test_walk_forward_purge_gap_indices_from_timestamps: every fold has
val_start - train_end == PURGE_BARS
- test_fold_generation_basic (gpu_walk_forward): updated to expect the
purge gap
Validation: cargo check --workspace clean (12.30s); 20/20 walk_forward
+ gpu_walk_forward tests pass; audit_sp18_consumers.sh --check exit 0.
Pre-existing test failures on local RTX 3050 Ti are unrelated SIGSEGVs
(VRAM exhaustion in chunked Thompson tests, pre-existing on baseline).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
e140392f86 |
feat(audit): per-regime val WR instrumentation (T/R/V buckets)
Adds 6 output slots to compute_backtest_metrics_kernel — per-bucket
(win_rate, trade_count) for the {Trending, Ranging, Volatile} regime
split — so the val backtest surfaces whether the long-running ~46%
aggregate WR hides regime-conditional edge. Trades are bucketed at
trade-OPEN by feature[40] (ADX-norm) per the structural thresholds
(T:ADX>0.4, R:ADX<0.2, V:otherwise), mirroring
gpu_walk_forward.rs::classify_regime_from_features. Block tree-reduce
only per feedback_no_atomicadd. Observability-only emission via new
HEALTH_DIAG[N]: val_regime [wr_T=... n_T=... wr_R=... n_R=... wr_V=...
n_V=...] line; thresholds remain kernel constants per
feedback_isv_for_adaptive_bounds (no controller consumer yet).
Implementation atomic (kernel + launcher + WindowMetrics + HEALTH_DIAG
emit + 3 GPU oracle tests + audit doc):
- backtest_metrics_kernel.cu: per-thread per-regime trade counters,
2-slot boundary buffer extension carrying open-bar regime through
block stitch, output stride 13 → 19, shmem 5 → 11 reduction tiles
- gpu_backtest_evaluator.rs: WindowMetrics +6 fields, metrics_buf
size 13 → 19, launcher passes features_buf + feature_dim, consume
populates per-regime fields
- metrics.rs: val_regime HEALTH_DIAG line in consume_validation_loss
- regime_wr_oracle_tests.rs (NEW): 1 CPU sanity + 3 GPU oracle tests
(bit-exact match ε=1e-5 vs CPU oracle on stratified 30/40/30 batch)
Validation: cargo check --workspace clean; 17/17 gpu_backtest_evaluator
unit tests pass; 1+3/4 regime WR tests pass on local RTX 3050 Ti
(1.89s); audit_sp18_consumers.sh --check exit 0.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
5ee5a1b655 |
chore(sp18): refresh audit fingerprint after Phase 4 docstring update
Off-by-one count drift in the td_lambda_kernel launch-sites bucket (19→20) from a single additional reference in sp18_td_lambda_q_next_oracle_tests.rs's docstring (the docstring now mentions td_lambda_kernel one extra time per the T_SKEL.3 atomic-refactor guard test). Pure documentation update; no source-code change. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
834eaec4bd |
feat(sp18 v2 P4.T1+T2): B-leg next_states hoist + compute_q_next_target_bootstrap skeleton (additive dead code)
Phase 4 lands the additive infrastructure for the B-leg target-Q DDQN
bootstrap that Phase 5 will swap into td_lambda_kernel's q_next argument
at gpu_experience_collector.rs:~4313 (post-Task-4.1 hoist).
Tasks landed:
- **Task 4.1**: build_next_states_f32 invocation hoisted to BEFORE
the TD(λ) launch block. Pure ordering change — verified via the
SP18 audit grep flagging 0 post-TD(λ) consumers of next_states.
- **Task 4.2**: compute_q_next_target_bootstrap skeleton method on
GpuExperienceCollector with full plan-aligned signature
(next_states, target_params_ptr, online_params_ptr, batch_size →
Result<CudaSlice<f32>, MLError>). Body is a hard-error early-return
per feedback_no_stubs — no production caller in Phase 4 (the
caller is wired by Phase 5 Task 5.1, replacing the self-bootstrap
clone). Returning Err satisfies Invariant 9 (no deferred-work
markers); any accidental pre-Phase-5 call hard-errors with a
clear pointer to the open sub-tasks.
Tasks 4.3 (online forward + DDQN per-branch argmax), 4.4 (target
forward + compute_expected_q gather), 4.5 (replace error early-return
with full implementation + GPU oracle test gates) are deferred to a
follow-up subagent dispatch per their plan-defined per-task TDD cycle
scope. Each requires substantial new infrastructure on the collector
(target-net trunk forward chain duplicating gpu_dqn_trainer.rs Pass 2's
VSN+BN+OFI+trunk+branch-head sequence).
Tests:
- crates/ml/tests/sp18_td_lambda_q_next_oracle_tests.rs (NEW):
3 introspection tests (no GPU required):
* compute_q_next_target_bootstrap_method_exists
* next_states_built_before_td_lambda (Task 4.1 ordering invariant)
* td_lambda_still_consumes_self_bootstrap_q_next_in_phase4
(Phase 4→Phase 5 atomic-refactor guard)
Atomic-refactor invariant (HARD — feedback_no_partial_refactor):
NO L40S DISPATCH between this Phase 4 close-out commit and the
Phase 5 Task 5.1 commit that replaces the q_next clone with
self.compute_q_next_target_bootstrap(&next_states, ...).
Verification:
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check --workspace → clean
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 \
cargo test -p ml --test sp18_td_lambda_q_next_oracle_tests
→ 3/3 pass
bash scripts/audit_sp18_consumers.sh --check → exit 0
Files:
crates/ml/src/cuda_pipeline/gpu_experience_collector.rs (Task 4.1+4.2)
crates/ml/tests/sp18_td_lambda_q_next_oracle_tests.rs (NEW; 3 tests)
docs/sp18-wireup-audit.md (Phase 4 close-out section + fingerprint)
docs/dqn-wire-up-audit.md (Invariant 7 entry)
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
5ea5aa9b8e |
feat(sp18 v2 P3.T1-T5): adaptive HOLD_REWARD_POS/NEG_CAP producer kernel
Lifts the Phase 2 caps from sentinel-driven cold-start (5.0/-10.0
fixed) to p99(|step_ret|) over Long/Short trade closes × 1.5 safety
factor with Wiener-optimal alpha blend. The Phase 2 consumer
(compute_sp18_hold_opportunity_cost) sees producer-driven slots
[483]/[484] from epoch 1 onward; sentinel branch in the consumer
remains as the cold-start fallback path (zero-trade-close epoch ⇒
producer early-returns ⇒ slots stay at sentinel ⇒ consumer falls
through to the +5/-10 macro defaults — bit-identical to pre-Phase-3).
Mirrors SP14 P0-A reward_cap_update_kernel structural template with
three differences: (1) filter (is_close && step_ret != 0) — both
winners AND losers (the consumer needs the magnitude scale of *all*
Long/Short closes); (2) Welford-derived Wiener-α (slots [487..493))
replaces fixed α=0.01, with floor at WELFORD_ALPHA_MIN=0.4 per
pearl_wiener_alpha_floor_for_nonstationary (the policy-realised
distribution is intrinsically non-stationary as the policy adapts);
(3) bounds [0.5, 50.0] (vs. position-side [1.0, 50.0]).
Atomic single-commit per feedback_no_partial_refactor:
- crates/ml/src/cuda_pipeline/hold_reward_cap_update_kernel.cu (NEW)
- crates/ml/build.rs cubin manifest entry
- HoldRewardCapUpdateOps in gpu_aux_trunk.rs (new struct + impl)
- HOLD_REWARD_CAP_UPDATE_CUBIN static + struct field +
launch_hold_reward_cap_update method + constructor instantiation +
field-init in gpu_dqn_trainer.rs (5 sites)
- Per-epoch boundary launch in training_loop.rs right AFTER
launch_reward_cap_update (shared step_ret/trade_close source buffers,
independent ISV slot pairs)
- HEALTH_DIAG[N]: hold_reward_cap [pos={:.4} neg={:.4} fire_rate={:.4}]
- 3 GPU oracle tests (T5 producer-drives-slots, Pearl-A REPLACE,
no-closes preserves-isv) — all pass on local RTX 3050 Ti
- Phase 3 close-out sections in docs/sp18-wireup-audit.md and
docs/dqn-wire-up-audit.md
Pearls applied: feedback_no_atomicadd, pearl_first_observation_bootstrap,
pearl_wiener_optimal_adaptive_alpha, pearl_wiener_alpha_floor_for_nonstationary,
pearl_no_host_branches_in_captured_graph, pearl_symmetric_clamp_audit,
pearl_audit_unboundedness_for_implicit_asymmetry (NEG = -2 × POS at
producer time, single source of truth), feedback_isv_for_adaptive_bounds,
pearl_fused_per_group_statistics_oracle.
Validation: cargo check --workspace clean; 3 GPU oracle tests pass on
local RTX 3050 Ti; scripts/audit_sp18_consumers.sh --check exits 0
(no fingerprint drift in tracked sections).
Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md
Phase 4-5 (B-leg target-net forward + q_next replacement) follows.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
1f4cc0f207 |
chore(sp18): refresh audit fingerprint for archaeology cleanup self-references
The cleanup commit's own additions to docs/dqn-wire-up-audit.md (the 2026-05-09 archaeology-cleanup-pass entry) cite HOLD_COST_SCALE_INDEX, hold_cost_scale_update_kernel, and HOLD_COST_INDEX by name in the prose describing what was stripped. The audit script's per-section grep counts those legitimately, producing a +1/+1/+4 hit-count drift on docs/dqn-wire-up-audit.md across the slot 380, slot 461, and hold_cost_scale_update_kernel sections respectively. Regenerate the fingerprint snapshot in docs/sp18-wireup-audit.md so `scripts/audit_sp18_consumers.sh --check` passes against the new post-cleanup baseline. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |