feat(sp5): Layer D Task D1 — PnL aggregation kernel (additive)
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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
|
||||
|
||||
@@ -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::<f32>() 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,
|
||||
|
||||
161
crates/ml/src/cuda_pipeline/pnl_aggregation_kernel.cu
Normal file
161
crates/ml/src/cuda_pipeline/pnl_aggregation_kernel.cu
Normal file
@@ -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<f32> 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();
|
||||
}
|
||||
}
|
||||
@@ -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<usize> = (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<usize> = (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);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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 }
|
||||
}
|
||||
|
||||
@@ -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 '{}'. \
|
||||
|
||||
@@ -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<CudaStream>) -> 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<CudaStream>) -> 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::<f32>() 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]);
|
||||
}
|
||||
|
||||
@@ -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`.
|
||||
|
||||
Reference in New Issue
Block a user