334b496647
Closes the deferred-consumer gap left by Phase 3.1
(r_quality_discipline_split_kernel, commit 5d36f3238), Phase 3.3
(dd_penalty_kernel), and Phase 3.5.2 (dd_asymmetric_reward_kernel) —
three SP15 reward-axis kernels that landed only as standalone scalar
producers awaiting deferred consumer wiring per
feedback_no_partial_refactor.md.
Wave 2 chooses Option β (layered, composable post-modifier) over
Option α (replace SP11 entirely): SP11 B1b stays canonical
"trader-quality" composer that writes out_rewards; the SP15 reward-axis
composition becomes a fused PER-(i,t) post-modifier read-modify-writing
the same buffer in place. This preserves SP11's z-score mag-ratio
contract (canary tests untouched) while landing all three deferred SP15
consumers atomically with zero parallel paths.
Architecture (Q1-Q4 user resolutions):
Q1: SP12 caps stay as state_layout.cuh macros (REWARD_NEG_CAP=-10,
REWARD_POS_CAP=+5), NOT lifted to ISV — spec'd constants per the
SP12 v3 design, not adaptive bounds.
Q2: Both on-policy + CF slots get the same DD-aware shaping. CF reward
= w_cf × r_cf (SP11 controller weight already applied) is composed
via the same helpers as on-policy. DD context is per-step, not
per-action.
Q3: New slot_completed_normally[N*L] flag preserves SP11's early-return
semantics (data-end at experience_kernels.cu:2142 → reward=0.0;
blown-account at :2203 → reward=-10.0). Fused kernel skips slots
where flag==0.
Q4: alpha_split_producer_kernel OWNS the warm-count increment (per-step
scalar producer; the fused kernel has 2*N*L threads and would
over-tick by that factor if it owned the increment).
Per-step launch sequence in gpu_experience_collector.rs (after Phase
1.3.b's dd_state launch): experience_env_step → alpha_split_producer
(reads grad-norm slots 418/419, writes ALPHA_SPLIT slot 417, increments
warm-count) → compute_sp15_final_reward (fused per-(i,t) over [N*2*L]
slots, applies α-blend → DD asymmetric → DD penalty → SP12 cap, writes
back to out_rewards in place).
experience_env_step signature change — 2 new output params:
r_discipline_out: float* [N*L] — per-step REGRET_EMA mirror, written
at end of normal reward composition.
slot_completed_normally_out: int* [N*L] — 0 default at entry, set to
1 only on the path that reaches out_rewards[out_off] = reward.
3 deleted kernel files:
- r_quality_discipline_split_kernel.cu (composer + producer; replaced
by renamed alpha_split_producer_kernel.cu keeping ONLY the producer
with the moved warm-count increment).
- dd_penalty_kernel.cu (replaced by sp15_dd_penalty __device__ helper).
- dd_asymmetric_reward_kernel.cu (replaced by sp15_dd_asymmetric_reward
helper).
3 new files:
- alpha_split_producer_kernel.cu (per-step scalar producer of α from
grad-norm ratio, with warm-count increment moved here per Q4).
- sp15_reward_axis_helpers.cuh (4 __device__ inline helpers:
sp15_alpha_blend, sp15_dd_asymmetric_reward, sp15_dd_penalty,
sp15_apply_sp12_cap).
- compute_sp15_final_reward_kernel.cu (fused per-(i,t) parallel over
[N*2*L] slots — α-blend + DD-asymmetric + DD-penalty + SP12 cap,
skips early-return sentinel slots).
3 deleted launchers + 3 deleted CUBIN statics in gpu_dqn_trainer.rs:
- launch_sp15_r_quality_discipline_split (composer scalar variant).
- launch_sp15_dd_penalty.
- launch_sp15_dd_asymmetric_reward.
- SP15_R_QUALITY_DISCIPLINE_SPLIT_CUBIN.
- SP15_DD_PENALTY_CUBIN.
- SP15_DD_ASYMMETRIC_REWARD_CUBIN.
2 new launchers + 2 new CUBIN statics:
- launch_sp15_final_reward + SP15_FINAL_REWARD_CUBIN.
- SP15_ALPHA_SPLIT_PRODUCER_CUBIN (the retained launch_sp15_alpha_split_producer
now loads this).
GpuExperienceCollector: 2 new CudaSlice fields
(r_discipline_per_sample, slot_completed_normally_per_sample) +
sp15_alpha_warm_count_dev_ptr field + set_sp15_alpha_warm_count_ptr
setter + Step 5b launch block.
State reset registry — no new entries: r_discipline +
slot_completed_normally are per-step ephemeral (defaulted at every
kernel entry); cross-fold leakage impossible. Existing Phase 3.1/3.3/
3.5.2 ISV slot entries cover the rest.
6 new oracle tests in sp15_phase1_oracle_tests.rs::mod gpu drive
compute_sp15_final_reward_kernel directly with hand-crafted buffers:
- final_reward_alpha_blend_at_cold_start (Stage 1)
- final_reward_dd_penalty_above_threshold (Stage 3)
- final_reward_dd_asymmetric_gain (Stage 2 + R_GAIN_DD_BOOST diag)
- final_reward_dd_asymmetric_loss (Stage 2 asymmetric guard)
- final_reward_sp12_cap_clamps_both_directions (Stage 4)
- final_reward_skips_early_return_slots (Q3 sentinel preservation)
9 deleted scalar-kernel oracle tests:
- r_split_uses_sentinel_alpha_at_cold_start
- r_quality_subtracts_explicit_cost
- dd_penalty_quadratic_above_threshold + dd_penalty_zero_below_threshold
- dd_asymmetric_reward_gain_amplified_by_dd_pct +
dd_asymmetric_reward_loss_unchanged + dd_asymmetric_reward_no_op_at_ath
(the new fused-kernel tests cover the same behavioral surface
end-to-end through the in-place RMW path).
Verified: SQLX_OFFLINE=true cargo check -p ml --features cuda clean (18
unrelated warnings); all 29 SP15 phase1 oracle tests pass on RTX 3050 Ti
(includes the 6 new fused-kernel tests + 17 retained tests + 6 old);
SP11 mag-ratio canary tests still untouched (no canary-test renames or
deletions); ml lib suite holds 946 pass / 13 fail = baseline.
Hard rules: feedback_no_partial_refactor (3 phases' deferred consumers +
3 deletions + 3 new files + 6 new tests + audit doc all in this commit;
no parallel paths, no feature flags), feedback_wire_everything_up
(closes 3 SP15 phase orphan launchers atomically), feedback_no_legacy_aliases
(deletions land in same commit as replacement; no compatibility shim),
feedback_no_atomicadd (fused kernel is per-(i,t) parallel, pure scalar
arithmetic), pearl_audit_unboundedness_for_implicit_asymmetry (gain-only
DD multiplier + asymmetric NEG/POS caps preserve loss aversion),
pearl_symmetric_clamp_audit (SP12 cap is bilateral via fmaxf/fminf even
though bounds are intentionally asymmetric per spec),
pearl_no_host_branches_in_captured_graph (new launches happen inside
collect_experiences_gpu::launch_timestep_loop per-step, OUTSIDE the
experience-fwd CUDA Graph capture region — same precedent as Phase
1.3.b's dd_state launch).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
ml
10-model ML ensemble for the Foxhunt HFT system, built on Candle v0.9.1.
Models
- DQN (Rainbow) — deep Q-network with prioritized replay, dueling heads, noisy nets
- PPO — proximal policy optimization with GAE, LSTM policies, clip-higher
- TFT — temporal fusion transformer for multi-horizon forecasting
- Mamba2 — state space model for sequence prediction
- Liquid Networks — biologically inspired networks for non-stationary data
- TLOB — transformer-based limit order book analysis
- KAN — Kolmogorov-Arnold networks
- xLSTM — extended LSTM architecture
- TGGN — temporal graph neural network
- Diffusion — diffusion-based generative model
Key Modules
ensemble— model ensemble coordination and confidence aggregationhyperopt— PSO-based hyperparameter optimization with per-model adapterstrainers— unified training loops (DQN, PPO, supervised)inference—InferenceAdaptertrait for predictioncheckpoint— model checkpointing and restorationevaluation— walk-forward evaluation pipeline
Usage
use ml::dqn::DQN;
use ml::ppo::PpoTrainer;