6657e56265
Final piece of the SP13 Layer B chain. B1.1a flipped the aux head from K=1
MSE regression to K=2 softmax CE classification but aux_nb_label_buf was
zero-init — model was training on "all bars are class 0 (down)". B1.1b
lands the producer kernel that fills i32 -1/0/1 labels from the 30-bar
price trajectory, the replay direct-path 8th gather that carries those
labels into the trainer, and the experience collector hoist that ensures
bar_indices_pinned is always populated (producer + hindsight relabel both
consume it). Aux head finally trains on real classification signal.
Recovery commit: this completes a B1.1b agent dispatch that crashed
mid-edit. The implementer had landed ~95% of the cascade (kernel file,
build.rs, replay buffer signature + direct path, fused_training getter,
trainer accessor, both training_loop.rs callers, kernel field + cubin
loader on the experience collector) before being killed. The missing
pieces (experience collector launch + bar_indices_pinned hoist + 6
producer tests + audit doc) were completed manually post-crash and
verified end-to-end.
Three contracts (atomic single commit per feedback_no_partial_refactor):
1. NEW aux_sign_label_kernel.cu producer — pure per-thread O(1) map
reading targets[bar*6+2] (raw_close column) at bar and bar+lookahead,
writing -1 (skip if bar+lookahead >= total_bars), 0 (down/flat under
strict greater-than tie-break), or 1 (up). Replaces B0 alloc_zeros.
2. Replay direct-path 8th gather — set_trainer_buffers gains 8th arg
trainer_aux_sign_labels_ptr; direct branch in sample_proportional
adds gather_i32_scalar into the trainer ptr; fallback gather wrapped
in if !direct_to_trainer (avoids wasted DtoD). Direct-mode
GpuBatchPtrs return points aux_sign_labels_ptr at trainer ptr.
3. Experience collector bar_indices_pinned cpu-fill hoist — moved out
of if hindsight_fraction > 0.0 so producer + hindsight share it.
Files (9 total):
- crates/ml/src/cuda_pipeline/aux_sign_label_kernel.cu (NEW)
- crates/ml/build.rs (cubin registration)
- crates/ml/src/cuda_pipeline/gpu_experience_collector.rs (kernel
field + cubin loader + struct init + hoist + producer launch)
- crates/ml-dqn/src/gpu_replay_buffer.rs (8th arg + direct gather +
fallback skip + GpuBatchPtrs return)
- crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs (aux_nb_label_buf_ptr
accessor mirrors 6 existing trainer-buf accessors)
- crates/ml/src/trainers/dqn/fused_training.rs
(trainer_aux_sign_labels_buf_ptr getter)
- crates/ml/src/trainers/dqn/trainer/training_loop.rs (both
set_trainer_buffers callers updated)
- crates/ml/tests/sp13_layer_b_oracle_tests.rs (6 NEW producer tests)
- docs/dqn-wire-up-audit.md (B1.1b section)
Hard rules upheld:
- feedback_no_partial_refactor: every consumer of the 3 contracts
migrates atomically
- feedback_no_atomicadd: producer is pure map; no reductions
- feedback_cpu_is_read_only: producer GPU-only; only host work is
pre-existing bar_indices_pinned cpu-fill (hoisted unchanged)
- feedback_no_stubs: kernel output flows through real chain — ring
buffer → direct gather → aux_nb_label_buf → CE consumer
- feedback_no_legacy_aliases: 8-arg setter gets
#[allow(clippy::too_many_arguments)] not an alias shim
- feedback_no_htod_htoh_only_mapped_pinned: targets_buf and
bar_indices_pinned both pre-existing mapped-pinned
Build + test:
- cargo check --workspace clean (only pre-existing warnings)
- cargo check --workspace --tests clean
- 17 tests in sp13_layer_b_oracle_tests.rs:
- 2 CPU-only (fingerprint bump + HEALTH_DIAG snap stable) pass
- 15 GPU on RTX 3050 Ti pass (9 B1.1a + 6 new B1.1b producer):
aux_sign_label_monotone_up_all_ones
aux_sign_label_monotone_down_all_zeros
aux_sign_label_flat_all_zeros_strict_gt
aux_sign_label_last_30_bars_skip
aux_sign_label_boundary_first_valid_last_skip
aux_sign_label_multi_episode_per_episode_skip
Producer tests cover every edge case in the kernel:
- Monotone trajectories (label=1 / label=0 across all valid bars)
- Flat tie-break (strict greater-than means flat → 0)
- Skip sentinel for last lookahead bars
- First/last bar boundary (bar=0 valid, bar=L-1 skip)
- Multi-episode global skip semantics
Next: Smoke A — L40S 5-epoch validation of full SP13 stack
(P0a + P0b + B0 + B0.1 + B1.0 + B1.1a + B1.1b). Expected: aux head
trains on real K=2 softmax CE labels; aux_dir_acc_short_ema rises above
0.5 within first epoch (vs B1.1a degraded baseline at 0.5);
HEALTH_DIAG aux_b1_diag emits per-epoch with n_down/n_up/n_skip/mask_frac.
If aux_dir_acc_short_ema > 0.55 by epoch 5, B1.1b is validated and the
chain merges to main.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
911 lines
53 KiB
Rust
911 lines
53 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",
|
||
// 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",
|
||
];
|
||
|
||
// 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,
|
||
}
|
||
}
|