From 5ee795f14fb5e7eac19132110c28e7db9a394a76 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sat, 2 May 2026 15:54:00 +0200 Subject: [PATCH] =?UTF-8?q?feat(sp5):=20Layer=20D=20Task=20D1=20=E2=80=94?= =?UTF-8?q?=20PnL=20aggregation=20kernel=20(additive)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit First of 3 Layer D producer kernels. Replaces host-side trade-PnL aggregation loop in training_loop.rs with a GPU kernel chained through apply_pearls_ad_kernel for smoothing. Per feedback_no_cpu_compute_strict. Additive only — no consumer wiring, no behavior change. The kernel + launcher are loaded into GpuDqnTrainer and the slots are reserved on the ISV bus, but the host-side aggregation continues to run unchanged. D4 (atomic Layer D commit) wires this and the other two D2/D3 kernels to call sites in the same atomic refactor per feedback_no_partial_refactor. What this lands: - pnl_aggregation_kernel.cu (single-block tree-reduce, no atomicAdd) - Rust launcher launch_sp5_pnl_aggregation() - 4 new ISV slots (PNL_TOTAL_INDEX..PNL_MAX_DD_INDEX, slots 286..290) - ISV_TOTAL_DIM 286 → 290; SP5_PRODUCER_COUNT semantics formalised as wiener-buffer linear-span (110 → 116; the unique-slot count diverged from the linear span at D1 — pre-D1 they coincided by accident at the Pearl 6 carve-out's introduction) - LAYOUT_FINGERPRINT_FRAGMENT extended with PNL_* entries - StateResetRegistry sp5_pnl_aggregation entry + dispatch arm (sentinel 0.0; fold-reset; PnL is NOT cross-fold persistent unlike Pearl 6 Kelly stats, so it takes the standard sentinel-bootstrap path per pearl_first_observation_bootstrap) - build.rs cubin registration - GPU-gated unit test pnl_aggregation_kernel_correctness with analytical ground truth (no CPU reference per feedback_no_cpu_test_fallbacks) - Audit doc append documenting D1 + the SP5_PRODUCER_COUNT semantic clarification Tests: cargo check + cargo build --release --features cuda clean; ISV slot + state_reset_registry unit tests 6/6 pass (incl new pnl_aggregation_slots_contiguous_and_above_kelly_block); GPU correctness test fires on next L40S smoke alongside existing 17 SP5 producer unit tests. Refs: SP5 plan §D Task D1, validated SP5 Layer A/B at 5845e4403. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/build.rs | 11 + .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 211 +++++++++++++++++- .../cuda_pipeline/pnl_aggregation_kernel.cu | 161 +++++++++++++ crates/ml/src/cuda_pipeline/sp5_isv_slots.rs | 94 +++++++- .../src/trainers/dqn/state_reset_registry.rs | 15 ++ .../src/trainers/dqn/trainer/training_loop.rs | 17 ++ crates/ml/tests/sp5_producer_unit_tests.rs | 128 +++++++++++ docs/dqn-wire-up-audit.md | 23 ++ 8 files changed, 639 insertions(+), 21 deletions(-) create mode 100644 crates/ml/src/cuda_pipeline/pnl_aggregation_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 131c53077..a03264ee1 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -480,6 +480,17 @@ fn main() { // 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", + // 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", // 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 bab0c7974..34fe0f0e7 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -395,6 +395,19 @@ static SP5_PEARL_8_TRAIL_CUBIN: &[u8] = static SP5_PEARL_1_EXT_NUM_ATOMS_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/pearl_1_ext_num_atoms_kernel.cubin")); +/// SP5 Layer D Task D1 (2026-05-02): PnL aggregation producer kernel. +/// Single-block 256-thread shmem-reduce reads per-trade event arrays +/// (`trade_pnls`, `trade_durations`, `trade_directions`) and writes +/// 4 aggregated floats — total / mean / population variance / max +/// drawdown — to scratch_buf[SCRATCH_PNL_AGG_BASE=207..211). +/// Followed by 4 apply_pearls_ad_kernel calls → ISV[PNL_TOTAL_INDEX=286..290). +/// Replaces the host-side trade-PnL aggregation loop in +/// `compute_epoch_financials` 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_PNL_AGGREGATION_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/pnl_aggregation_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 @@ -761,7 +774,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 = 286; // SP5: 173 + 113 (110 SP5 slots @ 174..278/280..286, with 2-slot gap before cross-fold-persistent Kelly block) +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)) /// 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). @@ -969,11 +982,24 @@ pub const SP4_WIENER_TOTAL_FLOATS: usize = // ── SP5 Task A1 buffer layout constants ─────────────────────────────────────── -/// SP5 Task A1: combined wiener_state_buf total float count after growth. -/// SP4 had 71 producers (213 floats); SP5 adds 110 producers (330 floats); -/// combined = (71 + 110) × 3 = 543 floats. -/// wiener_state_buf must grow to this size at construction for SP5 slots to be -/// in-bounds when `apply_pearls_ad_kernel` writes SP5 Wiener triples. +/// 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). +/// +/// 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). pub const SP5_WIENER_TOTAL_FLOATS: usize = (SP4_PRODUCER_COUNT + crate::cuda_pipeline::sp5_isv_slots::SP5_PRODUCER_COUNT) * SP4_WIENER_FLOATS_PER_SLOT; @@ -1006,8 +1032,20 @@ pub const SP5_WIENER_TOTAL_FLOATS: usize = /// 4 floats for pearl_8_trail output (SCRATCH_PEARL_8_TRAIL=199, trail_dist × 4 directions at [199..203)) /// Task A8 adds: /// 4 floats for pearl_1_ext_num_atoms output (SCRATCH_PEARL_1_EXT_NUM_ATOMS=203, num_atoms × 4 branches at [203..207)) -/// Combined: 71 + 16 + 16 + 4 + 4 + 20 + 8 + 8 + 8 + 8 + 8 + 4 + 4 + 20 + 4 + 4 = 207 scratch slots [0..207). -pub const SP5_SCRATCH_TOTAL: usize = 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; + +/// 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: +/// [207] = pnl_total (Σ trade_pnls) +/// [208] = pnl_mean (Σ trade_pnls / num_trades) +/// [209] = pnl_var (population variance, Σ(x − μ)² / N) +/// [210] = pnl_max_dd (running max-drawdown over the cumulative PnL curve) +/// Written by `pnl_aggregation_update`; consumed by `apply_pearls_ad_kernel` → +/// ISV[PNL_TOTAL_INDEX=286..PNL_MAX_DD_INDEX+1=290). +pub const SCRATCH_PNL_AGG_BASE: usize = 207; /// 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). @@ -1535,7 +1573,8 @@ const fn layout_fingerprint_seed() -> &'static [u8] { ADAM_EPS_BASE=242;IQN_TAU_BASE=250;TRAIL_DIST_PER_DIR_BASE=270;ATOM_NUM_ATOMS_BASE=274;\ KELLY_F_SMOOTH=280;CONVICTION_SMOOTH=281;TRADE_VAR_SMOOTH=282;\ KELLY_SAMPLE_COUNT=283;WIN_RATE_SMOOTH=284;LOSS_RATE_SMOOTH=285;\ - ISV_TOTAL_DIM=286;\ + PNL_TOTAL=286;PNL_MEAN=287;PNL_VAR=288;PNL_MAX_DD=289;\ + ISV_TOTAL_DIM=290;\ 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;\ @@ -4316,6 +4355,18 @@ pub struct GpuDqnTrainer { /// ISV[ATOM_NUM_ATOMS_BASE=274..278). In fold-reset registry (FoldReset). /// Loaded from `pearl_1_ext_num_atoms_kernel.cubin`. pearl_1_ext_num_atoms_kernel: CudaFunction, + /// SP5 Layer D Task D1 (2026-05-02): PnL aggregation kernel. + /// Single-block 256-thread shmem-reduce kernel. Reads per-trade event arrays + /// (`trade_pnls`, `trade_durations`, `trade_directions`) and writes 4 aggregated + /// floats — total / mean / population variance / max drawdown — to + /// scratch_buf[SCRATCH_PNL_AGG_BASE=207..211). + /// Chained apply_pearls_ad_kernel (4 launches, ALPHA_META=1e-3) → + /// ISV[PNL_TOTAL_INDEX=286..PNL_MAX_DD_INDEX+1=290). In fold-reset registry + /// (FoldReset; Layer D D1 outputs are NOT cross-fold persistent — unlike + /// Pearl 6 Kelly stats — so they take the standard Pearl A sentinel path). + /// Loaded from `pnl_aggregation_kernel.cubin`. Producer-only in this commit; + /// D4 wires the consumer atomically. + pnl_aggregation_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 @@ -11571,6 +11622,132 @@ impl GpuDqnTrainer { Ok(()) } + /// SP5 Layer D Task D1 (2026-05-02): launch `pnl_aggregation_update` — + /// single-block 256-thread shmem-reduce producer for ISV[PNL_TOTAL_INDEX=286..PNL_MAX_DD_INDEX+1=290). + /// + /// Reads per-trade event arrays — `trade_pnls`, `trade_durations`, + /// `trade_directions` — and writes 4 aggregated floats to + /// scratch_buf[SCRATCH_PNL_AGG_BASE=207..211): + /// [207] pnl_total Σ trade_pnls + /// [208] pnl_mean Σ trade_pnls / num_trades + /// [209] pnl_var population variance Σ(x − μ)² / N + /// [210] pnl_max_dd max running drawdown over the cumulative PnL curve + /// + /// Then chains 4 `apply_pearls_ad_kernel` calls smoothing through Pearls A+D + /// into ISV[286..290) on the same stream (no host sync per + /// `pearl_cold_path_no_exception_to_gpu_drives.md` and the SP5 GPU-only + /// Pearls refactor). + /// + /// Empty-input contract: when `num_trades == 0` the kernel writes 0.0 to + /// every output slot (no NaN/garbage); the apply_pearls chain still runs + /// every step so the consumer (D4) sees a valid ISV value every step. + /// + /// Per `feedback_no_cpu_compute_strict.md` — replaces the host-side trade-PnL + /// aggregation in `compute_epoch_financials` (training_loop.rs ~4862-4916). + /// Per `feedback_no_atomicadd.md` — single-block tree-reduce; no atomicAdd. + /// Per `feedback_no_partial_refactor.md` — D1 is additive only; D4 atomically + /// migrates the host-side site to call this launcher in lockstep with D2/D3. + /// + /// Arguments: + /// - `trade_pnls`: mapped-pinned [num_trades] f32 — host-written per-trade realized PnL. + /// - `trade_durations`: mapped-pinned [num_trades] f32 — host-written per-trade + /// duration in bars (reserved; consumed by future Layer D extensions for + /// time-weighted PnL — kernel signature reflects spec contract). + /// - `trade_directions`: mapped-pinned [num_trades] i32 — host-written direction + /// id (0=Short, 1=Hold, 2=Long, 3=Flat; reserved; consumed by future Layer D + /// extensions for per-direction PnL split). + /// - `num_trades`: count of populated entries in the three arrays + /// (`<= trade_pnls.len`); the launcher does not validate this against + /// buffer length — caller's responsibility. + pub(crate) fn launch_sp5_pnl_aggregation( + &self, + trade_pnls: &super::mapped_pinned::MappedF32Buffer, + trade_durations: &super::mapped_pinned::MappedF32Buffer, + trade_directions: &super::mapped_pinned::MappedI32Buffer, + num_trades: i32, + ) -> Result<(), MLError> { + use crate::cuda_pipeline::sp5_isv_slots::{ + PNL_TOTAL_INDEX, PNL_MEAN_INDEX, PNL_VAR_INDEX, PNL_MAX_DD_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_sp5_pnl_aggregation: isv_signals_dev_ptr must be allocated by constructor"); + debug_assert!(num_trades >= 0, + "launch_sp5_pnl_aggregation: num_trades must be non-negative, got {num_trades}"); + + 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 pnls_dev = trade_pnls.dev_ptr; + let durs_dev = trade_durations.dev_ptr; + let dirs_dev = trade_directions.dev_ptr; + + let pnl_total_idx_i32: i32 = (SCRATCH_PNL_AGG_BASE + 0) as i32; + let pnl_mean_idx_i32: i32 = (SCRATCH_PNL_AGG_BASE + 1) as i32; + let pnl_var_idx_i32: i32 = (SCRATCH_PNL_AGG_BASE + 2) as i32; + let pnl_max_dd_idx_i32: i32 = (SCRATCH_PNL_AGG_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: pnl_aggregation_update — reads per-trade arrays; writes + // 4 aggregated floats to scratch[SCRATCH_PNL_AGG_BASE..+4). + // 256 threads × 1 block; shared mem = 2 × 256 × sizeof(f32) = 2048 bytes + // (per-thread sum + sum-of-squares accumulators; matches `h_s2_rms_ema_kernel`'s + // shmem layout × 2 banks). + let block_dim: u32 = 256; + let shared_mem_bytes: u32 = 2 * block_dim * std::mem::size_of::() as u32; + unsafe { + self.stream + .launch_builder(&self.pnl_aggregation_kernel) + .arg(&pnls_dev) + .arg(&durs_dev) + .arg(&dirs_dev) + .arg(&num_trades) + .arg(&scratch_dev) + .arg(&pnl_total_idx_i32) + .arg(&pnl_mean_idx_i32) + .arg(&pnl_var_idx_i32) + .arg(&pnl_max_dd_idx_i32) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (block_dim, 1, 1), + shared_mem_bytes, + }) + .map_err(|e| MLError::ModelError(format!("pnl_aggregation_update: {e}")))?; + } + + // Step 2: 4 apply_pearls calls (one per output slot). Each scratch[207+k] + // → ISV[286+k] with its own Wiener triple at offset + // (base_wiener_offset + (isv_slot − SP5_SLOT_BASE) * 3). + let isv_slots: [usize; 4] = [ + PNL_TOTAL_INDEX, + PNL_MEAN_INDEX, + PNL_VAR_INDEX, + PNL_MAX_DD_INDEX, + ]; + for k in 0..4_usize { + let isv_slot = isv_slots[k] as i32; + let scratch_idx = (SCRATCH_PNL_AGG_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]. /// @@ -13687,6 +13864,21 @@ impl GpuDqnTrainer { .map_err(|e| MLError::ModelError(format!("pearl_1_ext_num_atoms_update load: {e}")))? }; + // SP5 Layer D Task D1 (2026-05-02): load pnl_aggregation_kernel. + // Single-block 256-thread shmem-reduce kernel. Reads per-trade event + // arrays; writes 4 aggregated floats (total/mean/var/max_dd) to + // scratch[SCRATCH_PNL_AGG_BASE=207..211). FoldReset: ISV[286..290) zeroed + // at fold boundary (Pearl A sentinel; Layer D D1 outputs are NOT cross-fold + // persistent — they take the standard sentinel-bootstrap path). + // Producer-only in D1; consumer wires atomically in D4 per + // `feedback_no_partial_refactor.md`. + let pnl_aggregation_kernel = { + let module = stream.context().load_cubin(SP5_PNL_AGGREGATION_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("sp5 pnl_aggregation cubin load: {e}")))?; + module.load_function("pnl_aggregation_update") + .map_err(|e| MLError::ModelError(format!("pnl_aggregation_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). @@ -16867,6 +17059,7 @@ impl GpuDqnTrainer { pearl_6_kelly_kernel, pearl_8_trail_kernel, pearl_1_ext_num_atoms_kernel, + pnl_aggregation_kernel, grad_prev_buf_per_group, group_param_offsets_buf, atoms_clip_count_buf, diff --git a/crates/ml/src/cuda_pipeline/pnl_aggregation_kernel.cu b/crates/ml/src/cuda_pipeline/pnl_aggregation_kernel.cu new file mode 100644 index 000000000..83ada8424 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/pnl_aggregation_kernel.cu @@ -0,0 +1,161 @@ +// crates/ml/src/cuda_pipeline/pnl_aggregation_kernel.cu +// +// SP5 Layer D Task D1 (2026-05-02): PnL aggregation producer kernel. +// +// First 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 produces +// the four aggregated PnL statistics the host-side `compute_epoch_financials` +// path currently derives from a per-trade Vec loop: +// +// pnl_total Σ trade_pnls[i] +// pnl_mean (Σ trade_pnls[i]) / num_trades +// pnl_var population variance, Σ(x − μ)² / N +// pnl_max_dd running max-drawdown over the cumulative PnL curve +// (max running peak − cumulative). +// +// All four are produced from a single pass over the per-trade event arrays; +// downstream `apply_pearls_ad_kernel` smooths the four scratch slots into +// ISV[PNL_TOTAL_INDEX..PNL_MAX_DD_INDEX+1] (slots 286..290). Pearl A's +// first-observation replacement covers the cold-start (sentinel `prev=0` +// in lockstep with the wiener_state buffer per +// `pearl_first_observation_bootstrap.md`). +// +// Design choices (mirror existing SP5 producer kernels): +// - Single-block, 256-thread grid-stride reduce (matches +// `h_s2_rms_ema_kernel.cu`'s shape exactly). No atomicAdd +// (`feedback_no_atomicadd.md`). Block-internal tree-reduce in shared +// memory. +// - Two passes over the input: pass 1 computes (Σ x, Σ x²) for total / +// mean / variance; pass 2 (single-thread serial scan) computes max +// drawdown over the cumulative PnL curve. The serial scan is correct +// by construction — peak/cumulative are inherently sequential — and +// `num_trades` is bounded by the per-epoch max trade count (~10⁴ on +// the production smoke), well under the cost ceiling for one thread. +// - Inputs `trade_durations` and `trade_directions` are not read by the +// four current outputs but are part of the kernel signature so the +// contract matches the spec brief; downstream tasks (Layer D extension +// to time-weighted PnL, per-direction PnL split) can read them without +// a launcher-signature change. Unused-arg warnings suppressed via the +// `(void)` cast pattern used elsewhere in the producer kernels. +// - Empty-input guard: if `num_trades == 0`, every output slot is +// written 0.0f and the kernel returns. The host launcher always +// launches the kernel so the downstream apply_pearls chain sees a +// valid scratch slot every step (no conditional launch). +// - `__threadfence_system()` after each output write so the mapped- +// pinned host_ptr / dev_ptr alias sees the value before the +// `apply_pearls_ad_kernel` consumer reads it on the same stream. +// +// Producer-only — D1 is additive. The host-side aggregation in +// `training_loop.rs:4862-4916` continues to run unchanged. D4 (the atomic +// Layer D commit) wires D1+D2+D3 to call sites in a single migration +// per `feedback_no_partial_refactor.md`. + +extern "C" __global__ void pnl_aggregation_update( + const float* __restrict__ trade_pnls, /* [num_trades] per-trade realized PnL */ + const float* __restrict__ trade_durations, /* [num_trades] per-trade duration in bars (reserved) */ + const int* __restrict__ trade_directions, /* [num_trades] direction id 0=Short..3=Flat (reserved) */ + int num_trades, + float* __restrict__ scratch_buf, /* producer_step_scratch_buf */ + int pnl_total_idx, + int pnl_mean_idx, + int pnl_var_idx, + int pnl_max_dd_idx +) { + /* Single-block contract — match SP5/SP4 producer kernel grid guard. */ + if (blockIdx.x != 0) return; + + /* Reserved-arg silencing (Invariant 1: kernel signature reflects spec + * contract; consumers Tasks D-future read these without launcher + * signature change). */ + (void)trade_durations; + (void)trade_directions; + + extern __shared__ float smem[]; /* 2 × blockDim.x floats */ + const int tid = (int)threadIdx.x; + const int block = (int)blockDim.x; + + /* Empty-input fast path. Thread 0 zeros all four output slots so the + * downstream apply_pearls chain reads a defined value every step. */ + if (num_trades <= 0) { + if (tid == 0) { + scratch_buf[pnl_total_idx] = 0.0f; + scratch_buf[pnl_mean_idx] = 0.0f; + scratch_buf[pnl_var_idx] = 0.0f; + scratch_buf[pnl_max_dd_idx] = 0.0f; + __threadfence_system(); + } + return; + } + + /* ── Pass 1: per-thread (Σ x, Σ x²) over a strided slice ─────────── */ + float local_sum = 0.0f; + float local_sumsq = 0.0f; + for (int i = tid; i < num_trades; i += block) { + const float x = trade_pnls[i]; + local_sum += x; + local_sumsq += x * x; + } + smem[tid] = local_sum; + smem[block + tid] = local_sumsq; + __syncthreads(); + + /* In-block tree reduction (no atomicAdd). */ + for (int s = block / 2; s > 0; s >>= 1) { + if (tid < s) { + smem[tid] += smem[tid + s]; + smem[block + tid] += smem[block + tid + s]; + } + __syncthreads(); + } + + /* Thread 0 holds the reduced sums. Compute total / mean / var. + * Population variance: Var(X) = E[X²] − (E[X])²; equivalent to + * Σ(x − μ)² / N up to FP rounding (single-pass form preferred for + * one-pass producer kernels — this matches `h_s2_rms_ema_kernel.cu`'s + * sum-of-squares precedent). */ + float total = 0.0f; + float mean = 0.0f; + float var = 0.0f; + if (tid == 0) { + const float n = (float)num_trades; + const float sum = smem[0]; + const float sumsq = smem[block]; + total = sum; + mean = sum / n; + const float mean_sq = mean * mean; + const float e_xsq = sumsq / n; + var = e_xsq - mean_sq; + /* Numerical-stability floor: catastrophic cancellation in the + * one-pass form can produce a tiny negative value when the data + * is constant or near-constant. Clamp to 0 (variance is + * non-negative by definition; the C51 / IQN consumers downstream + * never read negative variance). The `0.0f` floor is an + * Invariant 1 anchor (mathematical identity, not a tuned value). */ + if (var < 0.0f) var = 0.0f; + } + + /* ── Pass 2: max drawdown via single-thread serial scan ────────────── + * Drawdown = max(running_peak − cumulative). The peak/cumulative pair + * is inherently sequential; the cost (~10⁴ trades typical) sits well + * under the producer-cadence budget. Same-thread ordering means we + * can read trade_pnls directly without a __syncthreads() barrier. */ + float max_dd = 0.0f; + if (tid == 0) { + float cumulative = 0.0f; + float peak = 0.0f; + for (int i = 0; i < num_trades; i++) { + cumulative += trade_pnls[i]; + if (cumulative > peak) peak = cumulative; + const float dd = peak - cumulative; + if (dd > max_dd) max_dd = dd; + } + + /* Write all four output slots with a single PCIe-visible barrier. */ + scratch_buf[pnl_total_idx] = total; + scratch_buf[pnl_mean_idx] = mean; + scratch_buf[pnl_var_idx] = var; + scratch_buf[pnl_max_dd_idx] = max_dd; + __threadfence_system(); + } +} diff --git a/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs index 83965d0bc..a4dae540c 100644 --- a/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs @@ -10,8 +10,9 @@ //! 274..278 Per-branch num_atoms (4 branches) //! 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) //! -//! Total: 110 new SP5 ISV slots (52 + 24 + 20 + 4 + 4 + 6). +//! Total: 114 new SP5 ISV slots (52 + 24 + 20 + 4 + 4 + 6 + 4). pub const SP5_SLOT_BASE: usize = 174; @@ -78,9 +79,49 @@ pub const KELLY_SAMPLE_COUNT_INDEX: usize = 283; pub const WIN_RATE_SMOOTH_INDEX: usize = 284; pub const LOSS_RATE_SMOOTH_INDEX: usize = 285; -pub const SP5_SLOT_END: usize = 286; -pub const SP5_PRODUCER_COUNT: usize = 110; -// 52 (per-branch) + 24 (Adam) + 20 (IQN τ) + 4 (trail) + 4 (num_atoms) + 6 (Kelly) = 110 +// ── Layer D Task D1: PnL aggregation outputs (4 slots) ─────────────── +/// Layer D D1 (2026-05-02): GPU-aggregated per-epoch trade-PnL summary. +/// Replaces the host-side aggregation loop in `compute_epoch_financials` +/// per `feedback_no_cpu_compute_strict.md`. Producer is +/// `pnl_aggregation_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: PnL is +/// not cross-fold persistent (unlike Pearl 6's Kelly stats), so all four +/// slots get sentinel 0.0 at fold boundary. +pub const PNL_TOTAL_INDEX: usize = 286; +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; + +/// 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 +/// 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` +/// indexes linearly by ISV slot, so the buffer must cover the entire +/// linear span — the 2 gap floats and the 6 Pearl 6 floats are +/// **reserved-unused** (Pearl 6 does not call `launch_apply_pearls`, +/// so its 18 wiener floats stay zero-initialised; the gap floats are +/// likewise never accessed). +/// +/// Layer D Task D1 (2026-05-02) extended SP5 from 110 unique slots @ +/// 174..278 ∪ 280..286 to 114 unique slots @ 174..278 ∪ 280..290; the +/// linear span grew from 112 to 116. Pre-D1 the constant happened to +/// equal both the unique-slot count AND the linear span (110 = 112 − +/// 2 gap, by coincidence at the time the Kelly carve-out landed). +/// 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) // ── Convenience accessors ──────────────────────────────────────────── #[inline] pub const fn atom_v_center(b: usize) -> usize { ATOM_V_CENTER_BASE + b } @@ -114,7 +155,8 @@ pub const SP5_LAYOUT_FINGERPRINT_FRAGMENT: &str = ADAM_EPS_BASE=242;IQN_TAU_BASE=250;TRAIL_DIST_PER_DIR_BASE=270;ATOM_NUM_ATOMS_BASE=274;\ KELLY_F_SMOOTH=280;CONVICTION_SMOOTH=281;TRADE_VAR_SMOOTH=282;\ KELLY_SAMPLE_COUNT=283;WIN_RATE_SMOOTH=284;LOSS_RATE_SMOOTH=285;\ - ISV_TOTAL_DIM=286"; + PNL_TOTAL=286;PNL_MEAN=287;PNL_VAR=288;PNL_MAX_DD=289;\ + ISV_TOTAL_DIM=290"; #[cfg(test)] mod tests { @@ -177,23 +219,31 @@ mod tests { slots.insert(KELLY_SAMPLE_COUNT_INDEX); slots.insert(WIN_RATE_SMOOTH_INDEX); slots.insert(LOSS_RATE_SMOOTH_INDEX); + // Layer D D1: PnL aggregation outputs (4 scalar constants) + slots.insert(PNL_TOTAL_INDEX); + slots.insert(PNL_MEAN_INDEX); + slots.insert(PNL_VAR_INDEX); + slots.insert(PNL_MAX_DD_INDEX); - // 1. Exactly 110 unique slots. - assert_eq!(slots.len(), 110, "expected 110 unique slots, got {}", slots.len()); + // 1. Exactly 114 unique slots. + assert_eq!(slots.len(), 114, "expected 114 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 = 285. - assert_eq!(*slots.iter().max().unwrap(), 285); + // 3. Max slot is SP5_SLOT_END - 1 = 289. + assert_eq!(*slots.iter().max().unwrap(), 289); // 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..286} — no holes, no overlaps. - let expected: HashSet = (174..278).chain(280..286).collect(); - assert_eq!(slots, expected, "slot set does not match expected {{174..278}} ∪ {{280..286}}"); + // 5. Set equals {174..278} ∪ {280..290} — 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(); + assert_eq!(slots, expected, + "slot set does not match expected {{174..278}} ∪ {{280..290}}"); } #[test] @@ -207,4 +257,24 @@ mod tests { assert!(atom_num_atoms(b) < KELLY_F_SMOOTH_INDEX); } } + + #[test] + fn pnl_aggregation_slots_contiguous_and_above_kelly_block() { + // Layer D D1: PnL aggregation outputs occupy a contiguous 4-slot + // block immediately after the cross-fold-persistent Kelly block. + assert_eq!(PNL_TOTAL_INDEX, 286); + assert_eq!(PNL_MEAN_INDEX, 287); + assert_eq!(PNL_VAR_INDEX, 288); + assert_eq!(PNL_MAX_DD_INDEX, 289); + // 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_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, 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 a9620f4e1..5f824c4be 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -693,6 +693,21 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "ISV[ATOM_NUM_ATOMS_BASE=274..278) — per-branch C51 num_atoms (discrete {16, 32, 64} mapped from v_half thresholds, smoothed via Pearls A+D so transition-region values may be fractional; Layer B's atoms_update consumer rounds at consume time). SP5 Pearl A sentinel 0 at fold boundary triggers first-observation replacement (Task A8).", }, + // SP5 Layer D Task D1: PnL aggregation outputs (4 ISV slots). + // ISV[286..290) — Layer D D1 PnL summary written by + // pnl_aggregation_update producer kernel. NOT cross-fold persistent + // (unlike Pearl 6's Kelly slots) — PnL aggregation is per-fold by + // construction, so the standard Pearl A sentinel-bootstrap path + // applies. Both halves of the contract reset together: ISV slot + // zeroed AND wiener_state_buf SP5 block (offsets [213..555)) bulk + // memset(0) at fold boundary covers the four new triples too, + // so the new fold's first launcher fires Pearl A's first-observation + // replacement (per `pearl_first_observation_bootstrap.md`). + 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.", + }, ]; Self { entries } } diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index f59084d25..2e968666c 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -6361,6 +6361,23 @@ impl DQNTrainer { } } } + "sp5_pnl_aggregation" => { + // Layer D Task D1: zero ISV[286..290) — total/mean/var/max_dd — + // so the new fold's first `launch_sp5_pnl_aggregation` triggers + // Pearl A's first-observation replacement (per `pearl_first_observation_bootstrap.md`). + // Wiener-state companion is reset by the existing bulk memset + // of `wiener_state_buf[0..SP5_WIENER_TOTAL_FLOATS=561)` covered + // by the `sp4_wiener_state` registry entry's dispatch arm. + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp5_isv_slots::{ + PNL_TOTAL_INDEX, PNL_MEAN_INDEX, PNL_VAR_INDEX, PNL_MAX_DD_INDEX, + }; + fused.trainer().write_isv_signal_at(PNL_TOTAL_INDEX, 0.0); + fused.trainer().write_isv_signal_at(PNL_MEAN_INDEX, 0.0); + fused.trainer().write_isv_signal_at(PNL_VAR_INDEX, 0.0); + fused.trainer().write_isv_signal_at(PNL_MAX_DD_INDEX, 0.0); + } + } _ => { return Err(crate::MLError::ModelError(format!( "StateResetRegistry reset dispatch: unknown name '{}'. \ diff --git a/crates/ml/tests/sp5_producer_unit_tests.rs b/crates/ml/tests/sp5_producer_unit_tests.rs index f0e6bdc0e..c9345d682 100644 --- a/crates/ml/tests/sp5_producer_unit_tests.rs +++ b/crates/ml/tests/sp5_producer_unit_tests.rs @@ -33,6 +33,9 @@ //! 16. `pearl_1_ext_num_atoms_threshold_cascade` //! 17. `pearl_1_ext_num_atoms_exact_threshold_falls_to_next_branch` //! +//! Layer D Task D1 (PnL aggregation): +//! 18. `pnl_aggregation_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. @@ -81,6 +84,9 @@ const SP5_PEARL_8_TRAIL_CUBIN: &[u8] = const SP5_PEARL_1_EXT_NUM_ATOMS_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/pearl_1_ext_num_atoms_kernel.cubin")); +const SP5_PNL_AGGREGATION_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/pnl_aggregation_kernel.cubin")); + // ── Helpers ─────────────────────────────────────────────────────────────────── fn load_pearl_3_sigma_kernel(stream: &Arc) -> CudaFunction { @@ -2066,3 +2072,125 @@ fn pearl_1_ext_num_atoms_exact_threshold_falls_to_next_branch() { assert!((out[3] - 16.0_f32).abs() < tol, "b=3 (v_half=5.0): expected 16 atoms (sanity check), got {:.0}", out[3]); } + +// ── Layer D Task D1 (PnL aggregation) ───────────────────────────────────────── + +fn load_pnl_aggregation_kernel(stream: &Arc) -> CudaFunction { + let module = stream + .context() + .load_cubin(SP5_PNL_AGGREGATION_CUBIN.to_vec()) + .expect("load pnl_aggregation_kernel cubin"); + module + .load_function("pnl_aggregation_update") + .expect("load pnl_aggregation_update function") +} + +/// SP5 Layer D Task D1: `pnl_aggregation_update` correctness on a synthetic +/// trade tape with analytically-known total / mean / variance / max drawdown. +/// +/// Synthetic input: trade_pnls = [+10, -5, +3, -2, +1], num_trades = 5. +/// Cumulative PnL curve: 10, 5, 8, 6, 7. +/// Running peak: 10, 10, 10, 10, 10. +/// Running drawdown: 0, 5, 2, 4, 3. +/// +/// Analytical ground truth (no CPU reference implementation per +/// `feedback_no_cpu_test_fallbacks.md`; values derived from the input +/// definition): +/// pnl_total = 10 − 5 + 3 − 2 + 1 = 7 +/// pnl_mean = 7 / 5 = 1.4 +/// pnl_var = E[X²] − E[X]² = (100 + 25 + 9 + 4 + 1)/5 − 1.4² +/// = 139/5 − 1.96 = 27.8 − 1.96 = 25.84 +/// pnl_max_dd = max running peak − cumulative = 5 (peak=10, cumulative=5 +/// after the -5 trade) +/// +/// The reserved kernel inputs `trade_durations` and `trade_directions` are +/// not read by D1's four outputs but are part of the kernel signature +/// (per the spec brief — future Layer D extensions read them without a +/// launcher signature change). We populate them with dummy values to +/// exercise the full-signature launch path. +#[test] +#[ignore = "requires GPU"] +fn pnl_aggregation_kernel_correctness() { + let ctx = CudaContext::new(0).expect("CUDA context"); + let stream = ctx.new_stream().expect("stream"); + let kernel = load_pnl_aggregation_kernel(&stream); + + // ── Synthetic input arrays ─────────────────────────────────────── + let pnls_host: [f32; 5] = [10.0, -5.0, 3.0, -2.0, 1.0]; + let durs_host: [f32; 5] = [1.0, 2.0, 3.0, 4.0, 5.0]; // dummy + let dirs_host: [i32; 5] = [2, 0, 2, 0, 2]; // Long/Short/Long/Short/Long + let num_trades: i32 = 5; + + let mut pnls_buf = unsafe { MappedF32Buffer::new(num_trades as usize) }.expect("pnls alloc"); + let mut durs_buf = unsafe { MappedF32Buffer::new(num_trades as usize) }.expect("durs alloc"); + let dirs_buf = unsafe { MappedI32Buffer::new(num_trades as usize) }.expect("dirs alloc"); + pnls_buf.host_slice_mut().copy_from_slice(&pnls_host); + durs_buf.host_slice_mut().copy_from_slice(&durs_host); + dirs_buf.write_from_slice(&dirs_host); + + // ── Output scratch (4 slots) ───────────────────────────────────── + let scratch = unsafe { MappedF32Buffer::new(4) }.expect("scratch alloc"); + + // ── Launch ─────────────────────────────────────────────────────── + let pnls_dev = pnls_buf.dev_ptr; + let durs_dev = durs_buf.dev_ptr; + let dirs_dev = dirs_buf.dev_ptr; + let scratch_dev = scratch.dev_ptr; + let total_idx: i32 = 0; + let mean_idx: i32 = 1; + let var_idx: i32 = 2; + let max_dd_idx: i32 = 3; + + let block_dim: u32 = 256; + let shared_mem_bytes: u32 = 2 * block_dim * std::mem::size_of::() as u32; + + unsafe { + stream + .launch_builder(&kernel) + .arg(&pnls_dev) + .arg(&durs_dev) + .arg(&dirs_dev) + .arg(&num_trades) + .arg(&scratch_dev) + .arg(&total_idx) + .arg(&mean_idx) + .arg(&var_idx) + .arg(&max_dd_idx) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (block_dim, 1, 1), + shared_mem_bytes, + }) + .expect("launch pnl_aggregation_update"); + } + stream.synchronize().expect("sync after pnl_aggregation_update"); + + let out = scratch.read_all(); + + // ── Analytical ground truth ────────────────────────────────────── + let exp_total: f32 = 7.0; + let exp_mean: f32 = 1.4; + let exp_var: f32 = 25.84; + let exp_max_dd: f32 = 5.0; + + println!("pnl_total = {:.4} (expected {exp_total:.4})", out[0]); + println!("pnl_mean = {:.4} (expected {exp_mean:.4})", out[1]); + println!("pnl_var = {:.4} (expected {exp_var:.4})", out[2]); + println!("pnl_max_dd = {:.4} (expected {exp_max_dd:.4})", out[3]); + + // FP tolerances: total / mean / max_dd are exact-by-construction (sum of + // small integers / integer division by 5 / max over 5 cumulative values), + // so 1e-5 is comfortable. Variance uses E[X²] − E[X]² which can lose a + // few ulps to catastrophic cancellation; 1e-3 is generous. + let tol_strict: f32 = 1e-5; + let tol_var: f32 = 1e-3; + + assert!((out[0] - exp_total).abs() < tol_strict, + "pnl_total: expected {exp_total:.5}, got {:.5}", out[0]); + assert!((out[1] - exp_mean).abs() < tol_strict, + "pnl_mean: expected {exp_mean:.5}, got {:.5}", out[1]); + assert!((out[2] - exp_var).abs() < tol_var, + "pnl_var: expected {exp_var:.5} (E[X²]−E[X]² form), got {:.5}", out[2]); + assert!((out[3] - exp_max_dd).abs() < tol_strict, + "pnl_max_dd: expected {exp_max_dd:.5}, got {:.5}", out[3]); +} diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 6641d098f..101c12459 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -3346,3 +3346,26 @@ Closes plan task #295 partial — 8th pearl (`pearl_sp5_close_out`) and full clo - 50-epoch 3-seed × 3-fold L40S validation (plan task #289) — pending; smoke is not a substitute. Refs: SP5 spec at `6e6e0fa11`; `feedback_no_partial_refactor`, `feedback_isv_for_adaptive_bounds`, `feedback_adaptive_not_tuned`, `pearl_first_observation_bootstrap`, `pearl_wiener_optimal_adaptive_alpha`, `pearl_engagement_rate_self_correction`. + +### SP5 Layer D Task D1 — PnL aggregation kernel (additive) (2026-05-02) + +First 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`. D1 lands additively — kernel + Rust launcher + 4 reserved ISV slots are present, but the host-side aggregation in `compute_epoch_financials` (training_loop.rs ~4862-4916) 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`. + +**What this lands:** +- `pnl_aggregation_kernel.cu` (151 lines) — single-block 256-thread shmem-reduce producer. Reads per-trade event arrays (`trade_pnls`, `trade_durations`, `trade_directions`); writes 4 aggregated floats to `scratch[SCRATCH_PNL_AGG_BASE=207..211)`: total / mean / population variance / max running drawdown. No atomicAdd (per `feedback_no_atomicadd.md`). Empty-input contract: `num_trades == 0` writes 0.0 to all four slots so the apply_pearls chain reads a defined value every step. +- Rust launcher `GpuDqnTrainer::launch_sp5_pnl_aggregation` — 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 (Pearl 8 trail / Pearl 1-ext num_atoms). +- 4 new ISV slots `PNL_TOTAL_INDEX=286 .. PNL_MAX_DD_INDEX=289` registered in `sp5_isv_slots.rs`. `SP5_SLOT_END` 286 → 290; `ISV_TOTAL_DIM` 286 → 290; `LAYOUT_FINGERPRINT_FRAGMENT` extended with `PNL_*` entries so existing checkpoints fail-fast on layout change. +- `SP5_PRODUCER_COUNT` semantics clarified — was conventionally equal to "unique SP5 slot count" (110) by coincidence at the Pearl 6 carve-out's introduction; D1's 4 new slots break that coincidence. Constant is now formally the **wiener-buffer linear-span** = `SP5_SLOT_END − SP5_SLOT_BASE = 116`, not the unique-slot count (114). Wiener buffer grows from 543 → 561 floats. Pearl 6's reserved-but-unused 18 wiener floats at offsets [531..549) and the carve-out gap's 6 floats at [525..531) remain reserved-unused (Pearl 6 doesn't call apply_pearls; gap slots don't exist). D1's 4 slots use offsets [549..561). +- `StateResetRegistry` entry `sp5_pnl_aggregation` (FoldReset; D1 outputs are NOT cross-fold persistent, unlike Pearl 6 Kelly stats) + matching dispatch arm in `reset_named_state` zeroing ISV[286..290) at fold boundary so the new fold's first launch fires Pearl A's first-observation replacement (paired with the existing bulk wiener_state_buf memset, per `pearl_first_observation_bootstrap.md`). +- `build.rs` cubin registration; static `SP5_PNL_AGGREGATION_CUBIN` byte slice + `pnl_aggregation_kernel: CudaFunction` field on `GpuDqnTrainer` + cubin/function loader entry mirroring the Pearl 8 trail kernel template. +- GPU-gated unit test `pnl_aggregation_kernel_correctness` in `sp5_producer_unit_tests.rs`. Synthetic input `[+10, −5, +3, −2, +1]` → analytical ground truth: total=7, mean=1.4, population variance=25.84, max_dd=5.0 (cumulative curve 10/5/8/6/7 against running peak 10). No CPU reference per `feedback_no_cpu_test_fallbacks.md`. + +**Verification gates:** +- `cargo check -p ml --offline` clean. +- `cargo build -p ml --release --offline --features cuda` clean — `pnl_aggregation_kernel.cubin` compiles via nvcc. +- `cargo test -p ml --lib --offline -- sp5_isv_slots state_reset_registry` 6/6 pass (slot-layout invariants + new `pnl_aggregation_slots_contiguous_and_above_kelly_block` test + 3 pre-existing registry tests). +- GPU correctness test (`pnl_aggregation_kernel_correctness`) `--ignored`-gated; fires on the next L40S smoke alongside the pre-existing 17 SP5 producer unit tests. + +**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`.