996e61b6eff10b8a9d32a13e7e92fb8cc39fef29
5227 Commits
| Author | SHA1 | Message | Date | |
|---|---|---|---|---|
jgrusewski
|
996e61b6ef |
rename(ml-alpha): MultiHorizonAttention (no version suffix in identifiers)
Removes the version suffix from the trainer-state bundle introduced
in
|
||
|
jgrusewski
|
5aad1eb846 |
feat(ml-alpha): PerceptionV2State bundle (axes A+B+C+D+E) [V9]
Bundles all v2 components into one cohesive trainer field that
PerceptionTrainer wires through in V10. Owns:
(C) horizon_tokens + Q for the horizon-token attention pool
+ per-batch grad scratches + 2 AdamWs
(E) inverted attention pool + saved attn + d_scores scratch
(D) MoE gate weights + experts (W, b) + per-batch sparse-by-expert
grad scratches + 3 AdamWs + top_e / gate_probs / aux_loss
(A) log_sigma_h per-horizon Kendall scalar + AdamW (0.25× LR)
(B) AnchorController (Wiener-α host-side) + anchor_l2 kernel +
init snapshots of horizon_tokens, Q, experts_w, experts_b
for the L2 anchor
Init scale: 1/√HIDDEN_DIM Xavier for horizon_tokens/Q/experts_W;
zeros for experts_b and log_sigma_h. Anchor controller bootstrap
uses ‖horizon_tokens_init‖₂ + ‖Q_init‖₂ + their total numel to
derive the signal-driven floor (no tuned constants).
zero_grads uses memset_zeros only — capture-safe.
Compile-clean. V10 wires this state into perception.rs::step_batched
forward / backward / AdamW.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
55aeddaebd |
feat(ml-alpha): anchor_l2 kernel + Wiener-α controller (v2 B) [V8]
L2 anchor regularization toward initialization (axis B). Anchors
horizon_tokens + Q + MoE experts toward their init values to prevent
the calibration drift observed in v1 (where val_loss climbed as α
opened past epoch 1 in 2 of 3 folds).
KERNEL (`anchor_l2_fwd_bwd`):
loss_out = λ · Σ_i (p[i] − p_init[i])²
grad_p[i] += 2λ · (p[i] − p_init[i])
- Warp-shuffle reduce; one block per parameter group; strided thread
loop over n. Cross-warp reduce uses one __syncthreads.
- Coalesced grad write via stride loop.
- λ passed as device-side [1]-buffer (host writes scalar before launch
— capture-safe).
CONTROLLER (`trainer::anchor_controller::AnchorController`):
- Signal-driven λ floor: λ_floor = ‖p_init‖ / (100 · √numel).
Cross-fold-persistent per pearl_kelly_cap_signal_driven_floors.
- Wiener-α smoother (α = diff_var / (diff_var + sample_var + ε))
on val_loss change; α floored at 0.4 per
pearl_wiener_alpha_floor_for_nonstationary.
- λ blends toward target = |ema_change|·scale with α; floored at
λ_floor per pearl_blend_formulas_must_have_permanent_floor.
- First-observation bootstrap (sentinel state replaced directly on
first step) per pearl_first_observation_bootstrap.
- 4/4 unit tests PASS: signal-floor init, bootstrap returns floor,
floor protection across 1000 steps, λ_max cap.
NUMGRAD VERIFICATION (RTX 3050 sm_86):
anchor_l2_numgrad PASSES with closed-form parity (machine precision)
and central-difference parity (4 random positions) within 5e-2 rel.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
11b96dac6a |
feat(ml-alpha): regime_moe_gate kernel + numgrad (v2 D) [V6]
Top-1 Mixture-of-Experts gate per Switch Transformer. Three kernels
in one .cu file:
- regime_moe_gate_fwd: select top-1 expert from gate_logits, apply
its [H, H] linear to fused_ctx → routed_ctx [B, N_H, H].
- regime_moe_gate_bwd: chain-rule grads through the SELECTED
expert's W and bias (sparse-by-expert scratches), plus
d_fused_ctx accumulator. Inactive experts get zero contribution.
- regime_moe_gate_aux: softmax of gate_logits + load-balancing
auxiliary loss (frac · prob_mean × N_EXPERTS).
ARCHITECTURE:
- N_EXPERTS = 4. Each expert is a [H, H] linear with bias.
- Total expert params: 4 · 128 · 128 + 4 · 128 = 66 KB. Cheap.
- STE on gate: gate logit grad is zero from the expert path (top-1
is non-differentiable); the load-balance aux loss provides the
differentiable signal that pushes routing toward balanced usage.
PERFORMANCE:
- Grid (B, N_H, 1), block (HIDDEN_DIM). One block per (b, h).
- Forward: each thread computes one output channel via a dot
product over HIDDEN_DIM input dims (#pragma unroll 8).
- Backward d_fused_ctx: each thread accumulates over d_out
sequentially (HIDDEN_DIM iterations) since the weight matrix
column is naturally aligned to the thread's d_in index.
- Backward d_W/d_b scratches are sparse-by-expert; downstream
reduce_axis0 collapses over (B, N_H).
- Top-1 chosen by thread 0 per block, broadcast via shared mem.
NUMGRAD VERIFICATION (RTX 3050 sm_86):
forward_matches_host_reference_and_backward_matches_numgrad
PASSES 11 checks (4 on d_W, 3 on d_b, 4 on d_fused_ctx) within
5e-2 rel / 5e-3 abs.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
d94696620d |
feat(ml-alpha): inverted_attention_pool kernel + numgrad (v2 E) [V3]
iTransformer-style cross-variate attention pass: each of HIDDEN_DIM
features becomes a "variate token" with its K-trajectory as embedding.
FORWARD:
X_inv[h, k] = ln_out[k, h] # transpose
scores[h, j] = inv_scale · Σ_k X_inv[h, k] · X_inv[j, k]
attn[h, j] = softmax_j(scores)
pooled[h] = Σ_j attn[h, j] · mean_k_X_inv[j] # mean-K commutes out
BACKWARD: three independent chains into ln_out:
- value path: (1/K) · attn[h', my_h] · d_pooled[h'] (per-k constant)
- query path: inv_scale · Σ_j d_scores[my_h, j] · X_inv[j, k]
- key path: inv_scale · Σ_h' d_scores[h', my_h] · X_inv[h', k]
Softmax bwd: d_scores[h, j] = attn · (d_attn - Σ_l attn · d_attn)
IMPLEMENTATION NOTES:
- First attempt cached attn [H, H] = 64 KB in shared mem → tripped
the 48 KB dynamic-shared limit on sm_86 (CUDA_ERROR_INVALID_VALUE).
- Fixed by moving d_scores to a DRAM scratch buffer; shared mem
holds only X_inv [H, K] (≤ 16 KB at K = 32). One block-wide barrier
between the d_scores write and the value/query/key accumulation.
- All per-batch slice writes; no atomicAdd, no cross-block race.
- Pooled computation uses the mean-K commute (Σ_k attn · X_inv =
attn · mean_k_X_inv), saving an entire H×K accumulation pass.
LOCAL VERIFICATION (RTX 3050 sm_86):
forward_then_backward_matches_central_difference PASSES 6 numgrad
checks on ln_out positions within 5e-2 rel / 5e-3 abs.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
e8f6c4721f |
feat(ml-alpha): horizon_token_attention_pool kernel + numgrad (v2 C) [V2]
Replaces the falsified per-horizon Q_h pool with a single shared
query Q over an extended key sequence [horizon_tokens; LN_b_out],
producing per-horizon outputs via TFT-style horizon-token mixing.
FORWARD:
scores[i] = Σ_d Q[d] · ext[i, d] (i ∈ [0, N_H + K))
attn = softmax_i(scores)
S[d] = Σ_k attn[N_H + k] · ln_out[k, d] (shared time agg)
ctx[h, d] = attn[h] · horizon_tokens[h, d] + S[d] (per-horizon out)
BACKWARD: full chain rule with softmax-centring; gradients to
horizon_tokens, Q, and ln_out via the saved attn weights.
NVIDIA-grade implementation per feedback_nvidia_grade_perf_for_kernels:
- Warp-shuffle reduce (block_reduce_sum / block_reduce_max helpers)
for all per-d dot products and softmax aggregates.
- Cross-warp reduce uses exactly one __syncthreads.
- Non-divergent shuffles: inactive lanes contribute 0 via ternary.
- Block-per-batch + horizon-loop inside block → grad_ln_out += is
race-free without atomicAdd.
- Smem layout computed at launch: [s_attn(N_H+K); s_warp(N_WARPS);
s_d_S(H) on bwd]. No over-allocation.
LOCAL VERIFICATION (RTX 3050 sm_86):
forward_then_backward_matches_central_difference PASSES 12 numgrad
checks (4 each on horizon_tokens / Q / ln_out) at 5e-2 rel / 5e-3
abs envelope. First-try pass.
NOTE: .gitignore adjusted with narrow allow-rules for crates/ml-alpha/{
cuda,src,tests}/horizon_token_* paths — the broad "*token*" rule
intended for auth tokens was hiding these source files. Explicit
allow keeps the security rule intact while exempting these specific
files.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
679ab3f5eb |
feat(ml-alpha): Kendall sigma-weighted BCE kernel + numgrad (v2 A) [V7]
New `bce_multi_horizon_sigma_forward_backward` kernel implementing the
Kendall homoscedastic uncertainty weighting per spec axis A:
raw_bce_h = Σ_{i in h, m_i=1} L_i
count_h = #{i in h : m_i = 1}
mean_bce_h = raw_bce_h / count_h
w_h = base_weight_h / (2 · exp(2 · log_sigma_h))
total_loss = Σ_h [ w_h · mean_bce_h + log_sigma_h ]
d L / d p_i = m_i · (w_h / count_h) · (p − y) / (p (1 − p))
d L / d log_sigma_h = 1 − 2 · w_h · mean_bce_h
NVIDIA-grade implementation per feedback_nvidia_grade_perf_for_kernels:
- Warp-shuffle reduction (`__shfl_xor_sync`) for both per-horizon
sums and the global valid count, replacing block tree-reduce.
- One `__syncthreads` for the cross-warp aggregate; no inner-loop
barriers.
- Non-divergent shuffles: inactive lanes contribute 0 via ternary,
never via `if (tid < N) shuffle`.
- Coalesced strided access in both forward and gradient passes.
- Pre-compiled cubin via build.rs; no nvrtc.
Independent of the legacy `bce_loss_multi_horizon` kernel — that one
stays untouched so eval/smoke paths are unaffected. The v2 trainer
wires this kernel in via commit V10.
Standalone helper `bce_sigma_loss_and_grad_gpu` in `trainer::loss_sigma`
for numgrad parity tests. Three numgrad tests all PASS on RTX 3050
(sm_86) within 5e-2 rel / 5e-3 abs:
- d_log_sigma_h ↔ central-difference (numgrad on log_sigma)
- grad_probs ↔ central-difference (8 random positions)
- total_loss ↔ closed-form reconstructed from mean_bce_per_h
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
41292303dc |
refactor(ml-alpha): remove per-horizon Q_h path (C21-C25 falsified) [V1]
3-fold A/B sweep 2026-05-18 at commit
|
||
|
jgrusewski
|
4425a77844 |
plan(ml-alpha): v2 multi-horizon implementation plan (V1-V13)
Concrete TDD-driven commit map for the v2 spec (2026-05-18-ml-alpha-v2-multi-horizon-design.md). Thirteen commits ordered by dependency: delete falsified path, build new kernels with numgrad parity (V2-V8), trainer state bundle (V9), wire into PerceptionTrainer (V10), local smoke (V11), cluster smoke (V12), 3-fold A/B (V13). Per-commit verification gates and explicit stop conditions. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
1c2ee1d7c3 |
spec(ml-alpha): v2 multi-horizon redesign (A+B+C+D+E integrated)
Integrated design spec for the post-A/B redesign: Kendall sigma BCE (A), L2 anchor on horizon tokens + shared Q (B), horizon-token K-prepend replacing per-horizon Q_h (C), regime-aware MoE gate (D), and inverted-axis attention pass (E). Bundled per pearl_no_deferrals_for_complementary_fixes — all five axes have orthogonal architectural scope and independent kill criteria. Spec drops the C21-C25 per-horizon Q_h path (falsified by sweep 2026-05-18: mean_auc -0.019 vs single-Q baseline) and migrates the existing init buffers into the new horizon-tokens prefix. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
286ea26e2a |
perf(ml-alpha): warp-shuffle reduce in per-horizon kernels
Cluster A/B sweep with C25 wiring showed 86 s/epoch vs 17 s/epoch baseline = ~5x regression. Root cause: per-horizon attention pool + residual head used block tree-reduce with 8 __syncthreads per K-step in a serialised K-loop, repeated H=5 times in both fwd and bwd = ~3200 barriers/step. Plus the prob_blend bwd reduce kernel ran with a single thread per block, fully serialising over K*B. Replacements: - per_horizon_attention_pool fwd/bwd: introduce block_reduce_sum / block_reduce_max helpers using intra-warp __shfl_xor_sync + cross-warp shuffle (1 syncthread per K instead of 8). Smem shrinks to [K + N_WARPS] / [2K + N_WARPS]. - per_horizon_residual_head fwd: same warp-shuffle reduce pattern. - per_horizon_prob_blend_reduce_alpha_residual: 1 thread → 1 warp per horizon, lane-strided reduction over K*B via shfl_xor_sync. Launch config updated to block_dim=(32,1,1). Tricky bug found while implementing: the cross-warp reduce in the residual head originally guarded `__shfl_xor_sync(0xffffffff, ...)` with `if (tid < PHR_N_WARPS)`, leaving 28 of 32 lanes in warp 0 outside the call. Mask 0xffffffff requires all 32 lanes to participate — divergence is UB and hung the full-pipeline smoke on Ampere/Ada. Fix: read s_warp via ternary into all 32 lanes, then shuffle inside `if (tid < 32)`. Matches the pattern used in block_reduce_sum. Verified locally on RTX 3050 (sm_86): per_horizon_attention_pool numgrad, per_horizon_residual_head numgrad, and per_horizon_full_pipeline_smoke (zero-init identity + non-zero end-to-end) all PASS. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
83546b5c37 |
fix(ml-alpha): drop synchronize() in per-horizon pool + head bindings
After
|
||
|
jgrusewski
|
adaf275af3 |
fix(ml-alpha): C25 per-horizon path graph-capture safety
Four code paths in PerHorizonTrainState violated CUDA Graph capture invariants and tripped CUDA_ERROR_STREAM_CAPTURE_UNSUPPORTED on smoke alpha-perception-qk2p9: 1. stream.synchronize() inside forward_with_blend (illegal in capture) 2. stream.synchronize() inside backward_through_blend (idem) 3. reduce_per_batch_scratches_to_shared used host-side vec allocations + memcpy_dtoh + CPU summation + memcpy_htod (forbidden during capture per pearl_no_host_branches_in_captured_graph) 4. zero_grads allocated host zero vectors + memcpy_htod each step (idem) Fix: - Remove both synchronizes; same-stream issue order is sufficient. - Replace host-side reduction with three reduce_axis0 GPU kernel launches (q_h, w_res, bias_res). PerHorizonTrainState now owns its own reduce_axis0 cubin handle. - Replace host-zero memcpy with stream.memset_zeros for all nine gradient buffers plus d_alpha_reduced. Verified locally: per_horizon_full_pipeline_smoke (2/2) and both numgrad parity tests (attention pool + residual head) pass on RTX 3050. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
f50b466f77 |
feat(ml-alpha): PerceptionTrainer wires C21+C22 fwd+bwd+AdamW (C25)
The per-horizon attention pool from C21+C22 is now fully integrated
into step_batched's hot loop. Forward and backward both flow; AdamW
updates all four new parameter groups (Q_h, w_res, bias_res, α) every
training step. At α=0 init the contribution is byte-identical to
baseline; training discovers whether α should grow.
New kernel: cuda/per_horizon_prob_blend.cu
per_horizon_prob_blend_fwd
Reads logit_per_k_d (already stored by GRN forward) +
sigmoid(logit_baseline + tanh(α[h]) * residual[b, h]) →
overwrites probs_per_k_d in place. At α=0: r_contrib=0, output
== sigmoid(logit_baseline) == probs_baseline → bit-identical.
per_horizon_prob_blend_reduce_alpha_residual
Reads probs_per_k (= p_final, post-blend) + grad_probs_per_k
(= ∂L/∂p_final from BCE) and computes:
d_logit[k,b,h] = grad_probs[k,b,h] * p_final * (1 - p_final)
d_residual[b,h] = tanh(α[h]) * Σ_k d_logit[k,b,h]
d_alpha[h] = sech²(α[h]) * Σ_{k,b} d_logit[k,b,h] * residual[b,h]
No separate prob_blend_bwd needed — chain-rule equivalence
∂L/∂logit_baseline = ∂L/∂r_contrib (both flow through the same
sigmoid derivative) means the existing GRN backward is UNCHANGED.
trainer/per_horizon_state.rs extensions:
forward_with_blend(ln_b_out, logit_per_k, probs_per_k)
Pool fwd → context_h; head fwd → residual; prob_blend fwd
in-place rewrites probs_per_k.
backward_through_blend(probs_per_k, grad_probs_per_k, ln_b_out,
grad_ln_b_out_target)
Reduce kernel → d_residual + d_alpha. Then:
head bwd → d_w_res_scratch, d_bias_res_scratch, d_context.
pool bwd → d_q_h_scratch, += grad_h_enriched_seq_d.
Per-batch scratches reduced to shared grads host-side
(n_batch ≤ 64 → sub-millisecond on host).
adamw_step()
Steps the four optimizers using the shared grad buffers.
zero_grads()
Called once per step before forward to clear scratch.
trainer/perception.rs step_batched integration:
── 4.5 (after GRN K-loop, before BCE): zero_grads + forward_with_blend
overwrites probs_per_k_d with p_final.
── 5 (existing BCE consumes probs_per_k_d as today; grad_probs is
now ∂L/∂p_final automatically).
── 5a (after BCE, before ISV-lambda + heads bwd): backward_through_blend.
Existing GRN bwd path is UNTOUCHED — the chain rule absorbs
the bias.
── 9 (after existing 17 AdamW group steps): per_horizon.adamw_step
updates Q_h, w_res, bias_res, α.
Verification:
- 34 ml-alpha lib tests still green.
- Per-horizon kernel numgrad parity (C21, C22) still green.
- Per-horizon end-to-end pipeline smoke (C23, including the
alpha=0 byte-identity invariant) still green.
- Full workspace builds clean.
Closes the kernels+wiring portion of #203 (per-horizon attention pool
kernels + wiring). What remains (#204): 30-epoch × 3-fold A/B vs
single-Q baseline. The branch is ready for that sweep when GPU time
is budgeted; the implementation is structurally adoption-safe
(α=0 → identity to baseline) so it can be merged before the A/B if
desired.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
5eb6567c4c |
feat(ml-alpha): PerceptionTrainer per-horizon trainer state (C24)
Allocates the device buffers + AdamW optimizers + kernel bindings for
the per-horizon attention-pool training path from C21 + C22, bundled
into a single PerceptionTrainer.per_horizon field.
crates/ml-alpha/src/trainer/per_horizon_state.rs (NEW):
PerHorizonTrainState — owns:
Learnable params (zero-init for α + bias, Xavier-scale for Q_h,
0.1× scale for w_res):
q_h_d [N_HORIZONS, HIDDEN_DIM] — attention queries
w_res_d [N_HORIZONS, HIDDEN_DIM] — residual head weights
bias_res_d [N_HORIZONS] — residual bias
alpha_d [N_HORIZONS] — learnable gate (init 0)
Forward intermediates (per-batch):
context_d [B, N_HORIZONS, HIDDEN_DIM]
attn_weights_d [B, N_HORIZONS, K]
residual_d [B, N_HORIZONS]
Backward grad scratch + reduced grads + AdamW state.
Kernel bindings: PerHorizonAttentionPool + PerHorizonResidualHead.
PerHorizonTrainState::new(dev, n_batch, k_seq, lr, seed)
Allocates all buffers, runs Xavier-style init, constructs four
AdamW optimizers (q_h, w_res, bias_res at param-LR; α at 0.25× LR
per spec §5 open Q2 default — slow gate ramp). Captures the seed
via wrapping_add(0xA110C00A) from cfg.seed for determinism.
PerHorizonTrainState::zero_grads()
Clears all grad-scratch buffers between training steps.
trainer/mod.rs:
pub mod per_horizon_state — module export.
trainer/perception.rs:
PerceptionTrainer gains one field:
pub per_horizon: PerHorizonTrainState
Initialised in new() with cfg.n_batch + cfg.seq_len + cfg.lr_cfc.
α-gate init=0 ⇒ tanh(0)=0 ⇒ contribution to per-batch logits is exactly
0 at step 0 (per C23 alpha_zero_init_is_identity_to_baseline byte-equality
test). Adopting this commit produces bit-identical training behaviour
to the previous commit until C25 wires the forward+backward calls into
step_batched; even then the α=0 init means a one-epoch smoke against
existing baseline should match within FP rounding noise.
All 34 ml-alpha lib tests + 4 per-horizon GPU tests (numgrad pair +
end-to-end pipeline pair) green.
Next:
C25 — forward + backward integration into step_batched. The new path
runs as ADDITIONAL kernel launches before/after the existing
captured graph (not inside it) so the graph stays unchanged
and the integration risk is contained.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
69c64f2266 |
test(ml-alpha): per-horizon pipeline end-to-end smoke + α-gate (C23)
Composes the C21 + C22 kernels with a host-side learnable α-gate to
prove the full per-horizon contribution path works end-to-end without
yet doing the captured-graph integration in PerceptionTrainer.
Pipeline:
LNb [B, K, HIDDEN_DIM]
→ per_horizon_attention_pool_fwd → context_h [B, N_HORIZONS, HIDDEN_DIM]
→ per_horizon_residual_head_fwd → residual [B, N_HORIZONS]
→ final[b, h] = baseline[b, h] + tanh(α[h]) * residual[b, h]
Two tests cover the critical invariants for adoption-safety:
alpha_zero_init_is_identity_to_baseline
With α = [0, 0, 0, 0, 0] and any random Q_h / w_res / bias_res,
final_logit MUST be bit-identical to baseline_logit (because
tanh(0) = 0). Verified via to_bits() byte equality. Proves that
initialising the new variant with α=0 makes it a strict superset
of the existing path — switching to AttentionPoolVariant::PerHorizon
cannot regress before any training has happened.
alpha_nonzero_changes_output_and_grads_flow_end_to_end
With α = [0.5, -0.3, 0.2, -0.1, 0.4]:
* final ≠ baseline (residual contributing) ✓
* all final logits finite ✓
* full backward chain (residual_head_bwd → attention_pool_bwd)
produces finite d_Q_h_scratch + finite d_LNb with at least
one non-zero entry in each → gradients flow back to both the
attention queries and the LN_b input ✓
This closes the kernel-side correctness story. The remaining
integration commits (C24+) are operational:
C24: extend CheckpointV1 → V2 (add q_h, w_res, bias_res, alpha
fields; V1 files load as Variant::SharedQuery)
C25: PerceptionTrainer wiring — allocate the device buffers, fold
attention + residual + gate into the captured graph, plumb
gradients into AdamW's param list
C26: 1-epoch smoke (assert no NaN, loss decreases vs baseline) —
needs real training data + multi-GPU time
C27: 30-epoch × 3-fold A/B (task #204) — decision gate per
docs/superpowers/specs/2026-05-18-per-horizon-attention-pool-design.md
§0 falsifiable claim
C24-C25 are 1-2 day work even when carefully scoped; C26-C27 need
real GPU-hours + result analysis. C21-C23 land the validatable kernel
correctness piece without committing to that time investment yet.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
8c335caef7 |
feat(ml-alpha): per-horizon residual head kernel + numgrad parity (C22)
Companion kernel to C21's per_horizon_attention_pool. Computes a
per-horizon scalar residual from each horizon's context vector:
residual[b, h] = Σ_d w_res[h, d] * context_h[b, h, d] + bias_res[h]
Designed to be added (behind a learnable α-gate) to the existing
multi_horizon_heads logit output — keeps the existing GRN head kernel
completely unchanged. The per-horizon attention pool's contribution
flows through this lightweight projection without weight-shape
changes elsewhere or checkpoint-V2-bumping.
Path A integration sketch (deferred to follow-up commit C23):
alpha_logit_per_horizon = existing_head(h_K)[h] # from current path
+ tanh(α[h]) * residual_kernel(context_h)[h]
where α[h] is a learnable 5-vector init'd to 0 (no effect at start).
Training discovers per-horizon whether the residual contributes.
This is a strict superset of the existing path — α=0 → bit-identical
to today.
Backward kernel produces:
d_w_res — per-block scratch [B, N_HORIZONS, HIDDEN_DIM]
for host reduce_axis0 → shared [N_HORIZONS, HIDDEN_DIM]
d_bias_res — per-block scratch [B, N_HORIZONS], same reduction
d_context_h — per-batch indexed; += chained with attention bwd
Single-writer discipline preserved (no atomicAdd per
feedback_no_atomicadd.md); horizon loop inside the per-batch block.
Numgrad parity test:
- B=3, N_HORIZONS=5, HIDDEN_DIM=128 fixture.
- Loss = Σ residual_out (so d_residual = 1).
- Probes 8 random w_res indices, all 5 bias_res entries, 8 random
context_h indices via central-difference at ±eps=1e-2.
- All within 5e-2 rel-tol or 5e-3 abs-floor.
- Passes on RTX 3050.
Same scope discipline as C21: kernel + binding + numgrad first;
trainer wiring + α-gate + smoke training + A/B sweep follow once
both kernels are individually validated (now done).
Closes the second kernel-correctness portion of #203. Remaining:
C23: trainer wiring (capture attn_pool fwd into the graph; sum
residual into existing head output with α-gate)
C24: CheckpointV1 → V2 bump (add q_h, w_res, bias_res, alpha fields)
C25: 1-epoch smoke (assert no NaN, loss decreases vs baseline)
C26: 30-epoch × 3-fold A/B (#204) — decision gate per spec §0
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
edc449eecf |
feat(ml-alpha): per-horizon attention pool kernel + numgrad parity (C21)
First implementation slice of the per-horizon attention pool design
(docs/superpowers/specs/2026-05-18-per-horizon-attention-pool-design.md).
Lands the kernel + Rust binding + numgrad verification; downstream
wiring into PerceptionTrainer's captured graph + CheckpointV2 bump +
A/B sweep are follow-up commits gated on this proving correctness.
Kernel (cuda/per_horizon_attention_pool.cu):
per_horizon_attention_pool_fwd
Q_h[N_HORIZONS, HIDDEN_DIM] × LNb[B, K, HIDDEN_DIM]
→ context_h[B, N_HORIZONS, HIDDEN_DIM]
attn_h_weights[B, N_HORIZONS, K]
Per-block math identical to the single-Q variant, looped over
N_HORIZONS sequentially within each batch's block. Grid stays
(B, 1, 1) so backward grad_ln_out writes are race-free
(per feedback_no_atomicadd.md — no cross-block contention).
Per-batch shared mem ~k_seq + BLOCK + HIDDEN_DIM floats.
per_horizon_attention_pool_bwd
Same chain-rule pattern as attention_pool_bwd but with the horizon
loop inside the block: each (b, h) slice updates grad_ln_out in
place (sequential horizon accumulation), grad_Q_h is written as
per-block scratch [B, N_HORIZONS, HIDDEN_DIM] for host reduce.
Single-writer discipline preserved.
Rust binding (src/per_horizon_attention_pool.rs):
PerHorizonAttentionPool::{new, forward, backward}. Self-contained;
doesn't yet touch PerceptionTrainer or CfcTrunk. Loads the cubin
via the standard env!("OUT_DIR") path. Dynamic shared-mem byte
count computed per launch from k_seq.
Numgrad parity test (tests/per_horizon_attention_pool_numgrad.rs):
- B=2, K=8, HIDDEN_DIM=128, N_HORIZONS=5 fixture.
- Loss = Σ context_h (so d_context = 1 everywhere — clean analytical).
- Backward kernel produces analytical grads; central-difference of
forward kernel at ±eps=1e-2 across 8 random Q_h indices + 8 random
LNb indices verifies analytical matches CD within 5e-2 rel-tol or
5e-3 abs-floor.
- Passes on RTX 3050.
build.rs picks up the new .cu file automatically via the existing
KERNELS list; cubin compiles cleanly at sm_86 + sm_89.
Same scope discipline as Phase 2D.2 (VSN numgrad) — kernel correctness
first, integration second. The follow-up commit set per the spec §3
appendix:
C22: extend multi_horizon_heads.cu signature to accept per-horizon
context input + bump head_w shape to [N_HORIZONS, 2*HIDDEN_DIM]
C23: wire PerHorizonAttentionPool into PerceptionTrainer + CfcTrunk
captured graph behind AttentionPoolVariant config flag
C24: CheckpointV1 → V2 bump with discriminant + optional q_h field
C25: smoke training (one epoch, no NaN, loss decreases)
C26: 30-epoch × 3-fold A/B sweep (#204) — decision gate per spec §0
Closes the kernel-correctness portion of #203.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
f5632649ca |
spec(ml-alpha): per-horizon attention pool design (C20)
Captures the brainstormed "alternative attention pool variants"
follow-on from the original real-LOB integration brainstorm (Axis 1,
deferred from the LOB workstream as a separate model-side spec).
Design:
Replace shared learned query Q[HIDDEN_DIM] with per-horizon queries
Q_h[N_HORIZONS, HIDDEN_DIM]. Per-horizon softmax + context vectors
feed multi-horizon heads directly (PATH A) — each horizon attends
to a different part of the K=6000 LN_b output sequence. CfC k=0
state is initialised by the MEAN of per-horizon contexts so the
K-loop recurrence + state amplification (per
pearl_state_amplifies_short_horizon_into_long_horizon) survives.
Heads consume per-horizon context concat CfC h_K (residual) with a
default 75/25 weight split.
Falsifiable claim (§0): A/B-tested win means h6000 mean_auc lifts by
≥ +0.01 absolute OR per-horizon distribution shifts toward short
horizons (h1000, h300) with no net h6000 loss. The 3-fold variance
band on the current architecture (mean_auc 0.7749 ± 0.024) means a
+0.01 lift is within noise — a meaningful effect needs ≥ +0.024 or
qualitative distributional shift.
Two new kernels (per_horizon_attention_pool_fwd + _bwd) + signature
extension on multi_horizon_heads_{fwd,bwd}. Variant-toggle config flag
(SharedQuery vs PerHorizonQuery) keeps the existing path fully
functional; new variant is opt-in. CheckpointV1 → V2 with explicit
discriminant + optional q_h field; V1 files load as SharedQuery, new
V2 training writes the discriminant.
Three validation rings:
1. Per-(b,h) numgrad parity at K=16
2. One-epoch smoke (no NaN, loss decreases)
3. 30-epoch × 3-fold A/B (#204) — decision gate per §0 falsifiable claim
Implementation explicitly deferred. The decision to invest depends on
(a) GPU time budget (~3-6 hrs on L40S × 5 GPUs for the A/B), (b)
whether per-horizon cost-frontier sweeps (#202 follow-ups) surface
viable horizons beyond h6000 that would benefit from per-horizon
specialisation, and (c) the 3-fold variance noise floor making the
expected effect size visible.
Next step when ready: invoke superpowers:writing-plans against this
spec for the ~6-8 commit implementation plan.
Closes the "good to have" question from the recent brainstorm with a
concrete decision framework rather than ad-hoc implementation.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
62b1fc0965 |
infra(argo): lob-backtest-sweep workflow + argo-lob-sweep.sh (C19)
Cluster fan-out for the `fxt-backtest sweep` single-machine path.
Reads the same grid YAML format as the binary; runs each cell on a
dedicated GPU pod in parallel; aggregates at the end on a CPU pod.
infra/k8s/argo/lob-backtest-sweep-template.yaml:
WorkflowTemplate `lob-backtest-sweep` with three job templates:
ensure-binary — cache-or-compile fxt-backtest by short-SHA into
/mnt/training-data/bin/<sha>/. Mirrors the
train-multi-seed-template.yaml ensure-binary
shape but for a single binary.
run-cell — single GPU pod (ci-training-l40s default per
feedback_default_to_l40s_pool.md). Receives
cell-name + every Run arg via inputs.parameters.
Writes artifacts to <sweep-root>/<sweep-tag>/<cell>/.
aggregate — CPU pod runs `fxt-backtest aggregate <sweep-dir>`
producing aggregate.parquet + pareto_frontier.json
at the sweep root.
DAG marker `# __SWEEP_CELLS__` replaced at submission time with N
WorkflowTask stanzas (one per cell), and the aggregate's
`dependencies: [ensure-binary]` is rewritten to include every
run-cell-* dep — so aggregate waits for ALL cells.
scripts/argo-lob-sweep.sh:
Companion submission script following the argo-train.sh pattern.
Parses the grid YAML via python3 + PyYAML (no `yq` dependency —
yq isn't used elsewhere in foxhunt scripts; python3+PyYAML is
universal in our CI images). Emits per-cell WorkflowTask stanzas
+ aggregate dependency list, awks them into the template, then
`kubectl apply` + `argo submit`. Supports --dry-run for offline
rendering and --watch for live log following.
Defaults match the spec / pearl set:
- sm_89 / ci-training-l40s default (override via --gpu-pool
ci-training-h100 for sm_90 + 80 GB)
- data root /mnt/training-data/futures-baseline/ES.FUT
- sweep results under /mnt/training-data/sweeps/lob-backtest/<tag>/
- sweep-tag defaults to <basename of grid>-<short-sha>
Verified locally:
- bash -n syntax-check passes
- --help renders
- --dry-run against the existing
config/ml/sweep_decision_stride_example.yaml renders a valid
workflow with 4 cells + correct aggregate dependency list
Live submission is operational work that needs cluster access to
verify; the rendered YAML follows the same conventions as the
existing argo-train.sh workflows that ship in this repo.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
8d3efa450e |
test(ml-backtesting): integrated Ring 2 fuzz over full pipeline (C18)
Existing lob_sim_fuzz only exercised book_update + market orders. This new test suite drives the FULL integrated pipeline under random adversarial conditions: apply_snapshot (random-walk book) → step_resting_orders (random signed trade-flow signal) → broadcast_alpha (random per-horizon probs every 4th event) → step_decision_with_latency (mixed latency=0 + latency=50ms cells) → submit_market_immediate (immediate path) → pnl_track_step + isv_kelly_update_on_close Warm-starts every backtest with random-but-plausible Kelly state (at least h4 set credibly profitable so decisions actually open positions). Random target_annual_vol + annualisation_factor + max_lots per decision to vary the Kelly cap. Per-50-event invariants: • book monotonicity (bid_px[k] ≤ bid_px[k-1], ask_px[k] ≥ ask_px[k-1]) • no-crossed-book (ask[0] > bid[0]) • Pos.realized_pnl + Pos.vwap_entry finite (no NaN/Inf leaks) • Pos.position_lots in plausible range (|≤100|) • All 5 per-horizon IsvKellyState fields finite Three test sizes: integrated_fuzz_n1_short — N=1, 200 events integrated_fuzz_n8_medium — N=8, 500 events integrated_fuzz_n64_long — N=64, 1000 events All three pass on RTX 3050. The N=64 × 1000 case exercises 64,000 event-snapshots × 16,000 decisions × ~12,800 trade attempts without any invariant violation or NaN propagation. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
dd8d731dcf |
test(ml-backtesting): Ring 3 production-day replay validation (C17)
Closes the spec §8 Ring 3 deferral. Two integration tests against
FOXHUNT_TEST_DATA real MBP-10 day files:
buy_and_hold_full_day
Walk the first .dbn.zst from open to close; submit a market buy
of 1 lot at the first event, run the full step-loop
(apply_snapshot + step_resting_orders) over every subsequent
event, then close with a market sell at the last event. Assert
realized P&L matches (close_mid - open_mid) within ±2 ticks per
spec slippage budget. Proves the book-walking + position
accounting + step-loop orchestration are wire-level correct
against real data, not just hand-crafted JSON fixtures.
walk_book_one_full_day
Same fixture, no trades — just iterates every event through
apply_snapshot + step_resting_orders and asserts the book
invariants (bid/ask monotonicity, no-crossed-book) hold at every
10_000-event checkpoint across the full day. Stress test of
the kernel against real market microstructure (gaps, locked
markets, regime shifts).
Both tests skip gracefully when FOXHUNT_TEST_DATA is absent OR the
predecoded sidecars are empty (the current 40-byte placeholder
fixtures in test_data/futures-baseline/ES.FUT/*.predecoded.bin will
trigger the skip path; populating those with a real Databento decode
makes both tests fire end-to-end). Ring 3 is operational, not
CI-mandatory.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
19986c8d96 |
feat(ml-alpha): tick-rule signed trade-flow inference at L1 (C16)
Replaces the placeholder `trade_signed_vol = trade_count_delta` (always
non-negative, sign-neutral) with a proper Databento-standard tick-rule
inference applied to L1 size + price deltas across consecutive MBP-10
snapshots:
• ask_px[0] unchanged AND ask_sz[0] decreased → aggressive buys
consumed ask depth; add (prev.ask_sz − cur.ask_sz).
• bid_px[0] unchanged AND bid_sz[0] decreased → aggressive sells hit
the bid; subtract (prev.bid_sz − cur.bid_sz).
• ask_px[0] moved up → previous best ask cleared; add prev.ask_sz.
• bid_px[0] moved down → previous best bid hit; subtract prev.bid_sz.
• Pure cancellations (size shrank but price moved AWAY from us) =
ambiguous; ignore.
Convention matches `Mbp10RawInput::trade_signed_vol`: positive =
buyer-initiated, negative = seller-initiated. This is a LOWER-BOUND
estimator — won't catch trades that cleared multiple levels (those
manifest only via deeper-level deltas) or trades against hidden /
off-book liquidity. Acceptable for v1 queue-decay signal; production
deployments can layer the trades-stream loader for ground-truth flow.
Wired into BacktestHarness::run() → sim.step_resting_orders(ts, vol)
so the queue-decay branch of resting_orders.cu finally fires with
non-zero input. Previously the harness passed 0.0 unconditionally,
which meant resting limits could only fill via the price-cross
marketability branch — same-price queue-decay was inert.
Six new unit tests cover each branch of the inference (pure cancel,
ask-shrank, bid-shrank, ask-px-up, bid-px-down, mixed both-sides).
34 ml-alpha lib tests + 33 ml-backtesting lib + 12 GPU fixtures + 3
fuzz still green.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
f9b57f4c18 |
feat(fxt-backtest): sweep subcommand + example grid YAML (C15)
New `fxt-backtest sweep --grid <yaml> --out <dir>` subcommand: iterates
over a grid of Run configs, writes each cell's artifacts to
<out>/<cell_name>/, then automatically invokes the existing aggregate
path to produce aggregate.parquet + pareto_frontier.json at the root.
All cells run sequentially on the same GPU (single-machine). For
cluster fan-out the underlying mechanism is the same — Argo can wrap
this binary in a workflow that runs each cell as a separate pod
(left as infra-side work for a follow-up commit; the binary's
contract is the same).
Sweep grid YAML schema:
base: # defaults applied to every cell unless overridden
data: ...
n_parallel: ...
decision_stride: ...
latency_ns: ...
target_annual_vol_units: ...
annualisation_factor: ...
max_lots: ...
max_events: ...
seed: ...
checkpoint: ... # optional — load real trained weights
cells:
- name: cell_a
decision_stride: 1 # override base
- name: cell_b
latency_ns: 250000000
...
Each SweepCell may override any subset of fields; unset fields fall
back to base defaults. --max-cells gates the run for smoke-testing
large grids.
Adds an example grid at config/ml/sweep_decision_stride_example.yaml
that re-runs the decision_stride ∈ {1, 2, 4, 8} sweep deferred from
the original plan (task #202).
Closes #201 (sweep tool). #202 (first decision_stride sweep) is now
executable end-to-end via the example grid.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
3fa215ad2e |
feat(ml-alpha): CfcTrunk save/load checkpoint + --checkpoint CLI (C14)
CfcTrunk::save_checkpoint(path) reads each device weight tensor back
via memcpy_dtoh and bincode-serialises into a CheckpointV1 envelope:
{ version, n_in, n_hid,
w_in, w_rec, b, tau,
heads_w, heads_b,
proj_w, proj_b, proj_g, proj_n }
Total ~22k-25k f32 = ~90 KB per trunk. Tiny.
CfcTrunk::load_checkpoint(dev, cfg, path) deserialises + validates
(version == 1, n_in/n_hid match the supplied CfcConfig — a model
trained for one arch can't silently load against another). Constructs
a fresh trunk via new_random (for kernel bindings + scratch buffers)
then overwrites every weight tensor via memcpy_htod. The random init
values are thrown away — marginally wasteful, but keeps the
construction code paths unified.
Roundtrip test (--ignored, CUDA-required): save trunk_A → load → read
back every device tensor and assert bit-equality between trunk_A and
the loaded trunk_B. Passed locally. Dim-mismatch rejection test runs
without CUDA (verifies bincode envelope serialise/deserialise).
bin/fxt-backtest --checkpoint <path>: when set, overrides --seed and
loads from disk. When absent, warns loudly that the trunk is
random-initialised and backtest results are noise. This makes the
binary genuinely useful as a deployment tool — point it at a trained
checkpoint and run real backtests.
Adds bincode workspace dep to ml-alpha (was already in workspace
dependencies, just not in ml-alpha's [dependencies] block). serde
features bumped to ["derive"] (was using workspace default which
omits derive macros).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
ed19985c95 |
feat(ml-backtesting): bytecode VM dispatcher for Strategy compositions (C13)
The hardcoded WeightedByRealizedSharpe path from C7 is now joined by
a stack-based bytecode interpreter that consumes Strategy::flatten()
output, unlocking RegimeSwitch / Portfolio / non-default Ensemble /
single-Leaf compositions specified via policy-grid YAML.
cuda/decision_policy.cu — new kernel `decision_policy_program`:
Stack-based VM with parallel (value, attribution_mask) stacks.
Opcodes mirror src/policy/mod.rs::OpCode exactly:
NoOp / PushScalar / EvalRegime / BranchIfRegime
EmitPerHorizonSize (computes sized intent from alpha[h] +
IsvKellyState[h] using same Kelly + ISV-cap formula as the
hardcoded default)
AggMean / AggWeightedSharpe / AggMaxConfidence (pop n values,
push aggregated, OR attribution masks)
ApplyConflict (v1 no-op, reserved for Portfolio)
WriteOrder (terminal — converts top-of-stack to market_target +
open_horizon_masks attribution if currently flat)
AggWeightedSharpe recovers the source horizon from a single-bit
attribution mask to look up recent_sharpe; multi-bit masks (nested
aggregators that collapsed horizons) fall back to uniform weight.
decision_policy_default extended with a program_lens param: skips any
backtest whose plen > 0 (the program kernel handled it). The two
kernels run sequentially in step_decision_with_latency with mutual
exclusivity on each backtest slot.
LobSimCuda gains:
upload_program(b, &Program) — uploads a Strategy::flatten() output
to backtest b's slot in program_table_d, updates program_lens_d.
set_regime(b, regime_id) — writes regimes_d for OP_EVAL_REGIME /
OP_BRANCH_IF_REGIME consumption.
BacktestHarnessConfig.strategies (Vec<Strategy>) — empty means every
cell uses the hardcoded default; non-empty len must equal n_parallel
and each strategy is flattened + uploaded at construction.
bin/fxt-backtest --policy-grid <yaml> path now actually plumbs through
to the kernel (was parsed-but-ignored in C9). Empty grid keeps the
default behaviour.
New fixture decision_program_h4_only: uploads Strategy::Leaf(h4_only)
flattened to (EmitPerHorizonSize, WriteOrder) — 2 instructions, 16
bytes — and verifies the bytecode kernel produces equivalent end-state
to the hardcoded default (3 lot buy at vwap=5500). Proves the VM
dispatcher works end-to-end.
12/12 GPU fixtures green. 33 lib unit tests + 3 Ring 2 fuzz tests
still green.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
344d7a67d7 |
feat(ml-backtesting): wire --latency-ns end-to-end + fixture (C12)
The in-flight machinery from C11 is now reachable from the harness +
CLI. New step_decision_with_latency on LobSimCuda:
latency_ns == 0 → existing immediate-match path (submit_market_immediate
kernel fills against current book).
latency_ns > 0 → host reads market_targets, converts each non-noop
into seed_limit_order(active=2, price=very-aggressive,
arrival_ts_ns = current + latency_ns). The
resting_orders kernel promotes + fills at arrival
time, so the order sees whatever book exists then —
not the book at decision time.
Aggressive-price heuristic: buy at 1e9 / sell at 0 — the marketability
check (price ≥ best ask for buy, ≤ best bid for sell) unconditionally
crosses at arrival; the kernel then walks book levels for the actual
fill price. This models "market order with latency" correctly because
slippage emerges from the book-walking at arrival, not from a limit
price boundary.
BacktestHarness gains a latency_ns field; harness::run() always calls
step_decision_with_latency now. Also calls sim.step_resting_orders
per event (with trade_signed_vol=0.0 — the Mbp10RawInput trade-flow
hookup is deferred to a trades-feed integration commit) so in-flight
orders get a chance to promote on every snapshot.
bin/fxt-backtest no longer ignores --latency-ns; the flag value
propagates through BacktestHarnessConfig.latency_ns into the kernel
path. Default stays at 100_000_000 (IBKR + Scaleway baseline per
spec §4). Setting --latency-ns 0 selects the legacy immediate path.
Adds latency_in_flight_miss GPU fixture: limit at 5510 submitted
active=2 at T=1s with arrival_ts=T+100ms. step_resting at T+50ms
sees no promotion (still in-flight). Book moves +5 against buyer
during the 100ms window. step_resting at T+150ms promotes the
limit and fills it at the WORSE post-move book (vwap=5505 vs the
original 5500 it would have hit without latency). Slot ends active=0.
11/11 GPU fixtures green on RTX 3050. 33 lib tests still green.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
1fb2decb79 |
feat(ml-backtesting): resting orders + stops + OCO + audit ring (C11)
Picks up the deferred work from C5/C7 trim notes. Adds:
cuda/lob_state.cuh — LimitSlot (32 B) + StopSlot (32 B) + Orders
(limits[32] + stops[16] = 1536 B/backtest). u64-aligned with
explicit _pad[6] on both slot structs so the Rust mirrors
(LimitSlotFlat / StopSlotFlat) bytemuck::Pod-derive without
panics about implicit padding.
cuda/resting_orders.cu — three kernels:
resting_orders_step (runs per event after book_update):
1. In-flight → resting promotion at arrival_ts_ns. Queue position
initialised pessimistically (full level depth ahead) per spec §9.
2. Queue decay against same-side trade_signed_vol at level: ahead-
of-us first then us; partial-fill emits OrderEvent + pos update.
3. Marketability check: book moved through our price → cross at
the just-arrived book (slippage-aware fill walks levels).
IOC cancels remainder; FOK does too (partial-fill not allowed).
4. Stop trigger detection: best ask up for buy-stop, best bid
down for sell-stop. StopMarket walks book; StopLimit converts
to a new LimitSlot (overflow → submission_overflow++).
5. OCO mutual exclusion: oco_pair byte (0..31 = limit, 32..47 =
stop) — when one leg fills/triggers, the paired slot is freed.
seed_limit_slot / seed_stop_slot — host-side single-slot seeders
for fixture testing + as a v1 entry point until the decision
kernel learns to emit resting orders. Both set an overflow_flag
if the target slot is already in use (gen_counter bumped on
successful seed for SlotTag freshness).
src/sim.rs — wires three new methods:
seed_limit_order(b, slot, side, kind, active_state, oco_pair,
price, size, queue_position_init, arrival_ts_ns)
seed_stop_order(b, slot, side, kind, active_state, oco_pair,
trigger_price, limit_price, size, arrival_ts_ns)
step_resting_orders(current_ts_ns, trade_signed_vol)
→ launches resting_orders_step kernel + chains step_pnl_track
so any fill that closes a position emits a TradeRecord.
Plus read_limit_slot / read_stop_slot / read_audit_records for
fixture inspection.
Audit-ring buffers (deferred from C2 originally) now allocated:
audit_d (n × 256 × 24 B) + audit_head_d (n × u32). Both
resting_orders.cu and the decision kernel's WriteOrder path
populate it via the inlined pack_slot_tag + emit_audit helpers.
Four new GPU fixtures:
limit_rest_marketable_fill — resting buy at 5500 with book at
ask 5500 fills immediately on step_resting (3 lots, vwap=5500).
stop_trigger — buy 4 lots @ 5500, seed sell-stop at trigger=5495,
book moves so bid=5494.50: stop triggers, walks book, position
closes flat, TradeRecord emitted, stop slot active=0.
oco_one_cancels_other — pair {buy@5495, sell@5505} with oco_pair
cross-linked: book moves so bid=5505.50, sell leg fills (pos=-2),
buy leg auto-cancelled. Both slots end active=0.
submission_overflow — host loop fills all 32 limit slots, then 33rd
seed_limit_order call must return Err (slot 0 already in use).
All 10 GPU fixtures green on RTX 3050 (6 original + 4 new). Ring 2
fuzz at N ∈ {1, 16, 256} still green. 33 lib unit tests still green.
The decision kernel (C7) and submit_market_immediate (C5) paths
are unchanged — they continue to use the "immediate fill" route.
A follow-up commit can teach the decision kernel to emit resting
limits via submit_limit_alloc (already defined in resting_orders.cu).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
771faac723 |
test(ml-backtesting): Ring 2 invariant fuzz at N ∈ {1, 16, 256} (C10)
Property-based fuzz tests with random-walk MBP-10 sequences applied to
the LOB simulator at three backtest counts; assert per-snapshot
invariants that must hold regardless of input or block scheduling:
fuzz_n1_book_only (200 events, no orders)
Pure book-update kernel — verifies mid-drift random walks preserve
book monotonicity and never produce a crossed book.
fuzz_n16_with_orders (200 events, market orders every 8th event)
16 backtests in parallel, each submitting random buy/sell market
orders 1-3 lots at every 8th event. Asserts book invariants on
each backtest's state PLUS:
- position_lots stays within ±30 (plausible given fixture book depth)
- realized_pnl + vwap_entry finite (no NaN/Inf leaks)
fuzz_n256_with_orders (100 events, market orders)
Production-scale parallelism. Same invariants as N=16. Each block
has its own per-backtest Pos + OpenTradeState + TradeLog, exercising
the per-block isolation discipline established in C5-C7.
All 3 pass on RTX 3050. Spec §8 Ring 2 confidence gate hit.
Adds rand + rand_chacha to ml-backtesting dev-dependencies.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
63ed6d0217 |
feat(ml-backtesting): artifacts + sweep aggregate + fxt-backtest CLI (C9)
artifacts.rs:
- Summary struct (total_pnl_usd, sharpe_ann, sortino_ann,
max_drawdown_usd, calmar, n_trades, win_rate, avg_win/avg_loss,
profit_factor, total_fees_usd, exposure_pct,
kelly_cap_history_sample).
- compute_summary(records, pnl_curve_usd) — non-overlapping
annualisation × √825 per pearl_phase1d4_backtest_cost_edge_frontier
(K=6000 holding × 250 trading days ≈ 825 trades/year).
Sharpe + Sortino + max drawdown + Calmar.
- write_summary (JSON pretty-printed), write_trades_csv (with USD
conversion from fp ×100), write_pnl_curve_bin (bytemuck-cast f32
slice). 5 unit tests with tempdir.
aggregate.rs:
- aggregate_sweep_dir walks <root>/<cell>/summary.json, builds an
arrow RecordBatch (cell name + 9 stats columns), writes
SNAPPY-compressed aggregate.parquet.
- pareto_frontier: cells are kept unless another cell weakly
dominates on all three of (sharpe_ann maxed, max_drawdown_usd
minimised, total_fees_usd minimised) AND strictly improves on one.
Written to pareto_frontier.json (Vec<cell-name>).
- 2 unit tests (3-cell mutual-non-dominance; B-dominates-A).
harness.rs:
- run() now samples Pos.realized_pnl × $50/index-pt per event into
self.pnl_curves[b], so the per-cell P&L curve is ready for
write_artifacts() without an extra sim pass.
- write_artifacts(out_dir) — per-cell <out>/cell_NNNN/{summary.json,
trades.csv, pnl_curve.bin}.
bin/fxt-backtest:
- clap-derive CLI with two subcommands:
run --data <dir> [--predecoded-dir <dir>] [--policy-grid <yaml>]
[--n-parallel N] [--decision-stride S] [--latency-ns N]
[--target-annual-vol-units F] [--annualisation-factor F]
[--max-lots N] [--max-events N] [--seed N] --out <dir>
aggregate <sweep_dir>
- Constructs MlDevice::cuda(0) + CfcTrunk::new_random for the trunk
(v1 — ml-alpha has no checkpoint format yet; --seed gates init).
- Parses --policy-grid YAML if given but doesn't yet plumb to the
LobSimCuda decision kernel (the v1 kernel hardcodes the
Strategy::default_for path; bytecode VM is C7's deferred follow-up).
Parse step kept end-to-end so the YAML format is validated now.
- --latency-ns parsed but not consumed — reserved for follow-up
resting-order in-flight promotion (deferred from C5).
Adds parquet + arrow + arrow-array + arrow-schema + serde_yaml to
ml-backtesting deps; bin/fxt-backtest added to workspace members.
All 33 lib tests + 6 GPU fixture tests green. CLI --help renders both
subcommands correctly.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
a4127935a8 |
feat(ml-backtesting): BacktestHarness orchestrator + Ring 1b parity (C8)
BacktestHarness::new(cfg, &dev, trunk) constructs a MultiHorizonLoader
(inference_only=true), captures the trunk's perception Graph A using
loader.peek_first() as the template, then allocates a LobSimCuda.
run() walks the chronological snapshot stream, calls apply_snapshot on
every event, and at decision-stride boundaries:
trunk.update_input_buffers(raw)
→ trunk.perception_forward_captured() → (probs[N_HORIZONS], proj)
→ sim.broadcast_alpha(&probs)
→ sim.step_decision(ts, target_vol, ann_factor, max_lots)
Returns RunStats { events_processed, decisions_taken }. The harness
deliberately accepts an externally-built CfcTrunk (random-init in v1)
because a checkpoint format isn't pinned in ml-alpha yet; when one
lands, the binary CLI (C9) can switch from new_random to load_checkpoint.
trainer_parity.rs Ring 1b — two ignored tests:
peek_first_byte_equal_across_modes
Verifies the Mbp10RawInput produced by the loader path used by the
backtest harness (inference_only=true) is BYTE-EQUAL to what the
trainer's loader (inference_only=false) produces from the same
source. Guards against any future refactor accidentally diverging
the two paths (e.g. someone special-casing inference path to skip
regime feature computation). All 50 f32 fields + scalars compared
via .to_bits() equality.
inference_iteration_matches_chronological_snapshots
Verifies next_inference_input() yields monotone-ts snapshots with
correct cur==prev semantics on the first read and prev_ts==prior_ts
afterwards.
Both tests skip gracefully when FOXHUNT_TEST_DATA fixtures lack
populated sidecars (the placeholder-empty bins committed in the
test_data/ tree).
GPU-side inference parity (probs[N_HORIZONS] bit-equal across loader
modes) is deferred until ml-alpha pins a checkpoint format and we have
a small checkpoint fixture to gate it on FOXHUNT_TEST_CKPT.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
25f43a4268 |
feat(ml-backtesting): decision_policy + per-horizon ISV-Kelly (C7)
Two new kernels in cuda/decision_policy.cu:
decision_policy_default — alpha[N_HORIZONS] × per-horizon IsvKellyState
→ market target. Per-horizon Kelly fraction × signal magnitude × ISV-cap
(target_annual_vol / sqrt(realised_return_var × annualisation_factor)),
aggregated via WeightedByRealizedSharpe (weights = max(0, recent_sharpe)
/ Σ; auto-shifts capital toward empirically winning horizons per spec §5
+ §6). No floor — sub-1-lot intents become no-op. Round-to-nearest on
the final lot count to dodge an f32-truncation off-by-one where
0.8(f32) * 0.75 * 5.0 ≈ 2.99999976 → trunc-int 2.
isv_kelly_update_on_close — for every horizon flagged in
closed_horizon_mask[b], updates pnl_ema_{win,loss} via Wiener-α (floor
0.4 per pearl_wiener_alpha_floor_for_nonstationary), win_rate_ema,
Welford-ish realised_return_var, and the recent_sharpe composite.
First-observation bootstrap (per pearl_first_observation_bootstrap):
sentinel n_trades_seen=0 → direct EMA replacement, no zero-bias warmup.
The full bytecode VM from spec §6 is NOT in this commit — the default
policy is hardcoded in the kernel as the path Strategy::default_for()
already produces. Bytecode plumbing in src/policy/mod.rs stays put for
v2 expansion (custom RegimeSwitch / Portfolio compositions).
IsvKellyState struct added to lob_state.cuh (24 bytes per horizon × 5
per backtest); host mirror IsvKellyStateHost from C3 cast-compatible
via bytemuck::Pod. LobSimCuda gains broadcast_alpha + step_decision +
read_isv_kelly + write_isv_kelly (warm-start). step_decision chains:
decision_policy → merge_open_mask → submit_market_immediate
→ pnl_track_step → host close-detect → isv_kelly_update_on_close.
PRE-submit pos/pnl/mask snapshots feed the host close-detection;
captured via three small DtoH copies (cold path, 24 bytes × n_backtests).
decision_alpha_buy_close fixture: warm-start h4 with positive Kelly
state (n=50, recent_sharpe=1.0), broadcast alpha[4]=0.9 → buy 3 lots
@ ask top 5500.00. Snapshot moves to bid 5505.00, broadcast alpha[4]=0.1
→ sell 3 lots → close at +15 P&L. Verify ISV-Kelly h4: n_trades 50→51
exactly, others unchanged. PASS — all 6 Ring 1 fixtures green on RTX 3050.
Out-of-scope for this commit (defer to follow-ups, per plan trim notes):
- stop_trigger / oco_one_cancels_other / submission_overflow fixtures
(need resting-order LimitSlot[32] machinery deferred from C5)
- Bytecode VM dispatch (RegimeSwitch / Portfolio compositions)
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
1b679b5e40 |
feat(ml-backtesting): pnl_track kernel + segment_complete TradeRecord emission
pnl_track_step runs after each matching pass, compares per-block
position-state-now against persisted OpenTradeState (24 B scratch)
and either:
- records entry context (entry_ts_ns, entry_px_x100, entry_size,
realised_at_open) on open transition (prev==0, now!=0); or
- emits a 40-byte TradeRecord into the per-block trade-log buffer
on close transition (prev!=0, now==0), reconstructing implied
exit_px from the realized P&L delta and converting to USD ×100
fixed-point ($50/index-point × 100 = ×5000 multiplier).
Multi-fill averaging (scale-in then partial close) deferred to v2 —
v1 covers the clean open→close case the spec calls out as primary.
LobSimCuda owns three new buffers: open_trade_state_d (n × 24),
trade_log_d (n × TRADE_LOG_CAP × 40), trade_log_head_d (n × u32).
submit_market now takes current_ts_ns and chains pnl_track_step
internally; step_pnl_track() exposed for caller-driven orchestration.
read_trade_records(backtest_idx) drains the per-block ring as
Vec<TradeRecord>; LSP-pinned 40-byte Pod struct from C2 lines up
1:1 with the kernel's hand-rolled byte writes.
pnl_accounting_buy_close fixture: buy 4 lots @ ask top (5500.00),
book moves +5 to bid top 5505.00, sell 4 to close. Expected
realized_pnl = (5505 − 5500.00) × 4 = $20 in price-units, which is
$20 × $50/contract × 100 = 100000 USD ×100 fixed-point. PASS within
$1 fixed-point tolerance.
All 5 Ring 1 fixtures green on RTX 3050.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
78f0618294 |
feat(ml-backtesting): order_match immediate market-fill kernel + Pos accounting
submit_market_immediate walks the ask/bid book for a buy/sell market
order, computes notional cost + filled lots across up to 10 levels,
and updates per-backtest Pos { position_lots, vwap_entry, realized_pnl }
with correct VWAP scale-in + counter-direction realized-P&L accounting.
Single-writer (thread 0) per block — no atomics per
feedback_no_atomicadd.md. Each backtest gets its own target slot
{ side, size } in the device-global targets array (side=2 = no-op).
Pos struct added to lob_state.cuh (24 bytes); host mirror PosFlat in
src/lob/mod.rs is repr(C) bytemuck::Pod for direct host↔device cast.
LobSimCuda gains submit_market(backtest_idx, side, size) +
read_pos(backtest_idx) -> PosFlat. Fixture harness extended with the
SubmitMarket event variant + ExpectedPos check (vwap_tol + realized_pnl_tol
for FP tolerance).
market_order_consumes_top fixture verifies a buy of 5 lots into
ask[3@5500.00, 10@5500.25] fills 3+2 → position_lots=5,
vwap_entry=5500.10 within 1e-3 tolerance. All 4 Ring 1 fixtures
(book_update × 3 + market_order × 1) green on RTX 3050.
Scope deliberately trimmed from plan C5: this commit lands market-fill
matching only. Queue decay on resting limits, in-flight latency
promotion, stop-trigger detection, and OCO leg-cancellation will land
in follow-up commits — each in its own kernel addition. This keeps the
incremental delta reviewable.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
33636d250b |
feat(ml-backtesting): build.rs + book_update CUDA kernel + sim skeleton
build.rs compiles every .cu under crates/ml-backtesting/cuda/ to \$OUT_DIR/<name>.cubin via nvcc (no nvrtc per feedback_no_nvrtc.md); pairs every env::var with rerun-if-env-changed per pearl_build_rs_rerun_if_env_changed.md. Default arch sm_86 (RTX 3050 local); production sets FOXHUNT_CUDA_ARCH=sm_89 (L40S) or sm_90 (H100). First kernel: book_update_apply_snapshot — block-per-backtest replace of the 10-level book from a broadcast MBP-10 snapshot. Single-writer per level (thread tid = level idx), no atomics per feedback_no_atomicadd.md. lob_state.cuh defines the canonical per-block shared-mem layout that all subsequent kernels share (Book struct in this commit; Orders, Pos, IsvKellyState[5] added in C5-C7). LobSimCuda host struct (in src/sim.rs) constructed from an &MlDevice; owns per-backtest device book state + mapped-pinned input snapshots. apply_snapshot launches the kernel and synchronises; read_books drains device state for tests. Three Ring 1 fixture tests (book_update_replace, _two_step, _n_backtests=4) — N≤4 bit-exact GPU oracle assertions loaded from JSON. All three pass on RTX 3050 locally. Verifies the build-script → cubin → cudarc.load_cubin → kernel launch → DtoH read pipeline end-to-end. cudarc added as direct path dep (../../vendor/cudarc) since it's not in [workspace.dependencies] — same vendored fork ml-alpha uses. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
1ca034e22f |
feat(ml-backtesting): policy tree + bytecode flatten
Strategy (recursive Leaf | Ensemble | Portfolio | RegimeSwitch) +
StrategyConfig (horizon_idx + sizing_policy + sl_tp_rules +
max_concurrent_lots). Strategy::default_for(max_lots) returns the v1
default: per-horizon adaptive Ensemble with WeightedByRealizedSharpe
aggregator. No static horizon mask — capital auto-shifts via per-horizon
ISV-Kelly recent_sharpe weights (see spec §5 + §6).
Strategy::flatten() walks the tree → linear bytecode Program
(repr(C) Pod Instruction { op, arg0, arg1, arg2 }, 8 bytes/instruction)
for upload to device-global memory at slot blockIdx.x. RegimeSwitch
emits BranchIfRegime chains with arg2 patched to the post-child jump
offset.
Tests cover: 5-horizon default tree shape, flatten output for default
+ single-leaf + 2-leaf Portfolio (verifying ApplyConflict op emission),
Instruction layout size pin (8 bytes), LatencyConfig default 100ms
(IBKR+Scaleway baseline), Empirical wrap-around + empty case,
Decomposed summation.
IsvKellyStateHost host-mirror of cuda IsvKellyState (24 bytes, Pod)
defined in policy::sizing for warm-start file I/O in later commits.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
0b8dfa946f |
feat(ml-backtesting): host-side order types + SlotTag
OrderIntent (ergonomic host-side enum for policy YAML config + audit
reconstruction), OrderEvent (24-byte repr(C) Pod for the device-side
audit ring), TradeRecord (40-byte repr(C) Pod for the device-side
trade-log buffer), SlotTag (u32-packed {SlotKind, index, generation}
for cancel/modify against slot reuse), plus Side / LimitKind /
SlotKind / LimitParams.
Size pinning tests (24/40 bytes) ensure any future field-add forces
the kernel-side struct to migrate in lockstep. Serde roundtrips cover
OCO + Cancel variants. Renamed `.gen()` accessor to `.generation()`
to dodge the Rust 2024 reserved-keyword clash.
OCO is two limit legs linked by oco_pair byte on-device (not recursive
in OrderIntent) per spec §3 — keeps device-side layout flat for the
matching kernel coming in C5.
Adds ml-alpha + bytemuck + serde_json + tracing + thiserror to deps.
cudarc + parquet + arrow added later when CUDA + sweep aggregator
land (C4 / C9). No CUDA dependency in this commit.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
e9c4acfb12 |
feat(ml-alpha): MultiHorizonLoader inference-only mode
Add inference_only flag to MultiHorizonLoaderConfig that skips per-file forward-label precomputation (~half the file-load cost), plus peek_first() and next_inference_input() chronological-streaming methods for the ml-backtesting LOB harness. - min_size relaxed to cfg.seq_len when inference_only=true (training still requires seq_len + max_horizon + 1 for label generation) - New cursor fields (inference_file_idx, inference_snap_idx) walk every loaded snapshot in chronological order; reset() zeros both - peek_first() seeds CfcTrunk::capture_graph_a with cur==prev semantics (prev_ts_ns==ts_ns, trade_signed_vol=0) — natural stream-start - next_inference_input() errors if cfg.inference_only=false (guard against accidental mixing of training/inference paths) - All trainer call-sites (alpha_train example + multi_horizon_loader tests) updated with inference_only: false (zero behaviour change) - Inline test module exercises both modes; tests skip gracefully when fixture data isn't populated rather than panicking See docs/superpowers/specs/2026-05-18-real-lob-integration-design.md §1 (trainer parity) + §7 (orchestrator). Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
fab2d70c92 |
plan(ml-backtesting): real-LOB integration implementation plan
10 atomic commits, ~40 sub-tasks, TDD per superpowers:writing-plans:
C1 ml-alpha loader inference_only mode
C2 ml-backtesting order types + SlotTag
C3 policy tree + bytecode flatten
C4 build.rs + book_update kernel + sim skeleton (3 fixtures)
C5 order_match kernel + latency-aware fills (4 fixtures)
C6 pnl_track kernel (1 fixture)
C7 decision_policy kernel + per-horizon ISV-Kelly (3 fixtures)
C8 BacktestHarness orchestrator + Ring 1b trainer-parity test
C9 artifacts + aggregate + fxt-backtest binary
C10 Ring 2 invariant fuzz (N in {1, 16, 256})
Plan ends with green Ring 1 (11 fixtures) + Ring 1b + Ring 2 (3 fuzz)
and a working fxt-backtest run|aggregate CLI. Ring 3 (production-day
replay), per-horizon cost-frontier sweep, model-checkpoint sweeps,
IBKR live adapter, and multi-fill P&L explicitly deferred.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
bfbaea2661 |
spec(ml-backtesting): real-LOB integration design
Greenfield LOB simulator + execution policy + backtest harness inside existing ml-backtesting crate. Turns the AUC predictor into a P&L producer at IBKR+Scaleway latency profile (100ms baseline). Key design decisions: - Reuse trainer's MultiHorizonLoader, Mbp10RawInput, snap_feature_assemble cubin, CfcTrunk captured graph verbatim (zero train-vs-deploy skew). - Per-horizon ISV-Kelly + adaptive WeightedByRealizedSharpe aggregator (no static horizon mask; policy self-shifts capital). - Block-per-backtest CUDA (32 threads/warp/block), ~1.4 KB shared mem. - Three captured graphs (perception, step-event, decision); host branches only between captures. - Mapped-pinned for all CPU/GPU buffers (hard requirement per feedback_no_htod_htoh_only_mapped_pinned). - Three validation rings: N=1 bit-exact fixtures, trainer-parity check, N>1 invariant fuzz; production-day replay as Ring 3. Single binary fxt-backtest with run + aggregate subcommands. No new crates; extends existing ml-backtesting alongside barrier_backtest. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
73925b15d3 |
fix(ml-alpha): eval-path cfc_step launch config — block-per-batch (Phase B fix-up)
Phase B commit 1 (cfc_step block-per-batch refactor) updated the TRAINING-path cfg_cfc to grid=(B,1,1) but missed the same change in evaluate_batched. The eval path silently kept the legacy grid=(1,1,1) launch config, which against the refactored kernel meant only block 0 ran — batches 1..B-1 got GARBAGE h_new (whatever was in scratch memory), which propagated through CfC + GRN to produce garbage probs, which BCE-eval'd to chance-level AUC. Caught by t6z89-vs-txftz cluster A/B at L40S: baseline (t6z89, pre-Phase-B): mean_auc=0.7428 / h6000=0.7211 @ E0 broken (txftz, Phase B): mean_auc=0.4973 / h6000=0.5136 @ E0 train_loss matched within noise (0.6232 vs 0.6258) — the smoking gun for "training works, eval is broken". Why the perception_overfit smokes didn't catch it: those tests check that loss converges on a synthetic up-ramp signal, exercising only the training path. Eval is exercised by `evaluate_works_after_*` smokes but those use n_batch=1, where grid=(1,1,1) and grid=(B,1,1)=(1,1,1) are bit-identical — the bug only manifests at B>1. Fix: eval cfg_cfc → grid=(b_sz, 1, 1), matching training. Plus an explanatory comment so future eyes don't repeat the mistake. Phase B perf gains are unchanged (training path was correct). Only eval's wall time may grow slightly because of the now-correct per-batch parallelism doing the work it should have done. Follow-up: re-submit cluster A/B vs t6z89 to confirm AUC trajectory recovers to ±0.005 of baseline. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
a478ba3d84 |
perf(ml-alpha): block-per-batch attention pool bwd refactor (Phase B commit 4)
attention_pool_bwd refactored from grid=(1,1,1) to grid=(B,1,1). The existing per-batch grad_ln_out writes were already uniquely indexed; only grad_Q needed scratch+reducer (1 scratch, 1 reducer launch). Adds 1 per-batch grad scratch buffer + 1 reduce_axis0 launch: attn_grad_q_scratch_d [B, HIDDEN_DIM] ~16 KB scratch at B=32 — trivially small. attn_pool bwd runs 1×/step (not in K-loop) so the absolute wall-time win here is tiny. With this commit every single-SM bwd kernel in the trainer has been refactored to block-per-batch + scratch+reducer. Phase B kernel work complete. Next: local + cluster A/B perf benchmark to verify acceptance gates 6, 7, 8 from the spec. All 9 perception_overfit smokes pass. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
9607f33518 |
perf(ml-alpha): block-per-row VSN bwd refactor (Phase B commit 3)
variable_selection_bwd refactored from grid=(1,1,1) to grid=(B*K,1,1). VSN's n_rows = B*K positions (one row per (batch, K-position) pair); block-per-row matches the existing fwd kernel's layout. Adds 2 per-row grad scratch buffers + 2 reduce_axis0 launches: vsn_grad_w_scratch_d [B*K, FEATURE_DIM, FEATURE_DIM] vsn_grad_b_scratch_d [B*K, FEATURE_DIM] ~210 KB scratch at B=32, K=64. VSN bwd runs 1×/step (not K×) so the absolute wall-time win here is small versus commits 1+2. Done for pattern uniformity — every per-batch or per-row bwd in the trainer now uses scratch+reducer. All 9 perception_overfit smokes pass. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
5c2c3b65a8 |
perf(ml-alpha): block-per-batch GRN bwd refactor (Phase B commit 2)
multi_horizon_heads_grn_bwd_batched refactored from grid=(1,1,1) to grid=(B,1,1). Removes the single-SM bottleneck on the second-most-called K-loop kernel (64×/step like cfc_bwd). Adds 10 per-batch grad scratch buffers (one per GRN param tensor) + 10 reduce_axis0 launches collapsing B → final grad after the K-loop: grn_grad_w1_scratch_d [B, 5, HEAD_MID, HIDDEN] grn_grad_b1_scratch_d [B, 5, HEAD_MID] grn_grad_w2_scratch_d [B, 5, HEAD_MID, HEAD_MID] grn_grad_b2_scratch_d [B, 5, HEAD_MID] grn_grad_w_gate_scratch_d [B, 5, HEAD_MID] grn_grad_b_gate_scratch_d [B, 5] grn_grad_w_main_scratch_d [B, 5, HEAD_MID] grn_grad_b_main_scratch_d [B, 5] grn_grad_w_skip_scratch_d [B, 5, HIDDEN] grn_grad_b_skip_scratch_d [B, 5] Total: ~8 MB scratch at B=32. All 9 perception_overfit smokes pass (including stacked_trainer_loss_ shrinks_at_batch_32 which exercises the cross-batch reducer path on both cfc and GRN grads). Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
494a2e4827 |
perf(ml-alpha): block-per-batch cfc_step + reduce_axis0 reducer (Phase B commit 1)
Per docs/superpowers/specs/2026-05-17-kloop-parallelization-design.md. cfc_step_batched (fwd + bwd) refactored from grid=(1,1,1) with internal n_batch loop to grid=(B,1,1) — each block handles one batch. Removes the single-SM bottleneck on the K-loop's most-called kernel (64×/step). Param-grad accumulation moves to per-batch scratch: cfc_grad_w_in_scratch_d [B, n_hid, n_in] cfc_grad_w_rec_scratch_d [B, n_hid, n_hid] cfc_grad_b_scratch_d [B, n_hid] cfc_grad_tau_scratch_d [B, n_hid] Zeroed once per training step, K-loop's 64 bwd calls += into them, then 4 reduce_axis0 launches collapse B → final grad buffers (OVERWRITE) before AdamW. New AdamW-after-reducer invariant: final grads are meaningful only after the reducer has run in the current step. New reduce_axis0 kernel: single parameterised reducer [B, N] → [N] via block tree-reduce (no atomicAdd per feedback_no_atomicadd.md). Same pattern as layer_norm_reduce_param_grads — CUDA-Graph-safe. cfc_step_backward_batched shared-mem dropped from (B+1)*n_hid*4 to 2*n_hid*4 bytes per block (only one row of sd_pre needed per block bi). Tests: - New stacked_trainer_loss_shrinks_at_batch_32: FIRST test that actually exercises the cross-batch reduction code path; existing perception_overfit suite was all B=1. Initial 0.24 → final 0.00. - Scratch-clears test removed (explanatory comment kept): structurally hard to assert directly due to begin_capture/end_capture not executing kernels; the B=32 convergence smoke implicitly validates scratch zeroing since divergence would otherwise be immediate. All 9 perception_overfit smokes + 4 backward_finite_diff tests pass. build.rs: - KERNELS list adds "reduce_axis0" - Cache-bust → v11 Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
c508d101c1 |
plan(perf): K-loop block-per-batch parallelization implementation plan
Bite-sized 22-task plan implementing the design at
docs/superpowers/specs/2026-05-17-kloop-parallelization-design.md.
Four atomic commits per the spec's Rollout section:
Commit 1: reduce_axis0 kernel + cfc_step refactor (Tasks 1-10)
Commit 2: GRN bwd refactor (Tasks 11-14)
Commit 3: VSN bwd refactor (Tasks 15-17)
Commit 4: attention_pool bwd refactor (Tasks 18-20)
Each commit covers kernel rewrite + per-batch scratch buffers +
reducer launches + memset_zeros wiring + smoke validation.
Tests added across commits:
- cfc_bwd_b1_oracle.rs (Task 6): B=1 oracle vs single-sample helper
within relative_eq 1e-5 (FP-tolerant, not bit-exact)
- stacked_trainer_loss_shrinks_at_batch_32 (Task 7): first test
that exercises the cross-batch reducer path
- cfc_bwd_scratch_clears_between_steps (Task 8): scratch-zero
regression guard
Acceptance gates 1-8 from the spec are mapped to Tasks 6, 7, 8, 9,
21, 22 (local + cluster A/B perf). Reference baseline for gate #8
is t6z89's per-epoch AUC trajectory.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
ea1d8703b7 |
spec(perf): revise K-loop parallelization design — full critical pass
Revises all 12 issues from the self-review pass:
1. (HIGH) Soften bit-equivalence claim — single-sample helper and
batched kernel at B=1 are different CUDA kernels; FP order may
differ. Acceptance is relative_eq ≤ 1e-6, not bit-exact.
2. (HIGH) Explicit asymmetry: only cfc_step has a single-sample GPU
oracle. GRN/VSN/attn bwd rely on smoke + chain-rule preservation.
feedback_no_cpu_test_fallbacks.md forbids a CPU reference oracle.
3. (HIGH) Realistic targets — 3× floor, 5× stretch. Drops 15× claim
which was Amdahl-bounded under any reasonable assumption.
4. (MED) AdamW-after-reducer invariant stated explicitly: final grad
buffer is OVERWRITE by reducer, meaningful only after reducer ran.
5. (MED) New B=32 smoke test (stacked_trainer_loss_shrinks_at_batch_32)
actually exercises the cross-batch reduction path; existing
perception_overfit suite is all B=1.
6. (MED) Rollout commit 1 bundles reduce_axis0 + first consumer
(cfc_step refactor) to avoid feedback_wire_everything_up.md
orphan-kernel anti-pattern.
7. (LOW) Drop "merge to ml-alpha-phase-a" — user already chose direct
commits to that branch; clarify in Rollout.
8. (LOW) Add explicit scratch-sizing formula:
scratch_bytes ≈ B × Σ(param tensor sizes per kernel).
9. (LOW) Remove resolved open question (memset_zeros ordering).
10. (STRUCT) Post-refactor bottleneck analysis section — names the
next L40S floor (Mamba2 scan, launch latency, cuBLAS).
11. (STRUCT) cudaFuncSetAttribute note — refactored cfc_bwd's per-
block shared mem drops to 1 KiB; no attribute change needed.
12. (STRUCT) Explicit Rollback section — atomic-commit-per-kernel +
revert strategy; rollback baseline is the spec commit
(
|
||
|
jgrusewski
|
54aa69c108 |
spec(perf): K-loop block-per-batch parallelization design
Documents the design for fixing the single-SM bottleneck in five backward/K-loop kernels: cfc_step (fwd+bwd), GRN bwd, VSN bwd, and attention_pool bwd. All currently use grid=(1,1,1) with an internal n_batch loop — on L40S (142 SMs) with B=32 this puts <1% of the GPU to work in the K-loop critical path. Architecture: block-per-batch (grid=(B,1,1)) for the kernel body, plus per-batch grad scratch buffers reduced via a single parameterised reduce_axis0 kernel (block tree-reduce, no atomicAdd per feedback_no_atomicadd.md). Same pattern as the existing LayerNorm bwd reducer — CUDA-Graph-safe, debuggable, and consistent with foxhunt's no-cooperative-groups discipline. Target: ≥3× epoch wall speedup (stretch 8-15×). Makes 3-fold CV tractable (10.5h → 2-3h) and unblocks decision-stride / state-dim sweeps that compound the gain. Acceptance gates: (a) all 8 perception_overfit smokes still converge, (b) new B=1 bit-equivalence test asserts the refactored batched bwd kernel at B=1 matches the existing single-sample helper byte-for-byte, (c) cluster A/B vs t6z89 baseline shows AUC trajectory within ±0.005 and epoch wall ≥3× faster. Atomic refactor per kernel — one commit per kernel covering the kernel rewrite, scratch buffer alloc, reducer launch wiring, and smoke. No "_legacy" parallel kernels per feedback_no_legacy_aliases.md + feedback_no_partial_refactor.md. Open implementation-plan decisions: exact memset_zeros ordering inside the captured graph, batch-vs-per-tensor reducer launches, optimal block_dim for reduce_axis0 itself. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
7a558b88b7 |
feat(ml-alpha): wire attention pool into PerceptionTrainer (Phase 3.2)
Replaces CfC's zero-initialised h_old at k=0 with the attention-pooled
context vector — a learned content-addressable summary over all K LN_b
output positions. The K-loop's recurrent semantics (h_old at k+1 = h_new
at k) are preserved; only the INITIAL state at k=0 changes from zero to
the pooled context.
Forward chain change:
... → m2 → LN_b → ln_out_d [B, K, HIDDEN_DIM]
→ attention_pool_fwd(Q, ln_out_d)
→ attn_context_d [B, HIDDEN_DIM] (used as h_old@k=0)
→ attn_weights_d [B, K] (saved for bwd)
→ K-loop CfC (h_old@k=0 = attn_context, not zero)
Backward chain change:
K-loop bwd ends with grad_h_carry_d holding the gradient that would
have flowed into the initial h_old = grad on attn_context.
attention_pool_bwd consumes:
Q, ln_out_d, attn_weights_d (forward state)
grad_h_carry_d = grad_context
Writes (BOTH `+=`):
grad_attn_q_d (accumulates Q gradient — pre-zeroed at step start)
grad_h_enriched_seq_d (ADDS attn-path contribution onto LN_b output
gradient — chains with K-loop contribution)
LN_b bwd then consumes the now-summed grad_h_enriched_seq_d.
Trainer state additions (8 fields):
attn_q_d, attn_context_d, attn_weights_d, grad_attn_q_d,
attn_fwd_fn, attn_bwd_fn, _attn_module, opt_attn_q
Q is tiny (HIDDEN_DIM=128 floats); initialised near zero so initial
attention ≈ uniform 1/K (context ≈ mean of LN_b output). Model learns
content addressing from a near-uniform starting point.
Eval path mirrors training: attn_pool_fwd runs after LN_b fwd,
attn_context_d feeds the eval K-loop at k=0.
Trainer now manages 22 AdamW: CfC×4 + GRN heads×10 + LN×2 + LN_a×2 +
VSN×2 + Attn Q×1 + Mamba2×2 grouped.
Synthetic overfit smoke: stride=1 0.29 → 0.0000 in 50 steps (faster
than pre-attn 0.30), stride=4 0.30 → 0.0000. All 8 perception_overfit
tests PASS. Demonstrates the full Phase 1+2+3 stack (VSN → m1 → LN_a →
m2 → LN_b → attn pool → CfC + GRN heads) is wired forward + backward
end-to-end with every gradient flowing through every learned param.
Phase 1+2+3 capacity scale-up complete. The cumulative architectural
lift over the 3-fix-stack baseline (496q7 mean_auc=0.716 / h6000=0.704)
will be measured by deploying this stack head-to-head against bsml6.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
f76437c0c7 |
feat(ml-alpha): attention pool forward+backward CUDA kernels (Phase 3.1)
Single-head attention pool over Mamba2 K-positions, designed to replace
the CfC's zero-initialised `h_old` at k=0 with a learned content-
addressable summary over all K LN_b output positions. Forward math:
scores[k] = Q · keys[b, k, :] # [K]
attn[k] = softmax_k(scores) # [K]
context[h] = sum_k attn[k] * values[b, k, h] # [HIDDEN_DIM]
For our attention pool, keys == values == LN_b output [B, K, HIDDEN_DIM].
Single learned param: Q [HIDDEN_DIM]. Tiny (128 params).
Forward layout: grid = (B, 1, 1), block = HIDDEN_DIM=128 threads. Three
passes: (1) K dot-products with tree-reduce over HIDDEN_DIM, (2)
softmax over K with max-subtract+sum, (3) weighted sum into context.
Backward chain rule:
d_attn[k] = sum_h grad_context[h] * values[b, k, h]
d_scores[k] = attn[k] * (d_attn[k] - sum_kp attn[kp] * d_attn[kp])
d_Q[h] += sum_{b, k} d_scores[k] * values[b, k, h]
d_values[b, k, h] += grad_context[h] * attn[k] + d_scores[k] * Q[h]
Both `d_Q` and `d_values` use += semantics:
- d_Q: accumulates across batch (single block, internal n_batch loop).
- d_values: writes ADD onto whatever grad_ln_out already holds, so the
trainer can chain it on top of the K-loop's contribution to the LN_b
output gradient (no separate add-kernel needed).
Single-writer (no atomicAdd): one block per launch, thread h owns
column h of grad_ln_out for ALL (b, k). Internal n_batch loop matches
the GRN / VSN bwd pattern.
build.rs:
- "attention_pool" added to KERNELS
- Cache bust → v10
Wiring into PerceptionTrainer (Phase 3.2) is the follow-up commit:
add attn_q_d learned param + per-batch context + attn_weights buffers,
run attn_pool_fwd between LN_b fwd and the K-loop, use attn_context as
the K-loop's k=0 h_old (instead of zero_h_d), and chain attn_pool_bwd
after the K-loop reverse pass.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|