use std::path::{Path, PathBuf};
use std::process::Command;

fn main() {
    println!("cargo:rerun-if-changed=build.rs");

    // ── Feature schema fingerprint ────────────────────────────────────────────
    // Stamped into every fxcache file so caches built against a different
    // feature-extractor / state-layout / fxcache-format version fail
    // header-validation at load time and trigger automatic regeneration via
    // the Argo `ensure-fxcache` step. Removes the manual "remember to bump
    // FXCACHE_VERSION on schema change" ritual that broke the L40S deploy.
    //
    // Hashing strategy: FNV-1a 64-bit over the raw bytes of these source
    // files. Stable across rust versions and machines (unlike
    // `std::hash::DefaultHasher`). Cost: harmless cosmetic edits to these
    // files (whitespace / comments) trigger one cache regen on next deploy
    // (~5min). That cost is acceptable; missed schema drift is not.
    emit_feature_schema_hash();

    // Only compile CUDA kernels when the cuda feature is enabled
    if std::env::var("CARGO_FEATURE_CUDA").is_err() {
        return;
    }

    let out_dir = PathBuf::from(std::env::var("OUT_DIR").unwrap());
    let kernel_dir = Path::new("src/cuda_pipeline");
    let common_header = kernel_dir.join("common_device_functions.cuh");
    let trade_physics_header = kernel_dir.join("trade_physics.cuh");
    // SP4 Task A4: header-only `__device__` linear-histogram p99 estimator
    // included by `sp4_histogram_p99_test_kernel.cu` and (in subsequent
    // tasks A5-A9) by every SP4 magnitude-bound producer kernel.
    let sp4_histogram_p99_header = kernel_dir.join("sp4_histogram_p99.cuh");

    // Rebuild cubins when shared headers change
    println!("cargo:rerun-if-changed={}", trade_physics_header.display());
    println!("cargo:rerun-if-changed={}", sp4_histogram_p99_header.display());

    // Detect GPU architecture from env or default to sm_80.
    // `rerun-if-env-changed` is mandatory: the cargo-target PVC is shared
    // across H100 (sm_90) and L40S (sm_89) jobs, and without this trigger
    // cargo serves stale cubins compiled for the previous arch on every
    // subsequent build, producing CUDA_ERROR_NO_BINARY_FOR_GPU at runtime.
    println!("cargo:rerun-if-env-changed=CUDA_COMPUTE_CAP");
    let cuda_compute_cap = std::env::var("CUDA_COMPUTE_CAP").unwrap_or_else(|_| "80".to_string());
    let arch = format!("sm_{cuda_compute_cap}");

    // Check if nvcc is available -- gracefully skip if not (non-CUDA builds)
    let nvcc_path = find_nvcc();
    let nvcc = match nvcc_path {
        Some(p) => p,
        None => {
            eprintln!("  warning: nvcc not found, skipping CUDA kernel precompilation");
            eprintln!("           Install CUDA toolkit or set CUDA_HOME for GPU builds");
            return;
        }
    };

    // Read common header once
    let common_src = std::fs::read_to_string(&common_header)
        .unwrap_or_else(|e| panic!("Failed to read {}: {e}", common_header.display()));
    println!("cargo:rerun-if-changed={}", common_header.display());

    // All kernels to precompile.
    // Kernels marked "standalone" have their own helpers and don't need common header.
    // All others get common_device_functions.cuh prepended.
    let kernels_with_common = [
        // Original 24 kernels
        "epsilon_greedy_kernel.cu",
        "backtest_env_kernel.cu",
        "backtest_forward_ppo_kernel.cu",
        "backtest_forward_supervised_kernel.cu",
        "backtest_metrics_kernel.cu",
        "dt_kernels.cu",
        "ensemble_kernels.cu",
        "her_episode_kernel.cu",
        "her_relabel_kernel.cu",
        "signal_adapter_kernel.cu",
        "statistics_kernel.cu",
        "training_guard_kernel.cu",
        "c51_loss_kernel.cu",
        "mse_loss_kernel.cu",
        "curiosity_training_kernel.cu",
        "curiosity_inference_kernel.cu",
        "dqn_utility_kernels.cu",
        "attention_kernel.cu",
        "attention_backward_kernel.cu",
        "iql_value_kernel.cu",
        "iqn_dual_head_kernel.cu",
        "monitoring_kernel.cu",
        "nstep_kernel.cu",
        "reward_shaping_kernel.cu",
        "ppo_experience_kernel.cu",
        "bias_kernels.cu",
        // Formerly standalone — now need common header for BF16 types
        "experience_kernels.cu",
        "ema_kernel.cu",
        "relu_mask_kernel.cu",
        "c51_grad_kernel.cu",
        "mse_grad_kernel.cu",
        "q_stats_kernel.cu",
        "cql_grad_kernel.cu",
        "trade_stats_kernel.cu",
        "backward_kernels.cu",
        "iqn_cvar_kernel.cu",
        "mamba2_temporal_kernel.cu",
        "graph_utility_kernels.cu",
        "grad_decomp_kernel.cu",
        // SP7 Path A (2026-05-03): raw CQL norm reading `cql_grad_scratch`
        // directly so the SP7 controller has a budget-independent CQL
        // reference signal. Replaces the never-populated `cql` slot at
        // element offset 3 in `grad_decomp_result_pinned`. See
        // `cql_raw_norm_kernel.cu` and audit doc Fix 31 SP7 Path A.
        "cql_raw_norm_kernel.cu",
        "branch_grad_balance_kernel.cu",
        "backtest_plan_kernel.cu",
        "tau_update_kernel.cu",
        "epsilon_update_kernel.cu",
        "per_branch_gamma_update_kernel.cu",
        "kelly_cap_update_kernel.cu",
        "atoms_update_kernel.cu",
        "q_quantile_kernel.cu",
        // D.8 Plan 2 Task 6C: TLOB attention kernels (SDP + state scatter/read)
        "tlob_kernel.cu",
        // C.2 Plan 3 Task 1: per-component |reward| EMA into ISV[63..69)
        "reward_component_ema_kernel.cu",
        // B.2 Plan 3 Task 3: Flat→Positioned transition rate EMA into ISV[71]
        "trade_rate_ema_kernel.cu",
        // B.4 Plan 3 Task 4: per-batch readiness EMA + derived plan_threshold
        // (producer for ISV[PLAN_THRESHOLD_INDEX=49] + ISV[READINESS_EMA_INDEX=75])
        "plan_threshold_update_kernel.cu",
        // C.3 Plan 3 Task 7: train-vs-val state-distribution KL EMA + Flat-trap
        // escape amplifier (producer for ISV[78]+ISV[79]). Adaptive amp
        // multiplies B.1 opp_cost and B.2 bonus consumers.
        "state_kl_divergence_kernel.cu",
        // B.3 Plan 3 Task 8: GPU-only seeded warm-start scripted policies.
        // Replaces network-Q action source during the seed phase. 4 policies
        // (uniform/momentum/mean-rev/vwap-deviation) deterministically mixed
        // 40/20/20/20 by `i % 5`. Driven by ISV[SEED_STEPS_DONE/TARGET]
        // dispatch at the launcher boundary.
        "scripted_policy_kernel.cu",
        // B.3 Plan 3 Task 8: per-collect_experiences seed-phase progress
        // counter. Increments ISV[SEED_STEPS_DONE], EMAs the derived
        // seed-fraction into ISV[SEED_FRAC_EMA] (consumed by Task 9).
        "seed_step_counter_update_kernel.cu",
        // C.5 Plan 3 Task 9: CQL α ramp coupled to ISV[SEED_FRAC_EMA].
        // EMAs ISV[CQL_ALPHA_INDEX=48] toward `final × (1 - seed_frac)`.
        // Producer-only upgrade for slot 48; CQL gradient kernel consumer
        // (already reading ISV[48] per Plan 1 Task 12) unchanged.
        "cql_alpha_seed_update_kernel.cu",
        // E.5 Plan 4 Task 5 Mode A: attention-focus interpretability EMAs.
        // Multi-block GPU reductions — 3 blocks reducing VSN mag/dir param
        // slices + Mamba2 enriched-hidden buffer. EMA-updates ISV[87..90).
        // Diagnostic only; no consumer reads these slots in Mode A.
        "attention_focus_ema_kernel.cu",
        // Plan 4 follow-up: target-drift EMA, replaces legacy CPU-DtoH
        // per_branch_target_drift(). 2-block kernel computing
        // RMS(target - online) for mag + dir branches → ISV[92,93].
        "target_drift_kernel.cu",
        // Plan 4 Task 2c.1: Gated Residual Network kernels (forward + backward).
        // Module is additive — no production callers in this commit; Task 2c.3+4
        // wires it into the trunk encoder.
        "grn_kernel.cu",
        // Plan 4 Task 2c.3c.5: per-batch RMS EMA of `save_h_s2` into ISV[96].
        // Single-block shmem-reduce kernel (256 threads, no atomicAdd) launched
        // alongside `reward_component_ema` from `training_loop.rs`. Producer-only
        // in this commit — 2c.3c.6 wires the consumer in `mag_concat_qdir`'s
        // adaptive-scale path.
        "h_s2_rms_ema_kernel.cu",
        // Plan 4 Task 3 (E.3): IQN multi-quantile diagnostic EMAs into
        // ISV[99..103) (Q_p05/Q_p25/Q_p75/Q_p95; median τ=0.50 is the existing
        // greedy-Q diagnostic and not duplicated). 4-block kernel, one block
        // per off-median quantile, shmem-reduce over (B × TBA) of
        // `save_q_online`. No atomicAdd. Producer-only — diagnostic only.
        "iqn_quantile_ema_kernel.cu",
        // Plan 4 Task 1B-i: VSN feature-selection kernels (forward + backward).
        // Per-sample softmax-over-6-groups + gate-multiply. Caller-side cuBLAS GEMMs
        // for the per-group MLPs (Linear_1 + ReLU + Linear_2). Module is additive
        // in this commit — ZERO production callers; consumers wired in 1B-iii/iv.
        "vsn_feature_selection_kernel.cu",
        // Plan 4 Task 1B-i: VSN per-group mask EMA producer (single-block shmem
        // reduction, no atomicAdd). Reads vsn_mask saved by the forward kernel,
        // EMA-updates 6 ISV slots. Producer-only; consumer-side ISV slot
        // allocation lands in 1B-ii.
        "vsn_mask_ema_kernel.cu",
        // Plan 4 Task 6 Commit A: multi-task auxiliary heads (E.6).
        // Two `Linear(SH2 → 32) → ELU → Linear(32 → K)` MLPs branching off
        // h_s2 — next-bar return regression (K=1, MSE) + 5-class regime
        // classification (K=5, cross-entropy). Forward + backward + loss-
        // reduce + label-builder + per-tensor batch-reduce kernels. Module
        // is additive — ZERO production callers in this commit; Commit B
        // wires the forward/backward + training-loop loss accumulation.
        "aux_heads_kernel.cu",
        // Plan 4 Task 6 Commit A: aux-head loss EMA producer (single-thread
        // single-block kernel mirroring h_s2_rms_ema_kernel's footprint).
        // Producer-only; ISV[AUX_NEXT_BAR_MSE_EMA_INDEX] +
        // ISV[AUX_REGIME_CE_EMA_INDEX] consumer wires in Commit B.
        "aux_heads_loss_ema_kernel.cu",
        // Plan A Phase 0 (Thompson sampling spec 2026-04-26): standalone test
        // kernel exercised only by `distributional_q_tests.rs`. Implements the
        // Thompson direction sampling math (inverse-CDF over C51 atoms,
        // uniform-τ over IQN quantiles, argmax of E[Q]) that Phase 2 will
        // integrate into `experience_action_select`. Test-only; no production
        // callers in this commit.
        "thompson_test_kernel.cu",
        // MoE Regime Redesign Phase 2 Task 2.1: mixture-forward kernel.
        // Single-thread-per-(b,c), no atomicAdd, capture-friendly.
        // h_s2[b,c] = Σ_k g[b,k] · expert_outputs[k,b,c].
        // Subsequent tasks extend this cubin with moe_mixture_backward,
        // moe_load_balance_loss, moe_expert_util_ema_update.
        "moe_kernels.cu",
        // MoE adaptive load-balance λ controller (2026-04-27): single-block
        // single-thread cold-path producer kernel matching kelly_cap_update /
        // cql_alpha_seed_update precedent. Reads ISV[MOE_GATE_ENTROPY_EMA_INDEX=126]
        // (entropy EMA from moe_expert_util_ema_update), writes
        // ISV[MOE_LAMBDA_EFF_INDEX=128] = λ_floor + λ_max_extra × deficit.
        // Consumer (moe_load_balance_loss in moe_kernels.cu) reads slot 128
        // at runtime — no DtoH per feedback_isv_for_adaptive_bounds.md.
        "moe_lambda_eff_kernel.cu",
        // Plan C Phase 2 follow-up A.2 (2026-04-29): adaptive Polyak-tau
        // dampening signal. Single-block single-thread cold-path producer
        // mirroring moe_lambda_eff_kernel / kelly_cap_update precedent. Reads
        // ISV[Q_ABS_REF_INDEX=16] + ISV[Q_DIR_ABS_REF_INDEX=21] and the host-
        // passed q_mean_curr/prev scalars; writes
        // ISV[Q_DRIFT_RATE_INDEX=129] = clip(|q_curr-q_prev|/denom, 0, 4).
        // Consumer (tau_update_kernel.cu) multiplies tau_eff by
        // 1/(1+ISV[129]) so tau ∈ [tau_base/5, tau_base] — monotone
        // dampening under drift; healthy runs unaffected.
        "q_drift_rate_ema_kernel.cu",
        // Plan C Phase 2 follow-up K (2026-04-29): adaptive fold-boundary
        // warmup signal. Single-block single-thread cold-path producer
        // mirroring q_drift_rate_ema_kernel / moe_lambda_eff_kernel
        // precedent. Reads two grad-norm EMAs (fast α=0.1, slow α=0.001)
        // and writes ISV[FOLD_WARMUP_FACTOR_INDEX=130] = clamp(fast/slow, 0, 1).
        // Two consumers, both monotone (only dampen): lr_eff = lr_base × max(0.05, factor)
        // and clip_eff = clip_base × (0.1 + 0.9 × factor). After fold reset
        // factor starts at 0 → heavy damping; rises to 1 as gradients
        // stabilise → consumers return to baseline (healthy runs unaffected).
        "fold_warmup_factor_kernel.cu",
        // SP4 Layer C close-out C1 redesigned (2026-05-01): GPU-only
        // fast/slow grad-norm EMA update. Single-thread, single-block —
        // replaces the host-side arithmetic block in `update_adaptive_clip`
        // per `feedback_no_cpu_compute_strict`. Reads `grad_norm_buf[0]`
        // (populated earlier in the same stream by `launch_grad_norm_finalize`),
        // updates the fast (α=0.1, ~10-step horizon) and slow (α=0.001,
        // ~1000-step horizon) EMAs in mapped-pinned scalars consumed by
        // `fold_warmup_factor_kernel.cu` (the actual ISV[130] producer).
        // Cold-start sentinel (`prev ≤ 0.0` ⇒ assign obs directly) preserves
        // the deleted host-side formula's semantics exactly.
        "update_grad_norm_emas_kernel.cu",
        // SP4 Layer C close-out C2 (2026-05-01): GPU-only IQN readiness
        // gauge update. Single-thread single-block — replaces the host-side
        // arithmetic block in `update_iqn_readiness` per
        // `feedback_no_cpu_compute_strict`. Reads host-passed `iqn_loss`
        // scalar (already from `read_total_loss` mapped-pinned readback),
        // updates the three coupled scalars (loss_initial anchor, loss_ema
        // streaming smoothed loss, readiness gauge ∈ [0,1]) all backed by
        // mapped-pinned device-mapped storage. Cold-start sentinel
        // (`loss_initial < 1e-12` ⇒ bootstrap) preserves the deleted host-
        // side formula's semantics exactly. Same adaptive-α formula
        // (clamp(|err|/(|err|+0.01), 0.01, 0.30)) and improvement gauge
        // ((initial - ema) / max(initial, 1e-6)) clamped to [0, 1].
        // Consumer (c51_loss_kernel CVaR α via `iqn_readiness_dev_ptr`) is
        // unchanged — same dev_ptr, kernel writes the gauge through it.
        "update_iqn_readiness_kernel.cu",
        // SP4 Layer C close-out C3 (2026-05-01): GPU-only utilization EMA
        // update. Single-thread, single-block — replaces the host-side
        // arithmetic block in `reduce_current_q_stats` per
        // `feedback_no_cpu_compute_strict`. Reads host-passed `atom_util`
        // scalar (already a mapped-pinned readback at host[6]); updates
        // `utilization_ema_pinned` + `homeostatic_obs_pinned[1]` mirror in
        // lockstep. Cold-start sentinel `prev > 0.99 ⇒ assign obs directly`
        // (constructor init=1.0 marker) preserves the deleted host formula.
        "update_utilization_ema_kernel.cu",
        // SP4 Layer C close-out C4 (2026-05-01): GPU-only winsorized
        // adaptive grad-clip update. Single-thread, single-block — replaces
        // the host-side scalar-reduction + EMA arithmetic chain in
        // `update_adaptive_clip` per `feedback_no_cpu_compute_strict`. The
        // chain is winsor (1) + cold-start sentinel + EMA (3) + scalar
        // reduction (4) + ISV upper-bound clamp (5) + mapped-pinned write
        // (6); all six steps now run in a single kernel that reads
        // host-passed `observed_grad_norm` + the three mapped-pinned
        // scalars (`adaptive_clip_pinned`, `grad_norm_ema_pinned`,
        // `outlier_diag_pinned`) + `isv[GRAD_CLIP_BOUND_INDEX]`.
        "update_adaptive_clip_kernel.cu",
        // SP4 Layer C close-out follow-up (2026-05-01): GPU-only readiness-
        // driven homeostatic-target EMA. Single-block, six threads — one per
        // homeostatic slot — replaces the host-side EMA loop in
        // `calibrate_homeostatic_targets` per `feedback_no_cpu_compute_strict`.
        // Reads host-passed `readiness` scalar (already-clamped IQN gauge),
        // observations from `homeostatic_obs_pinned[6]` (mapped-pinned,
        // already GPU-readable via `homeostatic_obs_dev_ptr`), updates
        // `homeostatic_targets_pinned[k]` for k=1..5 with adaptive α
        // (`0.3 × (1 - readiness) + 0.01 × readiness`); thread 0 forces the
        // Q-mean invariant `targets[0] = 0.0`. Constructor-initialised target
        // defaults (`[0, 0.85, 0.1, 0, 1.0, 0.5]` at gpu_dqn_trainer.rs:14583-
        // 14588) make a separate cold-start sentinel unnecessary — the first
        // call's EMA blends those defaults with the first observation, same
        // algebraic shape the deleted host loop relied on.
        "calibrate_homeostatic_kernel.cu",
        // HEALTH_DIAG GPU port — Phase 2A landing (2026-04-28). Single-block
        // single-thread `health_diag_isv_mirror` kernel copies a curated set
        // of ISV signal-bus slots into the mapped-pinned `HealthDiagSnapshot`
        // (see `cuda_pipeline/health_diag.rs` + `gpu_health_diag.rs`).
        // Phases 2B-2E append further entry points to this same .cu file
        // (eval-histogram, q-mag-reduce, per-sample-reduce, finalise);
        // Phase 3 wires the launch chain into the captured graph; Phase 4
        // deletes the CPU-side reductions and HEALTH_DIAG emit-path Vecs.
        "health_diag_kernel.cu",
        // SP4 Layer A Task A4 (2026-04-30): standalone test kernel exercised
        // only by `tests/sp4_producer_unit_tests.rs`. Wraps the header-only
        // `sp4_histogram_p99<BLOCK_SIZE>` device function (in
        // `sp4_histogram_p99.cuh`) so the Rust unit test can drive the
        // three-pass single-block linear-histogram p99 estimator on a
        // controlled |N(0,1)| sample. Producer kernels (Tasks A5-A9) include
        // the same .cuh header directly; the test wrapper has no production
        // callers.
        "sp4_histogram_p99_test_kernel.cu",
        // SP12 v3 (2026-05-04): standalone test kernel exercised only by
        // `tests/sp12_reward_math_tests.rs`. Wraps the three reward
        // composition device functions added to `trade_physics.cuh`
        // (compute_asymmetric_capped_pnl, compute_min_hold_penalty,
        // compute_lump_sum_opp_cost) so the Rust GPU oracle tests can
        // drive each formula on synthetic inputs with analytically-known
        // expected outputs per `feedback_no_cpu_test_fallbacks`. Producer
        // call site is in `experience_kernels.cu::experience_env_step`
        // segment_complete branch; the test wrapper has no production
        // callers.
        "sp12_reward_math_test_kernel.cu",
        // SP4 Layer A Task A5 (2026-04-30): first end-to-end SP4 producer.
        // Single-block kernel reads `denoise_target_q_buf`, computes
        // p99(|target_q|) via `sp4_histogram_p99<256>`, writes the step
        // observation to `producer_step_scratch_buf[0]`. The host launcher
        // (`GpuDqnTrainer::launch_sp4_target_q_p99`) syncs, applies Pearls
        // A+D, writes new x_mean to ISV[TARGET_Q_BOUND_INDEX=131] + Wiener
        // state back to `wiener_state_buf`. Cold-path launch in this task;
        // Task A10 may relocate to the captured graph. Tasks A6-A9 mirror
        // the launcher template and slot-index discipline established here.
        "target_q_p99_kernel.cu",
        // SP4 Layer A Task A6 (2026-04-30): atom_pos_p99 × 4 branches.
        // Single parameterised single-block kernel reads ONE branch's slice
        // of `atom_positions_buf` (`[4 * num_atoms]`, branch `b` at offset
        // `b * num_atoms` per `atoms_update_kernel.cu`'s canonical layout),
        // computes p99(|atom_positions[branch]|) via `sp4_histogram_p99<256>`,
        // writes the step observation to
        // `producer_step_scratch_buf[scratch_idx]`. The launcher
        // (`launch_sp4_atom_pos_p99_all_branches`) loops `branch ∈ 0..4`,
        // launching this kernel four times with distinct `(branch_buf,
        // scratch_idx ∈ 1..5)` pairs and applying Pearls A+D host-side per
        // branch (writes back to ISV[ATOM_POS_BOUND_BASE + branch]).
        "atom_pos_p99_kernel.cu",
        // SP4 Layer A Task A7 (2026-04-30): Pearl B fused per-param-group
        // statistics oracle. Single-block kernel reads
        // `(params, grads, adam_m, adam_v)` once and runs 4 sequential
        // passes (5 for trunk):
        //   Pass A: WEIGHT_BOUND[group] = p99(|params|)  via sp4_histogram_p99
        //   Pass B: ADAM_M_BOUND[group] = p99(|adam_m|)  via sp4_histogram_p99
        //   Pass C: ADAM_V_BOUND[group] = p99(|adam_v|)  via sp4_histogram_p99
        //   Pass D: WD_RATE[group]      = |Σ w·g| / max(Σ w², EPS_DIV)
        //   Pass E: L1_LAMBDA[trunk]    = (mean|g|/mean|w|) × entropy_deficit
        //                                 (group 0 only, gated by
        //                                  l1_lambda_scratch_idx >= 0).
        // Pearl B 4× memory-bandwidth reduction vs four naive single-stat
        // producer kernels: each buffer read once for all outputs. The host
        // launcher (`launch_sp4_param_group_oracles_all_groups`) loops
        // `group ∈ 0..SP4_PARAM_GROUP_COUNT=8`, launching this kernel 8
        // times with distinct `(params, grads, m, v, count, scratch slots)`
        // tuples, then applies Pearls A+D host-side per output via
        // `pearls_ad_update`.
        "param_group_oracle_kernel.cu",
        // SP4 Layer A Task A8 (2026-04-30): grad_norm_p99 producer
        // (degenerate single-element). `grad_norm_buf` is a single-element
        // f32 device buffer; p99 over 1 element IS the element itself, so
        // the kernel reduces to a single-thread copy of `grad_norm_buf[0]`
        // into `producer_step_scratch_buf[scratch_idx=38]` with
        // `__threadfence_system()`. The host launcher
        // (`GpuDqnTrainer::launch_sp4_grad_norm_p99`) syncs, applies Pearls
        // A+D via `pearls_ad_update`, writes the new x_mean to
        // ISV[GRAD_CLIP_BOUND_INDEX=168] + Wiener state. Establishes the
        // launcher template that Mech 6's eventual Layer-B ISV read consumes.
        "grad_norm_p99_kernel.cu",
        // SP4 Layer A Task A9 (2026-04-30): h_s2_p99 producer for
        // ISV[H_S2_BOUND=169]. Single-block 256-thread kernel reading
        // `save_h_s2 [B × SH2]` (the trunk-output activation surface
        // populated each step by the online forward) via
        // `sp4_histogram_p99<256>`; output → `producer_step_scratch_buf[39]`
        // via host-side Pearls A+D applied at slot 39. Mirrors Task A5's
        // pattern exactly. NB: distinct from existing `h_s2_rms_ema_update`
        // (slot 96) which tracks RMS — different signal. Task A13 will
        // retrofit that existing kernel to also use Pearls A+D, separately.
        "h_s2_p99_kernel.cu",
        // SP4 Layer C #260 (2026-05-01): bw_d_h_s2_p99 producer for
        // ISV[BW_D_H_S2_BOUND_INDEX=171]. Single-block 256-thread kernel
        // reading `bw_d_h_s2 [B × SH2]` (the IQN trunk-backward d_h_s2 buffer
        // post-DtoD copy from `iqn_d_h_s2_ptr` and post slot-35 NaN check)
        // via `sp4_histogram_p99<256>`; output → `producer_step_scratch_buf[69]`
        // with `__threadfence_system()`. The host launcher
        // (`GpuDqnTrainer::launch_sp4_bw_d_h_s2_p99`) chains the GPU
        // `apply_pearls_ad_kernel` on the same stream — writes new x_mean to
        // ISV[171] + Wiener state at slot 69 (offset 120). Migrated from
        // SP1-Phase-C `1e6 × ISV[H_S2_RMS_EMA].max(1.0)` per
        // `feedback_isv_for_adaptive_bounds`. Mirrors `h_s2_p99_kernel.cu`'s
        // shape exactly — only differs in the consumed buffer + ISV slot.
        "bw_d_h_s2_p99_kernel.cu",
        // SP4 Layer C #260 (2026-05-01): q_dir_grad_p99 producer for
        // ISV[Q_DIR_GRAD_BOUND_INDEX=172]. Single-block 256-thread kernel
        // reading the union of `d_value_logits` ∪ `d_adv_logits` (both
        // dueling-head logits gradient buffers) via `sp4_histogram_p99_multi<256>`
        // with `n_sub=2`; output → `producer_step_scratch_buf[70]` with
        // `__threadfence_system()`. The host launcher
        // (`GpuDqnTrainer::launch_sp4_q_dir_grad_p99`) builds a 2-element
        // device-pointer table + 2-element count table in mapped-pinned memory
        // (reusing the same `oracle_subbuf_table_buf` allocated for
        // `param_group_oracle_kernel.cu`), then chains the GPU
        // `apply_pearls_ad_kernel` on the same stream — writes new x_mean to
        // ISV[172] + Wiener state at slot 70 (offset 123). Migrated from
        // SP1-Phase-C `1e6 × ISV[Q_DIR_ABS_REF].max(1.0)` per
        // `feedback_isv_for_adaptive_bounds`. The same bound was previously
        // applied to BOTH d_value_logits and d_adv_logits, so the union p99
        // is the natural ISV-driven replacement.
        "q_dir_grad_p99_kernel.cu",
        // SP5 Task A1 (2026-05-01): q_branch_stats producer.
        // Single-block 4-thread kernel reads q_out_buf [B, 13] row-major
        // and computes per-branch (mean, max_abs_dev, var, entropy) for
        // branches Direction(4)/Magnitude(3)/Order(3)/Urgency(3).
        // Writes 16 floats to producer_step_scratch_buf[71..87).
        // Shared signal source for Pearls 1, 2, 3, 5.
        "q_branch_stats_kernel.cu",
        // SP5 Task A1 (2026-05-01): Pearl 1 atom-span controller.
        // Single-block 4-thread kernel reads q_branch_stats scratch outputs
        // and per-branch atoms_clip_count, computes per-branch (v_center,
        // v_half, headroom, clip_rate) via the headroom controller
        // (TARGET_CLIP_RATE=0.01, HEADROOM_LR=0.01 — Invariant 1 anchors).
        // Writes 16 floats to producer_step_scratch_buf[87..103).
        // apply_pearls_ad_kernel then smooths via Pearls A+D into ISV[174..190).
        "pearl_1_atom_kernel.cu",
        // SP5 Task A2 (2026-05-01): Pearl 3 per-branch NoisyNet sigma controller.
        // Single-block 4-thread kernel reads ATOM_V_HALF (Pearl 1 output,
        // ISV[178..182)) + BRANCH_ENTROPY (q_branch_stats, ISV[218..222)) +
        // current SIGMA_FRACTION (ISV[214..218)); computes entropy-deficit
        // controller (target_entropy = log(n_actions[b]) * 0.7, SF_LR=0.005)
        // and new_sigma = new_sf * max(v_half, 1.0). Writes (sigma[4], SF[4])
        // to producer_step_scratch_buf[103..111).
        // apply_pearls_ad_kernel then smooths via Pearls A+D into ISV[210..218).
        "pearl_3_sigma_kernel.cu",
        // SP5 Task A3 (2026-05-01): Pearl 2 per-branch loss budget controller.
        // Single-block 4-thread kernel reads Q_VAR_PER_BRANCH (Pearl 1's
        // q_branch_stats, ISV[222..226)) + NOISY_SIGMA (Pearl 3 output,
        // ISV[210..214)); computes per-branch flatness = clamp(var_q/(σ²+EPS_DIV),
        // 0, 1) and derives (budget_c51, budget_iqn, budget_cql=0, budget_ens,
        // flatness) per branch. Writes 20 floats to scratch[111..131).
        // apply_pearls_ad_kernel then smooths via Pearls A+D into ISV[190..210).
        // Must run AFTER pearl_1_atom (Q_VAR) AND pearl_3_sigma (NOISY_SIGMA).
        "pearl_2_budget_kernel.cu",
        // SP5 Task A4 (2026-05-01): per-group gradient cosine similarity + L2 norm.
        // Single-block 8-thread kernel (one per param group: DqnTrunk/Value/Branches/
        // IQN/IqlHigh/IqlLow/Attn/Curiosity). Reads grad_buf + grad_prev_buf;
        // writes cosine_sim[8]→scratch[131..139) and l2_norm[8]→scratch[139..147).
        // Also writebacks grad_curr → grad_prev_buf for the next step's comparison.
        "grad_cosine_sim_kernel.cu",
        // SP5 Task A4 (2026-05-01): Pearl 4 per-param-group Adam β1, β2, ε.
        // Single-block 8-thread kernel reads cosine_sim + l2_norm from scratch;
        // writes β1[8]→scratch[147..155), β2[8]→scratch[155..163), ε[8]→scratch[163..171).
        // apply_pearls_ad_kernel then smooths via Pearls A+D (ALPHA_META=5e-4)
        // into ISV[226..250). Must run AFTER grad_cosine_sim_kernel.
        "pearl_4_adam_hparams_kernel.cu",
        // SP5 Task A5 (2026-05-01): auxiliary per-branch Q-distribution skew +
        // kurtosis reducer. Single-block 4-thread kernel reads save_q_online [B × 13];
        // computes per-branch (skew, excess kurtosis) via two-pass central moments;
        // writes skew[4]→scratch[171..175), ex_kurt[4]→scratch[175..179).
        // EPS_DIV=1e-12 Invariant 1 anchor. Structural envelopes [-3,+3]/[-3,+30].
        "q_skew_kurtosis_kernel.cu",
        // SP5 Task A5 (2026-05-01): Pearl 5 per-branch IQN τ schedule controller.
        // Single-block 4-thread kernel reads skew from scratch[skew_idx_base..+4);
        // computes τ = default + skew × SKEW_SHIFT=0.05 (Invariant 1 anchor),
        // clamped to [0.01, 0.99]; writes τ[4 branches × 5 quantiles]→scratch[179..199).
        // apply_pearls_ad_kernel then smooths via Pearls A+D (ALPHA_META=1e-3)
        // into ISV[IQN_TAU_BASE=250..270). Must run AFTER q_skew_kurtosis_kernel.
        "pearl_5_iqn_tau_kernel.cu",
        // SP5 Task A6 (2026-05-01): Pearl 6 cross-fold-persistent Kelly cap signals.
        // Single-block 6-thread kernel (one per ISV output slot [280..286)).
        // Reads portfolio_state [n_envs * PS_STRIDE]; writes directly to ISV[280..286).
        // CROSS-FOLD-PERSISTENT: NOT in fold-reset registry. In-kernel EWMA alpha=0.01
        // (Invariant 1 anchor). Cumulative-via-max() for sample_count slot (s==3).
        // No apply_pearls call (sentinel incompatible with cross-fold persistence).
        // Bayesian priors match kelly_cap_update_kernel.cu: prior_wins=2, prior_losses=2,
        // prior_sum_wins=0.01, prior_sum_losses=0.01 (Invariant 1 anchors).
        "pearl_6_kelly_kernel.cu",
        // SP5 Task A7 (2026-05-01): Pearl 8 per-direction trail-stop distance.
        // 4-thread single-block kernel (one per direction: Short=0, Hold=1, Long=2, Flat=3).
        // Reads features[bar_idx * market_dim + 9] (ATR_NORM) and denormalizes via
        // the canonical fxcache scheme (atr_abs = max(0.01, exp(atr_norm*16-7))).
        // Short and Long receive 2× ATR; Hold and Flat receive EPS_CLAMP_FLOOR=1.0.
        // Layer A simplification: Short and Long share the same ATR (direction-conditional
        // ATR tracking deferred to Layer C). Writes 4 floats to scratch_buf[199..203).
        // Followed by 4 apply_pearls_ad_kernel calls → ISV[TRAIL_DIST_PER_DIR_BASE=270..274).
        "pearl_8_trail_kernel.cu",
        // SP5 Task A8 (2026-05-01): Pearl 1-ext per-branch C51 num_atoms from per-branch v_half.
        // 4-thread single-block kernel (one per branch). Reads ISV[ATOM_V_HALF_BASE=178..182)
        // (populated by Pearl 1 in Task A1). Threshold cascade: v_half < 0.1 → 64 atoms,
        // v_half < 1.0 → 32 atoms, v_half >= 1.0 → 16 atoms.
        // Writes 4 floats to scratch_buf[203..207) = scratch[SCRATCH_PEARL_1_EXT_NUM_ATOMS].
        // Followed by 4 apply_pearls_ad_kernel calls → ISV[ATOM_NUM_ATOMS_BASE=274..278).
        "pearl_1_ext_num_atoms_kernel.cu",
        // SP7 Task 4 (2026-05-03): per-branch loss-balance controller for
        // CQL and C51 budgets. Single-block 8-thread kernel reads
        // grad_decomp pinned slots (IQN/CQL_SX/C51 at offsets 0/24/36 bytes),
        // FLATNESS_BASE (206..210), prior budgets (190..194 + 198..202),
        // and prior Wiener state (297..313). Writes 24 floats to scratch
        // (budgets[8] + diff_var[8] + sample_var[8]). apply_pearls_ad_kernel
        // then smooths via Pearls A+D (Wiener-α derived per-step from the
        // state slots themselves) into ISV[BUDGET_{CQL,C51}_BASE] and
        // ISV[LB_*_VAR_*_BASE]. Replaces Pearl 2's hardcoded-zero CQL
        // budget output and floor-pinned C51 budget.
        "loss_balance_controller_kernel.cu",
        // SP8 Fix 36 (2026-05-03): GPU-only train_active_frac canary
        // producer. Single-block, single-thread kernel reads
        // `monitoring_summary[5..17)` (the 12-bin action_counts block from
        // monitoring_reduce) and computes `(Short + Long) / total`. Writes
        // 1 float to `producer_step_scratch_buf[SCRATCH_TRAIN_ACTIVE_FRAC=250]`
        // and chains 1 apply_pearls_ad_kernel call → ISV[TRAIN_ACTIVE_FRAC_INDEX=321].
        // Replaces the host-side compute previously at
        // `training_loop.rs:3392-3396` per `feedback_no_cpu_compute_strict.md`.
        "train_active_frac_compute_kernel.cu",
        // SP8 Fix 36 (2026-05-03): ISV-driven adaptive MAX_BUDGET producer
        // for the loss-balance controller. Single-block, 8-thread kernel
        // (2 heads × 4 branches) reads `ISV[TRAIN_ACTIVE_FRAC_INDEX]`
        // (Pearls A+D smoothed canary) and computes per-(head, branch) cap
        // via linear interpolation `FLOOR + active_frac × (CEIL − FLOOR)`
        // (FLOOR=0.01, CEIL=1.0 — Invariant 1 numerical-stability anchors
        // per `feedback_isv_for_adaptive_bounds.md`). Writes 8 floats to
        // `producer_step_scratch_buf[SCRATCH_LB_MAX_BUDGET_*=251..259)` and
        // chains 8 apply_pearls_ad_kernel calls → ISV[LB_MAX_BUDGET_{CQL,C51}_BASE..+4).
        // Replaces the hardcoded `MAX_BUDGET=1.0f` constant in
        // `loss_balance_controller_kernel.cu` per
        // `pearl_controller_anchors_isv_driven.md`.
        "loss_balance_max_budget_compute_kernel.cu",
        // SP9 Fix 37 (2026-05-03): Kelly cold-start warmup floor — fully
        // ISV-driven. Four producer kernels per
        // `pearl_cold_start_exit_signal_or.md` and
        // `pearl_controller_anchors_isv_driven.md`:
        //   1. eval_intent_dist_compute — replaces DtoH read_eval_intent_magnitude_distribution
        //   2. intent_eval_divergence_compute — behavioral-axis numerator
        //   3. q_var_mag_ema_compute — base_floor self-relative baseline
        //   4. kelly_warmup_floor_compute — main producer (combines all axes)
        // All 4 are single-block kernels writing to producer_step_scratch_buf,
        // chained through apply_pearls_ad_kernel for Pearl A sentinel-bootstrap
        // + Pearl D Wiener-α steady-state smoothing. Per Fix 37.1
        // (2026-05-03): the 3 EMA target updaters that previously sat between
        // q_var_mag_ema and kelly_warmup_floor are deleted — the targets at
        // ISV[333..336) are Invariant-1 numerical anchors (sample-sufficiency
        // count, healthy divergence ratio, fold-temporal threshold) written
        // ONCE at constructor time per `feedback_isv_for_adaptive_bounds.md`,
        // not EMA-tracked. Pearl A sentinel-bootstrap immediately saturated
        // the EMAs to the first observation, which made conf=1.0 in epoch 1
        // (`floor=0`) — defeating the cold-start mechanism. Constructor-write
        // gives a fixed denominator so the ratio yields meaningful confidence.
        "eval_intent_dist_compute_kernel.cu",
        "intent_eval_divergence_compute_kernel.cu",
        "q_var_mag_ema_compute_kernel.cu",
        "kelly_warmup_floor_compute_kernel.cu",
        // SP11 Fix 39 (2026-05-04, Task A1): three canary producer kernels
        // for the reward-as-controlled-subsystem chain. Each is a single-
        // block producer chained through `apply_pearls_ad_kernel` for
        // Pearl A sentinel-bootstrap + Pearl D Wiener-α steady-state
        // smoothing. All three write to slots in [350..360) which no
        // consumer reads yet (Layer A is additive; consumer migration
        // lands atomically in Layer B). See spec §3.3 (canary slots) and
        // §5 (producer kernel listing).
        //   1. val_sharpe_delta_compute      → ISV[VAL_SHARPE_DELTA_EMA_INDEX=350,
        //                                          VAL_SHARPE_VAR_EMA_INDEX=351]
        //   2. saboteur_engagement_compute   → ISV[SABOTEUR_ENGAGEMENT_RATE_INDEX=358]
        //   3. reward_component_mag_ratio    → ISV[REWARD_COMPONENT_MAG_RATIO_BASE=352..358)
        //                                          ∪ ISV[PNL_REWARD_MAGNITUDE_EMA_INDEX=359]
        "val_sharpe_delta_compute_kernel.cu",
        "saboteur_engagement_compute_kernel.cu",
        "reward_component_mag_ratio_compute_kernel.cu",
        // SP11 Fix 39 B1b fix-up (2026-05-04): popart-component-specific
        // magnitude EMA producer. Reads the per-bar `r_popart` buffer
        // populated by `experience_env_step` and block-tree-reduces
        // mean(|x|) into ISV[POPART_COMPONENT_MAG_EMA_INDEX=360] via
        // chained Pearls A+D. Resolves the slot-63 PopArt-input vs
        // popart-component-magnitude overload that B1b smoke surfaced
        // (contamination → controller over-weighed popart → 10× sharpe
        // drop). Spec §4 amendment "Why slot 360".
        "popart_component_ema_kernel.cu",
        // SP11 Fix 39 (2026-05-04, Task A2): reward-subsystem controller +
        // SimHash novelty buffer. Three new kernels in this layer:
        //   1. reward_subsystem_controller — main producer (5 canaries → 10
        //      outputs in scratch). Chained `apply_pearls_ad_kernel` smooths
        //      into ISV[340..350) per spec §3.4.1 (output smoothing).
        //   2. novelty_simhash_proj_init   — one-shot Philox-driven init of
        //      the 42×16 projection matrix at trainer construct (CPU is
        //      read-only per `feedback_no_cpu_forwards`).
        //   3. novelty_simhash             — lookup + update (race-tolerated
        //      per `feedback_no_atomicadd`; under-counts bias novelty UP,
        //      the safe direction). Lookup/update launchers exist in A2 but
        //      are wired into the replay path in B1 — A2 just builds the
        //      kernels and reset registry contract end-to-end.
        "reward_subsystem_controller_kernel.cu",
        "novelty_simhash_proj_init_kernel.cu",
        "novelty_simhash_kernel.cu",
        // SP7 (2026-05-03): GPU dispatch for cql/c51 budget consumer.
        // Eliminates capture-time host-branch freeze in
        // `compute_adaptive_budgets`. The dispatch kernel resolves
        // `(active >= 0.5) ? controller_budget : bootstrap` on-device every
        // step (8 threads — 2 heads × 4 branches), writing trunk_mean +
        // per-branch correction to a 10 f32 mapped-pinned buffer. The same
        // cubin houses two new SAXPY/scale variants that read `alpha` from
        // a device pointer (instead of scalar arg), so the captured
        // `apply_c51_budget_scale` / `apply_cql_saxpy` chain reads runtime
        // ISV-derived values instead of step-0 frozen scalars. See
        // `feedback_no_cpu_compute_strict.md` and audit Fix 33.
        "consume_lb_budget_kernel.cu",
        // SP5 Layer D Task D1 (2026-05-02): PnL aggregation producer.
        // First of 3 Layer D producers replacing the host-side trade-PnL
        // aggregation in `compute_epoch_financials` per
        // `feedback_no_cpu_compute_strict.md`. Single-block 256-thread
        // shmem-reduce kernel (no atomicAdd) reads per-trade event arrays;
        // outputs cumulative PnL, mean, population variance, and max
        // drawdown to scratch[SCRATCH_PNL_AGG_BASE=207..211). Chained
        // `apply_pearls_ad_kernel` smooths via Pearls A+D into ISV[286..290).
        // Producer-only — D4 wires the consumer in the atomic Layer D
        // commit per `feedback_no_partial_refactor.md`.
        "pnl_aggregation_kernel.cu",
        // SP5 Layer D Task D2 (2026-05-02): LearningHealth composition producer.
        // Second of 3 Layer D producers replacing the host-side
        // `NormalizedComponents::from_raw` + `compose` chain in
        // `learning_health.rs` (called from `training_loop.rs:2658-2745`)
        // per `feedback_no_cpu_compute_strict.md`. Single-block 1-thread
        // arithmetic kernel: 7 raw inputs by value → 7 in-register
        // smoothstep normalisations → weighted sum into a scalar health
        // score; writes 4 scratch slots (composed score + q_gap_norm +
        // q_var_norm + grad_norm_norm) at scratch[SCRATCH_HEALTH_COMP_BASE
        // =211..215). Chained `apply_pearls_ad_kernel` smooths via Pearls
        // A+D into ISV[290..294). Producer-only — D4 wires the consumer
        // in the atomic Layer D commit per
        // `feedback_no_partial_refactor.md`.
        "health_composition_kernel.cu",
        // SP5 Layer D Task D3 (2026-05-02): training-metrics EMA producer.
        // Third (and final) Layer D producer replacing the host-side
        // adversarial-regime + adaptive-DSR EMA arithmetic block at
        // `training_loop.rs:5039-5113` per `feedback_no_cpu_compute_strict.md`.
        // Single-block 3-thread kernel — one thread per metric:
        //   tid=0 → training_sharpe_ema (adaptive α, host-tracked sentinel),
        //   tid=1 → max_dd_ema           (fixed α=0.1, sentinel<1e-12),
        //   tid=2 → low_dd_ratio         (fixed α=0.15, depends on tid=1's
        //                                  output via __syncthreads()).
        // Writes 3 scratch slots at scratch[SCRATCH_TRAINING_METRICS_EMA_BASE
        // =215..218). Chained `apply_pearls_ad_kernel` (3 launches,
        // ALPHA_META=1e-3) smooths via Pearls A+D into ISV[294..297).
        // Producer-only — D4 wires the consumer in the atomic Layer D
        // commit per `feedback_no_partial_refactor.md`.
        "training_metrics_ema_kernel.cu",
        // SP4 GPU-only Pearls A+D applicator (2026-05-01). Single-thread
        // device-side loop applies Pearls A+D to N consecutive ISV slots
        // (each with its own Wiener-state triple) on the producer's
        // stream. Replaces the pre-refactor host-side `apply_pearls_to_slot`
        // helper that broke CUDA Graph capture at `aux_label_scale_ema`
        // mid-step (smoke smoke-test-v9kjv: CUDA_ERROR_STREAM_CAPTURE_UNSUPPORTED).
        // Consumed by all 11 Pearls-using launchers + the inline
        // aux_heads_forward Step 2b call site. Constants (ALPHA_META,
        // EPS_DIV) mirror sp4_wiener_ema.rs's host-reference for the
        // GPU-vs-host oracle test.
        "apply_pearls_kernel.cu",
        // SP13 Phase 0a (2026-05-04): aux-head directional-accuracy
        // reducer. Single-block tree-reduce kernel scoring the aux next-
        // bar regression head's prediction sign against the next-bar
        // return label sign. Four shared-memory int arrays
        // (correct/pos_pred/pos_label/valid) reduce in lockstep — no
        // atomicAdd per `feedback_no_atomicadd.md`. Outputs three
        // batch-level fractions (dir_acc, pos_pred_frac, pos_label_frac)
        // into a 3-element mapped-pinned buffer. P0a reads regression-
        // mode aux scalar; Layer B swaps the source pointer to
        // `softmax[1] - softmax[0]` after the head becomes 2-class
        // softmax. Consumed by `launch_aux_dir_acc_reduce` in
        // `gpu_dqn_trainer.rs`.
        "aux_dir_acc_reduce_kernel.cu",
        // SP13 v3 Phase 0a P0a.T3 (2026-05-04): Hold-rate observer.
        // Single-block tree-reduce kernel counting per-bar Hold-action
        // picks (decoded from packed `batch_actions[i]` via
        // `dir = action_idx / (NUM_MAGNITUDES * NUM_ORD * NUM_URG)` —
        // P0a.T4 ABI alignment, 2026-05-04) and writing the fraction
        // `count(Hold) / batch_size` to a 1-element mapped-pinned buffer.
        // One shared-memory int array reduces in lockstep — no atomicAdd
        // per `feedback_no_atomicadd.md`. The fixed-α EMA applicator
        // smooths the observation into ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382];
        // the host-side controller reads ISV[382] + ISV[381] and writes the
        // priced Hold cost back to ISV[HOLD_COST_INDEX=380]. The reward
        // composition site in `experience_kernels.cu` then subtracts the
        // cost on every Hold-action bar — pricing the action so the
        // policy uses Hold deliberately rather than as a free CQL-bias-
        // anchored default. Replaces v2's atomic Hold elimination per
        // the v3 reframe (preserves the 4-way action space — no cross-
        // crate cascade). Consumed by `launch_hold_rate_observer` in
        // `gpu_dqn_trainer.rs`.
        "hold_rate_observer_kernel.cu",
        // SP13 Phase 0a P0a.T4 (2026-05-04): fixed-α EMA applicator.
        // Sibling of `apply_pearls_kernel` — needed because Pearls A+D
        // collapse two EMAs of the same signal to identical values
        // (variance ratio is α-independent), which would deadlock the
        // SP13 stagnation detector that compares slots 373 (α=0.3) vs
        // 374 (α=0.05). Per-thread first-observation-sentinel bootstrap
        // (caller passes the slot's registry sentinel — 0.5 for dir-acc,
        // 0.0 for hold-rate) plus a fixed-α blend. No atomicAdd per
        // `feedback_no_atomicadd.md`; pure GPU compute per
        // `feedback_no_cpu_compute_strict.md`. Consumed by
        // `launch_apply_fixed_alpha_ema` in `gpu_dqn_trainer.rs`.
        "apply_fixed_alpha_ema_kernel.cu",
        // SP13 Phase 0a P0a.T4 (2026-05-04): aux-head per-bar prediction
        // → ISV[AUX_DIR_PREDICTION_INDEX=375] tanh-bounded scalar
        // producer. Single-block tree-reduce reads the captured-graph
        // `aux_nb_pred_buf [B]` tile and writes
        // `mean(tanh(aux_pred[i]))` ∈ [-1, +1] to the SHARED ISV slot 375
        // (per-step state, not an EMA — slot is overwritten each launch).
        // tanh squash bounds the scalar so the downstream Q-head
        // consumer sees a well-conditioned signal regardless of
        // `label_scale` magnitude. No atomicAdd per
        // `feedback_no_atomicadd.md`; pure GPU compute per
        // `feedback_no_cpu_compute_strict.md`. Consumed by
        // `launch_aux_pred_to_isv_tanh` in `gpu_dqn_trainer.rs`.
        "aux_pred_to_isv_tanh_kernel.cu",
        // SP13 Layer B Commit B1.1b (2026-05-05): producer kernel for
        // the aux next-bar sign label `[total] i32` ring column. Replaces
        // the B0 `alloc_zeros::<i32>(total)` placeholder in
        // `gpu_experience_collector.rs::collect_experiences_gpu` with a
        // real classifier reading `targets[bar*6+2]` (raw_close column)
        // at `bar` and `bar+lookahead`, writing -1 (skip if window runs
        // past `total_bars`), 0 (down/flat: `p_fut <= p_now`), or 1
        // (up: `p_fut > p_now`) per experience. Pure per-thread O(1)
        // map — no atomicAdd / no reduction per
        // `feedback_no_atomicadd.md`. Consumer is the K=2 softmax CE
        // head wired in B1.1a (`aux_next_bar_loss_reduce` +
        // `aux_next_bar_backward` +`aux_dir_acc_reduce_kernel` +
        // `aux_pred_to_isv_tanh_kernel`); without B1.1b every label is
        // 0, so the head trains on "everything is class 0" — the
        // known degraded behavior between B1.1a and B1.1b that this
        // kernel resolves.
        "aux_sign_label_kernel.cu",
        // SP14 Layer B Task B.3 (2026-05-05): Earned Gradient Flow
        // producer kernel. Per-step computes the K=4↔K=2 mapped argmax
        // mismatch between the Q-head's 4-way direction action
        // (DIR_SHORT/DIR_HOLD/DIR_LONG/DIR_FLAT) and the aux head's
        // 2-way next-bar prediction (down/up), masking Hold/Flat (no
        // new directional commitment). Updates fast + slow EMAs of
        // the disagreement rate (ISV[383], ISV[384]) plus a Welford-
        // style variance EMA (ISV[389]) in a single launch — the
        // variance feeds the adaptive k_q sigmoid steepness consumed
        // by `alpha_grad_compute_kernel` in B.4. Block tree-reduce on
        // numerator + count separately; no atomicAdd per
        // `feedback_no_atomicadd.md`. Pearl-A first-observation
        // bootstrap on the EMAs (sentinel 0.5 → batch_mean direct
        // replacement) per `pearl_first_observation_bootstrap.md`.
        // Slot indices hardcoded inside the kernel via `#define` —
        // must match `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs`
        // (currently 383, 384, 389; the layout fingerprint regression
        // test catches drift).
        "q_disagreement_update_kernel.cu",
        // SP14 Layer B Task B.4 (2026-05-05): Earned Gradient Flow
        // consumer kernel — the heart of the EGF pearl. Single-thread
        // state-machine kernel that reads driver signals (target dir-
        // acc, aux dir-acc EMA, Q-disagreement EMA, three Welford
        // variance EMAs, persistent Gate-1 state) from ISV, runs a
        // Schmitt-trigger Gate 1 (open at target+0.03, close at
        // target-0.03), computes adaptive k_aux/k_q/β from the
        // variance EMAs, evaluates two sigmoids (Gate 1 = aux
        // competence, Gate 2 = Q-head agreement vs random-alignment
        // baseline 0.5), multiplies with the host-supplied warmup
        // gate to produce α_grad_raw, applies a β-rate-limiter to
        // produce α_grad_smoothed, and writes 7 outputs (k_aux, k_q,
        // β, var_α, gate1_state, α_raw, α_smoothed) to ISV slots
        // 385-393. Sigmoid arguments clipped to [-30, 30] for fp32
        // overflow guard. Per pearl_bounded_modifier_outputs_require_
        // structural_activation: sigmoid composition is structurally
        // bounded to [0, 1], no runtime clamp. KNOWN LIMITATION: the
        // var_aux producer (slot 388) is not yet wired as of B.4 —
        // until B.11 / fix-up adds it, k_aux is degenerate-but-non-
        // fatal at K_BASE_AUX (constant). Filed in the B.4 status
        // report. var_q (slot 389) is written by
        // q_disagreement_update_kernel (B.3), so adaptive k_q is
        // fully functional.
        "alpha_grad_compute_kernel.cu",
        // SP14 Layer B Task B.5 (2026-05-05): Earned Gradient Flow
        // anti-mesa-optimization circuit breaker. Single-thread state-
        // machine kernel that runs at the end of each epoch. Detects
        // suspected gradient hacking: gate1 open AND aux_dir_acc post-
        // open-minimum drops > LOCKOUT_TRIGGER_DROP (0.05) below the
        // Schmitt open-threshold (target + 0.03) AND q_disagreement rises
        // > LOCKOUT_TRIGGER_DIS_RISE (0.10) above the analytic baseline
        // 0.5, simultaneously. Action: force gate1_open_state = 0 and
        // set lockout_remaining = 2 epochs. During lockout, gate1 stays
        // force-closed. Tracks AUX_DIR_ACC_POST_OPEN_MIN (ISV[394]):
        // lowest aux_dir_acc since gate last opened; resets to 1.0 when
        // gate closes. Slot indices must match sp14_isv_slots.rs (shifted
        // +2 from the plan: 383, 391, 394, 395). Depends on B.4
        // alpha_grad_compute_kernel writing GATE1_OPEN_STATE (ISV[391])
        // before this kernel runs per epoch.
        "gradient_hack_detect_kernel.cu",
        // 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 B Task B.10 (2026-05-05): Earned Gradient Flow
        // backward gate. Scales the wire column of `dL/dx_concat
        // [B, SH2+1]` (gradient flowing FROM the direction Q-head's
        // first FC SGEMM TO `aux_softmax_diff`) by ISV[393] =
        // ALPHA_GRAD_SMOOTHED, computed by `alpha_grad_compute_kernel`
        // (B.4) and orchestrated per-step by B.11. Touches ONLY column
        // SH2 of the [B, SH2+1] gradient buffer; dW (Q-head's own
        // weight gradient computed by an independent SGEMM) stays
        // unscaled per the pearl design — Q-head learns to USE the
        // wire freely; only the back-flow to aux is gated. Pre-B.11
        // (unwired producer) ISV[393] sentinels at 0.0 → wire force-
        // closed for safety (gradient zeroed). Pure per-thread map:
        // one thread per batch row's wire element. No atomicAdd.
        "sp14_scale_wire_col_kernel.cu",
        // 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): drawdown reporting. Single-thread,
        // single-block per-step state machine. Reads existing
        // PS_PEAK_EQUITY (slot 7) and PS_PREV_EQUITY (slot 9) from the
        // position state buffer (canonical equity tracking — NO new ISV
        // slot, invariant verified during plan v2 critical review).
        // Writes 6 ISV slots: DD_CURRENT (401), DD_MAX (402),
        // DD_RECOVERY_BARS (403), DD_PERSISTENCE (404), CALMAR (405;
        // floored max_dd for host composer = mean_pnl / value-here),
        // DD_PCT (406). The 1e-4 calmar floor eliminates the saturation-
        // at-100 artifact previously seen in the train-dd4xl HEALTH_DIAG.
        "dd_state_kernel.cu",
        // SP15 Phase 1.4 partial (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 sharpe_net to its own ISV slot (409 buyhold,
        // 410 hold_only, 412 naive_momentum, 416 naive_reversion).
        // 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 Phase 1.5 (2026-05-06): dd_pct foundational state input
        // concat. LAYOUT FINGERPRINT BREAK — pre-SP15 checkpoints WILL
        // NOT LOAD after this commit. Greenfield OK per spec Q1. Per
        // spec §6.5: produces a [B, state_dim_padded + 1] concat buffer
        // whose last column is `isv[DD_PCT_INDEX=406]` broadcast across
        // batch. Eval-time policy SEES drawdown context on every forward
        // pass; Phase 3 teachings can condition on dd_pct directly via
        // state, not just reward modulation. Phase 1.5 lands kernel +
        // launcher + layout fingerprint marker only; trunk consumer
        // migration (re-pointing forward_online to consume the concat
        // buffer + bumping s1_input_dim) lands as a follow-up atomic
        // commit per the established Phase 1.1-1.4 precedent.
        "dd_pct_concat_kernel.cu",
        // SP15 Phase 3.1 (2026-05-06): r_quality + r_discipline split
        // kernel + ISV-driven α producer kernel (single source file, two
        // `extern "C" __global__` symbols sharing one cubin per the
        // 1:1-source-to-cubin / multi-kernel-per-file convention used by
        // `baseline_kernels.cu` and `novelty_simhash_kernel.cu`).
        // Composer (`r_quality_discipline_split_kernel`) reads
        // ALPHA_SPLIT_INDEX=417 + alpha_warm_count scratch buffer, writes
        // r_total + increments warm count. Producer
        // (`alpha_split_producer_kernel`) reads grad-norm slots 418/419
        // + warm count, writes ALPHA_SPLIT_INDEX clamped to [0.05, 0.95]
        // (silent no-op until `*alpha_warm_count >= N_WARM=100`).
        "r_quality_discipline_split_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,
    }
}
