main
2611 Commits
| Author | SHA1 | Message | Date | |
|---|---|---|---|---|
jgrusewski
|
a5de50503f |
feat(alpha): Poisson regression fill model scaffold (coeffs fitted in Task 5)
Phase E Task 4. Medium-tier fill simulator per the design memo.
For each (level ∈ {L1,L2,L3}, side ∈ {Bid,Ask}):
λ(features) = exp(β · [1, spread_bps, L1_imb, OFI_5, log(τ+1)])
Per-snapshot Bernoulli fill probability for a posted limit order:
p = 1 − exp(−λ)
Market orders fill immediately at the opposite-side L1 quote (no slippage
modeled at medium tier). Closing actions (Side::None) return λ=0 — they
are handled separately in the env.
Scaffolding only. Task 5 fits the 30 coefficients (6 distributions × 5
features) from the 5.2M-trade historical tape. The skeleton constructor
uses β=0 → λ=1 → p≈0.632, a stable sanity default for early smokes.
Numeric guard: `linear.exp().min(50.0)` caps λ to prevent f32 overflow
under outlier features before fitting lands. Fitted models should stay
well below this cap in practice.
4 unit tests:
- skeleton_has_uniform_fill_prob (β=0 → p≈0.632 within 0.01)
- fill_prob_bounded_under_outlier_features
- lambda_cap_prevents_f32_overflow (β=100 outlier path)
- side_none_returns_zero_rate
`cargo test -p ml --lib env::fill_model`: 4 passed.
|
||
|
jgrusewski
|
94f8f65571 |
feat(alpha): 9-action discrete execution action space + legality gating
Phase E Task 3. Introduces crates/ml/src/env/ for the execution-policy
environment. This commit lands the action space only; fill_model (Task 4)
and execution_env (Task 6) append to mod.rs in their respective commits
per feedback_wire_everything_up (no orphan declarations).
- 9-action discrete space: Wait + {Buy,Sell,Flat} × {Market,L1,L2 / L1}
- Soft action masking: illegal actions degrade to Wait in env::step
(not by masking Q-output) per design memo — preserves clean C51
categorical targets and Munchausen term (Task 10)
- `is_legal(position)` with position ∈ {-1, 0, +1} (sign only;
magnitude is decoupled into the Kelly layer in Task 11)
- repr(u8) discriminants round-trip with from_u8 for replay-buffer
storage by the existing GPU DQN trainer
4 unit tests:
- n_actions_matches_enum_cardinality
- legality_gating_by_position (covers flat/long/short × all 9 actions)
- decode_closing_flag_is_only_flat
- repr_u8_round_trip (replay-buffer contract)
`cargo test -p ml --lib env::action_space`: 4 passed.
|
||
|
jgrusewski
|
8bf8bdf874 |
feat(alpha): state_reset_registry entries + dispatch arms for slots 539..548
Per feedback_registry_entries_need_dispatch_arms — both must land in the same commit; the test `every_fold_and_soft_reset_entry_has_dispatch_arm` walks training_loop.rs::reset_named_state source and asserts coverage. - 10 RegistryEntry rows appended after SP22 block (full producer/consumer rationale per existing dense-description pattern) - 7 dispatch arms in reset_named_state (the 3 TrainingPersist anchors — slots 544/547/548 — are not dispatched per the existing convention) - All sentinels 0.0 per pearl_first_observation_bootstrap All 10 registry tests pass. Audit doc docs/isv-slots.md updated per Invariant-7. |
||
|
jgrusewski
|
aa5908aa53 |
feat(alpha): reserve ISV slot block 539..550 for alpha trading system
12 contiguous slots reserved for durable alpha-system infrastructure: - 539-542: diagnostics (Q-spread, action entropy, return-vs-random, early-Q-movement EMAs) — read by Phase E kill criteria - 543-546: stacker-threshold controller (threshold, target, observed-rate, Kelly attenuation) — engagement-rate self-correction - 547-548: random-uniform baseline (mean, std) — anchor for kill criterion 541 - 549-550: reserved spare (absorb growth without bumping ISV_TOTAL_DIM) Named `alpha_isv_slots` rather than `phase_e_isv_slots` because these slots are intended to outlive any individual Phase E/F/G milestone — the SP4..SP22 naming was iteration-scoped, but this block is system-scoped infrastructure. ISV_TOTAL_DIM bumped to 551. All indices validated in 4 unit tests (bounds, membership, uniqueness, total-dim coverage). Audit doc docs/isv-slots.md updated per Invariant-7. |
||
|
jgrusewski
|
db874b1841 |
feat(foxhuntq): Phase 1c snapshot-resolution alpha + leakage fix + variable-dim fxcache
Three things landing atomically because they're load-bearing for each other: 1. **Trend-scanning leakage fix** — trend_scanning.rs was emitting OLS slope+t-stat over a *forward* window [t, t+L]. With the Phase 1a label = sign(price[t+60] − price[t]), the forward feature window overlaps the label window, contaminating it. Purged walk-forward only sterilizes forward-looking *labels* that cross the train/val split, not forward-looking *features* that peek inside the same horizon the label measures. The leak inflated MLP accuracy from 0.49 (legacy 74-dim baseline) to 0.75 — vanished to 0.50 after switching to a trailing window. Bounded the perfect-fit t-stat sentinel from ±1e6 → ±20 (p<1e-30 is already meaningless); eliminated the 16k corruption-cap drops. 2. **Variable-dim alpha column** — fxcache schema now carries the alpha-feature width via metadata (`alpha_feature_dim`), not a compile-time constant. Same on-disk format hosts the 134-dim bar-level stack OR the 81-dim snapshot stack. Reader + auto-detect honor the metadata-declared dim; downstream MLP auto-sizes `in_dim`. Single schema, no forks. 3. **Snapshot pipeline (Phase 1c falsification)** — `snapshot_pipeline.rs`: 81-dim per-MBP10-snapshot extractor reusing 10 snapshot-native alpha blocks + 6 new snapshot-specific features (time-since-trade, time-since-snap, event-rate, spread-bps, L1-imbalance, microprice-mid drift). `precompute_features` gets `--row-unit snapshot` flag; emits one fxcache row per LOB update (1.97M rows from MBP-10 data vs 206K for bar mode). **Smoke verdict on real data** (ES.FUT, 1.97M snapshots, 384K val): - Bar-level honest alpha: accuracy=0.5005, AUC=0.5043 (no signal) - **Snapshot-level alpha**: accuracy=0.5241, AUC=0.6849 (real signal, 384K val) - GBM corroboration: accuracy=0.5401 (non-linear partitioning sees more) - Horizon decay: alpha peaks at K=20-50 snapshots (~5-25ms), gone by K=500 - Regime-conditional: spread-Q4 quintile hits 0.752 accuracy on 76k samples Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
19bb3bc3c8 |
fix(sp22-vnext): aux_outcome CF half mirror on-policy (NOT all -1)
Smoke train-k95mj epoch 0 + epoch 1 printed
`trade_outcome_ce=0.000e0` — K=3 CE EMA stuck at Pearl A sentinel
across all training steps.
Root cause chain:
1. insert_batch is called with bs = total = 4_096_000 (cf-mult-
expanded), replay cap = 300_000 (L40S GpuProfile). Tail-clip:
off = 3_796_000, slice(off..) takes the last 300K elements.
2. Prior fix (
|
||
|
jgrusewski
|
878cc9ba72 |
feat(sp22-vnext): F-3c follow-up — K=3 CE EMA in stdout HEALTH_DIAG aux line
The F-3c commit (
|
||
|
jgrusewski
|
256a5fa5ab |
fix(sp22-vnext): K=3 aux_outcome_labels CF-half buffer underrun
Phase B4b-2 introduced `exp_aux_to_label_per_sample` sized
`[alloc_episodes × alloc_timesteps]` (no cf-mult expansion — the
B4b-1 per-step kernel only writes the on-policy half). The batch
finalisation cloned into the emitted `aux_outcome_labels` with size
`total = base_total × 2` (cf-mult expanded), so
`dtod_clone_i32(src, total, ...)` invoked `src.slice(..total)` on a
half-sized source → `CudaSlice::try_slice` returned `None` → the
internal `unwrap()` panicked at cudarc safe/core.rs:1648.
Repro: workflow train-xzv56 panicked after rollout completed
(timestep=999) on fold 0; rollout itself ran clean, the OOM from
the prior commit is gone.
Fix: replace the single `dtod_clone_i32` with a 3-step build:
1. alloc_zeros::<i32>(total)
2. cuMemsetD32Async(ptr, 0xFFFFFFFFu32, total, stream) — fills
all `total` slots with i32 mask sentinel -1 (byte pattern
0xFFFFFFFF reinterprets as i32(-1))
3. memcpy_dtod the first `base_total` real labels from
exp_aux_to_label_per_sample into the on-policy half
CF half remains at -1. The K=3 sparse-CE loss masks `label == -1`
out of the mean and B_valid count, so CF samples contribute zero
gradient — exactly the semantic we want (CF actions have no
observed trade-close outcome to predict against).
Lib suite: 1016/0 green. Audit doc updated.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
0acf77e656 |
config(dqn): strip H100-tuned VRAM overrides from dqn-production.toml
dqn-production.toml hard-coded three H100-tuned values that shadow
GpuProfile auto-detection: batch_size=16384, buffer_size=500K,
gpu_n_episodes=4096. After the L40S-default flip lands (prior commit
|
||
|
jgrusewski
|
bbd52c3aa7 |
feat(sp22-vnext): FoldReset registry entries for B5b/C collector buffers
Two collector-side device buffers used by the K=3 trade-outcome head
were missing FoldReset registry coverage:
- prev_aux_outcome_probs [alloc_episodes × 3]
TRUE stale-read risk. Producer writes end-of-step; consumer
(experience_state_gather) reads start-of-next-step into
state[121..124). Without FoldReset the new fold's step-0 state
gather would inject the previous fold's last-step softmax probs
into the first batch's state slots.
- exp_aux_to_input_buf [alloc_episodes × 262]
Cleanliness-only. Concat kernel overwrites all 262 columns every
step before the K=3 forward reads them, so no steady-state stale-
read risk. Registered for parity with the rest of the K=3
pipeline + to satisfy feedback_registry_entries_need_dispatch_
arms (the pin test asserts every registry entry has a matching
dispatch arm in reset_named_state).
Both fields promoted to pub(crate) on GpuExperienceCollector so
reset_named_state can reach them. Matching dispatch arms added with
the standard memset_zeros pattern (is_win_per_env / hold_baseline_
buffer style).
Tests: All 10 state_reset_registry tests pass, including the critical
every_fold_and_soft_reset_entry_has_dispatch_arm pin test that walks
the dispatch body and validates parity with registry entries. Full
lib suite 1015/1 (the failing test is the pre-existing
test_dqn_checkpoint_round_trip NoisyLinear flake — pred1/pred2 sign
mismatch surfacing ~30-50% of full-suite runs, documented in
project_sp22_h6_vnext_resume memory as unrelated to this work).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
4b40710b7c |
feat(sp22-vnext): Phase B5b-2 collector trade plan forward — resolves K=3 asymmetry
Phase B5b's K=3 input concat passed plan_params=NULL because the collector had no trade plan launch. Trainer-side K=3 forward trained on real plan_params while the collector queried at plan_params=0 — documented train/inference asymmetry on the plan-conditioning surface. Phase B5b-2 mirrors the (now-corrected) trainer `launch_trade_plan_forward` chain on the collector inside the rollout step: 1. SGEMM: hidden[N, AH] = h_s2_q[N, SH2] @ W_fc[AH, SH2]^T 2. bias+relu in-place on hidden 3. SGEMM: pre_out[N, 6] = hidden[N, AH] @ W_out[6, AH]^T 4. trade_plan_activate → exp_plan_params[N, 6] Weight resolution uses the same `aux_w_ptrs` array the K=3 forward already consumes (`f32_weight_ptrs_from_base`); indices 91-94 match the corrected trainer-side reads. The `trade_plan_activate` kernel is loaded from `EXPERIENCE_KERNELS_CUBIN` (same cubin the rest of `exp_module_extra` uses; the trainer loads it from there too). The K=3 concat now takes `exp_plan_params.raw_ptr()` instead of NULL — both sides see f(h_s2; W_plan_*_init), symmetry restored. Plan tensors at [91..94] still have no backward (no Adam updates), so the plan-head weights stay at Xavier cold-start forever. This commit delivers the symmetry the K=3 head requires, not a learned plan signal — adding a real plan-head backward is a follow-up project. New struct fields on GpuExperienceCollector: - exp_trade_plan_hidden_buf [alloc_episodes × adv_h] - exp_trade_plan_pre_out_buf [alloc_episodes × 6] - exp_plan_params [alloc_episodes × 6] - exp_trade_plan_activate_kernel: CudaFunction Lib test suite: 1016/0 green maintained. Audit doc updated. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
fb9b62a1f9 |
fix(plan-head): trainer trade plan forward reads from wrong param indices
`launch_trade_plan_forward` read `w_fc/b_fc/w_out/b_out` from `padded_byte_offset(¶m_sizes, [82..86))` — i.e., the ISV-conditioning + recursive-confidence tensors (`b_isv_gate`, `w_isv_gamma`, `b_isv_gamma`, `w_conf_fc`). The actual plan tensors live at `[91..95)` per `compute_param_sizes` and the matching Xavier init block. No backward exists for the plan head, so the plan tensors at `[91..94]` sat at cold-start Xavier values forever while `plan_params` was being driven by whichever ISV/conf weights happened to occupy the wrong offsets. Every downstream consumer (`backtest_plan_kernel`, `plan_isv` slots in `experience_kernels`, `compute_plan_params` in `q_value_provider`, regime gating, and the SP22 H6 vNext B5b concat path) was therefore conditioning on noise correlated with ISV optimisation, not on a learned plan. Discovered while preparing Phase B5b-2 (collector trade plan launch). Two-line index fix + diagnostic comment + audit doc entry. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
cdd3dc6edb |
feat(sp22-vnext): Phase F-3c — HEALTH_DIAG snap + console line for K=3 CE EMA
Completes the F-3 observability story. Smoke runs now print
aux_trade_outcome_ce_ema every epoch in the standard HEALTH_DIAG output
— no ISV inspector required.
Changes:
health_diag.rs:
- New field aux_trade_outcome_ce: f32 appended at END of
HealthDiagSnapshot per "Field order is stable; adding fields
appends to the end" doc rule. Existing fields' byte offsets
preserved → no kernel-side word offset re-validation.
- snapshot_size_is_stable pin: 149 × 4 → 150 × 4 = 600 bytes.
health_diag_kernel.cu:
- New WORD_AUX_TRADE_OUTCOME_CE = 149 (appended at end).
- WORD_TOTAL = 150 (was 149); static_assert bumped.
- New kernel arg aux_trade_outcome_ce_idx appended after
moe_lambda_eff_idx.
- New mirror write after the existing MoE mirrors. Stream-implicit
ordering: aux_outcome_ce_ema_update (F-3b) fires before
health_diag_isv_mirror, so the read picks up the just-updated EMA.
gpu_health_diag.rs:
- launch_isv_mirror gets new aux_trade_outcome_ce_idx: i32 arg.
gpu_dqn_trainer.rs:
- launch_health_diag_isv_mirror passes
AUX_TRADE_OUTCOME_CE_EMA_INDEX as the new arg.
training_loop.rs:
- Per-epoch metrics push appends ("aux_trade_outcome_ce_ema",
ISV[538]) to the standard out vec. Console / CSV automatically
includes the new column.
End-to-end F-3 chain now closed:
K=3 fwd → aux_to_loss_scalar_buf → aux_outcome_ce_ema_update
→ ISV[538] → health_diag_isv_mirror → snap.aux_trade_outcome_ce
→ training_loop metrics → console.
Operator sees CE every epoch:
- Cold-start: 0.000
- After bootstrap: ~1.098 (= ln(3))
- After training: ideally 0.5-0.7 (head learning)
Phase F end-to-end ready. The vNext stack has:
- Full GPU kernel chain (A2-A5 + D + plan-conditioning)
- Full Rust wireup (B0-B4, B4b-1/2, C-1/C-2, B5b)
- Real labels reaching trainer
- K=3 → policy via state slots + atom-shift
- End-to-end CE observability
Verification:
- cargo check -p ml clean.
- cargo test -p ml --lib → 1016/0 green (incl. bumped
snapshot_size_is_stable byte-size pin test).
Audit: docs/dqn-wire-up-audit.md Phase F-3c section.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
02479c885d |
feat(sp22-vnext): Phase F-3b — K=3 CE EMA launcher wireup + reset registry
Completes the Phase F-3 observability chain. Kernel + ISV slot landed
in F-3a; this commit wires the launcher into the trainer's per-step
training graph and registers the FoldReset entry.
Changes:
- gpu_dqn_trainer.rs:
- New pub(crate) static AUX_OUTCOME_CE_EMA_CUBIN embed
- New aux_outcome_ce_ema_kernel: CudaFunction field on GpuDqnTrainer
- Constructor loads kernel handle + struct-init
- New launch_aux_outcome_ce_ema() method (single-thread launch,
ISV slot 538, α=0.05)
- training_loop.rs:
- Launch call appended after launch_aux_heads_loss_ema()
- New dispatch arm "aux_trade_outcome_ce_ema" in reset_named_state
- state_reset_registry.rs:
- New RegistryEntry for "aux_trade_outcome_ce_ema" (FoldReset
sentinel 0.0)
End-to-end observability now live: K=3 head's batch-mean sparse-CE
flows into ISV[538] every step via Pearl A-bootstrapped EMA.
Smoke runs can read this slot to verify learning:
- Cold-start: 0.0
- After first step with B_valid > 0: bootstrap to first observation
(~1.098 = ln(3) for uniform K=3 prediction)
- Healthy learning: monotonic decrease toward 0.5-0.7 over epochs
- Falsification: pinned at ~1.098 for many epochs
Phase F prep complete. The trade-outcome aux head's:
- Forward chain (A2-A5 + B0-B4)
- Real labels reach trainer (B4b-1/2)
- K=3 → policy state (C-1/C-2)
- K=3 → Q-target atom-shift (D)
- Plan-conditioning (B5b)
- Smoke observability (F-3 + F-3b)
are all wired. Phase F deployment (argo-train.sh smoke) is next.
Verification:
- cargo check -p ml clean.
- cargo test -p ml --lib → 1016/0 green (incl. every_fold_and_soft_
reset_entry_has_dispatch_arm pin test).
Audit: docs/dqn-wire-up-audit.md Phase F-3b section.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
de64935b78 |
feat(sp22-vnext): Phase F-3 — K=3 CE EMA producer + ISV slot
Adds observability infrastructure for Phase F smoke validation: new ISV slot AUX_TRADE_OUTCOME_CE_EMA_INDEX = 538 tracks the K=3 head's batch-mean sparse cross-entropy EMA. Producer kernel registered + cubin-built; launcher wireup follows in Phase F-3b commit. Changes: - sp22_isv_slots.rs: new pub const AUX_TRADE_OUTCOME_CE_EMA_INDEX = 538 - gpu_dqn_trainer.rs ISV_TOTAL_DIM: 538 → 539 - NEW kernel aux_outcome_ce_ema_kernel.cu: single-thread single-block direct-to-ISV EMA writer with Pearl A first-observation bootstrap. Fixed α=0.05 (slow-moving observability). NULL-tolerant + NaN-guarded. - build.rs: kernel registered; cubin compiles (3.2 KB) Why dedicated kernel (not extension of aux_heads_loss_ema_update): The K=2/K=5 EMA writes through Pearls A+D's 2-stage producer scratch. Extending it would require allocating a new scratch slot, threading Pearls A+D mapping, and touching apply_pearls_ad_kernel. The K=3 head's CE is OBSERVABILITY ONLY in this commit — no Pearls A+D adaptive α needed. Minimum-scope direct-to-ISV path. Re-routable through Pearls A+D if K=3 CE later becomes a controller anchor. Smoke validation signal interpretation: - Cold-start: ISV[538] = 0.0 (FoldReset sentinel) - First step with B_valid > 0: Pearl A bootstrap → ISV[538] = first CE - Typical uniform-K=3 cold-start CE: ln(3) ≈ 1.098 - Learning signal: CE drops from ~1.098 toward 0.5-0.7 over epochs - Falsification signal: CE pinned at ln(3) for many epochs → head can't learn (label noise / under-capacity / outcomes unconditional) Phase F-3b follow-up: - Trainer launcher call after aux_heads_loss_ema_update - HEALTH_DIAG snap layout extension + console column - Reset registry entry (FoldReset sentinel 0) Verification: - cargo check -p ml clean. - cargo test -p ml --lib → 1016/0 green. Audit: docs/dqn-wire-up-audit.md Phase F-3 section. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
93aa4cd6a2 |
feat(sp22-vnext): Phase D — 12-weight W atom-shift (4 actions × 3 outcomes)
THE K=3 HEAD NOW DIRECTLY MODULATES Q-TARGETS. Phase D extends the
Phase 3 atom-shift mechanism from per-action W[4] reading single state
slot 121 to per-(action, outcome) W[4, 3] = 12 weights reading 3 state
slots [121..124).
Mathematical change: shift[a, b] now sums Σ_k W[a*K+k] × state[121+k]
= W[a, 0]*p_Profit + W[a, 1]*p_Stop + W[a, 2]*p_Timeout.
5 atom-shift kernel sites updated (all coordinated):
1. experience_kernels.cu::compute_expected_q (replay path)
2. experience_kernels.cu::mag_concat_qdir (rollout path)
3. experience_kernels.cu::quantile_q_select
4. c51_loss_kernel.cu loss numerator (next_state CVaR side)
5. c51_loss_kernel.cu Bellman target (online + target combined)
Each site replaces `W[a] × state_121` with `Σ_k W[a*K+k] × state[121+k]`
unrolled 3 times. State hoist points lift 1 scalar → 3-element array.
5 supporting kernel updates:
- aux_w_prior_init_kernel.cu: writes 12 K=3 structural priors instead
of 4. Spec prior matrix:
Short × {Profit=+0.5, Stop=-0.5, Timeout=0}
Hold × {Profit= 0, Stop=+0.5, Timeout=0}
Long × {Profit=+0.5, Stop=-0.5, Timeout=0}
Flat × {Profit= 0, Stop=+0.5, Timeout=0}
Block dim bumped 4 → 12.
- c51_aux_dw_kernel.cu: grid bumped (4,1,1) → (b0_size×K=12,1,1).
blockIdx.x decoded as (a, k); reads state slot 121+k, writes
dw_aux[a*K + k]. New kernel arg aux_outcome_k=3.
- adam_w_aux_kernel.cu: W_AUX_DIM 4 → 12. Block dim 12 threads.
Trainer-side buffer resizes:
w_aux_to_q_dir [4] → [12]
adam_m_w_aux [4] → [12]
adam_v_w_aux [4] → [12]
dw_aux_buf [4] → [12]
Rust launcher updates:
- c51_aux_dw_kernel launch: grid (4,1,1) → (12,1,1) + new aux_kto arg
- adam_w_aux_kernel launch: block (4,1,1) → (12,1,1)
- aux_w_prior_init launch: block (4,1,1) → (12,1,1)
Cold-start gracefulness preserved: state[121..124] = 0.0 at step 0
(no K=3 prediction yet from C-1 producer). Σ_k W[a*K+k] × 0 = 0 →
zero atom-shift across all actions. After step 1+ when C-1 producer
fires, real softmax probs activate the prior W's structural bias and
Adam refines from there.
End-to-end K=3 → Q-target chain now active:
K=3 fwd (B3/B4) → softmax → C-1 producer → prev_aux_outcome_probs
→ C-2 state gather → state[121..124) → Phase D atom-shift
→ Q-target z_n + Σ_k W[a*K+k] × prob[k]
→ Bellman target + argmax + action_select all see aux's outcome
prediction.
The K=3 head now influences policy via TWO paths: state input (Phase
C-2) AND Q-target modulation (Phase D).
Verification:
- cargo check -p ml clean.
- cargo test -p ml --lib → 1016/0 green.
Remaining vNext work:
- Phase E: dW backward gradient validation tests
- Phase F: Validation smoke at structural prior — decisive spec test
- B5b-2 (deferred): collector trade plan launch
Audit: docs/dqn-wire-up-audit.md Phase D section.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
d2331f2f62 |
feat(sp22-vnext): Phase B5b — full plan-conditioning integration
The K=3 head's input is now [h_s2_aux (256) || plan_params (6)] = 262- dim, matching the spec's intended architecture. Implements the input concat in both the trainer's replay-batch path and the collector's rollout-step path, with appropriate handling of the backward-side stride mismatch. NEW kernel strided_row_saxpy_kernel.cu: row-truncating SAXPY that accumulates first n_cols_copy columns per row of src [B, src_cols] into dst [B, dst_cols] scaled by alpha. Handles stride mismatch (src_cols != dst_cols). Needed because backward emits dh_s2_aux_to_buf [B, 262] but dh_s2_aux_accum [B, 256] only consumes first 256 cols. PLAN_PARAM_DIM = 6 constant in gpu_aux_heads.rs. Ops struct updates: - AuxTradeOutcomeForwardOps gains concat_kernel + launch_concat() - AuxTradeOutcomeBackwardOps gains strided_saxpy_kernel + launch_strided_row_saxpy() - Both load new cubins in new() Weight tensor resize: - sizes[163] = H × (SH2 + PLAN_PARAM_DIM) = 128 × 262 = 33,536 floats - fan_dims[163] = (H, SH2 + PLAN_PARAM_DIM) Trainer changes: - New aux_to_input_buf [B × 262] field - aux_dh_s2_to_buf resized to [B × 262] - aux_partial_to_w1 resized to [B × H × 262] - max_aux_tensor_len bumped for param_grad_final scratch Trainer forward (aux_heads_forward): - Concat h_s2_aux + plan_params_buf → aux_to_input_buf - forward() with SH2_TOTAL=262 Trainer backward (aux_heads_backward): - backward() with SH2_TOTAL=262; reads aux_to_input_buf - saxpy_f32_kernel SAXPY for dh_s2_aux REPLACED by launch_strided_row_saxpy: copies only first SH2=256 cols per row; trailing 6 cols (plan_params gradient) are dropped — STOP-GRAD on trade plan head from aux loss. Collector forward (rollout): - New exp_aux_to_input_buf [N × 262] field - Concat with plan_params_ptr = 0 (NULL) → zero-fill trailing 6 cols - forward() with SH2_TOTAL=262 Train/inference asymmetry (documented): - Trainer: real plan_params from trade plan head output - Collector: zeros (no trade plan launch in collector) - The head is trained on real plan-conditional outcomes but queried at rollout time with plan_params=0. Phase B5b-2 follow-up would add a trade plan launch to the collector to resolve. Deferred — current state is functional, head still receives plan signal during training. Stop-grad on plan_params: backward writes full [B, 262] gradient, but strided SAXPY only copies first 256 cols. Trade plan head weights NOT trained by K=3 aux loss in this commit. Verification: - cargo check -p ml clean. - cargo test -p ml --lib → 1016/0 green. Phase D next: 12-weight W atom-shift (4 actions × 3 outcomes). Audit: docs/dqn-wire-up-audit.md Phase B5b section. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
53462a28d9 |
feat(sp22-vnext): Phase C-2 — state gather flip K=2 single-slot → K=3 3-slot
THE K=3 HEAD NOW REACHES THE POLICY. Phase C-2 is the consumer-side flip — slot 121's semantic changes from K=2's recentered p_up to K=3's p_Profit, and slots [122, 123] gain new meaning as (p_Stop, p_Timeout). Architectural constraint: aux_dir_prob_per_env (K=2 buffer) is ALSO consumed by experience_env_step's β reward at lines 2335 + 3724-3725. Cannot repurpose that pointer to point at the K=3 [N, 3] buffer — env_step β consumer would read wrong-stride memory. This commit adds SEPARATE arg aux_outcome_probs_per_env [N, 3] to state_gather kernels with NULL fallback. Branching semantic at state_gather read site: - aux_outcome_probs_per_env != NULL → K=3 active path: writes slots [121..124) via new assemble_state_outcome_k3 helper - aux_outcome_probs_per_env == NULL → K=2 fallback: writes slot 121 via existing assemble_state from the recentered scalar Changes: - state_layout.cuh: NEW __device__ helper assemble_state_outcome_k3 — mirrors assemble_state except padding slots [121..124) get p_Profit/p_Stop/p_Timeout (raw softmax probs [0, 1]), slots [124..128) zero for 8-alignment. - experience_kernels.cu: training-side experience_state_gather + eval-side backtest_state_gather both get new trailing arg aux_outcome_probs_per_env (NULL-tolerant). Read site branches: K=3 reads 3 floats / env → assemble_state_outcome_k3; K=2 fallback preserves legacy assemble_state call. - gpu_experience_collector.rs: training launcher passes self.prev_aux_outcome_probs.raw_ptr() → K=3 active in training. - gpu_backtest_evaluator.rs: eval launcher passes NULL → K=2 fallback in eval (eval has no aux producer infra yet). K=2 head still alive: - prev_aux_dir_prob still populated by aux_softmax_to_per_env_kernel - experience_env_step still reads it for β reward (independent consumer untouched) - EGF chain still reads exp_aux_nb_softmax_buf - Only K=2's slot 121 contribution to policy state is suppressed End-to-end K=3 chain now active: Label producer (A2) → per-(env, t) ring (B4b-1) → replay buffer scatter (B4b-2) → PER direct gather → trainer aux_to_label_buf → loss_reduce (B4) sparse CE on real labels → backward (B4) per-sample partials → Adam SAXPY (B1+B4) updates W1/b1/W2/b2 at [163..167) PARALLEL: Collector rollout K=3 forward (B3) → softmax tile → C-1 producer → prev_aux_outcome_probs [N, 3] → C-2 state gather → state[121..124] → policy reads in next step The K=3 head closes the loop: learns from real labels via replay, AND predictions reach policy via state assembly. Trainable + observable. Verification: - cargo check -p ml clean. - cargo test -p ml --lib → 1016/0 green. Remaining vNext work: - Phase D: 12-weight W atom-shift (4 actions × 3 outcomes) - Phase E: dW backward + Adam for W[4, 3] - Phase F: Validation smoke at structural prior - B5b (deferred): plan_params input concat Audit: docs/dqn-wire-up-audit.md Phase C-2 section. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
3286dc7dee |
feat(sp22-vnext): Phase C-1 — K=3 softmax → per-env 3-slot cache (producer)
First half of Phase C. Lands the producer side of the K=3 trade- outcome aux head's state bridge: new kernel populates a per-env 3-slot cache from the K=3 softmax tile every rollout step. The consumer side (state gather reading from this cache → state slots [121..124)) lands in Phase C-2. Mirrors the K=2 head's existing aux_softmax_to_per_env_kernel exactly at K=3: - K=2: prev_aux_dir_prob[env] = 2*softmax[env, 1] - 1 (recentered) - K=3: prev_aux_outcome_probs[env, k] = softmax[env, k] for k in [0, 3) Changes: - state_layout.rs: 3 new constants AUX_OUTCOME_PROFIT_INDEX = 121, AUX_OUTCOME_STOP_INDEX = 122, AUX_OUTCOME_TIMEOUT_INDEX = 123. PROFIT_INDEX aliases AUX_DIR_PROB_INDEX (same value, different semantic). Phase C-2 flips slot 121's meaning from K=2's recentered p_up to K=3's p_Profit. - aux_outcome_softmax_to_per_env_kernel.cu: new kernel + cubin. - gpu_dqn_trainer.rs: new SP22_AUX_OUTCOME_SOFTMAX_TO_PER_ENV_CUBIN embed. - gpu_experience_collector.rs: 2 new struct fields (cache buffer + kernel handle); cubin load + alloc in constructor; struct-init; per-step launch in rollout loop after K=3 forward. - build.rs: kernel registered. Encoding shift K=2 → K=3: K=2 used recentered [-1, +1] to match "no signal = 0" baseline of every other slot. K=3 keeps raw softmax probabilities [0, 1]. Cold-start sentinel 0.0 for all 3 slots = "no prediction yet" (mask). The 3-slot natural distribution is more informative than a scalar. Dead-code status: producer populates cache every step but experience_state_gather doesn't read from it yet — state slot 121 still receives K=2's prev_aux_dir_prob write. Phase C-2 swaps the state gather's source from K=2 cache to K=3 cache (3-slot write). Why split C into C-1 + C-2: experience_state_gather is a hot-path kernel with many consumers. Updating it touches training collector, eval-side backtest evaluator, Rust launcher arg list. C-2 lands that as an atomic state-semantic flip; C-1 lands the GPU-side scaffolding independently so the producer chain can be validated first. Verification: - cargo check -p ml clean. - cargo test -p ml --lib → 1016/0 green. Audit: docs/dqn-wire-up-audit.md Phase C-1 section. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
68f0481a9e |
feat(sp22-vnext): Phase B5a — input concat kernel scaffolding
Lands the plan-conditioning concat kernel for the K=3 trade-outcome forward as reusable scaffolding. Phase B5b (integration) is deferred with rationale: Phase C (state slots) is more critical for testing the K=3 head's effect on policy behavior, and can land independently of the plan-conditioning refinement. NEW kernel aux_to_input_concat_kernel.cu: - Writes [B, SH2+P] from h_s2_aux [B, 256] || plan_params [B, 6] - Pure GPU map; one thread per output element, no atomicAdd - NULL-tolerant on plan_params (zeros trailing P cols when source unavailable, e.g., collector cold-start where the trade plan head doesn't run) - Registered in build.rs; cubin compiles (5.7 KB). Dead code at this commit — no Rust launcher yet. Why B5 is split + B5b deferred: Full Phase B5 (integration) requires three coordinated changes: 1. Forward path: bump aux_to_fwd.forward() to SH2=262 + 262-dim input 2. Backward stride mismatch: backward emits dh_s2_aux_to_buf [B, 262], but dh_s2_aux_accum (input to aux trunk backward) is [B, 256]. A direct SAXPY mismatches row strides (262 vs 256) and corrupts the trunk's upstream gradient. Needs a strided-SAXPY kernel. 3. Collector-path plan_params unavailability: trade plan head only runs trainer-side. Workarounds: zero-fill, add trade plan to collector, or skip K=3 forward in collector. All have trade-offs. Phase B5b would need (1) strided-SAXPY kernel and (2) collector plan_params decision. Real work but NOT on the critical path for testing the K=3 head's effect on WR. Why Phase C should land first: The K=3 head currently trains on real labels (post-B4b) but doesn't influence policy behavior. Phase C wires the head's softmax into state slots [121..124) = (p_Profit, p_Stop, p_Timeout), replacing the K=2 single-slot 121 = 2*p_up - 1. WITH Phase C the policy reads aux's outcome predictions as state features → behavior changes → testable. Without Phase C, validation runs would show "K=3 head trains and converges" but predictions don't reach the policy → WR signal isn't a function of K=3 at all. We'd be testing nothing. Recommendation: skip the full B5 for now, do Phase C next, then Phase D (atom-shift). Phase B5b (plan-conditioning) is a refinement we add IF Phase C/D's no-plan-params version shows promise but plateaus below the WR ≥ 0.55 target. Audit: docs/dqn-wire-up-audit.md Phase B5a section. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
da5e564ccf |
feat(sp22-vnext): Phase B4b-2 — replay-buffer scatter + trainer setter
Second half of Phase B4b. Wires the trade-outcome label column through
the replay buffer (struct field, allocation, scatter on insert, gather
on sample, direct-to-trainer pointer + setter) and connects the
trainer's aux_to_label_buf to receive per-batch sampled labels.
Completes the end-to-end producer → ring → trainer i32 path the K=3
sparse CE consumer reads.
Replay buffer changes (crates/ml-dqn/src/gpu_replay_buffer.rs):
- New struct fields: aux_outcome_labels (capacity-sized ring),
sample_aux_outcome_labels (mbs-sized fallback gather), trainer_aux_
outcome_labels_ptr (direct-path destination)
- New aux_outcome_labels_ptr field on GpuBatchPtrs
- insert_batch signature: new aux_outcome arg between aux_conf_in and
bs. Scatters via existing K-generic scatter_insert_i32 (same kernel
the K=2 aux_sign_labels uses).
- sample_proportional direct gather when trainer_aux_outcome_labels_ptr
!= 0; fallback gather otherwise. Mirrors aux_sign_labels direct/
fallback semantic exactly.
- New setter set_trainer_aux_outcome_labels_ptr mirrors
set_trainer_aux_conf_ptr.
Collector emission:
- New aux_outcome_labels field on GpuExperienceBatch
- Populated at end of collect_experiences_gpu via dtod_clone_i32 from
Phase B4b-1's per-(env, t) producer scratch.
Trainer wireup:
- aux_to_label_buf_ptr() accessor on GpuDqnTrainer (mirrors
aux_nb_label_buf_ptr)
- trainer_aux_to_label_buf_ptr() delegating accessor on
FusedTrainingCtx
- New set_trainer_aux_outcome_labels_ptr call in training_loop at the
same site where set_trainer_buffers + set_trainer_aux_conf_ptr fire.
Two call sites updated (lines ~835 + ~2857).
- insert_batch call in training_loop passes &gpu_batch.aux_outcome
_labels as new arg.
Test fixtures updated: 4 smoke test files + 3 unit-test fixtures in
gpu_replay_buffer.rs alloc zero-init aux_outcome i32 arg.
End-to-end chain complete:
trade_outcome_label_kernel (A2)
→ collector per-step launch (B3)
→ collector emission (B4b-1)
→ replay-buffer insert + scatter (B4b-2)
→ PER sample + direct gather (B4b-2)
→ trainer aux_heads_forward.loss_reduce (B4) reads sparse {-1,0,1,2}
→ trainer aux_heads_backward (B4) computes per-sample partials
→ Adam SAXPY (B1+B4) updates W1, b1, W2, b2 at [163..167)
The "degraded predict-Profit-everywhere" cold-start from Phase B4 is
resolved. K=3 head trains on real sparse trade-outcome labels.
Verification:
- cargo check -p ml clean (21 warnings, none new).
- cargo test -p ml --lib → 1016/0 on clean runs; pre-existing
NoisyLinear flake still surfaces ~30-50% of runs (unrelated to
vNext work — see
|
||
|
jgrusewski
|
491bf7d3e6 |
feat(sp22-vnext): Phase B4b-1 — per-(env, t) label kernel output + collector buffer
First half of Phase B4b (replay-buffer label scatter chain). Amends the Phase A2 trade_outcome_label_kernel to emit a per-(env, t) output column alongside the existing per-env tile, and adds the collector-side buffer + per-step launch arg. Kernel amendment (trade_outcome_label_kernel.cu): - Added NULL-tolerant `out_labels_per_sample` arg after existing `out_labels`. When non-NULL, writes `out_labels_per_sample[env*L + t] = label` at the same offset as the `trade_close_per_sample[env*L + t]` read. NULL = no-op (preserves Phase A2/A3 contract for callers passing old signature). - Pattern mirrors the K=2 head's `aux_sign_labels` per-(i, t) ring column that threads through the replay buffer. Collector field + alloc + launch: - New struct field `exp_aux_to_label_per_sample: CudaSlice<i32>` sized `[alloc_episodes × alloc_timesteps]`. Sentinel -1 (mask) populated by alloc_zeros + per-step kernel writes — survives until a trade-close event overwrites the env's slot at that t. - Updated Phase B3 launcher in collect_experiences_gpu to pass `self.exp_aux_to_label_per_sample.raw_ptr()` as the new arg. Why split B4b into B4b-1 + B4b-2: the full replay-buffer wireup mirrors the K=2 head's aux_sign_labels pattern across ~8 distinct code sites (replay-buffer struct field, sample destination buffer, direct-to-trainer pointer, setter method, scatter on insert, gather direct, gather fallback, GpuBatchPtrs field). Splitting lets us validate the per-(i, t) producer in isolation before touching the consumer pipeline. B4b-1 (this commit) = producer chain complete. Per-(env, t) column populated correctly every rollout step. Consumer wiring (replay- buffer scatter + trainer setter) is B4b-2's scope. Verification: - cargo check -p ml clean (21 warnings, none new). - cargo test -p ml --lib → 1016/0 on clean runs; pre-existing test_dqn_checkpoint_round_trip NoisyLinear flake still surfaces ~50-70% of full-suite runs (flake predates Phase B4b, unrelated to trade-outcome head — disable_noise() zeros ε but leaves some other randomness source intact). Audit: docs/dqn-wire-up-audit.md Phase B4b-1 section. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
20e1aea27c |
feat(sp22-vnext): Phase B4 — trainer-side replay-batch chain wireup
Wires the trade-outcome head's forward + loss reduce + backward +
per-sample partial reduce + SAXPY into the trainer's aux_heads_forward
and aux_heads_backward methods. Adam SAXPY for the 4 new weight tensors
at [163..167) extends uniformly via the existing aux_param_specs array
iteration.
Changes to aux_heads_forward:
- Steps 7 + 8 appended after K=2 next-bar head's loss reduce
- K=3 forward reads weights at [163..167) (Phase B1) → writes to
aux_to_* save-for-backward tiles (Phase B2)
- K=3 loss_reduce writes aux_to_loss_scalar_buf + aux_to_valid_count_buf
Changes to aux_heads_backward:
- K=3 backward appended after K=5 regime backward, emits per-sample
partials (dW1, db1, dW2, db2) + per-sample dh_s2_aux
- aux_param_specs array extended from 8 → 12 entries. Per-tensor
reduce + SAXPY loop iterates uniformly, scaling each by aux_weight
- K=3 dh_s2_aux SAXPY appended after K=5's, all three heads'
gradients flow into aux trunk's dh_s2_aux_accum (encoder stop-grad
enforced structurally by aux_trunk_backward's missing dx_in output)
Label semantic (cold-start): aux_to_label_buf is alloc_zeros (all 0
= Profit) until Phase B4b lands replay-buffer label scatter. Model
trains on "predict Profit everywhere" — degraded but well-defined
(no NaN). Mirrors K=2 head's known-degraded state between B1.1a
(forward landed) and B1.1b (label producer wired).
Adam SAXPY: existing global SAXPY iterates 0..NUM_WEIGHT_TENSORS
(now 167) — 4 new weight slots get gradient SAXPYs followed by
Adam m/v updates uniformly. Architectural payoff of Phase B1's
NUM_WEIGHT_TENSORS bump.
Test flake mitigation: added bind_to_thread() to
ensemble::adapters::dqn::tests::shared_device() mirroring the
cuda_stream() test-helper pattern from the fix sweep at
|
||
|
jgrusewski
|
b28b349ac3 |
feat(sp22-vnext): Phase B3 — collector-side rollout buffers + forward chain wireup
Adds collector-side trade-outcome head: 5 struct fields + allocations
+ per-step forward + per-step label producer launches in the rollout
loop. Mirrors the K=2 next-bar head's collector wireup at K=3.
Collector struct additions:
- exp_aux_to_fwd: AuxTradeOutcomeForwardOps (3 kernel handles)
- exp_aux_to_hidden_buf [alloc_episodes × H=128] saved post-ELU
- exp_aux_to_logits_buf [alloc_episodes × K=3] saved logits
- exp_aux_to_softmax_buf [alloc_episodes × K=3] softmax tile
- exp_aux_to_label_buf [alloc_episodes] i32 sparse {-1, 0, 1, 2}
Per-step launches in collect_experiences_gpu rollout loop:
1. aux_trade_outcome_forward — launched immediately after the K=2
sibling's forward_next_bar, parallel on the same stream. Reads
exp_h_s2_aux + weights at flat-buffer indices [163..167) (Phase
B1 additions). Writes hidden/logits/softmax tiles. No consumer
yet — Phase C wires state assembly; Phase B4 wires trainer
scatter.
2. trade_outcome_label_kernel — launched immediately after
experience_env_step on the same stream, reading the save-for-
backward buffers (pnl_vs_target_at_close_per_env, pnl_vs_stop_at_
close_per_env) that env_step just wrote at segment_complete.
Stream-implicit producer→consumer ordering. Emits per-env
{-1, 0, 1, 2} labels — sparse, ~95-99% bars produce -1 (mask).
Dead-code discipline per feedback_wire_everything_up: every kernel arg
+ producer site is real wiring (not NULL placeholder) — only the
absence of consumers reading the produced tiles is "dead". The smoke
run produces softmax tiles + labels every step bit-identical to
pre-vNext baseline (no consumer = no effect on training behavior).
Phase B4 next: trainer-side replay-batch chain (forward + loss_reduce
+ backward + Adam SAXPY for the 4 new weight tensors).
Audit: docs/dqn-wire-up-audit.md Phase B3 section.
Verification:
- cargo check -p ml clean (21 warnings, none new on aux_to_*).
- cargo test -p ml --lib → 1016 passing / 0 failing (unchanged from
post-fix-sweep baseline at
|
||
|
jgrusewski
|
ebc1b15023 |
fix(tests): repair 14 pre-existing test failures across ml crate
Lib test suite was at 14 failures from accumulated layout/contract drift.
Fixed each by tracing root cause; lib suite now 1016 passing / 0 failing.
Failures fixed (test → root cause → fix):
1. sp14_isv_slots::sp20_isv_slots_reserved_510_to_520 — ISV_TOTAL_DIM
pin drifted from SP20-era 520 to current 538 via SP21/SP22 H6 bus
growth. Slot positions 510-519 still pinned. Fix: relax total-dim
check to >= 520; keep slot-position asserts tight.
2-3. gradient_budget::test_spectral_norm_all_heads_no_panic +
test_spectral_norm_constrains_operator_norm — test config used
cfg.state_dim=16 but trainer uses STATE_DIM=128 when bottleneck
off; also DuelingWeightBacking slices [2]/[3] used pre-GRN w_s2
shape, but post-GRN it's w_b_h_s1 [2*SH1, SH1] = 2048. Fix: add
s1_input_dim_for_test helper; correct slice sizes to GRN layout.
4-9. dqn::trainer::tests::test_feature_vector_to_state +
test_single_sample_batch + test_batched_action_selection +
test_batched_vs_sequential_action_selection_consistency +
test_batch_size_mismatch_larger/smaller_than_configured —
feature_vector_to_state wrapper falsely advertised "no OFI" by
signature but body required strict OFI. Contract violation —
broke 6 unit tests + hyperopt's public convert_to_state APIs.
Fix: wrapper now genuinely produces 45-dim no-OFI state; OFI-
strict callers use _with_ofi with Some(idx).
10. state_kl_monitor::observe_tracks_fire_rate_on_change — last_amp
initialized to 1.0 coincidentally equaled first observation value,
silently suppressing first fire. Test expected first observation
always fires (cold-start semantic). Fix: Option<f32> sentinel —
None means "no prior baseline" → first observe ALWAYS fires.
11. cuda_pipeline::tests::test_eval_action_select_thompson_picks_
proportionally — test launched experience_action_select kernel
with 1 missing arg (v_logits_dir from SP17 Commit C). Kernel
read garbage pointer → SIGSEGV. Also threshold 0.70 calibrated
for pre-SP17 raw-A distribution; post-SP17 sampler uses mean-zero
softmax(V + A_centered), Long still dominates ~4× but P(Long)
drops to ~0.66. Fix: add zero-filled v_logits_buf at correct
arg position; recalibrate threshold 0.70 → 0.60.
12. cuda_pipeline::tests::test_ppo_gpu_data_upload — flaked in parallel
test run with CUDA_ERROR_INVALID_CONTEXT from MappedF32Buffer's
cudaHostAlloc. cuda_stream() test helper used OnceLock to share
a CudaStream but didn't bind_to_thread on subsequent calls. CUDA
contexts are per-thread state. Fix: helper now calls bind_to_
thread() on every invocation (idempotent — mirrors trainer/
constructor.rs:111's pattern).
13. ppo::tests::test_reward_computation — test created hold action
with ExposureLevel::Flat but is_hold() matches only Hold (Flat
= close position = transaction with cost). So hold and buy got
same transaction-cost reduction; reward_hold > reward_buy assert
failed. Fix: use ExposureLevel::Hold for no-cost hold semantic.
14. training_profile::tests::test_production_profile_applies_all_
sections — pinned n_steps==5 and tau==0.005 (pre-TD(0)); production
flipped to n_steps=1, tau=0.01 per dqn-production.toml ("dense
micro-rewards cancel over n>1 bars; faster target tracking for
TD(0)"). Fix: update pins to current production values.
All 14 fixes target root causes — no #[ignore] masking. Verified:
- Lib-only run: cargo test -p ml --lib → 1016/0 (passed/failed)
- Stash-test pre-changes: 1001/15 (confirms 14 fixed atomically)
Remaining flake under `cargo test -p ml --tests`:
- ensemble::adapters::dqn::tests::test_dqn_checkpoint_round_trip —
CUDA-context-race-adjacent flake under parallel `--tests` mode;
passes in isolation and in `--lib` mode (single binary). Predates
this commit; deferred as separate triage.
Audit: docs/dqn-wire-up-audit.md "Fix sweep" section.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
b98dc2730d |
feat(sp22-vnext): Phase B2 — trade-outcome trainer saved-tensor + partial buffers
Adds 11 buffer fields + 2 orchestrator ops handles (fwd + bwd) to the
trainer struct, mirroring the existing aux_nb_* / aux_partial_nb_*
pattern at K=3 instead of K=2.
Trainer struct additions:
- aux_to_fwd: AuxTradeOutcomeForwardOps (Phase B0 scaffold)
- aux_to_bwd: AuxTradeOutcomeBackwardOps
- aux_to_hidden_buf [B, H=128] saved post-ELU
- aux_to_logits_buf [B, K=3] saved logits
- aux_to_softmax_buf [B, K=3] saved softmax (3 future consumers)
- aux_to_label_buf [B] i32 sparse {-1, 0, 1, 2}
- aux_to_loss_scalar_buf [1] mean CE
- aux_to_valid_count_buf [1] B_valid for backward
- aux_dh_s2_to_buf [B, SH2] SAXPYs into dh_s2_aux_accum
- aux_partial_to_w1 [B, H, SH2] per-sample dW1
- aux_partial_to_b1 [B, H] per-sample db1
- aux_partial_to_w2 [B, K=3, H] per-sample dW2
- aux_partial_to_b2 [B, K=3] per-sample db2
Memory: aux_partial_to_w1 = 256 MB at B=2048 — identical to K=2 head's
partial size (same SH2, same H). Total new aux-to footprint ≈ 260 MB.
The existing aux_param_grad_final_buf scratch is sized to the largest
tensor across all aux heads; trade-outcome head's largest is W1 [H, SH2]
= 32,768 floats — identical to K=2/K=5 W1s. No resize needed.
Cold-start label semantics: alloc_zeros yields label 0 (Profit) for
every sample. Until the producer wires in (B3), the trainer's CE loss
treats every sample as "should have predicted Profit" — degraded but
well-defined (no NaN). Mirrors the K=2 head's known-degraded state
between B1.1a and B1.1b.
No FoldReset registration: these buffers are overwritten every batch
— no stale-state-leak risk across folds (matches the existing aux_nb_*
pattern).
Phase B3 next: collector-side rollout buffers + forward chain wireup
into collect_experiences_gpu (per-env softmax → per-(i, t) fan-out
scatter for trainer's aux_to_softmax_buf population).
Audit: docs/dqn-wire-up-audit.md Phase B2 section.
Cargo check clean (21 warnings, none new on aux_to_* fields).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
205e46c171 |
feat(sp22-vnext): Phase B1 — trade-outcome head weight tensors + Xavier init
Adds the 4 weight tensors (W1, b1, W2, b2) for the SP22 H6 vNext
trade-outcome aux head into the trainer's flat params_buf at indices
[163..167). Adam machinery (m/v moment buffers, SAXPY iteration over
0..NUM_WEIGHT_TENSORS) picks up the new tensors uniformly — no
per-tensor wiring needed.
Changes:
- NUM_WEIGHT_TENSORS bumped 163 → 167. Most of the 54 references are
&[u64; NUM_WEIGHT_TENSORS] array-size generics that resize
uniformly with the constant.
- compute_param_sizes() adds 4 new size entries:
[163] aux_to_w1 [H=128, SH2=256] = 32,768 floats
[164] aux_to_b1 [H=128] = 128 floats
[165] aux_to_w2 [K=3, H=128] = 384 floats
[166] aux_to_b2 [K=3] = 3 floats
Total: 33,283 floats = ~133 KB params, ~266 KB Adam state.
- compute_param_sizes() debug_assert updated 163 → 167.
- Xavier fan_dims added: (H, SH2) for W1, (0, 0) for biases (zero-init),
(K=3, H) for W2. Cold-start: logits ≈ 0 → softmax ≈ uniform 1/3 → no
Profit/Stop/Timeout preference per pearl_first_observation_bootstrap.
SH2 stays at 256 in this commit (mirrors K=2 head exactly). The spec's
Phase B input concat (256 → 262 with plan_params 6-dim) will re-shape
slot [163] to 128 × 262 = 33,536 floats later — small touch-up vs the
full B1 commit.
Verification:
- cargo check -p ml clean (21 warnings, none new).
- 14 pre-existing test failures stashed-verified unrelated (OFI features
missing, ISV slot count drift — independent of NUM_WEIGHT_TENSORS).
Phase B2 next: saved-tensor + per-sample partial buffers (hidden_post,
logits, softmax, valid_count, dW*_partial, dh_s2_aux_out).
Audit: docs/dqn-wire-up-audit.md Phase B1 section.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
108a426c38 |
feat(sp22-vnext): Phase B0 — AuxTradeOutcome ops struct scaffolding
First Rust-side commit of Phase B for the SP22 H6 vNext trade-outcome aux head. Pure orchestrator scaffolding mirroring AuxHeadsForwardOps / AuxHeadsBackwardOps. Zero production callers — additive per the gpu_grn / aux_trunk commit-by-commit ordering convention (scaffold → trainer fields → collector wireup). Adds: - gpu_dqn_trainer.rs: 4 cubin embeds (TRADE_OUTCOME_LABEL_CUBIN, AUX_TRADE_OUTCOME_FORWARD_CUBIN, AUX_TRADE_OUTCOME_LOSS_REDUCE_CUBIN, AUX_TRADE_OUTCOME_BACKWARD_CUBIN). All #[allow(dead_code)] until B1+. - gpu_aux_heads.rs: AUX_OUTCOME_K = 3 constant (parallel to AUX_NEXT_BAR_K). - gpu_aux_heads.rs: AuxTradeOutcomeForwardOps struct holding 3 kernel handles (forward + loss_reduce + label producer). Launch methods: forward(), loss_reduce(), compute_label(). Per-env label kernel uses grid ceil(n_envs/256) — pure per-env map, no reduction. - gpu_aux_heads.rs: AuxTradeOutcomeBackwardOps struct holding 1 kernel handle (backward). Launch method: backward(). Reuses existing K-generic aux_param_grad_reduce from AuxHeadsBackwardOps — no separate reducer needed. Contract shapes mirror AuxHeadsForwardOps/AuxHeadsBackwardOps so subsequent wireup commits plug in with minimal contract drift. Phase B proper (input concat 256→262 with plan_params) becomes a small change touching only the buffer fill + W1 shape once B1-B4 land the wireup. Phase B1 next: trainer struct fields for W1/b1/W2/b2 + Adam state + Xavier init + reset registry entries. Audit: docs/dqn-wire-up-audit.md Phase B0 section. Cargo check clean (21 warnings, none new). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
9f4a25e623 |
feat(sp22-vnext): Phase A5 — aux_trade_outcome backward kernel (Phase A complete)
K=3 backward kernel that closes the forward → loss → backward chain for the trade-outcome aux head. Mirrors `aux_next_bar_backward` (K=2 sibling) line-for-line because gradient flow is K-independent: `d_logits = (softmax − one_hot)/B_valid` propagated through `Linear → ELU → Linear` chain via standard softmax-CE derivative. Per-sample partials (caller reduces via existing K-generic `aux_param_ grad_reduce` kernel): dW1_partial [B, H=128, SH2=256], db1_partial [B, H] dW2_partial [B, K=3, H], db2_partial [B, K=3] dh_s2_aux_out [B, SH2] Mask handling: labels[b] == -1 zeros the K-vector → all downstream partials zero (chain rule's multiplicative zero). All-skip batch produces valid_count=0 → d_logits=0 for every row → zero gradients across the board, no NaN. Sparse-label gradient amplification: B_valid is typically ~1-5% of nominal batch (trade-close events are rare), so inv_B = 1/B_valid is much larger than the K=2 sibling's inv_B = 1/(~B). Per-trade-close gradients have proportionally higher magnitude — correct credit assignment (rare signal speaks louder) but Phase E's Adam may need class-weighted CE or per-group LR tuning. ELU backward via post-activation identity: f'(x) = (h_post > 0) ? 1 : 1 + h_post — recovers derivative without re-evaluating x_pre. SP14 Phase C.5b separation preserved: reads h_s2_aux (aux trunk output), writes dh_s2_aux_out SAXPYing into dh_s2_aux_accum. Q's encoder structurally protected (aux_trunk_backward has no dx_in output). Phase A5 (this commit) is dead code — no Rust launcher. Phase B will land the full launcher chain (gpu_aux_heads.rs parallel ops struct, collector struct fields for W1/W2/b1/b2/Adam-state/saved-tensors/dW- partials, wireup in collect_experiences_gpu). Cubin: aux_trade_outcome_backward_kernel.cubin (24.8 KB). ═══ Phase A complete ═══ A1: ISV slots (none needed — reuses padding 121-123) A2: trade_outcome_label_kernel.cu (label producer) |
||
|
jgrusewski
|
3ddcfb8868 |
feat(sp22-vnext): Phase A4 — aux_trade_outcome loss reduce kernel
K=3 sparse cross-entropy reduce over the trade-outcome softmax tile produced by `aux_trade_outcome_forward` (Phase A3). Mirrors the K=2 sibling `aux_next_bar_loss_reduce` structurally: single-block shmem-tree reduce, two parallel partial strips (loss_numer + valid_count) reduced lockstep, fmaxf(p_tgt, 1e-30) numerical floor, fmaxf(valid, 1.0) all- skip-batch guard, valid_count_out[1] save-for-backward. Kept as SEPARATE kernel from the K=2 sibling: - Diagnostic isolation (distinct HEALTH_DIAG slot, distinct cubin in profiles for clean per-loss-source attribution) - Sparse-label semantic clarity (~95-99% mask=-1 vs ~50-100% valid for the K=2 next-bar head) - Future per-class weighting headroom (Profit/Stop/Timeout 3:1-10:1 imbalance will likely need class-weighted CE — surgical mod here without touching the K=2 head's contract) Phase A4 (this commit) is dead code — no Rust launcher yet. Phase A5 lands backward; Phase B wires the full forward→loss→backward chain. Discipline: feedback_no_atomicadd (single-block tree-reduce), feedback_ cpu_is_read_only (pure GPU), pearl_first_observation_bootstrap (sentinel 0 valid_count produces zero gradients gracefully on cold start). Audit: docs/dqn-wire-up-audit.md Phase A4 section. Cubin: aux_trade_outcome_loss_reduce_kernel.cubin (9.9 KB) compiles clean. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
07728f9efc |
feat(sp22-vnext): Phase A3 — aux_trade_outcome forward kernel + save-for-backward wireup
Phase A3 of the SP22 H6 vNext trade-outcome aux head (per
docs/plans/2026-05-14-sp22-h6-vNext-trade-outcome-aux.md). Two atomic pieces:
1. NEW kernel `aux_trade_outcome_forward_kernel.cu` — K=3 softmax aux head
forward (Linear → ELU → Linear → stable softmax). Mirrors
`aux_next_bar_forward` (K=2) but emits {Profit, Stop, Timeout} probs.
Saved tensors {hidden_out, logits_out, softmax_out} ready for A4 (loss
reduce) and A5 (backward). Dead code at this commit — no Rust launcher
yet. Registered in build.rs::kernels_with_common, cubin verified.
2. Save-for-backward buffers `pnl_vs_target_at_close_per_env` +
`pnl_vs_stop_at_close_per_env` ([alloc_episodes] f32 device-resident).
Producer: `experience_env_step::segment_complete` writes the trade's
realized P&L ratios vs profit_target / stop_loss at close (inline-
computed from `pre_trade_position × (raw_close − entry_price) /
(ps[PS_PLAN_PROFIT_TARGET] × prev_equity)`, symmetric-clamped to
[-2, +2] per pearl_symmetric_clamp_audit — same formula as the
sibling experience_state_gather's plan_isv[PLAN_ISV_PNL_VS_TARGET/_STOP]
slots). Consumer (eventual A4/A5 wireup): trade_outcome_label_kernel
classifies each close into {Profit, Stop, Timeout} via the ≥1.0
threshold-hit predicate.
Wireup discipline per feedback_registry_entries_need_dispatch_arms:
- New struct fields on GpuExperienceCollector
- stream.alloc_zeros at construct site
- Kernel-launch .arg() threading at experience_env_step launch
- StateResetRegistry entries (FoldReset sentinel 0.0)
- training_loop::reset_named_state dispatch arms
- All 10 registry pin tests pass including
every_fold_and_soft_reset_entry_has_dispatch_arm
Audit doc updated: docs/dqn-wire-up-audit.md Phase A3 section.
Next phases (per spec): A4 = aux_trade_outcome_loss_reduce (sparse CE,
mask=-1), A5 = aux_trade_outcome_backward, Phase B = 262-dim input
(h_s2_aux || plan_params), Phase C = 3-slot state assembly, Phase D =
12-weight W atom-shift, Phase E = dW + Adam, Phase F = validation smoke.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
26ce7ba690 |
feat(sp22-vNext): Phase A2 — trade-outcome label producer kernel
First foundation kernel for the H6 vNext trade-outcome aux head. Per-env classification at trade-close events into K=3 outcomes: - 0 = Profit (pnl_vs_target >= 1.0) - 1 = Stop (pnl_vs_stop >= 1.0) - 2 = Timeout (neither threshold hit) Sparse labels — most bars get -1 mask. Priority: Profit > Stop > Timeout. Pure per-env map; no atomicAdd, no reduction. Launch: grid(ceil(N/256)), block(256). Registered in build.rs. Cubin compiles clean. Currently dead code — launcher wireup comes in Phase A3+ commits. See docs/plans/2026-05-14-sp22-h6-vNext-trade-outcome-aux.md for the full vNext architecture. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
d0c037a3d2 |
feat(sp22): H6 Phase 3 FALSIFIED + vNext spec (trade-outcome aux head)
Decisive smoke train-xrkb7 @
|
||
|
jgrusewski
|
ebc7144434 |
feat(sp22): H6 Phase 3 RE-ACTIVATED with corrected aux head
Smoke train-8zwtf @
|
||
|
jgrusewski
|
465abc7e9b |
fix(sp14): dial back aux head/trunk enlargement after L40S OOM
Initial 8x head + 2x trunk-H2 (commit
|
||
|
jgrusewski
|
787ee7b86c |
feat(sp14/sp22): enlarge aux head + trunk capacity (8x head, 2x trunk-H2)
After Path C confirmed aux head's 28% accuracy at H=60 is the H6 Phase 3 bottleneck (not the mechanism itself), enlarge aux capacity: - AUX_HIDDEN_DIM: 32 -> 256 (8x, matches input dim, removes bottleneck) - AUX_TRUNK_H2: 128 -> 256 (2x, uniform trunk width) Architecture changes: - Aux head: 256 -> Linear -> 256 -> ELU -> Linear -> 2 (was 256->32->2) - Aux trunk: 256 -> 256 -> 256 -> 256 (was 256 -> 256 -> 128 -> 256) - +160K params total (mostly aux_nb_w1 + aux_rg_w1: [256, 256] each) Side effects: - Checkpoint fingerprint change (intentional) - Thread utilisation improves: AUX_BLOCK=256 threads x H=256 = 1:1 (vs 8:1 at H=32 — most threads idle previously) Phase 3 mechanism stays DORMANT (W=0, beta=0) for this validation smoke. Verdict criteria: aux_dir_acc improves from 0.28 toward 0.50+ with the larger capacity. If yes, re-activate Phase 3 priors. If no, the bottleneck is signal/horizon, not capacity. Cargo build clean (full nvcc rebuild). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
2c0911981a |
feat(sp22): H6 Phase 3 DORMANT - mechanism falsified, infrastructure preserved
Path C investigation revealed the aux head at epoch 1 is severely
anti-predictive at H=60 bars:
- Aux predicts UP 83% of the time
- Labels are 17% UP, 83% DOWN
- Accuracy = 28% (vs 50% random)
This means H6 Phase 3 hypothesis cannot help WR — atom-shift on an
anti-predictive signal produces no discriminative bias. Confirmed by
full-mechanism smoke at
|
||
|
jgrusewski
|
b4e26a3b45 |
feat(sp22): H6 Phase 3 α/β — RESTORE structural priors
Verification smoke train-5t6vb @ |
||
|
jgrusewski
|
79945987a5 |
fix(sp22): W-read guards across ALL atom-shift kernels
Verify smoke train-t5885 partial result: dir CLEAN, mag NaN. The isfinite(W) guard added in compute_expected_q didn't extend to the other atom-shift consumers. Adam-corrupted W propagates through them via 0*NaN=NaN. Adds isfinite(w_aux[a]) guards to all atom-shift kernels: - mag_concat_qdir - quantile_q_select - c51_loss_kernel (eq_per_action shift + effective_reward W[a0]/W[best_next_a]) Also strengthens c51_aux_dw_kernel: guard sp/isw/dz/gamma/done (not just dz<1e-7 which doesn't catch NaN). Any NaN/Inf input -> skip sample (no dW contribution). Defense-in-depth complete: NaN cannot propagate through any atom-shift path regardless of source. Cargo check clean. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
57cf5c1e08 |
fix(sp22): defensive NaN guards on Step 8/11 dW + Adam path
Bisect cut #1 (train-6zbcn) DEFINITIVELY identified Step 8/11 as the NaN source. With launches disabled, q_var/q_by_action/v_a_means all CLEAN. Forward atom-shift kernels confirmed innocent. Defensive fix (belt-and-suspenders): 1. c51_aux_dw_kernel: clamp NaN/inf total to 0 at dW write site. 2. adam_w_aux_kernel: early-return if any input non-finite. 3. compute_expected_q: guard w_aux[a] read with isfinite (in addition to state_121). Together these prevent any NaN propagation through atom-shift regardless of where the upstream NaN originated. Step 8/11 launches RESTORED in submit_adam_ops. The underlying NaN source is still unidentified but defensively neutralised. Smoke verification next. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
a1945de031 |
bisect(sp22): disable Step 8/11 launches (c51_aux_dw + adam_w_aux)
User-requested static-analysis-first bisect (Option 2): comment out both launch_c51_aux_dw and launch_adam_w_aux in submit_adam_ops to isolate forward vs backward as NaN source. With these disabled: - Forward atom-shift kernels still run (no-op with W=0) - W stays at [0,0,0,0] (no Adam update) - dW kernel doesn't run Hypothesis test: - Clean smoke -> backward is the bug - NaN persists -> forward is the bug Audit doc updated. Cargo check clean. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
a683c4bc52 |
fix(sp22): defensive NaN guard on state_121 reads in atom-shift kernels
Root cause: IEEE-754 rule 0 * NaN = NaN defeats W=0 safety. Two smokes proved the bug propagates regardless of W magnitude — pattern identical at W=[-0.5,0,+0.5,0] (train-th8pj) and W=[0,0,0,0] (train-gs4gx). The state_121 = batch_states[i * state_dim + 121] read in 4 atom-shift kernels can contain NaN (upstream source: aux head softmax producing NaN at some rollout step, stored in replay buffer, sampled into trainer's batch). 0 * NaN = NaN propagates through all atom-shift arithmetic. Defensive fix: guard each state_121 read with isfinite check, fallback to 0 if NaN/inf. Applies to: - compute_expected_q (experience_kernels.cu) - mag_concat_qdir (experience_kernels.cu) - quantile_q_select (experience_kernels.cu) - c51_loss_batched (c51_loss_kernel.cu) — state_121 AND next_state_121 - c51_aux_dw_kernel (s121 AND ns121) This unblocks Phase 3 mechanism validation. The actual state_121 NaN source (aux head or state assembly) is to be investigated separately. Cargo check clean. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
617eb61aab |
diagnostic(sp22): zero W prior + beta prior to isolate NaN source
First smoke at
|
||
|
jgrusewski
|
ff98edc774 |
fix(sp22): H6 Phase 3 beta - activate via 0.5 structural prior (B6-minimal)
Bug: state_reset_registry has FoldReset entries for sp22_reward_aux_align_ema
and sp22_aux_align_scale, but reset_named_state dispatch had no matching arms
-> unknown name error at first fold boundary with these registry entries.
Latent in
|
||
|
jgrusewski
|
5106e3b117 |
feat(sp22): H6 Phase 3 α Phase C1 — collector W ptr setter
Wires the trainer's `w_aux_to_q_dir [4]` device pointer into the GpuExperienceCollector so rollout-time action selection sees the active atom-shift for the direction branch. Changes: - gpu_experience_collector.rs: new field aux_w_to_q_dir_dev_ptr (u64, default 0) + setter set_aux_w_to_q_dir_ptr(). Both rollout launchers (compute_expected_q + quantile_q_select) now pass W ptr + exp_states_f32 + STATE_DIM_PADDED instead of NULL placeholders. NULL-safe via the kernels' existing aux_shift_active gating. - gpu_dqn_trainer.rs: w_aux_to_q_dir field promoted to pub(crate) for cross-module access via raw_ptr(). - trainers/dqn/trainer/training_loop.rs: new wire-up block after SP15 warm-count setter, mirroring the established setter pattern. NULL-safe on test scaffolds where fused_ctx or collector is absent. End-state — rollout activation: - compute_expected_q + quantile_q_select now return shifted E[Q] / quantile-blends per direction action during rollout. - With trainer's W trained each step (Step 8+11 Adam), the rollout policy's direction-action distribution actively reflects the learned aux→policy coupling. state_121's per-env value drives a per-(env, action) bias of magnitude W[a] (≤0.5 initial prior, learned thereafter). Verification: cargo check -p ml --lib clean (0 errors, 21 pre-existing warnings). Trainer + collector now smoke-ready end-to-end for the trainer/rollout side. Eval-side activation pending Phase D. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
5d163e0e1d |
feat(sp22): H6 Phase 3 α — adaptive W via dW backward + Adam (B9 Steps 8+11)
Step 8 — c51_aux_dw_kernel (new):
- Per-action block tree-reduce: grid=(4,1,1), block=(256,1,1). One block
per W index a, tree-reduces dW[a] across batch via warp shuffle + shmem.
Zero atomicAdd per pearl_no_atomicadd.
- Per-sample contributions:
a == a_d: dW[a] += inv_batch × isw × (SP_b/dz) × state_121
a == a*: dW[a] += inv_batch × isw × (-γ(1-done)) × (SP_b/dz) × next_state_121
c51_loss_kernel forward — new scratch outputs:
- aux_target_a_dir_buf[B] (i32): saves best_next_a for d==0 after Step c
sampling.
- aux_proj_logdiff_dir_buf[B] (f32): saves SP_b = Σ_n p_target_n ×
(current_lp[upper_n] - current_lp[lower_n]) after Step d's projection
via re-derivation of lower_n/upper_n (matching Huber compression +
clamp arithmetic of block_bellman_project_f).
Step 11 — adam_w_aux_kernel (new):
- Standard Adam with bias correction, grid=(1,1,1), block=(4,1,1).
- Graph-capture-safe: lr via self.lr_dev_ptr pointer arg; step via
self.ptrs.t_buf pointer arg (matches main Adam pattern). beta/eps
as value args from sp5_isv_slots constants.
- Bias-correction denominator floored at 1e-30 to avoid /0.
Trainer wiring (submit_adam_ops):
- launch_c51_aux_dw + launch_adam_w_aux added right after
launch_adam_update. Both inside the captured adam_child graph.
- New trainer fields: aux_target_a_dir_buf, aux_proj_logdiff_dir_buf,
c51_aux_dw_kernel, adam_w_aux_kernel. Cubin statics SP22_C51_AUX_DW_CUBIN
+ SP22_ADAM_W_AUX_CUBIN added.
NULL-safety:
- aux_shift_active=false in c51_loss_kernel forward → both scratch
buffers stay at alloc_zeros 0 → dW reads 0 → Adam W is a no-op.
- aux_target_a_dir_out / aux_proj_logdiff_dir_out are NULL-tolerant.
Deferred (deliberate scope):
- dL/dstate_121 backward (c51 → aux head): refinement, not correctness;
aux head trains via own supervised CE loss.
- Phase C1 collector W ptr setter.
- Phase D (eval-side aux infrastructure).
Verification:
- cargo build -p ml --lib: 0 errors, 21 pre-existing warnings.
- nvcc full recompile clean (1m05s for sm_89 target).
- All forward atom-shift consumers + adaptive W backward + Adam now wired.
End-state: adaptive W trains from structural prior [-0.5, 0, +0.5, 0]
via projection log-diff gradient. Aux head trains independently via
supervised CE. Together they form learned cross-coupling from aux
direction predictions to dir-branch Q distribution shifts. Smoke can
now measure adaptive W's effect on WR.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
a98f299823 |
feat(sp22): H6 Phase 3 α — atom-shift forward-side complete (B9 Steps 7+9+10)
Steps 7+9+10 land all FORWARD-path consumers of atom positions:
Step 7 — c51_loss_kernel atom-shift threading:
- 5 new args (w_aux, batch_states, next_batch_states, aux_dir_prob_index,
state_dim). NULL-safe — any NULL collapses to 0 shift, bit-identical
to pre-Phase-3-α.
- Per-sample state_121 + next_state_121 hoisted once at top of kernel.
- eq_per_action[a] += W[a]*next_state_121 in dir branch (d==0) before
Expected SARSA target-action sampling. Biases a* toward aux-aligned.
- Bellman projection: effective_reward = reward + γ*Δ_target*(1-done)
- Δ_online substituted for reward in block_bellman_project_f call.
Math derivation: see prior commit
|
||
|
jgrusewski
|
195c051e7a |
feat(sp22): H6 Phase 3 α — atom-shift wired through compute_expected_q (B9 Step 5+6)
W structural prior init kernel (Step 5): - aux_w_prior_init_kernel.cu: 4-thread one-time init writing W[a] = [-0.5, 0.0, +0.5, 0.0] (Short / Hold / Long / Flat). Hardcoded values in kernel — no HtoD per feedback_no_htod_htoh_only_mapped_pinned.md. - build.rs registration + gpu_dqn_trainer.rs cubin static + handle field + new() launch after alloc_zeros for w_aux_to_q_dir. compute_expected_q atom-shift (Step 6): - experience_kernels.cu signature grows 4 args (w_aux, batch_states, aux_dir_prob_index, state_dim). NULL-safe — collapses to 0 shift when either pointer is NULL, bit-identical to pre-Phase-3-α. - Per-action inner loop computes aux_atom_shift = w_aux[a] * state_121 once per (b, a) for d==0 only. Inner z-loop applies z_val += aux_atom_shift before all S/TZ/TZ²/TLM accumulators consume it. Launcher updates (3 trainer sites + 1 collector site): - populate_q_out: passes W + current states (online path). - replay_forward_for_q_values: passes W + current states (online replay). - compute_denoise_target_q: passes W + next_states (target on s'; W shared across online/target). - gpu_experience_collector.rs: passes NULL W + NULL states until Phase C1 wires the trainer's W ptr through a setter. Architectural notes: - Action selection (every compute_expected_q call) now uses shifted atom positions for direction branch (d==0). Other branches stay bit-identical (shift = 0). Step 7 (c51_loss_kernel) + Step 8 (c51_grad_kernel) will close the loop on training loss + W gradient in the same atomic commit to avoid the gradient mismatch trap per feedback_no_partial_refactor.md. - Adam wireup for w_aux_to_q_dir lands at Step 11; until then W stays at structural prior values (no Adam step modifies it). Verification: - cargo check -p ml --lib: 0 errors, 21 pre-existing warnings (Phase 3b baseline parity). - Audit doc updated with B9 Step 5+6 checkpoint entry. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
08fd5803c4 |
refactor(sp22): H6 Phase 3 α cleanup + runbook revision for atom-shift design
Same-session cleanup of Phase 3b scalar-bias residue (commit cb80b74ce's
additions that became architecturally redundant under the 2026-05-13
atom-shift revision). The scalar-bias α design is mathematically
ineffective in C51 distributional Q-learning (softmax-shift-invariance
→ W gradient = 0 → never trains). The revised design threads
`aux_atom_shift[b, a] = W[a] * state_121[b]` through compute_expected_q
+ c51_loss_kernel + c51_grad_kernel + mag_concat_qdir; dW + dstate
gradients integrate into c51_grad_kernel's projection backward.
Files
─────
- crates/ml/build.rs:
Removed `aux_to_q_dir_bias_kernel.cu` + `aux_to_q_dir_bias_backward_kernel.cu`
from kernels_with_common. Replaced with comment block documenting
C51 softmax-invariance reason + spec/audit pointers.
- crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs:
Removed `pub(crate) static SP22_AUX_TO_Q_DIR_BIAS_CUBIN` and
`SP22_AUX_TO_Q_DIR_BIAS_BWD_CUBIN` static declarations.
Removed 3 struct fields (aux_to_q_dir_bias_kernel,
aux_to_q_dir_bias_backward_dw_kernel,
aux_to_q_dir_bias_backward_dstate_kernel) and their new()
loading blocks + struct construction entries.
KEPT: w_aux_to_q_dir + adam_m_w_aux + adam_v_w_aux + dw_aux_buf
(still needed for atom-shift Adam-trained W).
Updated W doc-comment to describe atom-shift threading (was
scalar-bias) and structural-prior initialization `[-0.5, 0.0,
+0.5, 0.0]`.
- crates/ml/src/cuda_pipeline/aux_to_q_dir_bias_kernel.cu +
crates/ml/src/cuda_pipeline/aux_to_q_dir_bias_backward_kernel.cu:
Source files stay on disk as committed dead code (commit
|
||
|
jgrusewski
|
cb80b74ce9 |
feat(sp22): H6 Phase 3b — α infrastructure loaded (NOT yet wired)
Phase 3b builds on Phase 3a ( |