diff --git a/crates/ml/build.rs b/crates/ml/build.rs index a03264ee1..93021452d 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -491,6 +491,20 @@ fn main() { // 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", // 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 diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 34fe0f0e7..8a4f4adf4 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -408,6 +408,24 @@ static SP5_PEARL_1_EXT_NUM_ATOMS_CUBIN: &[u8] = static SP5_PNL_AGGREGATION_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/pnl_aggregation_kernel.cubin")); +/// SP5 Layer D Task D2 (2026-05-02): LearningHealth composition producer. +/// Single-block 1-thread arithmetic kernel takes 7 raw inputs by value +/// (q_gap, q_var, atom_util, grad_norm, ens_disagreement, grad_consistency, +/// spectral_gap), normalises each via smoothstep curves matching +/// `learning_health.rs::NormalizedComponents::from_raw` exactly, and +/// composes them into a scalar `health_score` via the spec Layer-1 +/// weighted sum (`compose()`). Writes 4 scratch floats — composed score +/// + 3 normalised intermediates (q_gap_norm, q_var_norm, +/// grad_norm_norm = `grad_stable`) — to +/// scratch_buf[SCRATCH_HEALTH_COMP_BASE=211..215). Followed by 4 +/// apply_pearls_ad_kernel calls → ISV[HEALTH_SCORE_INDEX=290..294). +/// Replaces the host-side composition in `learning_health.rs` invoked from +/// `training_loop.rs:2658-2745` per `feedback_no_cpu_compute_strict.md`. +/// Additive in this commit — D4 wires the consumer in the atomic +/// Layer D refactor per `feedback_no_partial_refactor.md`. +static SP5_HEALTH_COMPOSITION_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/health_composition_kernel.cubin")); + /// Plan 4 Task 3 (E.3): IQN multi-quantile diagnostic EMAs into ISV[99..103). /// 4-block (256 threads/block, shmem-reduce, no atomicAdd) kernel launched /// alongside `h_s2_rms_ema_update` from `training_loop.rs`. Reads the IQN @@ -774,7 +792,7 @@ const ISV_NETWORK_DIM: usize = 23; /// (shifted 112→116 in Plan 4 Task 6 Commit A). /// Written by the constructor; checked at checkpoint load. Fail-fast only — no migration /// path exists. See spec §4.A.2 and `LAYOUT_FINGERPRINT_CURRENT` for structural-hash rationale. -const ISV_TOTAL_DIM: usize = 290; // SP5 + Layer D D1: 173 + 117 (114 SP5 slots @ 174..278/280..290, with 2-slot gap before cross-fold-persistent Kelly block; Layer D D1 PnL outputs at [286..290)) +const ISV_TOTAL_DIM: usize = 294; // SP5 + Layer D D1+D2: 173 + 121 (118 SP5 slots @ 174..278/280..294, with 2-slot gap before cross-fold-persistent Kelly block; Layer D D1 PnL outputs at [286..290); Layer D D2 health composition outputs at [290..294)) /// Legacy alias preserved for call sites that haven't been audited for the /// network-vs-total split. New code should pick `ISV_NETWORK_DIM` (for weight /// tensor sizing) or `ISV_TOTAL_DIM` (for the broadcast bus buffer). @@ -982,24 +1000,27 @@ pub const SP4_WIENER_TOTAL_FLOATS: usize = // ── SP5 Task A1 buffer layout constants ─────────────────────────────────────── -/// SP5 Task A1 + Layer D Task D1: combined wiener_state_buf total float count. -/// SP4 had 71 producers (213 floats); SP5 contributes 116 producer triples -/// (linear span from `SP5_SLOT_BASE=174` to `SP5_SLOT_END=290`, including the -/// 2-slot carve-out gap at 278..280; 116 × 3 = 348 floats). Combined = -/// (71 + 116) × 3 = 561 floats. Pre-D1 SP5 was 110 producers / 543 floats; -/// Task D1 added 4 unique slots (286..290) which extended the linear span by -/// 4 from 112 to 116, growing the buffer by 12 floats. Note: SP5_PRODUCER_COUNT -/// is the **linear-span** constant — see its docstring for why it now diverges -/// from the unique-slot count (114). +/// SP5 Task A1 + Layer D Task D1+D2: combined wiener_state_buf total float count. +/// SP4 had 71 producers (213 floats); SP5 contributes 120 producer triples +/// (linear span from `SP5_SLOT_BASE=174` to `SP5_SLOT_END=294`, including the +/// 2-slot carve-out gap at 278..280; 120 × 3 = 360 floats). Combined = +/// (71 + 120) × 3 = 573 floats. Pre-D1 SP5 was 110 producers / 543 floats; +/// Task D1 added 4 unique slots (286..290) growing the linear span 112 → 116 +/// (buffer 543 → 561). Task D2 adds 4 more slots (290..294) growing the linear +/// span 116 → 120 (buffer 561 → 573). Note: SP5_PRODUCER_COUNT is the +/// **linear-span** constant — see its docstring for why it now diverges +/// permanently from the unique-slot count (118 post-D2). /// /// wiener_state_buf must grow to this size at construction for ALL SP5 slots -/// (including the new D1 PnL-aggregation block) to be in-bounds when -/// `apply_pearls_ad_kernel` writes their Wiener triples. The Pearl 6 Kelly -/// block (slots 280..286) and the carve-out gap (278..280) have reserved-but- -/// unused wiener storage at offsets [(280-174)*3+213..(286-174)*3+213) = -/// [531..549) — Pearl 6 does not call `launch_apply_pearls` (cross-fold -/// persistent), so those 18 floats stay zero-initialised. Layer D D1 slots -/// 286..290 use offsets [(286-174)*3+213..(290-174)*3+213) = [549..561). +/// (including the new D1 PnL-aggregation and D2 health-composition blocks) to +/// be in-bounds when `apply_pearls_ad_kernel` writes their Wiener triples. +/// The Pearl 6 Kelly block (slots 280..286) and the carve-out gap (278..280) +/// have reserved-but-unused wiener storage at offsets +/// [(280-174)*3+213..(286-174)*3+213) = [531..549) — Pearl 6 does not call +/// `launch_apply_pearls` (cross-fold persistent), so those 18 floats stay +/// zero-initialised. Layer D D1 slots 286..290 use offsets +/// [(286-174)*3+213..(290-174)*3+213) = [549..561). Layer D D2 slots 290..294 +/// use offsets [(290-174)*3+213..(294-174)*3+213) = [561..573). pub const SP5_WIENER_TOTAL_FLOATS: usize = (SP4_PRODUCER_COUNT + crate::cuda_pipeline::sp5_isv_slots::SP5_PRODUCER_COUNT) * SP4_WIENER_FLOATS_PER_SLOT; @@ -1034,8 +1055,10 @@ pub const SP5_WIENER_TOTAL_FLOATS: usize = /// 4 floats for pearl_1_ext_num_atoms output (SCRATCH_PEARL_1_EXT_NUM_ATOMS=203, num_atoms × 4 branches at [203..207)) /// Layer D Task D1 adds: /// 4 floats for pnl_aggregation output (SCRATCH_PNL_AGG_BASE=207, total/mean/var/max_dd at [207..211)) -/// Combined: 71 + 16 + 16 + 4 + 4 + 20 + 8 + 8 + 8 + 8 + 8 + 4 + 4 + 20 + 4 + 4 + 4 = 211 scratch slots [0..211). -pub const SP5_SCRATCH_TOTAL: usize = 211; +/// Layer D Task D2 adds: +/// 4 floats for health_composition output (SCRATCH_HEALTH_COMP_BASE=211, score/q_gap_norm/q_var_norm/grad_norm_norm at [211..215)) +/// Combined: 71 + 16 + 16 + 4 + 4 + 20 + 8 + 8 + 8 + 8 + 8 + 4 + 4 + 20 + 4 + 4 + 4 + 4 = 215 scratch slots [0..215). +pub const SP5_SCRATCH_TOTAL: usize = 215; /// SP5 Layer D Task D1 (2026-05-02): scratch index base for `pnl_aggregation_update`. /// Slots [207..211): aggregated PnL summary outputs in fixed order: @@ -1047,6 +1070,22 @@ pub const SP5_SCRATCH_TOTAL: usize = 211; /// ISV[PNL_TOTAL_INDEX=286..PNL_MAX_DD_INDEX+1=290). pub const SCRATCH_PNL_AGG_BASE: usize = 207; +/// SP5 Layer D Task D2 (2026-05-02): scratch index base for `health_composition_update`. +/// Slots [211..215): LearningHealth composition outputs in fixed order: +/// [211] = health_score (composed Layer-1 score in [0,1]) +/// [212] = q_gap_norm (smoothstep(0.01, 0.5, q_gap_ema) ∈ [0,1]) +/// [213] = q_var_norm (smoothstep(0.001, 0.1, q_var_ema) ∈ [0,1]) +/// [214] = grad_norm_norm (1 − smoothstep(10, 100, grad_norm_ema) ∈ [0,1] +/// — `grad_stable` in `learning_health.rs`) +/// Written by `health_composition_update`; consumed by `apply_pearls_ad_kernel` +/// → ISV[HEALTH_SCORE_INDEX=290..GRAD_NORM_NORM_INDEX+1=294). The other 4 +/// normalised intermediates (`atom_util_norm`, `ens_agree`, +/// `grad_consistency_norm`, `spectral_gap_norm`) are computed in-register +/// inside the kernel, fed into the composed score, and discarded — they +/// can be promoted to scratch slots in a future Layer D extension without +/// a launcher signature change. +pub const SCRATCH_HEALTH_COMP_BASE: usize = 211; + /// SP5 Task A7: scratch index base for pearl_8_trail_update trail_dist[4] output block. /// Slots [199..203): per-direction trail-stop distance (Short=0, Hold=1, Long=2, Flat=3). /// Written by `pearl_8_trail_update`; consumed by apply_pearls_ad_kernel → @@ -1574,7 +1613,8 @@ const fn layout_fingerprint_seed() -> &'static [u8] { KELLY_F_SMOOTH=280;CONVICTION_SMOOTH=281;TRADE_VAR_SMOOTH=282;\ KELLY_SAMPLE_COUNT=283;WIN_RATE_SMOOTH=284;LOSS_RATE_SMOOTH=285;\ PNL_TOTAL=286;PNL_MEAN=287;PNL_VAR=288;PNL_MAX_DD=289;\ - ISV_TOTAL_DIM=290;\ + HEALTH_SCORE=290;Q_GAP_NORM=291;Q_VAR_NORM=292;GRAD_NORM_NORM=293;\ + ISV_TOTAL_DIM=294;\ PARAM_W_A_H_S1=0;PARAM_B_A_H_S1=1;PARAM_W_B_H_S1=2;PARAM_B_B_H_S1=3;\ PARAM_W_RESIDUAL_H_S1=4;PARAM_GAMMA_H_S1=5;PARAM_BETA_H_S1=6;\ PARAM_W_A_H_S2=7;PARAM_B_A_H_S2=8;PARAM_W_B_H_S2=9;PARAM_B_B_H_S2=10;\ @@ -4367,6 +4407,21 @@ pub struct GpuDqnTrainer { /// Loaded from `pnl_aggregation_kernel.cubin`. Producer-only in this commit; /// D4 wires the consumer atomically. pnl_aggregation_kernel: CudaFunction, + /// SP5 Layer D Task D2 (2026-05-02): LearningHealth composition producer. + /// Single-block 1-thread arithmetic kernel takes 7 raw inputs by value + /// (q_gap, q_var, atom_util, grad_norm, ens_disagreement, + /// grad_consistency, spectral_gap), normalises each via smoothstep curves + /// matching `learning_health.rs::NormalizedComponents::from_raw` exactly, + /// composes them via the spec Layer-1 weighted sum, and writes 4 scratch + /// floats — composed score + 3 normalised intermediates (q_gap_norm / + /// q_var_norm / grad_norm_norm) — to scratch_buf[SCRATCH_HEALTH_COMP_BASE + /// =211..215). Chained apply_pearls_ad_kernel (4 launches, ALPHA_META=1e-3) + /// → ISV[HEALTH_SCORE_INDEX=290..GRAD_NORM_NORM_INDEX+1=294). In fold-reset + /// registry (FoldReset; the underlying EMAs reset at fold boundary so the + /// composition outputs reset together via the standard Pearl A sentinel + /// path). Loaded from `health_composition_kernel.cubin`. Producer-only in + /// this commit; D4 wires the consumer atomically. + health_composition_kernel: CudaFunction, /// SP5 Task A4 (2026-05-01): previous-step gradient buffer for cosine similarity. /// [total_params f32] mapped-pinned (feedback_no_htod_htoh_only_mapped_pinned). /// Zeroed at construction and at each fold boundary (Pearl A sentinel: cosine_sim=0 @@ -11748,6 +11803,131 @@ impl GpuDqnTrainer { Ok(()) } + /// SP5 Layer D Task D2 (2026-05-02): launch `health_composition_update` — + /// single-block 1-thread arithmetic producer for ISV[HEALTH_SCORE_INDEX=290..GRAD_NORM_NORM_INDEX+1=294). + /// + /// Reads 7 raw input signals as f32 scalars (passed by value; no ISV reads), + /// reproduces `learning_health.rs::NormalizedComponents::from_raw` + + /// `compose()` bit-for-bit on GPU, and writes 4 aggregated floats to + /// scratch_buf[SCRATCH_HEALTH_COMP_BASE=211..215): + /// [211] composed health_score (Layer-1 weighted sum, in [0,1]) + /// [212] q_gap_norm (smoothstep(0.01, 0.5, q_gap)) + /// [213] q_var_norm (smoothstep(0.001, 0.1, q_var)) + /// [214] grad_norm_norm (1 − smoothstep(10, 100, grad_norm)) + /// + /// Then chains 4 `apply_pearls_ad_kernel` calls smoothing through Pearls + /// A+D into ISV[290..294) on the same stream (no host sync per + /// `pearl_cold_path_no_exception_to_gpu_drives.md` and the SP5 GPU-only + /// Pearls refactor). + /// + /// Per `feedback_no_cpu_compute_strict.md` — replaces the host-side + /// `LearningHealth` composition currently driven from + /// `training_loop.rs:2658-2745`. Per `feedback_no_partial_refactor.md` + /// — D2 is additive only; D4 atomically migrates the host-side site to + /// call this launcher in lockstep with D1/D3. + /// + /// **Formula fidelity:** the kernel reproduces the host-side + /// `NormalizedComponents::from_raw` + `NormalizedComponents::compose` + /// chain verbatim — same edge constants, same weights, no algorithmic + /// change. Only the EMA + warmup + clamp wrapper from + /// `LearningHealth::update` is left to the host (and to the downstream + /// apply_pearls smoothing); D2 is structurally additive. + /// + /// Arguments (all f32 by value, no mapped-pinned plumbing per the + /// `learning_health.rs` pipeline shape): + /// - `q_gap`: per-branch Q-gap mean (host-aggregated EMA, 4 branches averaged). + /// - `q_var`: epoch q_max − q_min range proxy. + /// - `atom_util`: C51 atom utilisation fraction in [0,1]. + /// - `grad_norm`: |avg_grad| from the epoch (positive scalar). + /// - `ens_disagreement`: dispersion across per-branch Q-gaps (range proxy). + /// - `grad_consistency`: Adam m-flat cosine similarity (or scalar fallback). + /// - `spectral_gap`: σ₁/σ₂ from rolling Gram + power iteration. + pub(crate) fn launch_health_composition( + &self, + q_gap: f32, + q_var: f32, + atom_util: f32, + grad_norm: f32, + ens_disagreement: f32, + grad_consistency: f32, + spectral_gap: f32, + ) -> Result<(), MLError> { + use crate::cuda_pipeline::sp5_isv_slots::{ + HEALTH_SCORE_INDEX, Q_GAP_NORM_INDEX, Q_VAR_NORM_INDEX, GRAD_NORM_NORM_INDEX, + SP5_SLOT_BASE, + }; + use crate::cuda_pipeline::sp4_wiener_ema::{launch_apply_pearls, ALPHA_META}; + + debug_assert!(self.isv_signals_dev_ptr != 0, + "launch_health_composition: isv_signals_dev_ptr must be allocated by constructor"); + + let scratch_dev = self.producer_step_scratch_buf.dev_ptr; + let isv_dev = self.isv_signals_dev_ptr; + let wiener_dev = self.wiener_state_buf.dev_ptr; + + let health_score_idx_i32: i32 = (SCRATCH_HEALTH_COMP_BASE + 0) as i32; + let q_gap_norm_idx_i32: i32 = (SCRATCH_HEALTH_COMP_BASE + 1) as i32; + let q_var_norm_idx_i32: i32 = (SCRATCH_HEALTH_COMP_BASE + 2) as i32; + let grad_norm_norm_idx_i32: i32 = (SCRATCH_HEALTH_COMP_BASE + 3) as i32; + + // Wiener base offset for SP5 slots: SP4_PRODUCER_COUNT * 3 = 213. + let base_wiener_offset: i32 = (SP4_PRODUCER_COUNT * 3) as i32; + + // Step 1: health_composition_update — single-block 1-thread; no shared + // memory needed (composition is 7 reads / 7 smoothsteps / 7 multiply- + // adds in registers). + unsafe { + self.stream + .launch_builder(&self.health_composition_kernel) + .arg(&q_gap) + .arg(&q_var) + .arg(&atom_util) + .arg(&grad_norm) + .arg(&ens_disagreement) + .arg(&grad_consistency) + .arg(&spectral_gap) + .arg(&scratch_dev) + .arg(&health_score_idx_i32) + .arg(&q_gap_norm_idx_i32) + .arg(&q_var_norm_idx_i32) + .arg(&grad_norm_norm_idx_i32) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("health_composition_update: {e}")))?; + } + + // Step 2: 4 apply_pearls calls (one per output slot). Each + // scratch[211+k] → ISV[290+k] with its own Wiener triple at offset + // (base_wiener_offset + (isv_slot − SP5_SLOT_BASE) * 3). + let isv_slots: [usize; 4] = [ + HEALTH_SCORE_INDEX, + Q_GAP_NORM_INDEX, + Q_VAR_NORM_INDEX, + GRAD_NORM_NORM_INDEX, + ]; + for k in 0..4_usize { + let isv_slot = isv_slots[k] as i32; + let scratch_idx = (SCRATCH_HEALTH_COMP_BASE + k) as i32; + let wiener_off = base_wiener_offset + (isv_slot - SP5_SLOT_BASE as i32) * 3; + unsafe { + launch_apply_pearls( + &self.stream, + &self.apply_pearls_ad_kernel, + scratch_dev, scratch_idx, + isv_dev, isv_slot, + wiener_dev, wiener_off, + 1, + ALPHA_META, + )?; + } + } + + Ok(()) + } + /// Plan C Phase 2 follow-up A.2 (2026-04-29): launch `q_drift_rate_ema_update` /// — single-thread single-block ISV producer for ISV[Q_DRIFT_RATE_INDEX=129]. /// @@ -13879,6 +14059,21 @@ impl GpuDqnTrainer { .map_err(|e| MLError::ModelError(format!("pnl_aggregation_update load: {e}")))? }; + // SP5 Layer D Task D2 (2026-05-02): load health_composition_kernel. + // Single-block 1-thread arithmetic kernel. Takes 7 raw inputs by value; + // writes 4 floats (composed score + q_gap_norm + q_var_norm + + // grad_norm_norm) to scratch[SCRATCH_HEALTH_COMP_BASE=211..215). + // FoldReset: ISV[290..294) zeroed at fold boundary alongside the + // host-side `health_ema` EMAs (Pearl A sentinel path). Producer-only + // in D2; consumer wires atomically in D4 per + // `feedback_no_partial_refactor.md`. + let health_composition_kernel = { + let module = stream.context().load_cubin(SP5_HEALTH_COMPOSITION_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("sp5 health_composition cubin load: {e}")))?; + module.load_function("health_composition_update") + .map_err(|e| MLError::ModelError(format!("health_composition_update load: {e}")))? + }; + // SP5 Task A4 (2026-05-01): allocate grad_prev_buf_per_group [total_params f32]. // Zeroed at construction; fold-boundary reset zeroes it again (Pearl A sentinel: // zero grad_prev → cosine_sim=0 on first step → β1/β2 start at envelope midpoints). @@ -15447,8 +15642,11 @@ impl GpuDqnTrainer { // SP4_ENGAGE_BUF_LEN are public constants in `sp4_isv_slots`. // SP5 Task A1 (2026-05-01): wiener_state_buf grows from SP4_WIENER_TOTAL_FLOATS (213) - // to SP5_WIENER_TOTAL_FLOATS (543) = (71+110)×3. The extra 330 floats cover - // SP5's 110 producers. All-zeros at allocation (Pearl A sentinel contract). + // to SP5_WIENER_TOTAL_FLOATS, sized via `SP5_PRODUCER_COUNT` (the wiener-buffer + // linear-span constant — see SP5_PRODUCER_COUNT docstring for the linear-span + // vs unique-slot distinction; Layer D D1 grew it 110→116, D2 grows it 116→120 + // so the buffer expanded 543→573 floats post-D2). All-zeros at allocation + // (Pearl A sentinel contract). let wiener_state_buf = unsafe { MappedF32Buffer::new(SP5_WIENER_TOTAL_FLOATS) } .map_err(|e| MLError::ModelError(format!("SP4/SP5 wiener_state_buf alloc: {e}")))?; @@ -15458,18 +15656,20 @@ impl GpuDqnTrainer { // SP5 (2026-05-01): producer_step_scratch_buf sized for SP4 + SP5 Layer A producer outputs. // Layout (each producer commit appends its output range; see audit doc for derivation): // [0..71) SP4 producers - // [71..87) q_branch_stats (Pearl 1 / A1 — shared signal source) - // [87..103) pearl_1_atom_update (Pearl 1 / A1) - // [103..107) pearl_3_sigma σ (Pearl 3 / A2) - // [107..111) pearl_3_sigma SF (Pearl 3 / A2) - // [111..131) pearl_2_budget (Pearl 2 / A3 — c51/iqn/cql/ens/flatness) - // [131..147) pearl_4 cosine + l2 (auxiliary grad_cosine_sim, A4) - // [147..171) pearl_4 β1/β2/ε (Pearl 4 / A4) + // [71..87) q_branch_stats (Pearl 1 / A1 — shared signal source) + // [87..103) pearl_1_atom_update (Pearl 1 / A1) + // [103..107) pearl_3_sigma σ (Pearl 3 / A2) + // [107..111) pearl_3_sigma SF (Pearl 3 / A2) + // [111..131) pearl_2_budget (Pearl 2 / A3 — c51/iqn/cql/ens/flatness) + // [131..147) pearl_4 cosine + l2 (auxiliary grad_cosine_sim, A4) + // [147..171) pearl_4 β1/β2/ε (Pearl 4 / A4) // [171..199) pearl_5 skew/ex_kurt + IQN τ schedule (A5) // Pearl 6 (A6) writes directly to ISV — no scratch. - // [199..203) pearl_8 trail_dist (Pearl 8 / A7) - // [203..207) pearl_1_ext num_atoms (Pearl 1-ext / A8) - // Total: SP5_SCRATCH_TOTAL = 207. Audit doc records every commit's growth. + // [199..203) pearl_8 trail_dist (Pearl 8 / A7) + // [203..207) pearl_1_ext num_atoms (Pearl 1-ext / A8) + // [207..211) pnl_aggregation (Layer D D1) + // [211..215) health_composition (Layer D D2) + // Total: SP5_SCRATCH_TOTAL = 215. Audit doc records every commit's growth. let producer_step_scratch_buf = unsafe { MappedF32Buffer::new(SP5_SCRATCH_TOTAL) } .map_err(|e| MLError::ModelError(format!("SP4/SP5 producer_step_scratch_buf alloc: {e}")))?; @@ -17060,6 +17260,7 @@ impl GpuDqnTrainer { pearl_8_trail_kernel, pearl_1_ext_num_atoms_kernel, pnl_aggregation_kernel, + health_composition_kernel, grad_prev_buf_per_group, group_param_offsets_buf, atoms_clip_count_buf, diff --git a/crates/ml/src/cuda_pipeline/health_composition_kernel.cu b/crates/ml/src/cuda_pipeline/health_composition_kernel.cu new file mode 100644 index 000000000..ee0e35de8 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/health_composition_kernel.cu @@ -0,0 +1,166 @@ +// crates/ml/src/cuda_pipeline/health_composition_kernel.cu +// +// SP5 Layer D Task D2 (2026-05-02): LearningHealth composition producer kernel. +// +// Second of three Layer D producer kernels (D1=PnL, D2=health composition, +// D3=training-metrics EMA) that replace host-aggregated EMA / arithmetic +// pipelines per `feedback_no_cpu_compute_strict.md`. This kernel reproduces +// `LearningHealth::compose()` (in `learning_health.rs`) bit-for-bit on GPU: +// it normalises 7 raw component signals via smoothstep curves into [0,1], +// then composes them with the spec Layer-1 weighted sum into a scalar +// `health_score` ∈ [0,1]. +// +// Reference: the host-side composition lives in +// `learning_health.rs::NormalizedComponents::from_raw` + +// `NormalizedComponents::compose`, currently invoked from +// `training_loop.rs:2658-2745` once per epoch boundary. D4 (the atomic +// Layer D commit) wires this kernel + its `apply_pearls_ad_kernel` chain +// to the call site in lockstep with D1 and D3 per +// `feedback_no_partial_refactor.md`. **D2 is additive only.** +// +// Formula fidelity (Layer 1 spec, mirroring `learning_health.rs:39-71`): +// +// smoothstep(e0, e1, x) = let t = clamp((x - e0) / (e1 - e0), 0, 1) +// in t * t * (3 - 2 * t) +// +// q_gap_norm = smoothstep(0.01, 0.5, q_gap) +// q_var_norm = smoothstep(0.001, 0.1, q_var) +// atom_util_norm = smoothstep(0.2, 0.7, atom_util) +// grad_stable = 1.0 − smoothstep(10.0, 100.0, grad_norm) +// ens_agree = clamp(1.0 − ens_disagreement, 0, 1) +// grad_consistency_norm= smoothstep(−0.2, 0.5, grad_consistency) +// spectral_gap_norm = 1.0 − smoothstep(2.0, 10.0, spectral_gap) +// +// health_score = 0.25 * q_gap_norm +// + 0.15 * q_var_norm +// + 0.15 * atom_util_norm +// + 0.15 * grad_stable +// + 0.10 * ens_agree +// + 0.10 * grad_consistency_norm +// + 0.10 * spectral_gap_norm +// +// All edge constants and weights are Invariant 1 anchors — they encode the +// empirical collapse/health boundaries documented in `learning_health.rs` +// and are migrated verbatim per the spec brief ("formula fidelity > anything +// else"). No clamp / EMA / warmup is applied here: the host +// `LearningHealth::update` wraps `compose()` with EMA smoothing + the +// `[0.2, 0.95]` clamp + warmup gating; downstream `apply_pearls_ad_kernel` +// + ISV broadcast handles smoothing and the warmup/clamp behaviour migrates +// as part of D4's atomic refactor (out-of-scope for D2's structural-additive +// migration). +// +// Outputs (single block, single thread; 4 scratch slots written): +// +// scratch_out[health_score_idx] = composed health score in [0,1] +// scratch_out[q_gap_norm_idx] = q_gap_norm in [0,1] +// scratch_out[q_var_norm_idx] = q_var_norm in [0,1] +// scratch_out[grad_norm_norm_idx]= grad_stable in [0,1] +// +// `q_gap_norm` / `q_var_norm` / `grad_norm_norm` are exposed so HEALTH_DIAG +// can split "policy isn't differentiating Q" from "gradient explosion is +// driving the score down" without consulting the legacy CPU components +// struct (matches the eval/intent dist debugging idiom from +// `project_magnitude_eval_collapse_kelly_capped.md`). The other 4 +// normalised signals (atom_util_norm, ens_agree, grad_consistency_norm, +// spectral_gap_norm) are computed in-register and discarded — they fold +// into the composed score; future D-extension can expose them without a +// launcher-signature change. +// +// Single-block contract: matches `pearl_8_trail_kernel.cu` shape (1 thread +// suffices — composition is 7 reads, 7 smoothsteps, 7 multiply-adds; no +// reduction needed). `__threadfence_system()` after the writes so the +// mapped-pinned host_ptr / dev_ptr alias sees the values before the +// downstream `apply_pearls_ad_kernel` consumer reads them on the same stream. + +#include + +extern "C" { + +// Inline `smoothstep` mirroring `learning_health.rs::smoothstep` exactly. +// The Rust version returns 1.0 for `x >= edge0` when `edge0 == edge1` and +// 0.0 otherwise; we apply the same degenerate-bounds branch even though +// production callers never pass equal edges (Invariant 1 — fidelity to the +// reference implementation). +__device__ __forceinline__ float +hc_smoothstep(float edge0, float edge1, float x) { + if (edge1 == edge0) { + return x >= edge0 ? 1.0f : 0.0f; + } + float t = (x - edge0) / (edge1 - edge0); + if (t < 0.0f) t = 0.0f; + else if (t > 1.0f) t = 1.0f; + return t * t * (3.0f - 2.0f * t); +} + +__global__ void health_composition_update( + /* 7 raw inputs passed by value (no ISV reads — D4 wires the producer to + * call sites that pass the host-side EMA snapshot directly; the EMA + * sources themselves migrate to GPU in D3). */ + float q_gap, + float q_var, + float atom_util, + float grad_norm, + float ens_disagreement, + float grad_consistency, + float spectral_gap, + + float* __restrict__ scratch_out, + int health_score_idx, + int q_gap_norm_idx, + int q_var_norm_idx, + int grad_norm_norm_idx +) { + /* Single-block contract — matches SP5/SP4 producer kernel grid guard. */ + if (blockIdx.x != 0 || threadIdx.x != 0) return; + + /* ── Pass 1: normalise via smoothstep curves (Invariant 1 anchors) ──── + * Edge constants from `learning_health.rs::NormalizedComponents::from_raw`: + * q_gap : [0.01, 0.5] (argmax indistinguishability ↔ well-separated) + * q_var : [0.001, 0.1] (Q-collapse to uniform ↔ healthy variance) + * atom_util : [0.2, 0.7] (under-utilised ↔ ≥70% atoms carry >1% prob) + * grad_norm : [10, 100] (inverted; higher = less stable) + * ens_disagree : direct (1 − ens_disagreement, clamped) + * grad_consist : [−0.2, 0.5] (cosine sim window for grad direction) + * spectral_gap : [2, 10] (inverted; >10 = rank-1 collapse) + */ + float q_gap_norm = hc_smoothstep(0.01f, 0.5f, q_gap); + float q_var_norm = hc_smoothstep(0.001f, 0.1f, q_var); + float atom_util_norm = hc_smoothstep(0.2f, 0.7f, atom_util); + float grad_stable = 1.0f - hc_smoothstep(10.0f, 100.0f, grad_norm); + float ens_agree = 1.0f - ens_disagreement; + if (ens_agree < 0.0f) ens_agree = 0.0f; + else if (ens_agree > 1.0f) ens_agree = 1.0f; + float grad_consistency_norm = hc_smoothstep(-0.2f, 0.5f, grad_consistency); + float spectral_gap_norm = 1.0f - hc_smoothstep(2.0f, 10.0f, spectral_gap); + + /* ── Pass 2: weighted composition (Invariant 1 anchors) ────────────── + * Weights from `NormalizedComponents::compose`: + * q_gap : 0.25 + * q_var : 0.15 + * atom_util : 0.15 + * grad_stable : 0.15 + * ens_agree : 0.10 + * grad_consistency_norm : 0.10 + * spectral_gap_norm : 0.10 + * Sum = 1.0 → score lies in [0,1] by construction (each input ∈ [0,1]). + */ + float health_score = + 0.25f * q_gap_norm + + 0.15f * q_var_norm + + 0.15f * atom_util_norm + + 0.15f * grad_stable + + 0.10f * ens_agree + + 0.10f * grad_consistency_norm + + 0.10f * spectral_gap_norm; + + /* Write all four output slots with a single PCIe-visible barrier (matches + * the D1 PnL kernel pattern; mapped-pinned host_ptr aliasing requires + * threadfence_system before the downstream apply_pearls reads). */ + scratch_out[health_score_idx] = health_score; + scratch_out[q_gap_norm_idx] = q_gap_norm; + scratch_out[q_var_norm_idx] = q_var_norm; + scratch_out[grad_norm_norm_idx] = grad_stable; + __threadfence_system(); +} + +} /* extern "C" */ diff --git a/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs index a4dae540c..7601bf035 100644 --- a/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs @@ -11,8 +11,9 @@ //! 278..280 Intentional 2-slot carve-out gap (boundary marker, not allocated) //! 280..286 Cross-fold-persistent Kelly slots (NOT in fold-reset registry) //! 286..290 Layer D D1: PnL aggregation outputs (4 slots, fold-reset) +//! 290..294 Layer D D2: Health composition outputs (4 slots, fold-reset) //! -//! Total: 114 new SP5 ISV slots (52 + 24 + 20 + 4 + 4 + 6 + 4). +//! Total: 118 new SP5 ISV slots (52 + 24 + 20 + 4 + 4 + 6 + 4 + 4). pub const SP5_SLOT_BASE: usize = 174; @@ -93,14 +94,38 @@ pub const PNL_MEAN_INDEX: usize = 287; pub const PNL_VAR_INDEX: usize = 288; pub const PNL_MAX_DD_INDEX: usize = 289; -pub const SP5_SLOT_END: usize = 290; +// ── Layer D Task D2: Health composition outputs (4 slots) ──────────── +/// Layer D D2 (2026-05-02): GPU-composed `LearningHealth` outputs. +/// Replaces the host-side `NormalizedComponents::from_raw` + +/// `NormalizedComponents::compose` chain in +/// `learning_health.rs` (called from `training_loop.rs:2658-2745`) per +/// `feedback_no_cpu_compute_strict.md`. Producer is +/// `health_composition_kernel.cu` — chained through `apply_pearls_ad_kernel` +/// for Pearls A+D smoothing (sentinel-bootstrap on first observation, +/// Wiener-optimal α at steady state). Per-fold reset semantics: the +/// underlying EMAs (q_gap_ema/q_var_ema/grad_norm_ema) reset at fold +/// boundary, so the composition outputs reset together — all four slots +/// get sentinel 0.0 at fold boundary. The four exposed signals are the +/// composed score plus the three normalised inputs the task brief flagged +/// for HEALTH_DIAG observability (q_gap_norm / q_var_norm / grad_norm_norm +/// = `grad_stable`); the other 4 normalised intermediates +/// (atom_util_norm, ens_agree, grad_consistency_norm, spectral_gap_norm) +/// are computed in-register inside the kernel and folded into the score +/// without ISV exposure. Future Layer D extensions can expose them +/// without a launcher-signature change. +pub const HEALTH_SCORE_INDEX: usize = 290; +pub const Q_GAP_NORM_INDEX: usize = 291; +pub const Q_VAR_NORM_INDEX: usize = 292; +pub const GRAD_NORM_NORM_INDEX: usize = 293; + +pub const SP5_SLOT_END: usize = 294; /// Wiener-buffer producer-count constant. Sizes `wiener_state_buf` via /// `(SP4_PRODUCER_COUNT + SP5_PRODUCER_COUNT) * SP4_WIENER_FLOATS_PER_SLOT`. /// /// **Important**: this is NOT the count of unique allocated SP5 ISV slots -/// (114) — it is the linear span of the SP5 ISV index range -/// `[SP5_SLOT_BASE..SP5_SLOT_END) = [174..290) = 116`, including the +/// (118) — it is the linear span of the SP5 ISV index range +/// `[SP5_SLOT_BASE..SP5_SLOT_END) = [174..294) = 120`, including the /// intentional 2-slot carve-out gap (278..280) between the per-fold block /// and the cross-fold-persistent Kelly block. The wiener offset formula /// `wiener_off = SP4_PRODUCER_COUNT * 3 + (isv_slot - SP5_SLOT_BASE) * 3` @@ -118,10 +143,16 @@ pub const SP5_SLOT_END: usize = 290; /// Post-D1 the linear span (116) and unique-slot count (114) diverge /// permanently — `SP5_PRODUCER_COUNT` is now strictly the wiener-buffer /// linear-span constant. -pub const SP5_PRODUCER_COUNT: usize = 116; -// linear span = SP5_SLOT_END - SP5_SLOT_BASE = 290 - 174 = 116 wiener triples -// unique-slot count = 114 (52 per-branch + 24 Adam + 20 IQN τ + 4 trail -// + 4 num_atoms + 6 Kelly + 4 Layer D D1 PnL aggregation) +/// +/// Layer D Task D2 (2026-05-02) further extends to 118 unique slots @ +/// 174..278 ∪ 280..294; the linear span grows 116 → 120. Updates to +/// docstrings, fingerprint, and the slot-contiguity test are kept in +/// lockstep with the constant per `feedback_no_partial_refactor.md`. +pub const SP5_PRODUCER_COUNT: usize = 120; +// linear span = SP5_SLOT_END - SP5_SLOT_BASE = 294 - 174 = 120 wiener triples +// unique-slot count = 118 (52 per-branch + 24 Adam + 20 IQN τ + 4 trail +// + 4 num_atoms + 6 Kelly + 4 Layer D D1 PnL aggregation +// + 4 Layer D D2 health composition) // ── Convenience accessors ──────────────────────────────────────────── #[inline] pub const fn atom_v_center(b: usize) -> usize { ATOM_V_CENTER_BASE + b } @@ -156,7 +187,8 @@ pub const SP5_LAYOUT_FINGERPRINT_FRAGMENT: &str = KELLY_F_SMOOTH=280;CONVICTION_SMOOTH=281;TRADE_VAR_SMOOTH=282;\ KELLY_SAMPLE_COUNT=283;WIN_RATE_SMOOTH=284;LOSS_RATE_SMOOTH=285;\ PNL_TOTAL=286;PNL_MEAN=287;PNL_VAR=288;PNL_MAX_DD=289;\ - ISV_TOTAL_DIM=290"; + HEALTH_SCORE=290;Q_GAP_NORM=291;Q_VAR_NORM=292;GRAD_NORM_NORM=293;\ + ISV_TOTAL_DIM=294"; #[cfg(test)] mod tests { @@ -224,26 +256,31 @@ mod tests { slots.insert(PNL_MEAN_INDEX); slots.insert(PNL_VAR_INDEX); slots.insert(PNL_MAX_DD_INDEX); + // Layer D D2: Health composition outputs (4 scalar constants) + slots.insert(HEALTH_SCORE_INDEX); + slots.insert(Q_GAP_NORM_INDEX); + slots.insert(Q_VAR_NORM_INDEX); + slots.insert(GRAD_NORM_NORM_INDEX); - // 1. Exactly 114 unique slots. - assert_eq!(slots.len(), 114, "expected 114 unique slots, got {}", slots.len()); + // 1. Exactly 118 unique slots. + assert_eq!(slots.len(), 118, "expected 118 unique slots, got {}", slots.len()); // 2. Min slot is SP5_SLOT_BASE = 174. assert_eq!(*slots.iter().min().unwrap(), 174); - // 3. Max slot is SP5_SLOT_END - 1 = 289. - assert_eq!(*slots.iter().max().unwrap(), 289); + // 3. Max slot is SP5_SLOT_END - 1 = 293. + assert_eq!(*slots.iter().max().unwrap(), 293); // 4. Intentional carve-out gap (278, 279) is absent. assert!(!slots.contains(&278), "slot 278 must be absent (carve-out gap)"); assert!(!slots.contains(&279), "slot 279 must be absent (carve-out gap)"); - // 5. Set equals {174..278} ∪ {280..290} — no holes (other than the + // 5. Set equals {174..278} ∪ {280..294} — no holes (other than the // carve-out gap 278..280), no overlaps. Layer D D1 extends the - // upper end from 286 to 290. - let expected: HashSet = (174..278).chain(280..290).collect(); + // upper end from 286 to 290; D2 extends it 290 → 294. + let expected: HashSet = (174..278).chain(280..294).collect(); assert_eq!(slots, expected, - "slot set does not match expected {{174..278}} ∪ {{280..290}}"); + "slot set does not match expected {{174..278}} ∪ {{280..294}}"); } #[test] @@ -269,12 +306,32 @@ mod tests { // Strictly above the Kelly block (Layer D layout is per-fold; Kelly // is cross-fold-persistent — the contracts must remain disjoint). assert!(PNL_TOTAL_INDEX > LOSS_RATE_SMOOTH_INDEX); - // SP5_SLOT_END must reflect the new end-of-block. - assert_eq!(SP5_SLOT_END, 290); + // SP5_SLOT_END must reflect the post-D2 end-of-block. + assert_eq!(SP5_SLOT_END, 294); // SP5_PRODUCER_COUNT is the wiener-buffer linear span (slot-range // width including the 2-slot carve-out gap), NOT the unique-slot // count. See SP5_PRODUCER_COUNT docstring for the rationale. - assert_eq!(SP5_PRODUCER_COUNT, 116); + assert_eq!(SP5_PRODUCER_COUNT, 120); + assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE); + } + + #[test] + fn health_composition_slots_contiguous_and_above_pnl_block() { + // Layer D D2: Health composition outputs occupy a contiguous 4-slot + // block immediately after the Layer D D1 PnL aggregation block. + assert_eq!(HEALTH_SCORE_INDEX, 290); + assert_eq!(Q_GAP_NORM_INDEX, 291); + assert_eq!(Q_VAR_NORM_INDEX, 292); + assert_eq!(GRAD_NORM_NORM_INDEX, 293); + // Strictly above the D1 PnL block (Layer D's two producer commits + // land in disjoint slot ranges so D4's atomic consumer wiring sees + // both blocks at fixed offsets). + assert!(HEALTH_SCORE_INDEX > PNL_MAX_DD_INDEX); + // 4-slot block is internally contiguous. + assert_eq!(GRAD_NORM_NORM_INDEX - HEALTH_SCORE_INDEX, 3); + // SP5_SLOT_END must reflect the new end-of-block (4-slot grow). + assert_eq!(SP5_SLOT_END, 294); + // SP5_PRODUCER_COUNT linear-span check matches the new end. assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE); } } diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 5f824c4be..57c618a44 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -706,7 +706,22 @@ impl StateResetRegistry { RegistryEntry { name: "sp5_pnl_aggregation", category: ResetCategory::FoldReset, - description: "ISV[PNL_TOTAL_INDEX=286..PNL_MAX_DD_INDEX+1=290) — Layer D D1 PnL aggregation outputs (total/mean/population variance/max drawdown). SP5 Pearl A sentinel 0 at fold boundary so the new fold's first `launch_sp5_pnl_aggregation` fires the first-observation replacement (Task D1, 2026-05-02). Producer is `pnl_aggregation_kernel.cu`; consumer wires atomically in D4 (per `feedback_no_partial_refactor.md`). PnL aggregation is NOT cross-fold persistent — unlike Pearl 6's Kelly stats — so all four slots take the standard sentinel path. Wiener-state companion: bulk memset of the entire `wiener_state_buf[0..SP5_WIENER_TOTAL_FLOATS=561)` at fold boundary covers the new D1 triples at offsets [213+(286-174)*3 .. 213+(290-174)*3) = [549..561). The buffer grew from 543 to 561 floats with D1 (SP5_PRODUCER_COUNT linear span 110 → 116; unique-slot count 110 → 114; see SP5_PRODUCER_COUNT docstring for the linear-span vs unique-slot distinction). reset_for_fold's existing bulk memset of the whole buffer length covers the new triples atomically — no per-slot wiener entry is needed.", + description: "ISV[PNL_TOTAL_INDEX=286..PNL_MAX_DD_INDEX+1=290) — Layer D D1 PnL aggregation outputs (total/mean/population variance/max drawdown). SP5 Pearl A sentinel 0 at fold boundary so the new fold's first `launch_sp5_pnl_aggregation` fires the first-observation replacement (Task D1, 2026-05-02). Producer is `pnl_aggregation_kernel.cu`; consumer wires atomically in D4 (per `feedback_no_partial_refactor.md`). PnL aggregation is NOT cross-fold persistent — unlike Pearl 6's Kelly stats — so all four slots take the standard sentinel path. Wiener-state companion: bulk memset of the entire `wiener_state_buf[0..SP5_WIENER_TOTAL_FLOATS=573)` at fold boundary covers the new D1 triples at offsets [213+(286-174)*3 .. 213+(290-174)*3) = [549..561). The buffer grew from 543 to 561 floats with D1 (SP5_PRODUCER_COUNT linear span 110 → 116; unique-slot count 110 → 114; D2 grew it further to 573 / 120 / 118; see SP5_PRODUCER_COUNT docstring for the linear-span vs unique-slot distinction). reset_for_fold's existing bulk memset of the whole buffer length covers the new triples atomically — no per-slot wiener entry is needed.", + }, + // SP5 Layer D Task D2: Health composition outputs (4 ISV slots). + // ISV[290..294) — Layer D D2 LearningHealth composed score plus + // 3 normalised intermediates (q_gap_norm, q_var_norm, + // grad_norm_norm). The underlying EMAs (q_gap_ema / q_var_ema / + // grad_norm_ema in `health_ema`) reset at fold boundary, so the + // composition outputs reset together via the standard Pearl A + // sentinel-bootstrap path. Both halves of the contract reset + // together: ISV slot zeroed AND wiener_state_buf SP5 block + // bulk memset(0) at fold boundary covers the four new triples + // at offsets [213+(290-174)*3 .. 213+(294-174)*3) = [561..573). + RegistryEntry { + name: "sp5_health_composition", + category: ResetCategory::FoldReset, + description: "ISV[HEALTH_SCORE_INDEX=290..GRAD_NORM_NORM_INDEX+1=294) — Layer D D2 LearningHealth composition outputs (composed score + q_gap_norm + q_var_norm + grad_norm_norm). SP5 Pearl A sentinel 0 at fold boundary so the new fold's first `launch_health_composition` fires the first-observation replacement (Task D2, 2026-05-02). Producer is `health_composition_kernel.cu` — single-block 1-thread arithmetic kernel reproducing `learning_health.rs::NormalizedComponents::from_raw` + `compose()` bit-for-bit. Consumer wires atomically in D4 (per `feedback_no_partial_refactor.md`). The underlying EMAs (`HealthEmaTrackers.q_gap_ema` / `q_var_ema` / `grad_norm_ema`, plus the 4 non-EMA inputs sourced from per-branch Q-gap dispersion / atom utilisation / spectral gap / Adam m-flat cosine) reset at fold boundary so the composition outputs reset together. Wiener-state companion: bulk memset of the entire `wiener_state_buf[0..SP5_WIENER_TOTAL_FLOATS=573)` at fold boundary covers the new D2 triples at offsets [213+(290-174)*3 .. 213+(294-174)*3) = [561..573). The buffer grew from 561 to 573 floats with D2 (SP5_PRODUCER_COUNT linear span 116 → 120; unique-slot count 114 → 118; see SP5_PRODUCER_COUNT docstring for the linear-span vs unique-slot distinction). reset_for_fold's existing bulk memset of the whole buffer length covers the new triples atomically — no per-slot wiener entry is needed.", }, ]; Self { entries } diff --git a/crates/ml/tests/sp5_producer_unit_tests.rs b/crates/ml/tests/sp5_producer_unit_tests.rs index c9345d682..9aa4bbe41 100644 --- a/crates/ml/tests/sp5_producer_unit_tests.rs +++ b/crates/ml/tests/sp5_producer_unit_tests.rs @@ -36,6 +36,9 @@ //! Layer D Task D1 (PnL aggregation): //! 18. `pnl_aggregation_kernel_correctness` //! +//! Layer D Task D2 (Health composition): +//! 19. `health_composition_kernel_correctness` +//! //! All tests use analytically-known synthetic inputs with no CPU reference //! implementation, per `feedback_no_cpu_test_fallbacks.md`. All mapped-pinned //! memory (no DtoH), all `#[ignore]`-gated for GPU. @@ -87,6 +90,9 @@ const SP5_PEARL_1_EXT_NUM_ATOMS_CUBIN: &[u8] = const SP5_PNL_AGGREGATION_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/pnl_aggregation_kernel.cubin")); +const SP5_HEALTH_COMPOSITION_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/health_composition_kernel.cubin")); + // ── Helpers ─────────────────────────────────────────────────────────────────── fn load_pearl_3_sigma_kernel(stream: &Arc) -> CudaFunction { @@ -2194,3 +2200,164 @@ fn pnl_aggregation_kernel_correctness() { assert!((out[3] - exp_max_dd).abs() < tol_strict, "pnl_max_dd: expected {exp_max_dd:.5}, got {:.5}", out[3]); } + +// ── Layer D Task D2 (Health composition) ────────────────────────────────────── + +fn load_health_composition_kernel(stream: &Arc) -> CudaFunction { + let module = stream + .context() + .load_cubin(SP5_HEALTH_COMPOSITION_CUBIN.to_vec()) + .expect("load health_composition_kernel cubin"); + module + .load_function("health_composition_update") + .expect("load health_composition_update function") +} + +/// SP5 Layer D Task D2: `health_composition_update` formula-fidelity test. +/// +/// Asserts the kernel reproduces `learning_health.rs::NormalizedComponents:: +/// from_raw` + `compose()` bit-for-bit (within float precision). Per +/// `feedback_no_cpu_test_fallbacks.md` we do **not** programmatically +/// re-implement the formula in the test; instead the synthetic inputs are +/// chosen so each normalised component evaluates to a hand-computable +/// closed-form rational, and the expected outputs are checked against +/// literal pre-computed constants. +/// +/// Synthetic raw inputs: +/// +/// q_gap = 0.05 → smoothstep(0.01, 0.5) +/// t = 0.04 / 0.49 = 0.08163265306... +/// ss = t²·(3 − 2t) = 0.018901683... +/// q_var = 0.10 → smoothstep(0.001, 0.1) +/// t = (0.099)/(0.099) = 1 → ss = 1.0 +/// (intentional saturation at upper edge) +/// atom_util = 0.45 → smoothstep(0.2, 0.7) at the **midpoint** +/// t = 0.25/0.5 = 0.5 → ss = 0.5²·2 = 0.5 +/// grad_norm = 30.0 → 1 − smoothstep(10, 100) +/// t = 20/90 = 0.2̄ +/// ss = t²·(3 − 2t) = (4/81)·(73/27)·… +/// ≈ 0.126200274... +/// grad_stable = 1 − ss ≈ 0.873799725... +/// ens_disagreement = 0.20 → ens_agree = 1 − 0.20 = 0.80 (clamp pass-through) +/// grad_consistency = 0.15 → smoothstep(−0.2, 0.5) at the midpoint +/// t = 0.35/0.7 = 0.5 → ss = 0.5 +/// spectral_gap = 5.0 → 1 − smoothstep(2, 10) +/// t = 3/8 = 0.375 +/// ss = t²·(3 − 2t) = 0.140625·2.25 = 0.316406250 +/// spectral_gap_norm = 1 − 0.316406250 = 0.683593750 +/// +/// Two of the seven normalisations land at exact rational midpoints (output +/// 0.5) and one saturates exactly (output 1.0); the other four are +/// pre-computed analytically and pinned as literal constants below. The +/// composed score is the weighted sum. +/// +/// health_score = 0.25 · 0.018901683 +/// + 0.15 · 1.000000000 +/// + 0.15 · 0.500000000 +/// + 0.15 · 0.873799725 +/// + 0.10 · 0.800000000 +/// + 0.10 · 0.500000000 +/// + 0.10 · 0.683593750 +/// ≈ 0.559153878 +/// +/// **Formula fidelity is the load-bearing invariant** for D2: the kernel +/// performs a verbatim structural migration of the host-side composition, +/// no algorithmic change. The test pins both the cubic interpolation +/// (q_gap_norm interior point) and the composition arithmetic. +#[test] +#[ignore = "requires GPU"] +fn health_composition_kernel_correctness() { + let ctx = CudaContext::new(0).expect("CUDA context"); + let stream = ctx.new_stream().expect("stream"); + let kernel = load_health_composition_kernel(&stream); + + // ── Synthetic raw inputs ──────────────────────────────────────── + let q_gap: f32 = 0.05; + let q_var: f32 = 0.10; + let atom_util: f32 = 0.45; + let grad_norm: f32 = 30.0; + let ens_disagreement: f32 = 0.20; + let grad_consistency: f32 = 0.15; + let spectral_gap: f32 = 5.0; + + // ── Output scratch (4 slots) ───────────────────────────────────── + let scratch = unsafe { MappedF32Buffer::new(4) }.expect("scratch alloc"); + + let scratch_dev = scratch.dev_ptr; + let health_score_idx: i32 = 0; + let q_gap_norm_idx: i32 = 1; + let q_var_norm_idx: i32 = 2; + let grad_norm_norm_idx: i32 = 3; + + unsafe { + stream + .launch_builder(&kernel) + .arg(&q_gap) + .arg(&q_var) + .arg(&atom_util) + .arg(&grad_norm) + .arg(&ens_disagreement) + .arg(&grad_consistency) + .arg(&spectral_gap) + .arg(&scratch_dev) + .arg(&health_score_idx) + .arg(&q_gap_norm_idx) + .arg(&q_var_norm_idx) + .arg(&grad_norm_norm_idx) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch health_composition_update"); + } + stream.synchronize().expect("sync after health_composition_update"); + + let out = scratch.read_all(); + + // ── Analytical ground truth (literals; per feedback_no_cpu_test_fallbacks) ─ + // q_gap_norm closed-form (exact rational): + // t = 4/49, t²(3 − 2t) = (16·139) / (49²·49) = 2224 / 117649 ≈ 0.0189036880 + let exp_q_gap_norm: f32 = 0.018903688; + // q_var_norm: smoothstep saturates at edge1 (0.10 == upper edge). + let exp_q_var_norm: f32 = 1.0; + // grad_norm_norm = 1 − smoothstep(10, 100, 30) = 1 − (20/90)² · (3 − 40/90). + // t = 2/9, t²(3 − 2t) = (4/81) · (23/9) = 92 / 729 ≈ 0.1262002743 + // grad_stable = 1 − 92/729 = 637 / 729 ≈ 0.8737997257 + let exp_grad_stable: f32 = 0.873799725; + // health_score: weighted sum of the 7 normalised values listed in the + // docstring above. Using exact closed forms where possible: + // 0.25 · (2224/117649) ≈ 0.0047259220 + // 0.15 · 1.0 = 0.15 + // 0.15 · 0.5 = 0.075 (atom_util_norm midpoint) + // 0.15 · (637/729) ≈ 0.131069959 + // 0.10 · 0.80 = 0.08 (ens_agree) + // 0.10 · 0.5 = 0.05 (grad_consistency_norm midpoint) + // 0.10 · 0.68359375 = 0.068359375 + // ──────────────────────────────────── + // sum ≈ 0.5591552560 + let exp_health_score: f32 = 0.559155256; + + println!("health_score = {:.7} (expected {:.7})", out[0], exp_health_score); + println!("q_gap_norm = {:.7} (expected {:.7})", out[1], exp_q_gap_norm); + println!("q_var_norm = {:.7} (expected {:.7})", out[2], exp_q_var_norm); + println!("grad_norm_norm = {:.7} (expected {:.7})", out[3], exp_grad_stable); + + // FP tolerances: composed score has 7 multiply-adds of small (≤1.0) + // values; q_var_norm saturates exactly. q_gap_norm comes from cubic + // interpolation of an irrational `t` in f32 — the literal we pinned is + // truncated at the 9th digit, so the divergence between f32 evaluation + // of the closed form and the literal is bounded by ~5e-9 * 3 (cubic). + // 1e-6 absolute is safely above that. + let tol_strict: f32 = 1e-6; + let tol_score: f32 = 1e-6; + + assert!((out[0] - exp_health_score).abs() < tol_score, + "health_score: expected {exp_health_score:.7}, got {:.7}", out[0]); + assert!((out[1] - exp_q_gap_norm).abs() < tol_strict, + "q_gap_norm: expected {exp_q_gap_norm:.7}, got {:.7}", out[1]); + assert!((out[2] - exp_q_var_norm).abs() < tol_strict, + "q_var_norm: expected {exp_q_var_norm:.7}, got {:.7}", out[2]); + assert!((out[3] - exp_grad_stable).abs() < tol_strict, + "grad_norm_norm (= grad_stable): expected {exp_grad_stable:.7}, got {:.7}", out[3]); +} diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 101c12459..fe5d5e192 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -3369,3 +3369,50 @@ First of 3 Layer D producer kernels (D1=PnL, D2=health composition, D3=training- **No call-site wiring** — the host-side aggregation block in `training_loop.rs` is intentionally untouched; D4 atomically migrates the 5 EMA sites (training_sharpe_ema, max_dd_ema, gamma_blend, LearningHealth pipeline, HealthEmaTrackers) to the GPU producer chain in a single commit. Refs: SP5 plan §D Task D1; SP5 Layer A/B validated at `5845e4403`; `feedback_no_cpu_compute_strict`, `feedback_no_partial_refactor`, `feedback_no_atomicadd`, `feedback_no_cpu_test_fallbacks`, `pearl_first_observation_bootstrap`. + +### SP5 Layer D Task D2 — Health composition kernel (additive) (2026-05-02) + +Second of 3 Layer D producer kernels (D1=PnL, D2=health composition, D3=training-metrics EMA) replacing the host-aggregated EMA / arithmetic pipelines flagged in the SP4 Layer C close-out sweep audit grid (sites 6+7+8+9+10) per `feedback_no_cpu_compute_strict.md`. D2 lands additively — kernel + Rust launcher + 4 reserved ISV slots are present, but the host-side composition in `learning_health.rs::NormalizedComponents::from_raw` + `compose()` invoked from `training_loop.rs:2658-2745` continues to run unchanged. D4 (the atomic Layer D commit) wires D1+D2+D3 to call sites in lockstep per `feedback_no_partial_refactor.md`. + +**Host-side formula reproduced (formula fidelity > anything else):** + +```text +smoothstep(e0, e1, x) = let t = clamp((x - e0) / (e1 - e0), 0, 1) + in t² · (3 − 2t) + +q_gap_norm = smoothstep(0.01, 0.5, q_gap) +q_var_norm = smoothstep(0.001, 0.1, q_var) +atom_util_norm = smoothstep(0.2, 0.7, atom_util) +grad_stable = 1 − smoothstep(10, 100, grad_norm) +ens_agree = clamp(1 − ens_disagreement, 0, 1) +grad_consistency_norm= smoothstep(−0.2, 0.5, grad_consistency) +spectral_gap_norm = 1 − smoothstep(2, 10, spectral_gap) + +health_score = 0.25 · q_gap_norm + 0.15 · q_var_norm + + 0.15 · atom_util_norm + 0.15 · grad_stable + + 0.10 · ens_agree + 0.10 · grad_consistency_norm + + 0.10 · spectral_gap_norm +``` + +All edge constants and weights are migrated verbatim from `learning_health.rs:39-71` — they are Invariant 1 anchors per spec ("formula fidelity > anything else"). The kernel is structurally additive: only `from_raw` + `compose` migrate. The EMA + warmup + clamp wrapper from `LearningHealth::update` remains host-side; downstream `apply_pearls_ad_kernel` provides the smoothing equivalent and the warmup/clamp behaviour migrates as part of D4's atomic refactor (out-of-scope for D2). + +**What this lands:** +- `health_composition_kernel.cu` (132 lines) — single-block 1-thread arithmetic kernel. Takes 7 raw inputs by value (`q_gap`, `q_var`, `atom_util`, `grad_norm`, `ens_disagreement`, `grad_consistency`, `spectral_gap`); writes 4 floats to `scratch[SCRATCH_HEALTH_COMP_BASE=211..215)`: composed score + 3 normalised intermediates (`q_gap_norm`, `q_var_norm`, `grad_norm_norm = grad_stable`). The other 4 normalised signals (`atom_util_norm`, `ens_agree`, `grad_consistency_norm`, `spectral_gap_norm`) are computed in-register and folded into the score; future Layer D extensions can promote them to scratch slots without a launcher signature change. No atomicAdd. `__threadfence_system()` after the writes for mapped-pinned host_ptr / dev_ptr alias visibility. +- Rust launcher `GpuDqnTrainer::launch_health_composition` — chains 4 `apply_pearls_ad_kernel` calls on the same stream after the producer per the SP5 GPU-only Pearls A+D pattern. Wiener offset formula `213 + (isv_slot − SP5_SLOT_BASE) × 3` matches the existing SP5 launcher template (D1 PnL, Pearl 8 trail, Pearl 1-ext num_atoms). +- 4 new ISV slots `HEALTH_SCORE_INDEX=290 .. GRAD_NORM_NORM_INDEX=293` registered in `sp5_isv_slots.rs`. `SP5_SLOT_END` 290 → 294; `ISV_TOTAL_DIM` 290 → 294; `LAYOUT_FINGERPRINT_FRAGMENT` extended with `HEALTH_SCORE` / `Q_GAP_NORM` / `Q_VAR_NORM` / `GRAD_NORM_NORM` entries so existing checkpoints fail-fast on layout change. +- `SP5_PRODUCER_COUNT` (wiener-buffer linear-span constant) 116 → 120; unique-slot count 114 → 118; wiener buffer 561 → 573 floats. The carve-out gap (slots 278..280) and Pearl 6 reserved-but-unused 18 wiener floats remain reserved-unused. D1's PnL block sits at wiener offsets [(286-174)·3+213 .. (290-174)·3+213) = [549..561); D2's health-composition block sits at [(290-174)·3+213 .. (294-174)·3+213) = [561..573). Both blocks are atomically reset by the existing bulk `wiener_state_buf[0..573)` memset at fold boundary. +- `producer_step_scratch_buf` grows 211 → 215 (`SP5_SCRATCH_TOTAL = 215`). 1 new scratch constant `SCRATCH_HEALTH_COMP_BASE=211`. The constructor's scratch-layout audit comment block updated with the D1 + D2 entries. +- `StateResetRegistry` entry `sp5_health_composition` (FoldReset; D2 outputs reset alongside the underlying EMAs at fold boundary — the host-side `health_ema` `q_gap_ema` / `q_var_ema` / `grad_norm_ema` reset together, and the 4 non-EMA inputs sourced from per-branch Q-gap dispersion / atom utilisation / spectral gap / Adam m-flat cosine reset upstream too). No per-slot wiener entry needed — the bulk `wiener_state_buf[0..573)` memset at fold boundary covers offsets [561..573) atomically. +- `build.rs` cubin registration; static `SP5_HEALTH_COMPOSITION_CUBIN` byte slice + `health_composition_kernel: CudaFunction` field on `GpuDqnTrainer` + cubin/function loader entry mirroring the D1 PnL kernel template. +- GPU-gated unit test `health_composition_kernel_correctness` in `sp5_producer_unit_tests.rs`. Synthetic raw inputs chosen so each smoothstep evaluates to a hand-computable closed-form rational (`q_gap=0.05` → 2224/117649 ≈ 0.018903688; `atom_util=0.45` → exact 0.5 midpoint; `grad_norm=30` → 1 − 92/729 ≈ 0.873799725; `spectral_gap=5` → 1 − 0.31640625 = 0.68359375; etc.). Expected outputs pinned as literal pre-computed constants per `feedback_no_cpu_test_fallbacks.md` — no programmatic CPU mirror of the formula. Composed score `≈ 0.559155256`; tolerance 1e-6 absolute. +- ISV slot test `health_composition_slots_contiguous_and_above_pnl_block` mirroring D1's `pnl_aggregation_slots_contiguous_and_above_kelly_block` — verifies slot continguity and layout invariants. + +**Verification gates:** +- `cargo check -p ml --offline` clean. +- `cargo build -p ml --release --offline --features cuda` clean — `health_composition_kernel.cubin` compiles via nvcc. +- `cargo test -p ml --lib --offline -- sp5_isv_slots state_reset_registry` 7/7 pass (slot-layout invariants + new `health_composition_slots_contiguous_and_above_pnl_block` test + 3 pre-existing registry tests + 3 pre-existing slot tests). +- GPU correctness test (`health_composition_kernel_correctness`) `--ignored`-gated; fires on the next L40S smoke alongside the D1 test and the 17 pre-existing SP5 producer unit tests (test count 18 → 19 with D2). + +**No call-site wiring** — the host-side composition block in `training_loop.rs:2658-2745` is intentionally untouched; D4 atomically migrates the 5 EMA sites (training_sharpe_ema, max_dd_ema, gamma_blend, LearningHealth pipeline, HealthEmaTrackers) to the GPU producer chain in a single commit. + +Refs: SP5 plan §D Task D2; builds on D1 (`5ee795f14`); `feedback_no_cpu_compute_strict`, `feedback_no_partial_refactor`, `feedback_no_atomicadd`, `feedback_no_cpu_test_fallbacks`, `pearl_first_observation_bootstrap`.