foxhunt/crates/ml/build.rs

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",
        // SP20 Phase 2 Task 2.1 (2026-05-09): standalone test wrapper for
        // the 4-quadrant directional-correctness reward function defined
        // in `sp20_reward.cuh`. Wraps `sp20_compute_event_reward` so the
        // Rust GPU oracle tests in `tests/sp20_event_reward_test.rs` can
        // drive each quadrant of the truth table on synthetic inputs with
        // analytically-known expected outputs per
        // `feedback_no_cpu_test_fallbacks`. Producer call site lands in
        // Task 2.2 inside `experience_env_step::segment_complete` branch;
        // the test wrapper has no production callers.
        "sp20_event_reward_test_kernel.cu",
        // SP20 Phase 3 Task 3.2 (2026-05-10): standalone test wrappers
        // for the two device helpers defined in `sp20_hold_baseline.cuh`
        // (Component 2 — Hold opportunity-cost dual emission). Wraps
        // `sp20_compute_per_bar_aux_conf_k2`,
        // `sp20_sum_hold_baseline_over_trade`, and a fixture mirroring
        // the production per-bar emission site so the Rust GPU oracle
        // tests in `tests/sp20_hold_baseline_test.rs` drive the EXACT
        // same code paths the production caller
        // (`experience_env_step` per-bar branches + segment_complete
        // trade-close site) executes per
        // `feedback_no_cpu_test_fallbacks`. Producer call sites land in
        // Task 3.2 inside `experience_env_step` (per-bar branches +
        // segment_complete branch). Test wrapper has no production
        // callers.
        "sp20_hold_baseline_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",
        // SP22 H6 vNext (2026-05-13): trade-outcome label producer kernel.
        // Replaces the per-bar binary direction label with a per-trade-close
        // 3-way outcome label {Profit, Stop, Timeout}. Sparse labels
        // (only at trade-close events) but aligned with the system's
        // actual multi-bar trade decision horizon. See
        // `docs/plans/2026-05-14-sp22-h6-vNext-trade-outcome-aux.md` for
        // the full vNext architecture spec. This kernel is the FOUNDATION
        // (Phase A2) — additional kernels for aux forward / loss /
        // backward + 12-weight W atom-shift land in subsequent commits.
        "trade_outcome_label_kernel.cu",
        // SP22 H6 vNext Phase A3 (2026-05-13): K=3 softmax aux head forward
        // kernel for the trade-outcome aux head. Architecture mirrors
        // `aux_next_bar_forward` (K=2) but emits a 3-way distribution over
        // {Profit, Stop, Timeout}. Dead code until A4 (loss reduce) + A5
        // (backward) + Phase B (262-dim input concat with plan_params) +
        // Phase C (state slot 121-123 wireup) + Phase D (12-weight W
        // atom-shift) land in subsequent commits. See
        // `docs/plans/2026-05-14-sp22-h6-vNext-trade-outcome-aux.md` for
        // the full vNext architecture spec.
        "aux_trade_outcome_forward_kernel.cu",
        // SP22 H6 vNext Phase A4 (2026-05-13): sparse cross-entropy reduce
        // over the K=3 softmax tile from `aux_trade_outcome_forward`.
        // Mirrors `aux_next_bar_loss_reduce`'s single-block shmem-tree
        // pattern: mean over B_valid (skip label=-1 rows), fmaxf(p, 1e-30)
        // numerical floor, saved valid_count for backward. Dead code until
        // Phase A5 (backward) + Phase B (Rust launcher wireup) land. Kept
        // separate from the K=2 sibling for diagnostic isolation, sparse-
        // label semantic clarity, and future per-class weighting headroom.
        "aux_trade_outcome_loss_reduce_kernel.cu",
        // SP22 H6 vNext Phase A5 (2026-05-13): backward kernel for the
        // K=3 trade-outcome aux head. Mirrors `aux_next_bar_backward`
        // (K=2 sibling): per-sample partials (dW1, db1, dW2, db2,
        // dh_s2_aux), d_logits = (softmax − one_hot)/B_valid with mask
        // handling, ELU backward via post-activation identity. Sparse-
        // label arithmetic amplifies per-trade-close gradient magnitude
        // (smaller B_valid → larger inv_B); Phase E's Adam may need
        // class-weighted CE or per-group LR tuning. Reuses the existing
        // `aux_param_grad_reduce` kernel from aux_heads_kernel.cu for
        // the per-sample → final-grad batch-reduce step. Dead code
        // until Phase B wireup lands the Rust launcher chain.
        "aux_trade_outcome_backward_kernel.cu",
        // SP22 H6 vNext Phase B5a (2026-05-14): input concat kernel for
        // plan-conditioned K=3 trade-outcome forward. Writes [B, 262]
        // from h_s2_aux [B, 256] || plan_params [B, 6]. NULL-tolerant on
        // plan_params (zeros for the trailing P=6 cols when source
        // unavailable, e.g., collector-path cold-start). Dead code at
        // this commit — Phase B5b allocates buffers + wires launches +
        // flips SH2_TOTAL=262 in forward/backward.
        "aux_to_input_concat_kernel.cu",
        // SP22 H6 vNext Phase B5b (2026-05-14): row-truncating SAXPY for
        // backward dh_s2_aux accumulate. Adds first n_cols_copy columns
        // per row of [B, src_cols] into [B, dst_cols] with row stride
        // mismatch handling. Needed when B5b's plan-conditioned backward
        // emits dh_s2_aux_to_buf [B, 262] but the aux trunk's accumulator
        // dh_s2_aux_accum [B, 256] only consumes first 256 cols (plan_
        // params gradient is stop-grad).
        "strided_row_saxpy_kernel.cu",
        // SP22 H6 vNext Phase F-3 (2026-05-14): K=3 trade-outcome head's
        // sparse-CE EMA producer for observability. Single-thread
        // single-block direct-to-ISV writer with Pearl A bootstrap.
        // Writes ISV[AUX_TRADE_OUTCOME_CE_EMA_INDEX=538] every training
        // step alongside the existing K=2/K=5 aux EMA producers.
        // Dead code until launcher wired (next commit step in this
        // phase).
        "aux_outcome_ce_ema_kernel.cu",
        // SP22 H6 vNext Phase C (2026-05-14): per-env 3-slot softmax
        // extractor for the K=3 trade-outcome aux head. Copies
        // aux_outcome_softmax [n_envs, K=3] → prev_aux_outcome_probs
        // [n_envs, 3]. Consumed at the NEXT step's state assembly via
        // experience_state_gather (Phase C-2 follow-up). For Phase C
        // (this commit) the producer fires every step but no consumer
        // reads — dead-code scaffolding parallel to the K=2 head's
        // aux_softmax_to_per_env_kernel.
        "aux_outcome_softmax_to_per_env_kernel.cu",
        // SP22 H6 (2026-05-12): per-env p_up extractor that copies
        // `aux_softmax[env, 1]` → `prev_aux_dir_prob[env]` after each aux
        // forward. The cache buffer is read by `experience_state_gather`
        // and `backtest_state_gather*` at the next step's state assembly,
        // wiring the aux head's directional probability into policy state
        // (slot `SL_PADDING_START + 0 = 121`). One-step lag is intentional:
        // state is assembled BEFORE aux forward at step t, so step t+1
        // reads step t's aux output. Pure per-env gather; no reduction, no
        // atomicAdd per `feedback_no_atomicadd.md`; pure GPU compute per
        // `feedback_no_cpu_compute_strict.md`. Launched on the collector's
        // stream right after `aux_next_bar_forward` in
        // `gpu_experience_collector.rs`.
        "aux_softmax_to_per_env_kernel.cu",
        // SP22 H6 Phase 3 α scalar-bias kernels — REGISTRATIONS DELETED
        // 2026-05-13. The kernel `.cu` files
        // (aux_to_q_dir_bias_kernel.cu, aux_to_q_dir_bias_backward_kernel.cu)
        // stay on disk as committed dead code (commit 464bc5f7a) for git
        // history. They are mathematically ineffective in C51 distributional
        // Q-learning because adding a constant to Q is softmax-shift-
        // invariant — W gradient = 0, W never trains. The atom-position
        // shift replacement design threads `aux_atom_shift[b, a] = W[a] *
        // state_121[b]` through compute_expected_q + c51_loss_kernel +
        // c51_grad_kernel + mag_concat_qdir. The dW + dstate_121 gradient
        // computations integrate into c51_grad_kernel's existing
        // projection backward — no separate α kernel needed. See audit
        // doc Phase 3c entry + spec doc α section for the revised design.
        //
        // SP22 H6 Phase 3 α (atom-shift design 2026-05-13): structural-prior
        // init for w_aux_to_q_dir [4] — writes per-action prior
        // [-0.5, 0.0, +0.5, 0.0] for [Short, Hold, Long, Flat] at trainer
        // construction. One-time GPU init (NOT in captured graph). The
        // Adam-trained W refines from this prior via c51_grad_kernel's
        // projection backward (Phase 3-final B9 Step 8). Pure 4-thread
        // single-block; trivial cost.
        "aux_w_prior_init_kernel.cu",
        // SP22 H6 Phase 3 α Step 8 (2026-05-13): backward dW kernel for
        // w_aux_to_q_dir [4]. Reads scratch from c51_loss_kernel forward
        // (aux_target_a_dir + aux_proj_logdiff_dir) and computes per-action
        // gradient via block tree-reduce (grid=(4,1,1), block=(256,1,1))
        // — no atomicAdd per pearl_no_atomicadd. Per-block writes dw_aux[a].
        "c51_aux_dw_kernel.cu",
        // SP22 H6 Phase 3 α Step 11 (2026-05-13): Adam update kernel for
        // w_aux_to_q_dir [4]. Standard Adam with bias correction; 4
        // threads single block; launched once per training step OUTSIDE
        // the captured forward/backward graphs.
        "adam_w_aux_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",
        // Class A P0-A (2026-05-08): adaptive REWARD_POS/NEG_CAP producer.
        // Single-block kernel that sweeps the per-epoch
        // `step_ret_per_sample` + `trade_close_per_sample` buffers
        // (populated by `unified_env_step_core` in
        // `experience_kernels.cu`), computes p99(winning_returns) ×
        // safety_factor=1.5 → POS cap, NEG = −2 × POS (Kahneman 2:1
        // asymmetry preserved at producer-time per
        // `pearl_audit_unboundedness_for_implicit_asymmetry` — single
        // source of truth, no consumer applies the multiplier itself).
        // Pearl-A first-observation bootstrap (sentinels 5.0/-10.0
        // match pre-P0-A hardcoded constants for bit-identical
        // cold-start); α=0.01 slow EMA thereafter. Bounds POS∈[1, 50],
        // NEG∈[-100, -2] are Category-1 dimensional safety floors per
        // `feedback_isv_for_adaptive_bounds`. Replaces hardcoded
        // REWARD_POS_CAP=+5.0f / REWARD_NEG_CAP=-10.0f
        // (state_layout.cuh:266-267) — Class A audit
        // highest-suspected-impact fix for the months-long
        // WR-stuck-at-46-48% plateau across 11 superprojects. Per-epoch
        // boundary launch.
        "reward_cap_update_kernel.cu",
        // Class A P1-Producer (2026-05-08): adaptive Bayesian Kelly priors
        // producer. Single-block kernel that sweeps the per-fold-end
        // `portfolio_state[n_envs, PS_STRIDE]` buffer (same source the
        // `kelly_cap_update_kernel` reads from at the same boundary),
        // aggregates the realized PS_KELLY_{WIN_COUNT, LOSS_COUNT,
        // SUM_WINS, SUM_LOSSES} fields across envs, and slow-EMA-blends
        // into ISV[454..458). Pearl-A first-observation bootstrap
        // (sentinels 2.0/2.0/0.01/0.01 match pre-P1-Producer hardcoded
        // values for bit-identical cold-start); α=0.005 slow EMA
        // thereafter (per-fold cadence — the Bayesian prior is a *prior
        // belief* and should change slowly across folds; a fast
        // "prior" would just be a noisy Kelly estimate). Bounds
        // counts∈[0.5, 100], sums∈[0.001, 1.0] are Category-1 dimensional
        // safety floors per `feedback_isv_for_adaptive_bounds`. Replaces
        // hardcoded `prior_wins=2.0f / prior_losses=2.0f /
        // prior_sum_wins=0.01f / prior_sum_losses=0.01f` in
        // kelly_cap_update_kernel.cu:39-42 + trade_physics.cuh::
        // kelly_position_cap:304-307 — Class A audit P1-Producer batch.
        // Block-tree-reduce (no atomicAdd) per
        // `feedback_no_atomicadd.md`. Per-epoch boundary launch.
        "kelly_bayesian_priors_update_kernel.cu",
        // Class A audit-fix Batch 4-A (2026-05-08): adaptive DD saturation
        // floor producer. Single-block 256-thread kernel that sweeps the
        // per-env DD trajectory tile `dd_state_per_env[n_envs * 6]`
        // (populated by `dd_state_kernel`), aggregates DD_MAX (offset 1)
        // across envs via Welford `mean + Z_75 × sigma` p75 estimator
        // with `max(p75, mean)` robustness guard × 1.5 safety factor,
        // and slow-EMA blends into ISV[DD_SATURATION_FLOOR_ADAPTIVE_INDEX
        // =458]. Pearl-A first-observation bootstrap (sentinel 0.25
        // matches pre-fix hardcoded `0.25f` in
        // `trade_physics.cuh::apply_margin_cap` for bit-identical cold-
        // start); α=0.01 slow EMA thereafter (DD distribution is a slow-
        // moving fold-volatility property — same cadence as P0-A
        // REWARD_POS_CAP). Bounds [0.10, 0.50] are Category-1 dimensional
        // safety floors per `feedback_isv_for_adaptive_bounds`. Replaces
        // the hardcoded saturation floor (the upper end of the linear
        // position-size scaling ramp `dd_scale = max(0.05, 1.0 −
        // dd_frac/saturation_floor)`). Distinct from
        // `SP15_DD_THRESHOLD_INDEX=421` (the SP15 quadratic DD-penalty
        // *trigger* threshold, a *lower* bound). Atomically deleted
        // alongside this slot addition: legacy `compute_drawdown_penalty`
        // device function in trade_physics.cuh + its single caller at
        // experience_kernels.cu:3822 + the `w_dd` Rust config field +
        // the `dd_threshold` and `w_dd` kernel args (Item 2 Case A —
        // SP15's quadratic asymmetric DD penalty in
        // `compute_sp15_final_reward_kernel.cu:154` is the production-
        // grade replacement; layering both creates double-counting).
        // Block-tree-reduce (no atomicAdd) per `feedback_no_atomicadd.md`.
        // Per-epoch boundary launch.
        "dd_saturation_floor_update_kernel.cu",
        // min_hold_temperature_update_kernel.cu — reads
        // ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382] and
        // ISV[HOLD_RATE_TARGET_INDEX=381], maps overrun
        // `clamp(max(0, observed - target) / max(target, 0.01), 0, 1)`
        // to a [TEMP_MIN=5, TEMP_MAX=50] temperature, blends via
        // Pearl-A bootstrap + Welford EMA α=0.05, writes to
        // ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460]. Per SP16 Phase 1
        // (revised) — supersedes the prior slot-373 dir-acc-skill mapping;
        // see docs/dqn-wire-up-audit.md § "SP16 Phase 1 (revised)" for
        // rationale.
        "min_hold_temperature_update_kernel.cu",
        // SP18 v2 Phase 0 Task 0.1 (2026-05-08): D-leg per-action reward
        // decomposition diagnostic. Reads `reward_components_per_sample`
        // ([n_samples × 6]) + `actions` ([n_samples], factored-action
        // ints) from the experience collector, bins per direction
        // (DIR_HOLD / DIR_LONG / DIR_SHORT / DIR_FLAT — 4 bins), and
        // emits 5 stats per bin (mean r_micro, mean r_opp_cost, mean
        // r_popart, mean |reward|, fire_rate) into a 20-float
        // mapped-pinned `reward_decomp_diag_buf`. Block tree-reduce
        // (4 blocks × 256 threads, one block per bin); no atomicAdd
        // per `feedback_no_atomicadd.md`. Pure observability — does
        // NOT modify any production-path buffer; written values are
        // consumed exclusively by the per-epoch HEALTH_DIAG line emit.
        // Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md
        // § Phase 0 Task 0.1.
        "reward_decomp_diag_kernel.cu",
        // SP18 v2 Phase 0 Task 0.2 (2026-05-08): B-leg TD-error
        // magnitude EMA producer. Single-block 256-thread kernel
        // computing mean(|td_errors[b]|) over a B-element batch and
        // EMA-blending into ISV[TD_ERROR_MAG_EMA_INDEX=493] via
        // Pearl-A first-observation bootstrap (sentinel 0.0) + fixed
        // α=WELFORD_ALPHA_MIN=0.4 (per
        // `pearl_wiener_alpha_floor_for_nonstationary` — Welford-α
        // accumulators at [498..504) are reserved for Phase 4
        // q_next_target chain). Block tree-reduce (no atomicAdd).
        // Pure observability: pre-fix baseline for the B-DD9 ratio
        // gate. Plan: SP18 Phase 0 Task 0.2.
        "td_error_mag_ema_kernel.cu",
        // SP18 v2 Phase 3 (2026-05-09): D-leg adaptive HOLD_REWARD_POS/
        // NEG_CAP producer kernel. Single-block 256-thread kernel that
        // sweeps the per-sample `step_ret_per_sample` +
        // `trade_close_per_sample` buffers (populated by
        // `unified_env_step_core` in `experience_kernels.cu` line ~2855),
        // filters by `(is_close && step_ret != 0)` to extract the
        // Long/Short trade-close magnitude distribution (DD3=b), computes
        // a Welford `mean + Z_99 × sigma` p99 estimator with `max(p99,
        // max_observed)` conservative takeover × HOLD_REWARD_CAP_SAFETY_
        // FACTOR=1.5, applies Wiener-optimal α blend (per
        // `pearl_wiener_optimal_adaptive_alpha`; floor at WELFORD_ALPHA_
        // MIN=0.4 per `pearl_wiener_alpha_floor_for_nonstationary` since
        // the policy-realised distribution is non-stationary), and
        // writes both `ISV[HOLD_REWARD_POS_CAP_INDEX=483]` (POS cap) and
        // `ISV[HOLD_REWARD_NEG_CAP_INDEX=484]` (NEG = −2 × POS, DD5=b
        // mirrored asymmetry — same Kahneman 2:1 ratio as the position-
        // side `reward_cap_update_kernel` for ISV[452]/[453], single
        // source of truth at producer time per
        // `pearl_audit_unboundedness_for_implicit_asymmetry`). Welford
        // accumulators in slots [487..493). Pearl-A first-observation
        // bootstrap on POS (sentinel 5.0/-10.0 matches the SP14 P0-A
        // pattern for bit-identical cold-start). POS bounds [0.5, 50.0]
        // are Category-1 dimensional safety floors per
        // `feedback_isv_for_adaptive_bounds`; NEG ∈ [−100, −1] derived.
        // Block-tree-reduce in shmem (no atomicAdd per
        // `feedback_no_atomicadd.md`); `CudaFunction` pre-loaded at
        // construction per `pearl_no_host_branches_in_captured_graph`.
        // Loaded by `gpu_aux_trunk::HoldRewardCapUpdateOps`. Per-epoch
        // boundary launch right AFTER `launch_reward_cap_update` so the
        // two cap producers (position-side + Hold-side) share the same
        // step_ret/trade_close source buffers. Plan: SP18 Phase 3.
        "hold_reward_cap_update_kernel.cu",
        // SP18 v2 Phase 2 (2026-05-09): D-leg structural Hold
        // opportunity-cost device-fn GPU oracle test wrapper. Single
        // thread × single block; wraps `compute_sp18_hold_opportunity_cost`
        // from `trade_physics.cuh`. Mapped-pinned f32 input/output (no
        // htod/dtoh per `feedback_no_htod_htoh_only_mapped_pinned`); host
        // reads via `read_volatile` after stream sync (the
        // `__threadfence_system()` at thread exit guarantees PCIe
        // visibility before the sync returns). Same structural pattern
        // as `sp12_reward_math_test_kernel.cu`. Test-only; no production
        // callers. Drives `compute_sp18_hold_opportunity_cost`'s closed-
        // form math contract from the Rust GPU oracle in
        // `tests/sp18_hold_reward_oracle_tests.rs` against the CPU
        // oracle in the same file (ε=1e-5). Plan: SP18 Phase 2 Task 2.2.
        "sp18_hold_opp_test_kernel.cu",
        // SP14 Layer C Phase C.6 (2026-05-08): h_s2_aux RMS EMA producer.
        // Single-block 256-thread kernel computing RMS = sqrt(mean(x²))
        // over `h_s2_aux [B, SH2]` (aux trunk final output, no activation)
        // and EMA-blending into `ISV[H_S2_AUX_RMS_EMA_INDEX=449]`. Pearl-A
        // first-observation bootstrap embedded in kernel (sentinel 0.0 →
        // replace directly); fixed α=0.05 EMA blend thereafter. ISV slot 449
        // is outside the SP4/SP5 wiener buffer linear span so the shared
        // scratch+apply_pearls_ad_kernel path is unavailable — bootstrap
        // logic is self-contained per `avg_win_hold_time_update_kernel`
        // precedent. Block-tree-reduce (no atomicAdd) per
        // `feedback_no_atomicadd.md`. Per-collector-step launch cadence.
        "h_s2_aux_rms_ema_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",
        // SP17 Task PP.3 (2026-05-08): pre-centering V/A mean diagnostic.
        // Single-block tree-reduce kernel reading the existing
        // `on_v_logits_buf [B, NA]` and `on_b_logits_buf [B, (B0+B1+B2+B3)*NA]`
        // (BRANCH-MAJOR) and emitting 5 floats:
        //   [E[V], E[A_dir], E[A_mag], E[A_ord], E[A_urg]]
        // into a `MappedF32Buffer<5>`. Sequentially sweeps the 5 reduction
        // regions, reusing one shared-memory tile (1024 bytes for the
        // standard 256-thread launch) per `feedback_no_atomicadd.md`. Does
        // NOT modify any existing aggregation — pure observability over
        // the same buffers `compute_expected_q` reads from. Consumed by
        // `read_v_a_means_pre_centering()` in `gpu_dqn_trainer.rs` and
        // emitted into per-epoch HEALTH_DIAG by `training_loop.rs`. Sets
        // the diagnostic baseline for the SP17 Phase 0 kill criterion (V
        // dominance pre-centering).
        "v_a_means_diag_kernel.cu",
        // SP17 Phase 3.1 (2026-05-08): per-branch EMA of the over-actions
        // variance of `A_centered`. 4 blocks × 256 threads (one block per
        // branch — dir/mag/ord/urg). Reads the same BRANCH-MAJOR
        // `on_b_logits_buf` that `compute_expected_q` consumes; computes
        // per-atom mean of A across actions in the branch; sums squared
        // deviations from the per-atom mean across (B × n_branch × NA);
        // writes to ISV[A_VAR_EMA_DIR/MAG/ORD/URG_INDEX] with Pearl-A
        // first-observation bootstrap (sentinel SENTINEL_A_VAR_EMA=0.0)
        // + α=WELFORD_ALPHA_MIN=0.4 EMA blend per
        // `pearl_wiener_alpha_floor_for_nonstationary`. Block tree-reduce
        // (no atomicAdd) per `feedback_no_atomicadd.md`. Diagnostic
        // observability — does NOT modify Q-head logits or any consumer
        // path. Consumed by HEALTH_DIAG `dueling [a_var=...]` line via
        // `read_a_var_ema_per_branch()` in gpu_dqn_trainer.rs. See
        // docs/dqn-wire-up-audit.md § "SP17 Phase 3.1" for rationale.
        "sp17_a_var_ema_kernel.cu",
        // SP17 Phase 3.2 (2026-05-08): per-branch V_share EMA producer.
        // 4 blocks × 256 threads. Reads `on_v_logits_buf [B, NA]` and
        // BRANCH-MAJOR `on_b_logits_buf`; computes per-branch
        // `|E[V]| / (|E[V]| + |E[A_centered, picked]|)` where picked is
        // argmax_a Σ_z A_raw[i, a, z] (max-Q action — tractable
        // per-batch without depending on actions_history_buf). Writes
        // to ISV[V_SHARE_DIR/MAG/ORD/URG_INDEX] (slots 478..482) with
        // Pearl-A bootstrap (sentinel 0.5 = healthy 50/50 cold-start) +
        // α=WELFORD_ALPHA_MIN=0.4 EMA blend. Bilateral clamp [0, 1]
        // before write per `pearl_symmetric_clamp_audit`. Block
        // tree-reduce (no atomicAdd) per `feedback_no_atomicadd`.
        // Diagnostic observability — does NOT modify any consumer path.
        // Consumed by HEALTH_DIAG `dueling [v_share=...]` via
        // `read_v_share_per_branch()`. See
        // docs/dqn-wire-up-audit.md § "SP17 Phase 3.2" for rationale.
        "sp17_v_share_kernel.cu",
        // SP17 Phase 3.2 (2026-05-08): adaptive advantage_clip_bound
        // producer. Single block × 256 threads. Reads BRANCH-MAJOR
        // `on_b_logits_buf`; for each branch d, computes per-atom mean
        // of A across actions and writes |A_raw - mean[z]| into a
        // pre-allocated flat scratch buffer of size Σ_d B×b_d×NA;
        // calls `sp4_histogram_p99` over the full flat scratch (block
        // tree-reduce + per-warp tile binning, no atomicAdd per
        // `pearl_fused_per_group_statistics_oracle` and
        // `feedback_no_atomicadd`); blends `bound = p99 × 1.5` into
        // ISV[ADVANTAGE_CLIP_BOUND_INDEX=482] with Pearl-A bootstrap
        // (sentinel 1.0 = no effective clipping cold-start) + α=0.01
        // slow per-fold cadence + bilateral clamp [0.1, 100.0] per
        // `pearl_symmetric_clamp_audit`. Observability-only in this
        // commit — the actual clip wire-up is Phase 5 follow-up. See
        // docs/dqn-wire-up-audit.md § "SP17 Phase 3.2" for rationale.
        "sp17_advantage_clip_bound_kernel.cu",
        // SP18 v2 weight-drift diagnostic (2026-05-09): single block × 256
        // threads. Reads `params_flat [n]` and `best_params_snapshot [n]`
        // and writes [||params - best||₂, relative_drift] into a
        // `MappedF32Buffer<2>` for cold-path readback. Two block
        // tree-reductions (no atomicAdd per `feedback_no_atomicadd`)
        // share one shmem tile sequentially. Diagnostic-only — does NOT
        // modify any consumer path. Surfaces HD2→HD5 weight drift across
        // epochs so the post-deploy reviewer can correlate Sharpe-peak
        // decay with actual parameter trajectory. Consumed by
        // `read_weight_drift_diag()` in gpu_dqn_trainer.rs and emitted
        // into per-epoch HEALTH_DIAG `weight_drift` line by
        // training_loop.rs. See docs/dqn-wire-up-audit.md § "SP18 v2
        // weight-drift diagnostic".
        "weight_drift_diag_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",
        // SP20 Phase 1.1 (2026-05-09): aux confidence p50/std stats
        // producer. Single-block, BLOCK=256 kernel reading
        // `aux_logits [B, 3]` (the SP14-C aux head's 3-class direction
        // logits) and emitting `[p50, std]` of the per-row signal
        // `aux_conf[i] = max_c softmax(logits[i, *])[c] - 1/3` into a
        // `MappedF32Buffer<2>`. p50 uses the inlined `sp4_histogram_p99`
        // pattern with cumulative-from-bottom (target = ⌈B/2⌉)
        // substituted for cumulative-from-top, per
        // `pearl_fused_per_group_statistics_oracle` (one fused stream
        // for both stats) and `feedback_no_atomicadd` (per-warp tile
        // binning + block tree-reduce sums; no atomicAdd anywhere).
        // std uses two block tree-reductions sharing one shmem tile
        // sequentially. Component 5 / Kernel 3 of the SP20 fused-
        // producer chain (`sp20_emas_compute` + `sp20_controllers_
        // compute` land in Phase 1.2 + 1.3). Phase 1.4 wires the
        // production launch site; this commit adds kernel + Rust
        // launcher (`sp20_stats_compute.rs`) + GPU oracle tests
        // (`tests/sp20_stats_compute_test.rs`) atomically per
        // `feedback_no_partial_refactor`.
        "sp20_stats_compute_kernel.cu",
        // SP20 Phase 1.2 (2026-05-09): 8 Wiener-α EMA producer.
        // Component 5 / Kernel 1 of the SP20 fused-producer chain.
        // Single-block, single-thread kernel updating 4 ISV slots
        // (ALPHA_EMA, WR_EMA, HOLD_PCT_EMA, HOLD_REWARD_EMA) plus
        // 4 internal scratch slots (trade_duration_ema +
        // aux_conf_p50/std/dir_acc EMAs) per training step. Per-EMA
        // observation counters (i32 mapped-pinned) handle the
        // `pearl_first_observation_bootstrap` sentinel transition
        // (count==0 ⇒ replace, count>0 ⇒ Wiener-blend) so EMAs
        // that legitimately observe 0 (e.g., first-loss WR) don't
        // re-bootstrap. Wiener-α floor 0.4 hardcoded as a structural
        // property per `pearl_wiener_alpha_floor_for_nonstationary`.
        // No atomicAdd, no host branches — captureable in the per-
        // step CUDA Graph.
        // AMENDED 2026-05-09 (Path C, Phase 1.4): kernel signature
        // refactored from 9 scalar value-args to a single device
        // `SP20EmaInputs* ema_inputs` struct pointer so the upstream
        // `sp20_aggregate_inputs_kernel` can populate the 9 fields
        // from per-env GPU arrays without any host sync (forbidden
        // by `feedback_cpu_is_read_only` + `feedback_no_htod_htoh_only_
        // mapped_pinned`). Math semantics bit-identical. The kernel +
        // launcher signature change + tests + production wire-up +
        // aggregation kernel + audit/spec/plan amendments land
        // atomically per `feedback_no_partial_refactor`.
        "sp20_emas_compute_kernel.cu",
        // SP20 Phase 1.3 (2026-05-09): 6 derived ISV controller
        // outputs producer. Component 5 / Kernel 2 of the SP20 fused-
        // producer chain. Single-block, single-thread kernel reading
        // the ISV EMAs (WR_EMA, HOLD_PCT_EMA, prev HOLD_COST_SCALE)
        // and the 4 internal scratch EMAs (trade_duration_ema,
        // aux_conf_p50/std/dir_acc) Kernel 1 wrote, then deriving 6
        // controller outputs (LOSS_CAP, N_STEP, AUX_CONF_THRESHOLD,
        // AUX_GATE_TEMP, TARGET_HOLD_PCT, HOLD_COST_SCALE) per the
        // spec §4.5 formulas. HOLD_COST_SCALE is a two-sided
        // multiplicative ramp around TARGET_HOLD_PCT ±0.05 (deadband
        // path always writes the unchanged previous value to satisfy
        // `feedback_no_stubs`). All 6 controllers are anchored on
        // EMAs per `pearl_controller_anchors_isv_driven` — no
        // hardcoded magic constants drive the controllers, only
        // spec-frozen ramp / clamp parameters. No atomicAdd, no host
        // branches — captureable in the per-step CUDA Graph. Phase
        // 1.4 wires the production launch site atomically with
        // Kernel 1 + Kernel 3 per `feedback_no_partial_refactor`.
        "sp20_controllers_compute_kernel.cu",
        // SP20 Phase 1.4 (2026-05-09, Path C): per-env → SP20EmaInputs
        // aggregation kernel. Reads `trade_close_per_env`,
        // `step_ret_per_env`, `hold_at_exit_per_env` (all f32/i32
        // [N] arrays sliced to current_t from the experience
        // collector's per-sample `[N, L]` buffers) + packed factored
        // `actions_per_env` + upstream `sp20_stats_compute` outputs
        // (p50/std) + `aux_dir_acc_reduce_kernel` output [0]; tree-
        // reduces over envs and writes a single `SP20EmaInputs`
        // 36-byte struct that the post-Path-C `sp20_emas_compute_kernel`
        // reads via a device-pointer arg. Aggregation rules:
        // is_close = OR over envs; is_win = (≥0.5 of closed envs were
        // wins); trade_duration = round(mean over closed envs);
        // action_is_hold = (>n_envs/2 envs decoded to DIR_HOLD);
        // alpha / per_bar_hold_reward = 0.0 (Phase 2 / 3.2 forward
        // references). Single-block, BLOCK=64 tree-reduce; no
        // atomicAdd per `feedback_no_atomicadd`. Mandatory companion
        // to the Phase 1.2 buffer-arg kernel-signature refactor —
        // the two land atomically with the production wire-up + audit
        // / spec / plan amendments per `feedback_no_partial_refactor`.
        "sp20_aggregate_inputs_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,
    }
}