3b71d21834
Aux's original label was (p_{t+1} > p_t) — pure HFT-scale microstructure
noise that's unlearnable at our HFT-MFT trading frequency. Migrated to
(p_{t+H} > p_t) where H is read from ISV[AUX_PRED_HORIZON_BARS_INDEX=450].
Adaptive producer drives H from observed avg winning hold time:
- Pearl-A first-observation bootstrap: replace sentinel H=60 directly
on first valid observation
- Steady-state Wiener-α EMA blend, slow (α=0.01) for stable horizon
(no target-variance EMA available, fallback per
pearl_wiener_optimal_adaptive_alpha)
- "No winning trades yet" guard keeps sentinel until first valid observation
Lookahead truncation: labels at t where t+H >= total_bars are masked
(sentinel -1, loss-reduce skips). The existing aux_next_bar_loss_reduce
in aux_heads_kernel.cu already supports the -1 mask convention via the
B_valid count — no new valid_mask parameter needed.
Step 5b finding: Case B — existing per-sample buffers
(hold_at_exit_per_sample, trade_profitable_per_sample) populated by
unified_env_step_core, but no aggregate ISV slot. Added new aggregator
slot AVG_WIN_HOLD_TIME_BARS_INDEX=451 + new producer kernel
avg_win_hold_time_update_kernel.cu (block-tree-reduce, no atomicAdd).
ISV_TOTAL_DIM bumped 450 → 452.
ATOMIC migration per feedback_no_partial_refactor: both label kernels
(aux_sign_label_kernel.cu trajectory + aux_sign_label_per_step_kernel.cu
per-rollout-step) migrated together to the new
(targets, bar_indices, isv, isv_h_idx, out_labels, total, total_bars)
signature. The lookahead host-passed scalar argument is removed; H is
read from ISV inside the kernel (broadcast value, single read per
thread, on-device clamp [1, 240]).
Producer chain (per-epoch boundary): new
GpuDqnTrainer::launch_aux_horizon_chain orchestrates
avg_win_hold_time_update → aux_horizon_update sequentially alongside
launch_kelly_cap_update at the existing epoch-boundary slot in
training_loop.rs.
Trunk math (C.2/C.3/C.4) unchanged — separate aux trunk is label-
agnostic. Validation in C.10 will use H=60 cold-start; the adaptive
producer drives H from real winning-trade observations.
Tests (8 oracle, 5 new + 3 preserved):
- aux_trunk_forward_matches_numpy_reference (C.3) ✓
- aux_trunk_backward_gradient_check (C.4) ✓
- aux_trunk_backward_does_not_write_dx (C.4) ✓
- aux_sign_label_h_bar_horizon (NEW) ✓
- aux_sign_label_lookahead_mask (NEW) ✓
- aux_horizon_pearl_a_bootstrap (NEW) ✓
- aux_horizon_converges_to_steady_target (NEW) ✓
- aux_horizon_holds_sentinel_with_no_winning_trades (NEW) ✓
8/8 pass on RTX 3050 Ti.
Phase C.4b of SP14 Layer C separate-aux-trunk refactor.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
1290 lines
78 KiB
Rust
1290 lines
78 KiB
Rust
use std::path::{Path, PathBuf};
|
||
use std::process::Command;
|
||
|
||
fn main() {
|
||
println!("cargo:rerun-if-changed=build.rs");
|
||
|
||
// ── Feature schema fingerprint ────────────────────────────────────────────
|
||
// Stamped into every fxcache file so caches built against a different
|
||
// feature-extractor / state-layout / fxcache-format version fail
|
||
// header-validation at load time and trigger automatic regeneration via
|
||
// the Argo `ensure-fxcache` step. Removes the manual "remember to bump
|
||
// FXCACHE_VERSION on schema change" ritual that broke the L40S deploy.
|
||
//
|
||
// Hashing strategy: FNV-1a 64-bit over the raw bytes of these source
|
||
// files. Stable across rust versions and machines (unlike
|
||
// `std::hash::DefaultHasher`). Cost: harmless cosmetic edits to these
|
||
// files (whitespace / comments) trigger one cache regen on next deploy
|
||
// (~5min). That cost is acceptable; missed schema drift is not.
|
||
emit_feature_schema_hash();
|
||
|
||
// Only compile CUDA kernels when the cuda feature is enabled
|
||
if std::env::var("CARGO_FEATURE_CUDA").is_err() {
|
||
return;
|
||
}
|
||
|
||
let out_dir = PathBuf::from(std::env::var("OUT_DIR").unwrap());
|
||
let kernel_dir = Path::new("src/cuda_pipeline");
|
||
let common_header = kernel_dir.join("common_device_functions.cuh");
|
||
let trade_physics_header = kernel_dir.join("trade_physics.cuh");
|
||
// SP4 Task A4: header-only `__device__` linear-histogram p99 estimator
|
||
// included by `sp4_histogram_p99_test_kernel.cu` and (in subsequent
|
||
// tasks A5-A9) by every SP4 magnitude-bound producer kernel.
|
||
let sp4_histogram_p99_header = kernel_dir.join("sp4_histogram_p99.cuh");
|
||
|
||
// Rebuild cubins when shared headers change
|
||
println!("cargo:rerun-if-changed={}", trade_physics_header.display());
|
||
println!("cargo:rerun-if-changed={}", sp4_histogram_p99_header.display());
|
||
|
||
// Detect GPU architecture from env or default to sm_80.
|
||
// `rerun-if-env-changed` is mandatory: the cargo-target PVC is shared
|
||
// across H100 (sm_90) and L40S (sm_89) jobs, and without this trigger
|
||
// cargo serves stale cubins compiled for the previous arch on every
|
||
// subsequent build, producing CUDA_ERROR_NO_BINARY_FOR_GPU at runtime.
|
||
println!("cargo:rerun-if-env-changed=CUDA_COMPUTE_CAP");
|
||
let cuda_compute_cap = std::env::var("CUDA_COMPUTE_CAP").unwrap_or_else(|_| "80".to_string());
|
||
let arch = format!("sm_{cuda_compute_cap}");
|
||
|
||
// Check if nvcc is available -- gracefully skip if not (non-CUDA builds)
|
||
let nvcc_path = find_nvcc();
|
||
let nvcc = match nvcc_path {
|
||
Some(p) => p,
|
||
None => {
|
||
eprintln!(" warning: nvcc not found, skipping CUDA kernel precompilation");
|
||
eprintln!(" Install CUDA toolkit or set CUDA_HOME for GPU builds");
|
||
return;
|
||
}
|
||
};
|
||
|
||
// Read common header once
|
||
let common_src = std::fs::read_to_string(&common_header)
|
||
.unwrap_or_else(|e| panic!("Failed to read {}: {e}", common_header.display()));
|
||
println!("cargo:rerun-if-changed={}", common_header.display());
|
||
|
||
// All kernels to precompile.
|
||
// Kernels marked "standalone" have their own helpers and don't need common header.
|
||
// All others get common_device_functions.cuh prepended.
|
||
let kernels_with_common = [
|
||
// Original 24 kernels
|
||
"epsilon_greedy_kernel.cu",
|
||
"backtest_env_kernel.cu",
|
||
"backtest_forward_ppo_kernel.cu",
|
||
"backtest_forward_supervised_kernel.cu",
|
||
"backtest_metrics_kernel.cu",
|
||
"dt_kernels.cu",
|
||
"ensemble_kernels.cu",
|
||
"her_episode_kernel.cu",
|
||
"her_relabel_kernel.cu",
|
||
"signal_adapter_kernel.cu",
|
||
"statistics_kernel.cu",
|
||
"training_guard_kernel.cu",
|
||
"c51_loss_kernel.cu",
|
||
"mse_loss_kernel.cu",
|
||
"curiosity_training_kernel.cu",
|
||
"curiosity_inference_kernel.cu",
|
||
"dqn_utility_kernels.cu",
|
||
"attention_kernel.cu",
|
||
"attention_backward_kernel.cu",
|
||
"iql_value_kernel.cu",
|
||
"iqn_dual_head_kernel.cu",
|
||
"monitoring_kernel.cu",
|
||
"nstep_kernel.cu",
|
||
"reward_shaping_kernel.cu",
|
||
"ppo_experience_kernel.cu",
|
||
"bias_kernels.cu",
|
||
// Formerly standalone — now need common header for BF16 types
|
||
"experience_kernels.cu",
|
||
"ema_kernel.cu",
|
||
"relu_mask_kernel.cu",
|
||
"c51_grad_kernel.cu",
|
||
"mse_grad_kernel.cu",
|
||
"q_stats_kernel.cu",
|
||
"cql_grad_kernel.cu",
|
||
"trade_stats_kernel.cu",
|
||
"backward_kernels.cu",
|
||
"iqn_cvar_kernel.cu",
|
||
"mamba2_temporal_kernel.cu",
|
||
"graph_utility_kernels.cu",
|
||
"grad_decomp_kernel.cu",
|
||
// SP7 Path A (2026-05-03): raw CQL norm reading `cql_grad_scratch`
|
||
// directly so the SP7 controller has a budget-independent CQL
|
||
// reference signal. Replaces the never-populated `cql` slot at
|
||
// element offset 3 in `grad_decomp_result_pinned`. See
|
||
// `cql_raw_norm_kernel.cu` and audit doc Fix 31 SP7 Path A.
|
||
"cql_raw_norm_kernel.cu",
|
||
"branch_grad_balance_kernel.cu",
|
||
"backtest_plan_kernel.cu",
|
||
"tau_update_kernel.cu",
|
||
"epsilon_update_kernel.cu",
|
||
"per_branch_gamma_update_kernel.cu",
|
||
"kelly_cap_update_kernel.cu",
|
||
"atoms_update_kernel.cu",
|
||
"q_quantile_kernel.cu",
|
||
// D.8 Plan 2 Task 6C: TLOB attention kernels (SDP + state scatter/read)
|
||
"tlob_kernel.cu",
|
||
// C.2 Plan 3 Task 1: per-component |reward| EMA into ISV[63..69)
|
||
"reward_component_ema_kernel.cu",
|
||
// B.2 Plan 3 Task 3: Flat→Positioned transition rate EMA into ISV[71]
|
||
"trade_rate_ema_kernel.cu",
|
||
// B.4 Plan 3 Task 4: per-batch readiness EMA + derived plan_threshold
|
||
// (producer for ISV[PLAN_THRESHOLD_INDEX=49] + ISV[READINESS_EMA_INDEX=75])
|
||
"plan_threshold_update_kernel.cu",
|
||
// C.3 Plan 3 Task 7: train-vs-val state-distribution KL EMA + Flat-trap
|
||
// escape amplifier (producer for ISV[78]+ISV[79]). Adaptive amp
|
||
// multiplies B.1 opp_cost and B.2 bonus consumers.
|
||
"state_kl_divergence_kernel.cu",
|
||
// B.3 Plan 3 Task 8: GPU-only seeded warm-start scripted policies.
|
||
// Replaces network-Q action source during the seed phase. 4 policies
|
||
// (uniform/momentum/mean-rev/vwap-deviation) deterministically mixed
|
||
// 40/20/20/20 by `i % 5`. Driven by ISV[SEED_STEPS_DONE/TARGET]
|
||
// dispatch at the launcher boundary.
|
||
"scripted_policy_kernel.cu",
|
||
// B.3 Plan 3 Task 8: per-collect_experiences seed-phase progress
|
||
// counter. Increments ISV[SEED_STEPS_DONE], EMAs the derived
|
||
// seed-fraction into ISV[SEED_FRAC_EMA] (consumed by Task 9).
|
||
"seed_step_counter_update_kernel.cu",
|
||
// C.5 Plan 3 Task 9: CQL α ramp coupled to ISV[SEED_FRAC_EMA].
|
||
// EMAs ISV[CQL_ALPHA_INDEX=48] toward `final × (1 - seed_frac)`.
|
||
// Producer-only upgrade for slot 48; CQL gradient kernel consumer
|
||
// (already reading ISV[48] per Plan 1 Task 12) unchanged.
|
||
"cql_alpha_seed_update_kernel.cu",
|
||
// E.5 Plan 4 Task 5 Mode A: attention-focus interpretability EMAs.
|
||
// Multi-block GPU reductions — 3 blocks reducing VSN mag/dir param
|
||
// slices + Mamba2 enriched-hidden buffer. EMA-updates ISV[87..90).
|
||
// Diagnostic only; no consumer reads these slots in Mode A.
|
||
"attention_focus_ema_kernel.cu",
|
||
// Plan 4 follow-up: target-drift EMA, replaces legacy CPU-DtoH
|
||
// per_branch_target_drift(). 2-block kernel computing
|
||
// RMS(target - online) for mag + dir branches → ISV[92,93].
|
||
"target_drift_kernel.cu",
|
||
// Plan 4 Task 2c.1: Gated Residual Network kernels (forward + backward).
|
||
// Module is additive — no production callers in this commit; Task 2c.3+4
|
||
// wires it into the trunk encoder.
|
||
"grn_kernel.cu",
|
||
// Plan 4 Task 2c.3c.5: per-batch RMS EMA of `save_h_s2` into ISV[96].
|
||
// Single-block shmem-reduce kernel (256 threads, no atomicAdd) launched
|
||
// alongside `reward_component_ema` from `training_loop.rs`. Producer-only
|
||
// in this commit — 2c.3c.6 wires the consumer in `mag_concat_qdir`'s
|
||
// adaptive-scale path.
|
||
"h_s2_rms_ema_kernel.cu",
|
||
// Plan 4 Task 3 (E.3): IQN multi-quantile diagnostic EMAs into
|
||
// ISV[99..103) (Q_p05/Q_p25/Q_p75/Q_p95; median τ=0.50 is the existing
|
||
// greedy-Q diagnostic and not duplicated). 4-block kernel, one block
|
||
// per off-median quantile, shmem-reduce over (B × TBA) of
|
||
// `save_q_online`. No atomicAdd. Producer-only — diagnostic only.
|
||
"iqn_quantile_ema_kernel.cu",
|
||
// Plan 4 Task 1B-i: VSN feature-selection kernels (forward + backward).
|
||
// Per-sample softmax-over-6-groups + gate-multiply. Caller-side cuBLAS GEMMs
|
||
// for the per-group MLPs (Linear_1 + ReLU + Linear_2). Module is additive
|
||
// in this commit — ZERO production callers; consumers wired in 1B-iii/iv.
|
||
"vsn_feature_selection_kernel.cu",
|
||
// Plan 4 Task 1B-i: VSN per-group mask EMA producer (single-block shmem
|
||
// reduction, no atomicAdd). Reads vsn_mask saved by the forward kernel,
|
||
// EMA-updates 6 ISV slots. Producer-only; consumer-side ISV slot
|
||
// allocation lands in 1B-ii.
|
||
"vsn_mask_ema_kernel.cu",
|
||
// Plan 4 Task 6 Commit A: multi-task auxiliary heads (E.6).
|
||
// Two `Linear(SH2 → 32) → ELU → Linear(32 → K)` MLPs branching off
|
||
// h_s2 — next-bar return regression (K=1, MSE) + 5-class regime
|
||
// classification (K=5, cross-entropy). Forward + backward + loss-
|
||
// reduce + label-builder + per-tensor batch-reduce kernels. Module
|
||
// is additive — ZERO production callers in this commit; Commit B
|
||
// wires the forward/backward + training-loop loss accumulation.
|
||
"aux_heads_kernel.cu",
|
||
// Plan 4 Task 6 Commit A: aux-head loss EMA producer (single-thread
|
||
// single-block kernel mirroring h_s2_rms_ema_kernel's footprint).
|
||
// Producer-only; ISV[AUX_NEXT_BAR_MSE_EMA_INDEX] +
|
||
// ISV[AUX_REGIME_CE_EMA_INDEX] consumer wires in Commit B.
|
||
"aux_heads_loss_ema_kernel.cu",
|
||
// SP14 Layer C Phase C.3 (2026-05-08): aux trunk forward kernel —
|
||
// 3-layer Linear→ELU→Linear→ELU→Linear MLP parallel to Q's GRN
|
||
// trunk. Reads encoder output (`x_in [B, ENCODER_OUT_DIM]`),
|
||
// writes `h_s2_aux [B, AUX_HIDDEN_DIM]` plus saved-for-backward
|
||
// hidden activations `h_aux1` and `h_aux2`. Module is additive
|
||
// in this commit — wire-up into the collector chain lands in
|
||
// Phase C.5 (atomic). Backward kernel + Adam updates land in
|
||
// Phase C.4. See plan §C.3 in
|
||
// `2026-05-07-sp14-layer-c-separate-aux-trunk.md`.
|
||
"aux_trunk_forward_kernel.cu",
|
||
// SP14 Layer C Phase C.4 (2026-05-08): aux trunk backward kernel —
|
||
// three kernels in one cubin (`aux_trunk_bwd_dh_pre`,
|
||
// `aux_trunk_bwd_dW_reduce`, `aux_trunk_bwd_db_reduce`).
|
||
// Propagates `dh_s2_aux` through w3/w2/w1 to accumulate
|
||
// dW{1,2,3} + db{1,2,3} via block-tree-reduce (no atomicAdd per
|
||
// feedback_no_atomicadd). CRITICAL: kernel does NOT compute
|
||
// `dx_in` — the encoder boundary is the stop-gradient per
|
||
// locked design. Module is additive — wire-up into collector
|
||
// backward chain + Adam updates land in Phase C.5 (atomic).
|
||
// See plan §C.4 in
|
||
// `2026-05-07-sp14-layer-c-separate-aux-trunk.md`.
|
||
"aux_trunk_backward_kernel.cu",
|
||
// Plan A Phase 0 (Thompson sampling spec 2026-04-26): standalone test
|
||
// kernel exercised only by `distributional_q_tests.rs`. Implements the
|
||
// Thompson direction sampling math (inverse-CDF over C51 atoms,
|
||
// uniform-τ over IQN quantiles, argmax of E[Q]) that Phase 2 will
|
||
// integrate into `experience_action_select`. Test-only; no production
|
||
// callers in this commit.
|
||
"thompson_test_kernel.cu",
|
||
// MoE Regime Redesign Phase 2 Task 2.1: mixture-forward kernel.
|
||
// Single-thread-per-(b,c), no atomicAdd, capture-friendly.
|
||
// h_s2[b,c] = Σ_k g[b,k] · expert_outputs[k,b,c].
|
||
// Subsequent tasks extend this cubin with moe_mixture_backward,
|
||
// moe_load_balance_loss, moe_expert_util_ema_update.
|
||
"moe_kernels.cu",
|
||
// MoE adaptive load-balance λ controller (2026-04-27): single-block
|
||
// single-thread cold-path producer kernel matching kelly_cap_update /
|
||
// cql_alpha_seed_update precedent. Reads ISV[MOE_GATE_ENTROPY_EMA_INDEX=126]
|
||
// (entropy EMA from moe_expert_util_ema_update), writes
|
||
// ISV[MOE_LAMBDA_EFF_INDEX=128] = λ_floor + λ_max_extra × deficit.
|
||
// Consumer (moe_load_balance_loss in moe_kernels.cu) reads slot 128
|
||
// at runtime — no DtoH per feedback_isv_for_adaptive_bounds.md.
|
||
"moe_lambda_eff_kernel.cu",
|
||
// Plan C Phase 2 follow-up A.2 (2026-04-29): adaptive Polyak-tau
|
||
// dampening signal. Single-block single-thread cold-path producer
|
||
// mirroring moe_lambda_eff_kernel / kelly_cap_update precedent. Reads
|
||
// ISV[Q_ABS_REF_INDEX=16] + ISV[Q_DIR_ABS_REF_INDEX=21] and the host-
|
||
// passed q_mean_curr/prev scalars; writes
|
||
// ISV[Q_DRIFT_RATE_INDEX=129] = clip(|q_curr-q_prev|/denom, 0, 4).
|
||
// Consumer (tau_update_kernel.cu) multiplies tau_eff by
|
||
// 1/(1+ISV[129]) so tau ∈ [tau_base/5, tau_base] — monotone
|
||
// dampening under drift; healthy runs unaffected.
|
||
"q_drift_rate_ema_kernel.cu",
|
||
// Plan C Phase 2 follow-up K (2026-04-29): adaptive fold-boundary
|
||
// warmup signal. Single-block single-thread cold-path producer
|
||
// mirroring q_drift_rate_ema_kernel / moe_lambda_eff_kernel
|
||
// precedent. Reads two grad-norm EMAs (fast α=0.1, slow α=0.001)
|
||
// and writes ISV[FOLD_WARMUP_FACTOR_INDEX=130] = clamp(fast/slow, 0, 1).
|
||
// Two consumers, both monotone (only dampen): lr_eff = lr_base × max(0.05, factor)
|
||
// and clip_eff = clip_base × (0.1 + 0.9 × factor). After fold reset
|
||
// factor starts at 0 → heavy damping; rises to 1 as gradients
|
||
// stabilise → consumers return to baseline (healthy runs unaffected).
|
||
"fold_warmup_factor_kernel.cu",
|
||
// SP4 Layer C close-out C1 redesigned (2026-05-01): GPU-only
|
||
// fast/slow grad-norm EMA update. Single-thread, single-block —
|
||
// replaces the host-side arithmetic block in `update_adaptive_clip`
|
||
// per `feedback_no_cpu_compute_strict`. Reads `grad_norm_buf[0]`
|
||
// (populated earlier in the same stream by `launch_grad_norm_finalize`),
|
||
// updates the fast (α=0.1, ~10-step horizon) and slow (α=0.001,
|
||
// ~1000-step horizon) EMAs in mapped-pinned scalars consumed by
|
||
// `fold_warmup_factor_kernel.cu` (the actual ISV[130] producer).
|
||
// Cold-start sentinel (`prev ≤ 0.0` ⇒ assign obs directly) preserves
|
||
// the deleted host-side formula's semantics exactly.
|
||
"update_grad_norm_emas_kernel.cu",
|
||
// SP4 Layer C close-out C2 (2026-05-01): GPU-only IQN readiness
|
||
// gauge update. Single-thread single-block — replaces the host-side
|
||
// arithmetic block in `update_iqn_readiness` per
|
||
// `feedback_no_cpu_compute_strict`. Reads host-passed `iqn_loss`
|
||
// scalar (already from `read_total_loss` mapped-pinned readback),
|
||
// updates the three coupled scalars (loss_initial anchor, loss_ema
|
||
// streaming smoothed loss, readiness gauge ∈ [0,1]) all backed by
|
||
// mapped-pinned device-mapped storage. Cold-start sentinel
|
||
// (`loss_initial < 1e-12` ⇒ bootstrap) preserves the deleted host-
|
||
// side formula's semantics exactly. Same adaptive-α formula
|
||
// (clamp(|err|/(|err|+0.01), 0.01, 0.30)) and improvement gauge
|
||
// ((initial - ema) / max(initial, 1e-6)) clamped to [0, 1].
|
||
// Consumer (c51_loss_kernel CVaR α via `iqn_readiness_dev_ptr`) is
|
||
// unchanged — same dev_ptr, kernel writes the gauge through it.
|
||
"update_iqn_readiness_kernel.cu",
|
||
// SP4 Layer C close-out C3 (2026-05-01): GPU-only utilization EMA
|
||
// update. Single-thread, single-block — replaces the host-side
|
||
// arithmetic block in `reduce_current_q_stats` per
|
||
// `feedback_no_cpu_compute_strict`. Reads host-passed `atom_util`
|
||
// scalar (already a mapped-pinned readback at host[6]); updates
|
||
// `utilization_ema_pinned` + `homeostatic_obs_pinned[1]` mirror in
|
||
// lockstep. Cold-start sentinel `prev > 0.99 ⇒ assign obs directly`
|
||
// (constructor init=1.0 marker) preserves the deleted host formula.
|
||
"update_utilization_ema_kernel.cu",
|
||
// SP4 Layer C close-out C4 (2026-05-01): GPU-only winsorized
|
||
// adaptive grad-clip update. Single-thread, single-block — replaces
|
||
// the host-side scalar-reduction + EMA arithmetic chain in
|
||
// `update_adaptive_clip` per `feedback_no_cpu_compute_strict`. The
|
||
// chain is winsor (1) + cold-start sentinel + EMA (3) + scalar
|
||
// reduction (4) + ISV upper-bound clamp (5) + mapped-pinned write
|
||
// (6); all six steps now run in a single kernel that reads
|
||
// host-passed `observed_grad_norm` + the three mapped-pinned
|
||
// scalars (`adaptive_clip_pinned`, `grad_norm_ema_pinned`,
|
||
// `outlier_diag_pinned`) + `isv[GRAD_CLIP_BOUND_INDEX]`.
|
||
"update_adaptive_clip_kernel.cu",
|
||
// SP4 Layer C close-out follow-up (2026-05-01): GPU-only readiness-
|
||
// driven homeostatic-target EMA. Single-block, six threads — one per
|
||
// homeostatic slot — replaces the host-side EMA loop in
|
||
// `calibrate_homeostatic_targets` per `feedback_no_cpu_compute_strict`.
|
||
// Reads host-passed `readiness` scalar (already-clamped IQN gauge),
|
||
// observations from `homeostatic_obs_pinned[6]` (mapped-pinned,
|
||
// already GPU-readable via `homeostatic_obs_dev_ptr`), updates
|
||
// `homeostatic_targets_pinned[k]` for k=1..5 with adaptive α
|
||
// (`0.3 × (1 - readiness) + 0.01 × readiness`); thread 0 forces the
|
||
// Q-mean invariant `targets[0] = 0.0`. Constructor-initialised target
|
||
// defaults (`[0, 0.85, 0.1, 0, 1.0, 0.5]` at gpu_dqn_trainer.rs:14583-
|
||
// 14588) make a separate cold-start sentinel unnecessary — the first
|
||
// call's EMA blends those defaults with the first observation, same
|
||
// algebraic shape the deleted host loop relied on.
|
||
"calibrate_homeostatic_kernel.cu",
|
||
// HEALTH_DIAG GPU port — Phase 2A landing (2026-04-28). Single-block
|
||
// single-thread `health_diag_isv_mirror` kernel copies a curated set
|
||
// of ISV signal-bus slots into the mapped-pinned `HealthDiagSnapshot`
|
||
// (see `cuda_pipeline/health_diag.rs` + `gpu_health_diag.rs`).
|
||
// Phases 2B-2E append further entry points to this same .cu file
|
||
// (eval-histogram, q-mag-reduce, per-sample-reduce, finalise);
|
||
// Phase 3 wires the launch chain into the captured graph; Phase 4
|
||
// deletes the CPU-side reductions and HEALTH_DIAG emit-path Vecs.
|
||
"health_diag_kernel.cu",
|
||
// SP4 Layer A Task A4 (2026-04-30): standalone test kernel exercised
|
||
// only by `tests/sp4_producer_unit_tests.rs`. Wraps the header-only
|
||
// `sp4_histogram_p99<BLOCK_SIZE>` device function (in
|
||
// `sp4_histogram_p99.cuh`) so the Rust unit test can drive the
|
||
// three-pass single-block linear-histogram p99 estimator on a
|
||
// controlled |N(0,1)| sample. Producer kernels (Tasks A5-A9) include
|
||
// the same .cuh header directly; the test wrapper has no production
|
||
// callers.
|
||
"sp4_histogram_p99_test_kernel.cu",
|
||
// SP12 v3 (2026-05-04): standalone test kernel exercised only by
|
||
// `tests/sp12_reward_math_tests.rs`. Wraps the three reward
|
||
// composition device functions added to `trade_physics.cuh`
|
||
// (compute_asymmetric_capped_pnl, compute_min_hold_penalty,
|
||
// compute_lump_sum_opp_cost) so the Rust GPU oracle tests can
|
||
// drive each formula on synthetic inputs with analytically-known
|
||
// expected outputs per `feedback_no_cpu_test_fallbacks`. Producer
|
||
// call site is in `experience_kernels.cu::experience_env_step`
|
||
// segment_complete branch; the test wrapper has no production
|
||
// callers.
|
||
"sp12_reward_math_test_kernel.cu",
|
||
// SP4 Layer A Task A5 (2026-04-30): first end-to-end SP4 producer.
|
||
// Single-block kernel reads `denoise_target_q_buf`, computes
|
||
// p99(|target_q|) via `sp4_histogram_p99<256>`, writes the step
|
||
// observation to `producer_step_scratch_buf[0]`. The host launcher
|
||
// (`GpuDqnTrainer::launch_sp4_target_q_p99`) syncs, applies Pearls
|
||
// A+D, writes new x_mean to ISV[TARGET_Q_BOUND_INDEX=131] + Wiener
|
||
// state back to `wiener_state_buf`. Cold-path launch in this task;
|
||
// Task A10 may relocate to the captured graph. Tasks A6-A9 mirror
|
||
// the launcher template and slot-index discipline established here.
|
||
"target_q_p99_kernel.cu",
|
||
// SP4 Layer A Task A6 (2026-04-30): atom_pos_p99 × 4 branches.
|
||
// Single parameterised single-block kernel reads ONE branch's slice
|
||
// of `atom_positions_buf` (`[4 * num_atoms]`, branch `b` at offset
|
||
// `b * num_atoms` per `atoms_update_kernel.cu`'s canonical layout),
|
||
// computes p99(|atom_positions[branch]|) via `sp4_histogram_p99<256>`,
|
||
// writes the step observation to
|
||
// `producer_step_scratch_buf[scratch_idx]`. The launcher
|
||
// (`launch_sp4_atom_pos_p99_all_branches`) loops `branch ∈ 0..4`,
|
||
// launching this kernel four times with distinct `(branch_buf,
|
||
// scratch_idx ∈ 1..5)` pairs and applying Pearls A+D host-side per
|
||
// branch (writes back to ISV[ATOM_POS_BOUND_BASE + branch]).
|
||
"atom_pos_p99_kernel.cu",
|
||
// SP4 Layer A Task A7 (2026-04-30): Pearl B fused per-param-group
|
||
// statistics oracle. Single-block kernel reads
|
||
// `(params, grads, adam_m, adam_v)` once and runs 4 sequential
|
||
// passes (5 for trunk):
|
||
// Pass A: WEIGHT_BOUND[group] = p99(|params|) via sp4_histogram_p99
|
||
// Pass B: ADAM_M_BOUND[group] = p99(|adam_m|) via sp4_histogram_p99
|
||
// Pass C: ADAM_V_BOUND[group] = p99(|adam_v|) via sp4_histogram_p99
|
||
// Pass D: WD_RATE[group] = |Σ w·g| / max(Σ w², EPS_DIV)
|
||
// Pass E: L1_LAMBDA[trunk] = (mean|g|/mean|w|) × entropy_deficit
|
||
// (group 0 only, gated by
|
||
// l1_lambda_scratch_idx >= 0).
|
||
// Pearl B 4× memory-bandwidth reduction vs four naive single-stat
|
||
// producer kernels: each buffer read once for all outputs. The host
|
||
// launcher (`launch_sp4_param_group_oracles_all_groups`) loops
|
||
// `group ∈ 0..SP4_PARAM_GROUP_COUNT=8`, launching this kernel 8
|
||
// times with distinct `(params, grads, m, v, count, scratch slots)`
|
||
// tuples, then applies Pearls A+D host-side per output via
|
||
// `pearls_ad_update`.
|
||
"param_group_oracle_kernel.cu",
|
||
// SP4 Layer A Task A8 (2026-04-30): grad_norm_p99 producer
|
||
// (degenerate single-element). `grad_norm_buf` is a single-element
|
||
// f32 device buffer; p99 over 1 element IS the element itself, so
|
||
// the kernel reduces to a single-thread copy of `grad_norm_buf[0]`
|
||
// into `producer_step_scratch_buf[scratch_idx=38]` with
|
||
// `__threadfence_system()`. The host launcher
|
||
// (`GpuDqnTrainer::launch_sp4_grad_norm_p99`) syncs, applies Pearls
|
||
// A+D via `pearls_ad_update`, writes the new x_mean to
|
||
// ISV[GRAD_CLIP_BOUND_INDEX=168] + Wiener state. Establishes the
|
||
// launcher template that Mech 6's eventual Layer-B ISV read consumes.
|
||
"grad_norm_p99_kernel.cu",
|
||
// SP4 Layer A Task A9 (2026-04-30): h_s2_p99 producer for
|
||
// ISV[H_S2_BOUND=169]. Single-block 256-thread kernel reading
|
||
// `save_h_s2 [B × SH2]` (the trunk-output activation surface
|
||
// populated each step by the online forward) via
|
||
// `sp4_histogram_p99<256>`; output → `producer_step_scratch_buf[39]`
|
||
// via host-side Pearls A+D applied at slot 39. Mirrors Task A5's
|
||
// pattern exactly. NB: distinct from existing `h_s2_rms_ema_update`
|
||
// (slot 96) which tracks RMS — different signal. Task A13 will
|
||
// retrofit that existing kernel to also use Pearls A+D, separately.
|
||
"h_s2_p99_kernel.cu",
|
||
// SP4 Layer C #260 (2026-05-01): bw_d_h_s2_p99 producer for
|
||
// ISV[BW_D_H_S2_BOUND_INDEX=171]. Single-block 256-thread kernel
|
||
// reading `bw_d_h_s2 [B × SH2]` (the IQN trunk-backward d_h_s2 buffer
|
||
// post-DtoD copy from `iqn_d_h_s2_ptr` and post slot-35 NaN check)
|
||
// via `sp4_histogram_p99<256>`; output → `producer_step_scratch_buf[69]`
|
||
// with `__threadfence_system()`. The host launcher
|
||
// (`GpuDqnTrainer::launch_sp4_bw_d_h_s2_p99`) chains the GPU
|
||
// `apply_pearls_ad_kernel` on the same stream — writes new x_mean to
|
||
// ISV[171] + Wiener state at slot 69 (offset 120). Migrated from
|
||
// SP1-Phase-C `1e6 × ISV[H_S2_RMS_EMA].max(1.0)` per
|
||
// `feedback_isv_for_adaptive_bounds`. Mirrors `h_s2_p99_kernel.cu`'s
|
||
// shape exactly — only differs in the consumed buffer + ISV slot.
|
||
"bw_d_h_s2_p99_kernel.cu",
|
||
// SP4 Layer C #260 (2026-05-01): q_dir_grad_p99 producer for
|
||
// ISV[Q_DIR_GRAD_BOUND_INDEX=172]. Single-block 256-thread kernel
|
||
// reading the union of `d_value_logits` ∪ `d_adv_logits` (both
|
||
// dueling-head logits gradient buffers) via `sp4_histogram_p99_multi<256>`
|
||
// with `n_sub=2`; output → `producer_step_scratch_buf[70]` with
|
||
// `__threadfence_system()`. The host launcher
|
||
// (`GpuDqnTrainer::launch_sp4_q_dir_grad_p99`) builds a 2-element
|
||
// device-pointer table + 2-element count table in mapped-pinned memory
|
||
// (reusing the same `oracle_subbuf_table_buf` allocated for
|
||
// `param_group_oracle_kernel.cu`), then chains the GPU
|
||
// `apply_pearls_ad_kernel` on the same stream — writes new x_mean to
|
||
// ISV[172] + Wiener state at slot 70 (offset 123). Migrated from
|
||
// SP1-Phase-C `1e6 × ISV[Q_DIR_ABS_REF].max(1.0)` per
|
||
// `feedback_isv_for_adaptive_bounds`. The same bound was previously
|
||
// applied to BOTH d_value_logits and d_adv_logits, so the union p99
|
||
// is the natural ISV-driven replacement.
|
||
"q_dir_grad_p99_kernel.cu",
|
||
// SP5 Task A1 (2026-05-01): q_branch_stats producer.
|
||
// Single-block 4-thread kernel reads q_out_buf [B, 13] row-major
|
||
// and computes per-branch (mean, max_abs_dev, var, entropy) for
|
||
// branches Direction(4)/Magnitude(3)/Order(3)/Urgency(3).
|
||
// Writes 16 floats to producer_step_scratch_buf[71..87).
|
||
// Shared signal source for Pearls 1, 2, 3, 5.
|
||
"q_branch_stats_kernel.cu",
|
||
// SP5 Task A1 (2026-05-01): Pearl 1 atom-span controller.
|
||
// Single-block 4-thread kernel reads q_branch_stats scratch outputs
|
||
// and per-branch atoms_clip_count, computes per-branch (v_center,
|
||
// v_half, headroom, clip_rate) via the headroom controller
|
||
// (TARGET_CLIP_RATE=0.01, HEADROOM_LR=0.01 — Invariant 1 anchors).
|
||
// Writes 16 floats to producer_step_scratch_buf[87..103).
|
||
// apply_pearls_ad_kernel then smooths via Pearls A+D into ISV[174..190).
|
||
"pearl_1_atom_kernel.cu",
|
||
// SP5 Task A2 (2026-05-01): Pearl 3 per-branch NoisyNet sigma controller.
|
||
// Single-block 4-thread kernel reads ATOM_V_HALF (Pearl 1 output,
|
||
// ISV[178..182)) + BRANCH_ENTROPY (q_branch_stats, ISV[218..222)) +
|
||
// current SIGMA_FRACTION (ISV[214..218)); computes entropy-deficit
|
||
// controller (target_entropy = log(n_actions[b]) * 0.7, SF_LR=0.005)
|
||
// and new_sigma = new_sf * max(v_half, 1.0). Writes (sigma[4], SF[4])
|
||
// to producer_step_scratch_buf[103..111).
|
||
// apply_pearls_ad_kernel then smooths via Pearls A+D into ISV[210..218).
|
||
"pearl_3_sigma_kernel.cu",
|
||
// SP5 Task A3 (2026-05-01): Pearl 2 per-branch loss budget controller.
|
||
// Single-block 4-thread kernel reads Q_VAR_PER_BRANCH (Pearl 1's
|
||
// q_branch_stats, ISV[222..226)) + NOISY_SIGMA (Pearl 3 output,
|
||
// ISV[210..214)); computes per-branch flatness = clamp(var_q/(σ²+EPS_DIV),
|
||
// 0, 1) and derives (budget_c51, budget_iqn, budget_cql=0, budget_ens,
|
||
// flatness) per branch. Writes 20 floats to scratch[111..131).
|
||
// apply_pearls_ad_kernel then smooths via Pearls A+D into ISV[190..210).
|
||
// Must run AFTER pearl_1_atom (Q_VAR) AND pearl_3_sigma (NOISY_SIGMA).
|
||
"pearl_2_budget_kernel.cu",
|
||
// SP5 Task A4 (2026-05-01): per-group gradient cosine similarity + L2 norm.
|
||
// Single-block 8-thread kernel (one per param group: DqnTrunk/Value/Branches/
|
||
// IQN/IqlHigh/IqlLow/Attn/Curiosity). Reads grad_buf + grad_prev_buf;
|
||
// writes cosine_sim[8]→scratch[131..139) and l2_norm[8]→scratch[139..147).
|
||
// Also writebacks grad_curr → grad_prev_buf for the next step's comparison.
|
||
"grad_cosine_sim_kernel.cu",
|
||
// SP5 Task A4 (2026-05-01): Pearl 4 per-param-group Adam β1, β2, ε.
|
||
// Single-block 8-thread kernel reads cosine_sim + l2_norm from scratch;
|
||
// writes β1[8]→scratch[147..155), β2[8]→scratch[155..163), ε[8]→scratch[163..171).
|
||
// apply_pearls_ad_kernel then smooths via Pearls A+D (ALPHA_META=5e-4)
|
||
// into ISV[226..250). Must run AFTER grad_cosine_sim_kernel.
|
||
"pearl_4_adam_hparams_kernel.cu",
|
||
// SP5 Task A5 (2026-05-01): auxiliary per-branch Q-distribution skew +
|
||
// kurtosis reducer. Single-block 4-thread kernel reads save_q_online [B × 13];
|
||
// computes per-branch (skew, excess kurtosis) via two-pass central moments;
|
||
// writes skew[4]→scratch[171..175), ex_kurt[4]→scratch[175..179).
|
||
// EPS_DIV=1e-12 Invariant 1 anchor. Structural envelopes [-3,+3]/[-3,+30].
|
||
"q_skew_kurtosis_kernel.cu",
|
||
// SP5 Task A5 (2026-05-01): Pearl 5 per-branch IQN τ schedule controller.
|
||
// Single-block 4-thread kernel reads skew from scratch[skew_idx_base..+4);
|
||
// computes τ = default + skew × SKEW_SHIFT=0.05 (Invariant 1 anchor),
|
||
// clamped to [0.01, 0.99]; writes τ[4 branches × 5 quantiles]→scratch[179..199).
|
||
// apply_pearls_ad_kernel then smooths via Pearls A+D (ALPHA_META=1e-3)
|
||
// into ISV[IQN_TAU_BASE=250..270). Must run AFTER q_skew_kurtosis_kernel.
|
||
"pearl_5_iqn_tau_kernel.cu",
|
||
// SP5 Task A6 (2026-05-01): Pearl 6 cross-fold-persistent Kelly cap signals.
|
||
// Single-block 6-thread kernel (one per ISV output slot [280..286)).
|
||
// Reads portfolio_state [n_envs * PS_STRIDE]; writes directly to ISV[280..286).
|
||
// CROSS-FOLD-PERSISTENT: NOT in fold-reset registry. In-kernel EWMA alpha=0.01
|
||
// (Invariant 1 anchor). Cumulative-via-max() for sample_count slot (s==3).
|
||
// No apply_pearls call (sentinel incompatible with cross-fold persistence).
|
||
// Bayesian priors match kelly_cap_update_kernel.cu: prior_wins=2, prior_losses=2,
|
||
// prior_sum_wins=0.01, prior_sum_losses=0.01 (Invariant 1 anchors).
|
||
"pearl_6_kelly_kernel.cu",
|
||
// SP5 Task A7 (2026-05-01): Pearl 8 per-direction trail-stop distance.
|
||
// 4-thread single-block kernel (one per direction: Short=0, Hold=1, Long=2, Flat=3).
|
||
// Reads features[bar_idx * market_dim + 9] (ATR_NORM) and denormalizes via
|
||
// the canonical fxcache scheme (atr_abs = max(0.01, exp(atr_norm*16-7))).
|
||
// Short and Long receive 2× ATR; Hold and Flat receive EPS_CLAMP_FLOOR=1.0.
|
||
// Layer A simplification: Short and Long share the same ATR (direction-conditional
|
||
// ATR tracking deferred to Layer C). Writes 4 floats to scratch_buf[199..203).
|
||
// Followed by 4 apply_pearls_ad_kernel calls → ISV[TRAIL_DIST_PER_DIR_BASE=270..274).
|
||
"pearl_8_trail_kernel.cu",
|
||
// SP5 Task A8 (2026-05-01): Pearl 1-ext per-branch C51 num_atoms from per-branch v_half.
|
||
// 4-thread single-block kernel (one per branch). Reads ISV[ATOM_V_HALF_BASE=178..182)
|
||
// (populated by Pearl 1 in Task A1). Threshold cascade: v_half < 0.1 → 64 atoms,
|
||
// v_half < 1.0 → 32 atoms, v_half >= 1.0 → 16 atoms.
|
||
// Writes 4 floats to scratch_buf[203..207) = scratch[SCRATCH_PEARL_1_EXT_NUM_ATOMS].
|
||
// Followed by 4 apply_pearls_ad_kernel calls → ISV[ATOM_NUM_ATOMS_BASE=274..278).
|
||
"pearl_1_ext_num_atoms_kernel.cu",
|
||
// SP7 Task 4 (2026-05-03): per-branch loss-balance controller for
|
||
// CQL and C51 budgets. Single-block 8-thread kernel reads
|
||
// grad_decomp pinned slots (IQN/CQL_SX/C51 at offsets 0/24/36 bytes),
|
||
// FLATNESS_BASE (206..210), prior budgets (190..194 + 198..202),
|
||
// and prior Wiener state (297..313). Writes 24 floats to scratch
|
||
// (budgets[8] + diff_var[8] + sample_var[8]). apply_pearls_ad_kernel
|
||
// then smooths via Pearls A+D (Wiener-α derived per-step from the
|
||
// state slots themselves) into ISV[BUDGET_{CQL,C51}_BASE] and
|
||
// ISV[LB_*_VAR_*_BASE]. Replaces Pearl 2's hardcoded-zero CQL
|
||
// budget output and floor-pinned C51 budget.
|
||
"loss_balance_controller_kernel.cu",
|
||
// SP8 Fix 36 (2026-05-03): GPU-only train_active_frac canary
|
||
// producer. Single-block, single-thread kernel reads
|
||
// `monitoring_summary[5..17)` (the 12-bin action_counts block from
|
||
// monitoring_reduce) and computes `(Short + Long) / total`. Writes
|
||
// 1 float to `producer_step_scratch_buf[SCRATCH_TRAIN_ACTIVE_FRAC=250]`
|
||
// and chains 1 apply_pearls_ad_kernel call → ISV[TRAIN_ACTIVE_FRAC_INDEX=321].
|
||
// Replaces the host-side compute previously at
|
||
// `training_loop.rs:3392-3396` per `feedback_no_cpu_compute_strict.md`.
|
||
"train_active_frac_compute_kernel.cu",
|
||
// SP8 Fix 36 (2026-05-03): ISV-driven adaptive MAX_BUDGET producer
|
||
// for the loss-balance controller. Single-block, 8-thread kernel
|
||
// (2 heads × 4 branches) reads `ISV[TRAIN_ACTIVE_FRAC_INDEX]`
|
||
// (Pearls A+D smoothed canary) and computes per-(head, branch) cap
|
||
// via linear interpolation `FLOOR + active_frac × (CEIL − FLOOR)`
|
||
// (FLOOR=0.01, CEIL=1.0 — Invariant 1 numerical-stability anchors
|
||
// per `feedback_isv_for_adaptive_bounds.md`). Writes 8 floats to
|
||
// `producer_step_scratch_buf[SCRATCH_LB_MAX_BUDGET_*=251..259)` and
|
||
// chains 8 apply_pearls_ad_kernel calls → ISV[LB_MAX_BUDGET_{CQL,C51}_BASE..+4).
|
||
// Replaces the hardcoded `MAX_BUDGET=1.0f` constant in
|
||
// `loss_balance_controller_kernel.cu` per
|
||
// `pearl_controller_anchors_isv_driven.md`.
|
||
"loss_balance_max_budget_compute_kernel.cu",
|
||
// SP9 Fix 37 (2026-05-03): Kelly cold-start warmup floor — fully
|
||
// ISV-driven. Four producer kernels per
|
||
// `pearl_cold_start_exit_signal_or.md` and
|
||
// `pearl_controller_anchors_isv_driven.md`:
|
||
// 1. eval_intent_dist_compute — replaces DtoH read_eval_intent_magnitude_distribution
|
||
// 2. intent_eval_divergence_compute — behavioral-axis numerator
|
||
// 3. q_var_mag_ema_compute — base_floor self-relative baseline
|
||
// 4. kelly_warmup_floor_compute — main producer (combines all axes)
|
||
// All 4 are single-block kernels writing to producer_step_scratch_buf,
|
||
// chained through apply_pearls_ad_kernel for Pearl A sentinel-bootstrap
|
||
// + Pearl D Wiener-α steady-state smoothing. Per Fix 37.1
|
||
// (2026-05-03): the 3 EMA target updaters that previously sat between
|
||
// q_var_mag_ema and kelly_warmup_floor are deleted — the targets at
|
||
// ISV[333..336) are Invariant-1 numerical anchors (sample-sufficiency
|
||
// count, healthy divergence ratio, fold-temporal threshold) written
|
||
// ONCE at constructor time per `feedback_isv_for_adaptive_bounds.md`,
|
||
// not EMA-tracked. Pearl A sentinel-bootstrap immediately saturated
|
||
// the EMAs to the first observation, which made conf=1.0 in epoch 1
|
||
// (`floor=0`) — defeating the cold-start mechanism. Constructor-write
|
||
// gives a fixed denominator so the ratio yields meaningful confidence.
|
||
"eval_intent_dist_compute_kernel.cu",
|
||
"intent_eval_divergence_compute_kernel.cu",
|
||
"q_var_mag_ema_compute_kernel.cu",
|
||
"kelly_warmup_floor_compute_kernel.cu",
|
||
// SP11 Fix 39 (2026-05-04, Task A1): three canary producer kernels
|
||
// for the reward-as-controlled-subsystem chain. Each is a single-
|
||
// block producer chained through `apply_pearls_ad_kernel` for
|
||
// Pearl A sentinel-bootstrap + Pearl D Wiener-α steady-state
|
||
// smoothing. All three write to slots in [350..360) which no
|
||
// consumer reads yet (Layer A is additive; consumer migration
|
||
// lands atomically in Layer B). See spec §3.3 (canary slots) and
|
||
// §5 (producer kernel listing).
|
||
// 1. val_sharpe_delta_compute → ISV[VAL_SHARPE_DELTA_EMA_INDEX=350,
|
||
// VAL_SHARPE_VAR_EMA_INDEX=351]
|
||
// 2. saboteur_engagement_compute → ISV[SABOTEUR_ENGAGEMENT_RATE_INDEX=358]
|
||
// 3. reward_component_mag_ratio → ISV[REWARD_COMPONENT_MAG_RATIO_BASE=352..358)
|
||
// ∪ ISV[PNL_REWARD_MAGNITUDE_EMA_INDEX=359]
|
||
"val_sharpe_delta_compute_kernel.cu",
|
||
"saboteur_engagement_compute_kernel.cu",
|
||
"reward_component_mag_ratio_compute_kernel.cu",
|
||
// SP11 Fix 39 B1b fix-up (2026-05-04): popart-component-specific
|
||
// magnitude EMA producer. Reads the per-bar `r_popart` buffer
|
||
// populated by `experience_env_step` and block-tree-reduces
|
||
// mean(|x|) into ISV[POPART_COMPONENT_MAG_EMA_INDEX=360] via
|
||
// chained Pearls A+D. Resolves the slot-63 PopArt-input vs
|
||
// popart-component-magnitude overload that B1b smoke surfaced
|
||
// (contamination → controller over-weighed popart → 10× sharpe
|
||
// drop). Spec §4 amendment "Why slot 360".
|
||
"popart_component_ema_kernel.cu",
|
||
// SP11 Fix 39 (2026-05-04, Task A2): reward-subsystem controller +
|
||
// SimHash novelty buffer. Three new kernels in this layer:
|
||
// 1. reward_subsystem_controller — main producer (5 canaries → 10
|
||
// outputs in scratch). Chained `apply_pearls_ad_kernel` smooths
|
||
// into ISV[340..350) per spec §3.4.1 (output smoothing).
|
||
// 2. novelty_simhash_proj_init — one-shot Philox-driven init of
|
||
// the 42×16 projection matrix at trainer construct (CPU is
|
||
// read-only per `feedback_no_cpu_forwards`).
|
||
// 3. novelty_simhash — lookup + update (race-tolerated
|
||
// per `feedback_no_atomicadd`; under-counts bias novelty UP,
|
||
// the safe direction). Lookup/update launchers exist in A2 but
|
||
// are wired into the replay path in B1 — A2 just builds the
|
||
// kernels and reset registry contract end-to-end.
|
||
"reward_subsystem_controller_kernel.cu",
|
||
"novelty_simhash_proj_init_kernel.cu",
|
||
"novelty_simhash_kernel.cu",
|
||
// SP7 (2026-05-03): GPU dispatch for cql/c51 budget consumer.
|
||
// Eliminates capture-time host-branch freeze in
|
||
// `compute_adaptive_budgets`. The dispatch kernel resolves
|
||
// `(active >= 0.5) ? controller_budget : bootstrap` on-device every
|
||
// step (8 threads — 2 heads × 4 branches), writing trunk_mean +
|
||
// per-branch correction to a 10 f32 mapped-pinned buffer. The same
|
||
// cubin houses two new SAXPY/scale variants that read `alpha` from
|
||
// a device pointer (instead of scalar arg), so the captured
|
||
// `apply_c51_budget_scale` / `apply_cql_saxpy` chain reads runtime
|
||
// ISV-derived values instead of step-0 frozen scalars. See
|
||
// `feedback_no_cpu_compute_strict.md` and audit Fix 33.
|
||
"consume_lb_budget_kernel.cu",
|
||
// SP5 Layer D Task D1 (2026-05-02): PnL aggregation producer.
|
||
// First of 3 Layer D producers replacing the host-side trade-PnL
|
||
// aggregation in `compute_epoch_financials` per
|
||
// `feedback_no_cpu_compute_strict.md`. Single-block 256-thread
|
||
// shmem-reduce kernel (no atomicAdd) reads per-trade event arrays;
|
||
// outputs cumulative PnL, mean, population variance, and max
|
||
// drawdown to scratch[SCRATCH_PNL_AGG_BASE=207..211). Chained
|
||
// `apply_pearls_ad_kernel` smooths via Pearls A+D into ISV[286..290).
|
||
// Producer-only — D4 wires the consumer in the atomic Layer D
|
||
// commit per `feedback_no_partial_refactor.md`.
|
||
"pnl_aggregation_kernel.cu",
|
||
// SP5 Layer D Task D2 (2026-05-02): LearningHealth composition producer.
|
||
// Second of 3 Layer D producers replacing the host-side
|
||
// `NormalizedComponents::from_raw` + `compose` chain in
|
||
// `learning_health.rs` (called from `training_loop.rs:2658-2745`)
|
||
// per `feedback_no_cpu_compute_strict.md`. Single-block 1-thread
|
||
// arithmetic kernel: 7 raw inputs by value → 7 in-register
|
||
// smoothstep normalisations → weighted sum into a scalar health
|
||
// score; writes 4 scratch slots (composed score + q_gap_norm +
|
||
// q_var_norm + grad_norm_norm) at scratch[SCRATCH_HEALTH_COMP_BASE
|
||
// =211..215). Chained `apply_pearls_ad_kernel` smooths via Pearls
|
||
// A+D into ISV[290..294). Producer-only — D4 wires the consumer
|
||
// in the atomic Layer D commit per
|
||
// `feedback_no_partial_refactor.md`.
|
||
"health_composition_kernel.cu",
|
||
// SP5 Layer D Task D3 (2026-05-02): training-metrics EMA producer.
|
||
// Third (and final) Layer D producer replacing the host-side
|
||
// adversarial-regime + adaptive-DSR EMA arithmetic block at
|
||
// `training_loop.rs:5039-5113` per `feedback_no_cpu_compute_strict.md`.
|
||
// Single-block 3-thread kernel — one thread per metric:
|
||
// tid=0 → training_sharpe_ema (adaptive α, host-tracked sentinel),
|
||
// tid=1 → max_dd_ema (fixed α=0.1, sentinel<1e-12),
|
||
// tid=2 → low_dd_ratio (fixed α=0.15, depends on tid=1's
|
||
// output via __syncthreads()).
|
||
// Writes 3 scratch slots at scratch[SCRATCH_TRAINING_METRICS_EMA_BASE
|
||
// =215..218). Chained `apply_pearls_ad_kernel` (3 launches,
|
||
// ALPHA_META=1e-3) smooths via Pearls A+D into ISV[294..297).
|
||
// Producer-only — D4 wires the consumer in the atomic Layer D
|
||
// commit per `feedback_no_partial_refactor.md`.
|
||
"training_metrics_ema_kernel.cu",
|
||
// SP4 GPU-only Pearls A+D applicator (2026-05-01). Single-thread
|
||
// device-side loop applies Pearls A+D to N consecutive ISV slots
|
||
// (each with its own Wiener-state triple) on the producer's
|
||
// stream. Replaces the pre-refactor host-side `apply_pearls_to_slot`
|
||
// helper that broke CUDA Graph capture at `aux_label_scale_ema`
|
||
// mid-step (smoke smoke-test-v9kjv: CUDA_ERROR_STREAM_CAPTURE_UNSUPPORTED).
|
||
// Consumed by all 11 Pearls-using launchers + the inline
|
||
// aux_heads_forward Step 2b call site. Constants (ALPHA_META,
|
||
// EPS_DIV) mirror sp4_wiener_ema.rs's host-reference for the
|
||
// GPU-vs-host oracle test.
|
||
"apply_pearls_kernel.cu",
|
||
// SP13 Phase 0a (2026-05-04): aux-head directional-accuracy
|
||
// reducer. Single-block tree-reduce kernel scoring the aux next-
|
||
// bar regression head's prediction sign against the next-bar
|
||
// return label sign. Four shared-memory int arrays
|
||
// (correct/pos_pred/pos_label/valid) reduce in lockstep — no
|
||
// atomicAdd per `feedback_no_atomicadd.md`. Outputs three
|
||
// batch-level fractions (dir_acc, pos_pred_frac, pos_label_frac)
|
||
// into a 3-element mapped-pinned buffer. P0a reads regression-
|
||
// mode aux scalar; Layer B swaps the source pointer to
|
||
// `softmax[1] - softmax[0]` after the head becomes 2-class
|
||
// softmax. Consumed by `launch_aux_dir_acc_reduce` in
|
||
// `gpu_dqn_trainer.rs`.
|
||
"aux_dir_acc_reduce_kernel.cu",
|
||
// SP13 v3 Phase 0a P0a.T3 (2026-05-04): Hold-rate observer.
|
||
// Single-block tree-reduce kernel counting per-bar Hold-action
|
||
// picks (decoded from packed `batch_actions[i]` via
|
||
// `dir = action_idx / (NUM_MAGNITUDES * NUM_ORD * NUM_URG)` —
|
||
// P0a.T4 ABI alignment, 2026-05-04) and writing the fraction
|
||
// `count(Hold) / batch_size` to a 1-element mapped-pinned buffer.
|
||
// One shared-memory int array reduces in lockstep — no atomicAdd
|
||
// per `feedback_no_atomicadd.md`. The fixed-α EMA applicator
|
||
// smooths the observation into ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382];
|
||
// the host-side controller reads ISV[382] + ISV[381] and writes the
|
||
// priced Hold cost back to ISV[HOLD_COST_INDEX=380]. The reward
|
||
// composition site in `experience_kernels.cu` then subtracts the
|
||
// cost on every Hold-action bar — pricing the action so the
|
||
// policy uses Hold deliberately rather than as a free CQL-bias-
|
||
// anchored default. Replaces v2's atomic Hold elimination per
|
||
// the v3 reframe (preserves the 4-way action space — no cross-
|
||
// crate cascade). Consumed by `launch_hold_rate_observer` in
|
||
// `gpu_dqn_trainer.rs`.
|
||
"hold_rate_observer_kernel.cu",
|
||
// SP13 Phase 0a P0a.T4 (2026-05-04): fixed-α EMA applicator.
|
||
// Sibling of `apply_pearls_kernel` — needed because Pearls A+D
|
||
// collapse two EMAs of the same signal to identical values
|
||
// (variance ratio is α-independent), which would deadlock the
|
||
// SP13 stagnation detector that compares slots 373 (α=0.3) vs
|
||
// 374 (α=0.05). Per-thread first-observation-sentinel bootstrap
|
||
// (caller passes the slot's registry sentinel — 0.5 for dir-acc,
|
||
// 0.0 for hold-rate) plus a fixed-α blend. No atomicAdd per
|
||
// `feedback_no_atomicadd.md`; pure GPU compute per
|
||
// `feedback_no_cpu_compute_strict.md`. Consumed by
|
||
// `launch_apply_fixed_alpha_ema` in `gpu_dqn_trainer.rs`.
|
||
"apply_fixed_alpha_ema_kernel.cu",
|
||
// SP13 Phase 0a P0a.T4 (2026-05-04): aux-head per-bar prediction
|
||
// → ISV[AUX_DIR_PREDICTION_INDEX=375] tanh-bounded scalar
|
||
// producer. Single-block tree-reduce reads the captured-graph
|
||
// `aux_nb_pred_buf [B]` tile and writes
|
||
// `mean(tanh(aux_pred[i]))` ∈ [-1, +1] to the SHARED ISV slot 375
|
||
// (per-step state, not an EMA — slot is overwritten each launch).
|
||
// tanh squash bounds the scalar so the downstream Q-head
|
||
// consumer sees a well-conditioned signal regardless of
|
||
// `label_scale` magnitude. No atomicAdd per
|
||
// `feedback_no_atomicadd.md`; pure GPU compute per
|
||
// `feedback_no_cpu_compute_strict.md`. Consumed by
|
||
// `launch_aux_pred_to_isv_tanh` in `gpu_dqn_trainer.rs`.
|
||
"aux_pred_to_isv_tanh_kernel.cu",
|
||
// SP13 Layer B Commit B1.1b (2026-05-05): producer kernel for
|
||
// the aux next-bar sign label `[total] i32` ring column. Replaces
|
||
// the B0 `alloc_zeros::<i32>(total)` placeholder in
|
||
// `gpu_experience_collector.rs::collect_experiences_gpu` with a
|
||
// real classifier reading `targets[bar*6+2]` (raw_close column)
|
||
// at `bar` and `bar+lookahead`, writing -1 (skip if window runs
|
||
// past `total_bars`), 0 (down/flat: `p_fut <= p_now`), or 1
|
||
// (up: `p_fut > p_now`) per experience. Pure per-thread O(1)
|
||
// map — no atomicAdd / no reduction per
|
||
// `feedback_no_atomicadd.md`. Consumer is the K=2 softmax CE
|
||
// head wired in B1.1a (`aux_next_bar_loss_reduce` +
|
||
// `aux_next_bar_backward` +`aux_dir_acc_reduce_kernel` +
|
||
// `aux_pred_to_isv_tanh_kernel`); without B1.1b every label is
|
||
// 0, so the head trains on "everything is class 0" — the
|
||
// known degraded behavior between B1.1a and B1.1b that this
|
||
// kernel resolves.
|
||
"aux_sign_label_kernel.cu",
|
||
// SP14 β-migration step 3 (Option B, 2026-05-07): per-rollout-step
|
||
// variant of `aux_sign_label_kernel`. The trajectory kernel writes
|
||
// the entire `[total]` ring at the end of `collect_experiences_gpu`;
|
||
// this kernel writes the current step's `[n_episodes]` slice using
|
||
// `bar = episode_starts[ep] + t` so the collector-native EGF
|
||
// producer chain has fresh labels every rollout step. Same per-thread
|
||
// O(1) map; same lookahead window semantics; same skip sentinel.
|
||
"aux_sign_label_per_step_kernel.cu",
|
||
// SP14 Layer C Phase C.4b (2026-05-08): adaptive aux prediction
|
||
// horizon producer. Single-thread kernel writing
|
||
// `ISV[AUX_PRED_HORIZON_BARS_INDEX=450]` from
|
||
// `ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451]`. Pearl-A first-observation
|
||
// bootstrap (sentinel 60.0 → direct replacement); Wiener-α slow
|
||
// EMA blend thereafter (α=0.01 fixed; no target-variance EMA
|
||
// available in C.4b). Per-epoch boundary launch — H is
|
||
// slow-moving. Plan: docs/superpowers/plans/2026-05-07-sp14-layer-c-separate-aux-trunk.md §C.4b.
|
||
"aux_horizon_update_kernel.cu",
|
||
// SP14 Layer C Phase C.4b (2026-05-08): avg winning hold time
|
||
// EMA producer. Single-block kernel that sweeps the per-epoch
|
||
// `hold_at_exit_per_sample` + `trade_profitable_per_sample`
|
||
// buffers (populated by `unified_env_step_core` in
|
||
// `experience_kernels.cu`) and writes the EMA-blended mean to
|
||
// `ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451]` for downstream
|
||
// consumption by `aux_horizon_update_kernel`. Block-tree-reduce
|
||
// (no atomicAdd) per `feedback_no_atomicadd.md`. Pearl-A
|
||
// first-observation bootstrap; α=0.05 EMA thereafter. Plan: §C.4b.
|
||
"avg_win_hold_time_update_kernel.cu",
|
||
// SP14 Layer B Task B.3 (2026-05-05): Earned Gradient Flow
|
||
// producer kernel. Per-step computes the K=4↔K=2 mapped argmax
|
||
// mismatch between the Q-head's 4-way direction action
|
||
// (DIR_SHORT/DIR_HOLD/DIR_LONG/DIR_FLAT) and the aux head's
|
||
// 2-way next-bar prediction (down/up), masking Hold/Flat (no
|
||
// new directional commitment). Updates fast + slow EMAs of
|
||
// the disagreement rate (ISV[383], ISV[384]) plus a Welford-
|
||
// style variance EMA (ISV[389]) in a single launch — the
|
||
// variance feeds the adaptive k_q sigmoid steepness consumed
|
||
// by `alpha_grad_compute_kernel` in B.4. Block tree-reduce on
|
||
// numerator + count separately; no atomicAdd per
|
||
// `feedback_no_atomicadd.md`. Pearl-A first-observation
|
||
// bootstrap on the EMAs (sentinel 0.5 → batch_mean direct
|
||
// replacement) per `pearl_first_observation_bootstrap.md`.
|
||
// Slot indices hardcoded inside the kernel via `#define` —
|
||
// must match `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs`
|
||
// (currently 383, 384, 389; the layout fingerprint regression
|
||
// test catches drift).
|
||
"q_disagreement_update_kernel.cu",
|
||
// SP14 Layer C Phase C.1 (2026-05-08): cubin entries for
|
||
// alpha_grad_compute_kernel.cu, gradient_hack_detect_kernel.cu,
|
||
// and sp14_scale_wire_col_kernel.cu deleted atomically with
|
||
// those kernel files per `feedback_no_legacy_aliases` and
|
||
// `feedback_no_partial_refactor`. Coupling B (α-gated backward
|
||
// wire) eliminated; Coupling A (`dir_concat_qaux_kernel.cu`,
|
||
// forward feature wire) preserved below.
|
||
// SP14 Layer B Task B.6 (2026-05-05): Earned Gradient Flow —
|
||
// pre-SGEMM concat kernel. Forward-wires aux_softmax_diff into
|
||
// the direction Q-head input. Mirrors mag_concat_qdir precedent
|
||
// (experience_kernels.cu:4560). Concatenates [h_s2 ; softmax_diff]
|
||
// into scratch buffer [B, SH2+1]: positions 0..SH2 are a direct
|
||
// copy of h_s2[b, 0..SH2]; position SH2 is softmax[b,1]-softmax[b,0]
|
||
// (up_prob - down_prob). The diff is structurally bounded to [-1,+1]
|
||
// by softmax composition per pearl_bounded_modifier_outputs_require_
|
||
// structural_activation — no runtime clamp needed. Pure per-thread
|
||
// map: one thread per output element, no reduction, no atomicAdd
|
||
// (per feedback_no_atomicadd.md). Launch-order constraint: aux head
|
||
// forward MUST complete before this kernel reads aux_nb_softmax_buf
|
||
// (enforced by graph capture orchestrator in B.10/B.11). Does NOT
|
||
// read or write ISV slots — purely data-movement.
|
||
"dir_concat_qaux_kernel.cu",
|
||
// SP14 Layer C Phase C.1 (2026-05-08): sp14_scale_wire_col_kernel
|
||
// cubin entry deleted atomically with the kernel file —
|
||
// Coupling B (α-gated backward wire) eliminated.
|
||
// SP15 Phase 1.1 (2026-05-06): unified per-bar sharpe. Single-block,
|
||
// BLOCK=256 kernel computing mean/std/sharpe via 2-pass shared-memory
|
||
// tree-reduce — no atomicAdd. Used by both train and val paths with
|
||
// identical formula and identical window definition. Replaces the
|
||
// SP14-era sharpe_ema (per-batch EMA) vs sharpe_annualised
|
||
// (val × sqrt(525600)) split. Annualisation is the host-side caller's
|
||
// responsibility (sharpe_annualised = out[2] × sqrt(N_bars_per_year)).
|
||
"sharpe_per_bar_kernel.cu",
|
||
// SP15 Phase 1.2 (2026-05-06): cost-net sharpe. Single-block,
|
||
// BLOCK=256 kernel computing mean/std/sharpe of (pnl - cost) per
|
||
// spec §6.2 per-side semantics: commission_per_rt × rt_ind +
|
||
// half_spread × |pos| × side_ind + ofi_lambda × |pos| × |ofi| ×
|
||
// side_ind. Charges full spread per RT (entry + exit). Reads
|
||
// ofi_lambda from ISV[OFI_IMPACT_LAMBDA_INDEX=407]; emits
|
||
// mean cost-per-bar to ISV[COST_PER_BAR_AVG_INDEX=408]. Same
|
||
// kernel reads LobBar fields from synthetic markets and prod
|
||
// fxcache LOB (dev/prod parity per Q3).
|
||
"cost_net_sharpe_kernel.cu",
|
||
// SP15 Phase 1.3 (2026-05-06) + 1.3.b-followup (2026-05-07):
|
||
// per-env drawdown state tracking. Per-env grid `[n_envs, 1, 1]`,
|
||
// one thread per env: each thread reads PS_PEAK_EQUITY (slot 7)
|
||
// and PS_PREV_EQUITY (slot 9) from its env's row of the position
|
||
// state buffer and writes 6 scalars to the per-env tile output
|
||
// `dd_state_per_env[env*6 + k]` for k ∈ {DD_CURRENT, DD_MAX,
|
||
// DD_RECOVERY_BARS, DD_PERSISTENCE, CALMAR (floored max_dd
|
||
// denominator), DD_PCT}. The 6 ISV scalar slots [401..407) are
|
||
// populated by the separate `dd_state_reduce_kernel` below
|
||
// (mean-aggregate across envs for HEALTH_DIAG diagnostic) plus
|
||
// the new `DD_PERSISTENCE_MAX_INDEX = 443` slot (max-aggregate
|
||
// for the plasticity-injection trigger). Path A → Path B
|
||
// migration: the original Phase 1.3.b kernel chose env-0 as the
|
||
// canonical DD observable and wrote ISV scalars directly; the
|
||
// followup-A redesign threads each env's actual DD context
|
||
// through the reward shaping per spec §6.3.
|
||
"dd_state_kernel.cu",
|
||
// SP15 Phase 1.3.b-followup (2026-05-07): DD state reduction.
|
||
// Single-block tree-reduce (no atomicAdd) of the per-env tile
|
||
// `dd_state_per_env[n_envs * 6]` into 6 scalar ISV slots
|
||
// [401..407) (mean-aggregate across envs — smooth across-env
|
||
// summary for HEALTH_DIAG diagnostic) plus the new
|
||
// DD_PERSISTENCE_MAX_INDEX = 443 slot (max-aggregate, consumed
|
||
// by `plasticity_injection_kernel` so the global advantage-head
|
||
// reset fires when ANY env's persistence exceeds the threshold:
|
||
// one set of advantage weights → global firing semantics →
|
||
// max-aggregate is the only correct rule). BLOCK=256 with the
|
||
// strided initial accumulation pattern handles n_envs up to
|
||
// production sizes (32768 envs on H100). Layout fingerprint
|
||
// marker `DD_STATE_PER_ENV=sp15_phase_1_3_b_followup` enforces
|
||
// synchrony between the producer's tile layout (in
|
||
// `dd_state_kernel.cu`) and this consumer's reads. Single
|
||
// kernel per cubin (mirrors the SP15 1.3 / 1.4 / 3.X
|
||
// single-source-to-cubin convention).
|
||
"dd_state_reduce_kernel.cu",
|
||
// SP15 Phase 1.4 + Wave 3a (2026-05-06): 4 constant-policy
|
||
// counterfactual baselines per spec §6.4. Single source file with
|
||
// 4 `extern "C" __global__` symbols (baseline_buyhold_kernel,
|
||
// baseline_hold_only_kernel, baseline_naive_momentum_kernel,
|
||
// baseline_naive_reversion_kernel) sharing one cubin per the 1:1
|
||
// source-to-cubin convention — the four launchers in
|
||
// gpu_dqn_trainer.rs all load the same cubin module and resolve
|
||
// different `get_function()` symbols. Each kernel: single-block
|
||
// BLOCK=256 with a templated 2-pass shared-memory tree-reduce
|
||
// (no atomicAdd). Reads price/half_spread/ofi arrays + commission
|
||
// scalar; computes constant-policy or last-bar-return-based
|
||
// actions; writes per-window `[mean, std, raw_sharpe]` to a
|
||
// caller-provided output buffer (Wave 3a contract change — was
|
||
// ISV-scalar slots 409/410/412/416 pre-Wave-3a). Trunk-shared
|
||
// baselines (random_dir_kelly slot 411, aux_only slot 413,
|
||
// mag_quarter_fixed slot 414, trail_only slot 415) are out of
|
||
// scope for this commit — they need partial-policy-forward
|
||
// access from the main eval pass and are deferred to Task 1.4.b.
|
||
"baseline_kernels.cu",
|
||
// SP15 Wave 3a (2026-05-06): post-loop derivation of per-bar
|
||
// position state from `actions_history_buf` (factored ints
|
||
// recorded by env_step). Single-block-per-window launch with one
|
||
// thread per block (the position state machine is sequential per
|
||
// window — no in-window parallelism). Emits 3 f32 streams per
|
||
// (window, bar): position_history (-1/0/+1), side_ind (1.0 on
|
||
// position change), rt_ind (1.0 on transition-to-flat from
|
||
// non-flat). Consumed by Wave 3b's cost-net sharpe path; uses
|
||
// `action_decoding_helpers.cuh::factored_action_to_position`
|
||
// single-source-of-truth for the action→direction→position
|
||
// mapping (mirrors `trade_physics.cuh::decode_direction_4b`).
|
||
"position_history_derivation_kernel.cu",
|
||
// SP15 Wave 4.1a (2026-05-06): the standalone
|
||
// `dd_pct_concat_kernel.cu` from Phase 1.5 was DELETED. The
|
||
// standalone kernel was bottleneck-incompatible — it operated on
|
||
// raw `[B, state_dim_padded=128]` state but production trunk
|
||
// consumes `[B, s1_input_dim=102]` post-bottleneck. Wave 4.1a
|
||
// adds a bottleneck-aware extension `bn_tanh_concat_dd_kernel`
|
||
// INSIDE `dqn_utility_kernels.cu` that fuses the dd_pct append
|
||
// into the same launch as the bn_tanh + portfolio concat
|
||
// (output shape `[B, bn_dim + portfolio_dim + 1]`); Wave 4.1b
|
||
// lands the consumer migration (s1_input_dim 102→103, GRN w_s1
|
||
// reshape, 3 forward + 3 backward call-site updates).
|
||
// SP15 Wave 2 (2026-05-06): per-step ISV-driven α producer.
|
||
// Replaces the pre-Wave-2 `r_quality_discipline_split_kernel.cu`
|
||
// which contained both the scalar composer
|
||
// (`r_quality_discipline_split_kernel`) and the producer
|
||
// (`alpha_split_producer_kernel`). Wave 2 retains ONLY the
|
||
// producer; the composer responsibility moves into the fused
|
||
// post-modifier `compute_sp15_final_reward_kernel.cu` per-(i,t)
|
||
// kernel that reads ALPHA_SPLIT_INDEX=417 directly. The producer
|
||
// reads grad-norm slots 418/419 + the `sp15_alpha_warm_count`
|
||
// mapped-pinned scratch buffer, writes ALPHA_SPLIT_INDEX clamped
|
||
// to [0.05, 0.95] post-warmup, and OWNS the warm-count
|
||
// increment (Q4 resolution: the fused composer would over-tick
|
||
// by 2*N*L per step if it owned the increment).
|
||
"alpha_split_producer_kernel.cu",
|
||
// SP15 Wave 2 (2026-05-06): fused post-SP11 reward-axis composer.
|
||
// Single source file with one `extern "C" __global__` symbol
|
||
// (`compute_sp15_final_reward_kernel`) → one cubin → one launcher
|
||
// (`launch_sp15_final_reward`). Per-(i,t) parallel over [N*2*L]
|
||
// slots (on-policy + CF). Reads SP11's `out_rewards[idx]`
|
||
// in-place, applies the full SP15 §8.2 / §9.2 four-stage
|
||
// composition (α-blend split → DD asymmetric reward → quadratic
|
||
// DD penalty → SP12 v3 asymmetric cap via state_layout.cuh
|
||
// macros REWARD_NEG_CAP=-10 / REWARD_POS_CAP=+5) and writes the
|
||
// result back to `out_rewards[idx]`. Skips slots where
|
||
// `slot_completed_normally[idx % (N*L)] == 0` (data-end at
|
||
// experience_kernels.cu line 2142, blown-account at line 2203 —
|
||
// those sentinel rewards 0.0f / -10.0f pass through unmodified
|
||
// per Q3 resolution). Both on-policy and CF slots get the same
|
||
// DD-aware shaping per Q2 resolution. Subsumes the deleted
|
||
// `dd_penalty_kernel.cu` + `dd_asymmetric_reward_kernel.cu` +
|
||
// the deleted scalar `r_quality_discipline_split_kernel`
|
||
// composer; their bodies live as `__device__` helpers in
|
||
// `sp15_reward_axis_helpers.cuh`.
|
||
"compute_sp15_final_reward_kernel.cu",
|
||
// SP15 Phase 3.4 (2026-05-06): regret / opportunity-cost signal
|
||
// = r_discipline content. Single source file with one
|
||
// `extern "C" __global__` symbol (`regret_signal_kernel`) → one
|
||
// cubin → one launcher (`launch_sp15_regret_signal`). Mirrors
|
||
// the Phase 3.3 `dd_penalty_kernel.cu` single-kernel-per-cubin
|
||
// pattern. Reads ISV[REGRET_EMA_INDEX=423] (read-modify-write —
|
||
// EMA self-update via first-observation bootstrap per
|
||
// `pearl_first_observation_bootstrap` + simple exponential decay
|
||
// α=0.05; Wiener-α adaptive smoothing per
|
||
// `pearl_wiener_optimal_adaptive_alpha` is a documented
|
||
// follow-up). Computes `regret_bar = ideal_pnl_missed - cost_t`
|
||
// when `action == Hold (0)` AND
|
||
// `ideal_pnl_missed > cost_t + trail_dist`, else 0. Asymmetric
|
||
// gate: regret fires only when policy was Hold AND the ideal
|
||
// trade cleared the slippage+trail floor. Phase 3.4 lands kernel
|
||
// + launcher + 3 ISV anchor seeds + 3 registry entries + 3
|
||
// dispatch arms only; the host-side reward composer that
|
||
// composes `r_discipline -= LAMBDA_REGRET × REGRET_EMA` is
|
||
// deferred to a follow-up atomic commit per
|
||
// `feedback_no_partial_refactor.md` (matching Phase 1.1-1.3 +
|
||
// 3.1-3.3 precedent).
|
||
"regret_signal_kernel.cu",
|
||
// SP15 Phase 3.5 — `hold_floor_kernel.cu` REMOVED in Phase 3.5.b
|
||
// (2026-05-06): the bounded-sigmoid `hold_floor = α × σ(k ×
|
||
// (entropy − ε₀))` is now an inline `__device__` computation
|
||
// inside `experience_action_select` reading the same ISV slots
|
||
// 426/427/428/429 directly. The standalone kernel was scaffolding
|
||
// for the deferred consumer per `feedback_no_partial_refactor`;
|
||
// when the consumer landed inline, the kernel + launcher +
|
||
// cubin manifest entry became dead code per
|
||
// `feedback_wire_everything_up.md` + `feedback_no_legacy_aliases.md`.
|
||
// ISV slots + state_reset_registry entries remain (still
|
||
// consumed by the inline computation).
|
||
//
|
||
// SP15 Wave 2 (2026-05-06) — `dd_asymmetric_reward_kernel.cu`
|
||
// REMOVED: the Phase 3.5.2 standalone scalar kernel was
|
||
// scaffolding for a deferred consumer per
|
||
// `feedback_no_partial_refactor.md`. Wave 2's layered
|
||
// architecture replaces the deferred consumer with the
|
||
// `compute_sp15_final_reward_kernel.cu` fused post-modifier; the
|
||
// gain-only DD-recovery multiplier body now lives as a
|
||
// `__device__` inline `sp15_dd_asymmetric_reward` helper in
|
||
// `sp15_reward_axis_helpers.cuh` consumed by the fused composer.
|
||
// ISV slots 430 / 431 / 432 + state_reset_registry entries +
|
||
// dispatch arms remain (still read by the helper).
|
||
//
|
||
// SP15 Wave 2 (2026-05-06) — `dd_penalty_kernel.cu` REMOVED:
|
||
// the Phase 3.3 standalone scalar kernel was scaffolding for a
|
||
// deferred consumer per `feedback_no_partial_refactor.md`.
|
||
// Wave 2's layered architecture replaces the deferred consumer
|
||
// with the `compute_sp15_final_reward_kernel.cu` fused
|
||
// post-modifier; the quadratic DD penalty body now lives as a
|
||
// `__device__` inline `sp15_dd_penalty` helper in
|
||
// `sp15_reward_axis_helpers.cuh` consumed by the fused composer.
|
||
// ISV slots 420 / 421 / 422 + state_reset_registry entries +
|
||
// dispatch arms remain (still read by the helper).
|
||
// SP15 Phase 3.5.3 (2026-05-06): cooldown gate — after K
|
||
// consecutive losing trades, force Hold for M bars. K and M
|
||
// ISV-tracked (slots 433/434); initial sentinels K=5, M=20
|
||
// hardcoded in the trainer constructor (ISV-driven K from a
|
||
// running median of consecutive-loss streak lengths via
|
||
// `MEDIAN_STREAK_LENGTH=442` is a documented follow-up
|
||
// sub-task — two-heap algorithm; ISV-driven M from
|
||
// vol_normalizer time-to-mean-reversion is a parallel
|
||
// follow-up). Single source file with one
|
||
// `extern "C" __global__` symbol (`cooldown_kernel`) → one
|
||
// cubin → one launcher (`launch_sp15_cooldown`). Mirrors the
|
||
// Phase 3.3 `dd_penalty_kernel.cu` / 3.4
|
||
// `regret_signal_kernel.cu` / 3.5 `hold_floor_kernel.cu` /
|
||
// 3.5.2 `dd_asymmetric_reward_kernel.cu` single-kernel-per-cubin
|
||
// pattern. Reads ISV[COOLDOWN_K_THRESHOLD_INDEX=433] +
|
||
// ISV[COOLDOWN_M_BARS_INDEX=434], read-modify-writes
|
||
// ISV[COOLDOWN_BARS_REMAINING_INDEX=435] + a separate
|
||
// mapped-pinned `consecutive_losses` scratch buffer
|
||
// (`sp15_cooldown_consecutive_losses`, [1] f32, persistent
|
||
// streak state). Per spec §9.2 (3.5.3) post-amendment-2 fix:
|
||
// streak counter only updates on trade-close events; per-bar
|
||
// non-close calls just decrement the cooldown counter. Phase
|
||
// 3.5.3 lands kernel + launcher + 4 ISV anchor seeds + 1
|
||
// scratch buffer + 5 registry entries + 5 dispatch arms only;
|
||
// the action-selection wiring (force Hold while
|
||
// `cooldown_remaining > 0`) is deferred to a follow-up atomic
|
||
// commit per `feedback_no_partial_refactor.md` (matching Phase
|
||
// 1.1-1.5 + 3.1-3.5.2 precedent — kernel + launcher verify in
|
||
// isolation first via the GPU oracle tests below).
|
||
"cooldown_kernel.cu",
|
||
// SP15 Phase 3.5.4 (2026-05-06): plasticity injection — TWO-STEP
|
||
// recovery (Flat + cooldown) with Kaiming-He reset of last 10%
|
||
// of advantage-head weights. K=DD_PERSISTENCE threshold is
|
||
// ISV-tracked (slot 437); initial sentinel 100.0 hardcoded in
|
||
// the trainer constructor (ISV-driven from running mean of
|
||
// dd_persistence_bars is a documented follow-up sub-task per
|
||
// `feedback_isv_for_adaptive_bounds.md`). Single source file
|
||
// with one `extern "C" __global__` symbol
|
||
// (`plasticity_injection_kernel`) → one cubin → one launcher
|
||
// (`launch_sp15_plasticity_injection`). Mirrors the Phase 3.3
|
||
// `dd_penalty_kernel.cu` / 3.4 `regret_signal_kernel.cu` / 3.5
|
||
// `hold_floor_kernel.cu` / 3.5.2 `dd_asymmetric_reward_kernel.cu`
|
||
// / 3.5.3 `cooldown_kernel.cu` single-kernel-per-cubin pattern.
|
||
// Reads ISV[DD_PERSISTENCE_INDEX=404] +
|
||
// ISV[PLASTICITY_PERSISTENCE_THRESHOLD_INDEX=437]; read-modify-
|
||
// writes ISV[PLASTICITY_FIRED_THIS_FOLD_INDEX=436] +
|
||
// ISV[PLASTICITY_WARM_BARS_REMAINING_INDEX=438]. Per spec §9.2
|
||
// (3.5.4) post-amendment-2 fix: fire when DD_PERSISTENCE
|
||
// exceeds threshold AND PLASTICITY_FIRED_THIS_FOLD == 0 →
|
||
// set fired flag, set warm-bars counter to M_warm (default
|
||
// 200), AND reset last 10% of advantage-head weights to
|
||
// Kaiming-He init (`Normal(0, sqrt(2/fan_in))` sampled via
|
||
// cuRAND `curand_normal()`); per-bar warm-bars decrement;
|
||
// consumer-wiring follow-up reads max(COOLDOWN_BARS_REMAINING,
|
||
// PLASTICITY_WARM_BARS_REMAINING) and forces Hold while > 0.
|
||
// Phase 3.5.4 (2026-05-06) landed kernel + launcher + 3 ISV
|
||
// anchor seeds + 3 registry entries + 3 dispatch arms with
|
||
// the weight-reset deferred (kernel accepted
|
||
// advantage_head_weights + n_weights but no-op'd via (void)
|
||
// cast); Wave 4.2 / Phase 3.5.4.b (2026-05-07) lands the
|
||
// actual cuRAND Kaiming-He reset, extending the launcher
|
||
// signature with `fan_in: i32` and `seed: u64`. The kernel
|
||
// includes `<curand_kernel.h>` for device-side
|
||
// `curand_init`/`curand_normal` (inlined into the cubin by
|
||
// nvcc — no host-side cuRAND linker dependency required).
|
||
// The action-selection two-step Flat + cooldown consumer
|
||
// wiring remains a separate follow-up (Phase 3.5.4.c) per
|
||
// `feedback_no_partial_refactor.md`.
|
||
"plasticity_injection_kernel.cu",
|
||
// SP15 Phase 3.5.5 (2026-05-06): recovery curriculum in PER —
|
||
// per-step DD_TRAJECTORY_DECREASING proxy. Replaces non-existent
|
||
// episode-level metadata with a per-bar boolean computed from
|
||
// the dd_pct trajectory (`dd_pct(t) < dd_pct(t-1) AND
|
||
// dd_pct(t-1) > DD_TRAJECTORY_FLOOR`). True when transition is
|
||
// part of a recovery (DD shrinking from a non-trivial
|
||
// drawdown). Single source file with one `extern "C" __global__`
|
||
// symbol (`dd_trajectory_decreasing_kernel`) → one cubin → one
|
||
// launcher (`launch_sp15_dd_trajectory_decreasing`). Mirrors
|
||
// the Phase 3.3 `dd_penalty_kernel.cu` / 3.4
|
||
// `regret_signal_kernel.cu` / 3.5 `hold_floor_kernel.cu` / 3.5.2
|
||
// `dd_asymmetric_reward_kernel.cu` / 3.5.3 `cooldown_kernel.cu`
|
||
// / 3.5.4 `plasticity_injection_kernel.cu`
|
||
// single-kernel-per-cubin pattern. Reads ISV[DD_PCT_INDEX=406] +
|
||
// ISV[DD_TRAJECTORY_FLOOR_INDEX=441]; writes
|
||
// ISV[DD_TRAJECTORY_DECREASING_INDEX=439]; read-modify-writes a
|
||
// separate mapped-pinned `prev_dd_pct` scratch buffer
|
||
// (`sp15_dd_trajectory_prev_dd`, [1] f32, persistent
|
||
// previous-bar dd_pct between kernel calls — mirrors Phase
|
||
// 3.5.3 `sp15_cooldown_consecutive_losses` non-ISV
|
||
// mapped-pinned pattern). Per spec §9.2 (3.5.5)
|
||
// post-amendment-2 fix: PER sampling consumer
|
||
// (`sampling_weight = base × (1 + RECOVERY_OVERSAMPLE_WEIGHT ×
|
||
// DD_TRAJECTORY_DECREASING)`) AND the ISV-driven
|
||
// 25th-percentile DD_TRAJECTORY_FLOOR producer are deferred to
|
||
// follow-up atomic commits per `feedback_no_partial_refactor.md`
|
||
// (matching Phase 1.1-1.5 + 3.1-3.5.4 precedent — kernel +
|
||
// launcher verify in isolation first via the GPU oracle tests
|
||
// below). Phase 3.5.5 lands kernel + launcher + 3 ISV anchor
|
||
// seeds + 1 scratch buffer + 4 registry entries + 4 dispatch
|
||
// arms only.
|
||
"dd_trajectory_kernel.cu",
|
||
];
|
||
|
||
// ALL kernels get common header (BF16 types + wrappers)
|
||
let mut failed: Vec<&str> = Vec::new();
|
||
for kernel_name in &kernels_with_common {
|
||
if !try_compile_kernel(
|
||
&nvcc, kernel_dir, kernel_name, &arch, &out_dir, Some(&common_src),
|
||
) {
|
||
failed.push(kernel_name);
|
||
}
|
||
}
|
||
let passed = kernels_with_common.len() - failed.len();
|
||
eprintln!(" Precompiled {passed}/{} CUDA kernels ({arch}) — f32/TF32",
|
||
kernels_with_common.len());
|
||
if !failed.is_empty() {
|
||
eprintln!(" FAILED: {}", failed.join(", "));
|
||
panic!("nvcc failed to compile {} kernel(s): {}", failed.len(), failed.join(", "));
|
||
}
|
||
}
|
||
|
||
/// Compile a single .cu kernel file to a .cubin via nvcc.
|
||
///
|
||
/// If `common_header` is Some, it is prepended to the kernel source.
|
||
fn try_compile_kernel(
|
||
nvcc: &Path,
|
||
kernel_dir: &Path,
|
||
kernel_name: &str,
|
||
arch: &str,
|
||
out_dir: &Path,
|
||
common_header: Option<&str>,
|
||
) -> bool {
|
||
let kernel_path = kernel_dir.join(kernel_name);
|
||
let cubin_name = kernel_name.replace(".cu", ".cubin");
|
||
let cubin_path = out_dir.join(&cubin_name);
|
||
|
||
println!("cargo:rerun-if-changed={}", kernel_path.display());
|
||
|
||
let kernel_src = std::fs::read_to_string(&kernel_path)
|
||
.unwrap_or_else(|e| panic!("Failed to read {}: {e}", kernel_path.display()));
|
||
|
||
// Compose source: optional common header + kernel
|
||
let full_source = match common_header {
|
||
Some(header) => format!("{header}\n{kernel_src}"),
|
||
None => kernel_src,
|
||
};
|
||
|
||
// Write composed source to temp file
|
||
let tmp_src = out_dir.join(format!("_{kernel_name}"));
|
||
std::fs::write(&tmp_src, &full_source).unwrap();
|
||
|
||
// Compile with nvcc (-I kernel_dir for shared .cuh includes)
|
||
let include_flag = format!("-I{}", kernel_dir.display());
|
||
let status = Command::new(nvcc)
|
||
.args([
|
||
"-cubin",
|
||
&format!("-arch={arch}"),
|
||
"-O3",
|
||
"--ftz=true",
|
||
"--fmad=true",
|
||
"--prec-div=true",
|
||
"--prec-sqrt=true",
|
||
&include_flag,
|
||
"-o", cubin_path.to_str().unwrap(),
|
||
tmp_src.to_str().unwrap(),
|
||
])
|
||
.status();
|
||
|
||
match status {
|
||
Ok(s) if s.success() => {
|
||
eprintln!(" Compiled {kernel_name} -> {cubin_name} ({arch})");
|
||
true
|
||
}
|
||
Ok(s) => {
|
||
eprintln!(" FAILED: {kernel_name} (exit={})", s.code().unwrap_or(-1));
|
||
false
|
||
}
|
||
Err(e) => {
|
||
eprintln!(" FAILED: {kernel_name} (nvcc error: {e})");
|
||
false
|
||
}
|
||
}
|
||
}
|
||
|
||
/// Compute and emit the feature-schema fingerprint as `FEATURE_SCHEMA_HASH`
|
||
/// (decimal u64 string) so `fxcache.rs` can pick it up via `env!`.
|
||
///
|
||
/// Hashes the raw bytes of every source file that defines the on-disk fxcache
|
||
/// schema. Any edit (including whitespace / comments) bumps the hash. The
|
||
/// fxcache header records this value at write time; `validate()` rejects
|
||
/// caches whose recorded hash mismatches the compiled-in const, and
|
||
/// `precompute_features` regenerates them.
|
||
fn emit_feature_schema_hash() {
|
||
// Source files whose bytes determine the on-disk schema.
|
||
// Touching ANY of these files implies "fxcache regen on next deploy".
|
||
//
|
||
// 2026-04-27: added `examples/precompute_features.rs` after diagnosing the
|
||
// L40S `label_scale=5420` (raw-price magnitude) vs smoke `label_scale=0.05`
|
||
// (z-score) divergence. Root cause: the z-score normalisation step landed
|
||
// in `precompute_features.rs:625-630` on 2026-04-03 but that file was not
|
||
// tracked here, so the hash didn't change. PVC fxcaches written by an
|
||
// earlier build (no normalisation) still validated against the current
|
||
// binary, silently feeding raw features through column 0 of next_states
|
||
// and leaking future-bar prices into the policy via the aux head's
|
||
// shared trunk → impossible Sharpe (141 at epoch 0, 0.32% max DD over
|
||
// 214k bars). Adding the precompute pipeline to the hash forces every
|
||
// fxcache to be regenerated when its writer changes — including any
|
||
// future tweaks to feature normalisation, target ordering, or per-feature
|
||
// post-processing.
|
||
let schema_sources: &[&str] = &[
|
||
"src/features/extraction.rs",
|
||
"src/fxcache.rs",
|
||
"examples/precompute_features.rs",
|
||
"../ml-core/src/state_layout.rs",
|
||
];
|
||
|
||
// FNV-1a 64-bit init constants (RFC-style).
|
||
const FNV_OFFSET: u64 = 0xcbf2_9ce4_8422_2325;
|
||
const FNV_PRIME: u64 = 0x0000_0100_0000_01b3;
|
||
|
||
let mut hash: u64 = FNV_OFFSET;
|
||
for rel in schema_sources {
|
||
let path = PathBuf::from(rel);
|
||
let bytes = std::fs::read(&path)
|
||
.unwrap_or_else(|e| panic!("Failed to read schema source {}: {e}", path.display()));
|
||
// Mix file path bytes too so reordering / renaming bumps the hash even
|
||
// if the bytes happen to match.
|
||
for b in rel.as_bytes() {
|
||
hash ^= *b as u64;
|
||
hash = hash.wrapping_mul(FNV_PRIME);
|
||
}
|
||
// Length separator so concatenation boundaries can't collide.
|
||
for b in (bytes.len() as u64).to_le_bytes() {
|
||
hash ^= b as u64;
|
||
hash = hash.wrapping_mul(FNV_PRIME);
|
||
}
|
||
for b in &bytes {
|
||
hash ^= *b as u64;
|
||
hash = hash.wrapping_mul(FNV_PRIME);
|
||
}
|
||
println!("cargo:rerun-if-changed={}", path.display());
|
||
}
|
||
|
||
// Decimal so the consumer can use the const-stable `u64::from_str_radix`
|
||
// with radix 10 without any hex-prefix stripping.
|
||
println!("cargo:rustc-env=FEATURE_SCHEMA_HASH={hash}");
|
||
}
|
||
|
||
/// Find nvcc: prefer $CUDA_HOME/bin/nvcc, then check PATH
|
||
fn find_nvcc() -> Option<PathBuf> {
|
||
// Try $CUDA_HOME/bin/nvcc first
|
||
if let Ok(home) = std::env::var("CUDA_HOME") {
|
||
let nvcc = PathBuf::from(home).join("bin/nvcc");
|
||
if nvcc.exists() {
|
||
return Some(nvcc);
|
||
}
|
||
}
|
||
|
||
// Try common CUDA paths
|
||
for path in &["/usr/local/cuda/bin/nvcc", "/usr/bin/nvcc"] {
|
||
let p = PathBuf::from(path);
|
||
if p.exists() {
|
||
return Some(p);
|
||
}
|
||
}
|
||
|
||
// Check if nvcc is in PATH
|
||
match Command::new("nvcc").arg("--version").output() {
|
||
Ok(output) if output.status.success() => Some(PathBuf::from("nvcc")),
|
||
_ => None,
|
||
}
|
||
}
|