From a92ff28a9865a0b40ad4c8398697d8f453ecb381 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Wed, 6 May 2026 13:46:03 +0200 Subject: [PATCH] =?UTF-8?q?feat(sp15-p1.2):=20cost-net=20sharpe=20kernel?= =?UTF-8?q?=20=E2=80=94=20commission=20+=20spread=20+=20OFI-impact?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Per spec §6.2 per-side semantics: cost_t = commission_per_rt × rt_ind[t] + half_spread[t] × |pos[t]| × side_ind[t] + ofi_lambda × |pos[t]| × |ofi[t]| × side_ind[t] Commission charged at close; half-spread × |pos| at entry AND exit (sums to one full spread per RT); OFI impact same per-side. Initial λ=2.0e-4 in ISV[OFI_IMPACT_LAMBDA_INDEX=407] as Invariant-1 anchor (constructor-write + FoldReset rewrite per feedback_isv_for_adaptive_bounds); per-fold ISV refit may overwrite in later phases. Mean cost-per-bar emitted to ISV[COST_PER_BAR_AVG_INDEX=408] (stateful kernel output; FoldReset sentinel 0 + Pearl A first-observation bootstrap). Same kernel reads LobBar fields from synthetic markets (Phase 2A) and real fxcache LOB (prod) — dev/prod parity per Q3. Phase 1.2 lands kernel + launcher + anchor seed + 2 registry entries + 2 dispatch arms only; consumer migration deferred to a follow-up commit per feedback_no_partial_refactor (mirrors Phase 1.1 atomic pattern). Oracle test cost_net_sharpe_round_trip_charges_full_spread validates single round-trip cost = 2.00 / 10 bars = 0.20/bar with mean_pnl_net = 0.80 on 1.69s RTX 3050 Ti (sm_86). Zero regressions introduced (946 passed / 13 pre-existing failures, same as Task 1.1). Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/build.rs | 10 ++ .../cuda_pipeline/cost_net_sharpe_kernel.cu | 111 +++++++++++++++ .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 91 +++++++++++++ .../src/trainers/dqn/state_reset_registry.rs | 25 ++++ .../src/trainers/dqn/trainer/training_loop.rs | 27 ++++ crates/ml/tests/sp15_phase1_oracle_tests.rs | 128 +++++++++++++++++- docs/dqn-wire-up-audit.md | 2 + 7 files changed, 392 insertions(+), 2 deletions(-) create mode 100644 crates/ml/src/cuda_pipeline/cost_net_sharpe_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index bfe92fe59..7f3ae7ed1 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -833,6 +833,16 @@ fn main() { // (val × sqrt(525600)) split. Annualisation is the host-side caller's // responsibility (sharpe_annualised = out[2] × sqrt(N_bars_per_year)). "sharpe_per_bar_kernel.cu", + // SP15 Phase 1.2 (2026-05-06): cost-net sharpe. Single-block, + // BLOCK=256 kernel computing mean/std/sharpe of (pnl - cost) per + // spec §6.2 per-side semantics: commission_per_rt × rt_ind + + // half_spread × |pos| × side_ind + ofi_lambda × |pos| × |ofi| × + // side_ind. Charges full spread per RT (entry + exit). Reads + // ofi_lambda from ISV[OFI_IMPACT_LAMBDA_INDEX=407]; emits + // mean cost-per-bar to ISV[COST_PER_BAR_AVG_INDEX=408]. Same + // kernel reads LobBar fields from synthetic markets and prod + // fxcache LOB (dev/prod parity per Q3). + "cost_net_sharpe_kernel.cu", ]; // ALL kernels get common header (BF16 types + wrappers) diff --git a/crates/ml/src/cuda_pipeline/cost_net_sharpe_kernel.cu b/crates/ml/src/cuda_pipeline/cost_net_sharpe_kernel.cu new file mode 100644 index 000000000..b1dcd5192 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/cost_net_sharpe_kernel.cu @@ -0,0 +1,111 @@ +// crates/ml/src/cuda_pipeline/cost_net_sharpe_kernel.cu +// +// SP15 Phase 1.2 — cost-net sharpe. +// +// Per spec §6.2 per-side semantics: +// cost_t = commission_per_rt × rt_ind[t] +// + half_spread[t] × |position[t]| × side_ind[t] +// + ofi_lambda × |position[t]| × |ofi[t]| × side_ind[t] +// +// Commission is charged at close (rt_ind=1). Half-spread × |position| is +// charged on side-events at entry AND exit (sums to one full spread per +// round-trip). OFI impact same per-side. ofi_lambda is read from +// ISV[OFI_IMPACT_LAMBDA_INDEX=407] (constant across the window; +// constructor-seeded to 2e-4, FoldReset-rewritten to 2e-4 since this is an +// Invariant-1 anchor, not a stateful EMA). +// +// Outputs: +// isv[COST_PER_BAR_AVG_INDEX=408] = sum(cost_t) / n +// out[0] = mean(pnl - cost_t) +// out[1] = std(pnl - cost_t) +// out[2] = sharpe_net = out[0] / max(out[1], 1e-12) +// +// Same kernel reads LobBar fields from synthetic markets (Phase 2A) and +// real fxcache LOB (prod) — dev/prod parity per Q3. +// +// Single-block launch, BLOCK=256. No atomicAdd (per feedback_no_atomicadd): +// reduction uses shared-memory tree-reduce only. `__threadfence_system()` +// flushes mapped-pinned writes before the host reads via `read_volatile`. + +extern "C" __global__ void cost_net_sharpe_kernel( + const float* __restrict__ pnl, + const float* __restrict__ half_spread, + const float* __restrict__ ofi, + const unsigned int* __restrict__ side_ind, + const unsigned int* __restrict__ rt_ind, + const float* __restrict__ position, + int n, + float commission_per_rt, + float* __restrict__ isv, + float* __restrict__ out +) { + if (blockIdx.x != 0) return; + + const int tid = (int)threadIdx.x; + const int BLOCK = (int)blockDim.x; + __shared__ float reduce_a[256]; + __shared__ float reduce_b[256]; + __shared__ float mean_pnl_shared; + __shared__ float mean_cost_shared; + + // Read OFI lambda once (slot 407 = OFI_IMPACT_LAMBDA_INDEX). + const float ofi_lambda = isv[407]; + + // Pass 1: parallel sums of pnl_t and cost_t. + float local_pnl = 0.0f; + float local_cost = 0.0f; + for (int i = tid; i < n; i += BLOCK) { + local_pnl += pnl[i]; + float c_comm = commission_per_rt * (float)rt_ind[i]; + float pos_abs = fabsf(position[i]); + float c_spread = half_spread[i] * pos_abs * (float)side_ind[i]; + float c_ofi = ofi_lambda * pos_abs * fabsf(ofi[i]) * (float)side_ind[i]; + local_cost += c_comm + c_spread + c_ofi; + } + reduce_a[tid] = local_pnl; + reduce_b[tid] = local_cost; + __syncthreads(); + for (int s = BLOCK / 2; s > 0; s >>= 1) { + if (tid < s) { + reduce_a[tid] += reduce_a[tid + s]; + reduce_b[tid] += reduce_b[tid + s]; + } + __syncthreads(); + } + if (tid == 0) { + mean_pnl_shared = (n > 0) ? reduce_a[0] / (float)n : 0.0f; + mean_cost_shared = (n > 0) ? reduce_b[0] / (float)n : 0.0f; + // ISV slot 408 = COST_PER_BAR_AVG_INDEX (per sp15_isv_slots.rs). + isv[408] = mean_cost_shared; + __threadfence_system(); + } + __syncthreads(); + + const float mean_net = mean_pnl_shared - mean_cost_shared; + + // Pass 2: variance of (pnl_t - cost_t). + float local_sq = 0.0f; + for (int i = tid; i < n; i += BLOCK) { + float c_comm = commission_per_rt * (float)rt_ind[i]; + float pos_abs = fabsf(position[i]); + float c_spread = half_spread[i] * pos_abs * (float)side_ind[i]; + float c_ofi = ofi_lambda * pos_abs * fabsf(ofi[i]) * (float)side_ind[i]; + float net = pnl[i] - (c_comm + c_spread + c_ofi); + float d = net - mean_net; + local_sq += d * d; + } + reduce_a[tid] = local_sq; + __syncthreads(); + for (int s = BLOCK / 2; s > 0; s >>= 1) { + if (tid < s) reduce_a[tid] += reduce_a[tid + s]; + __syncthreads(); + } + if (tid == 0) { + float var = (n > 0) ? reduce_a[0] / (float)n : 0.0f; + float std = sqrtf(fmaxf(var, 1e-12f)); + out[0] = mean_net; + out[1] = std; + out[2] = (std > 0.0f) ? mean_net / std : 0.0f; + __threadfence_system(); + } +} diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index a33ebdedd..bd0a7c840 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -754,6 +754,82 @@ pub fn launch_sp15_sharpe_per_bar( Ok(()) } +/// SP15 Phase 1.2 (2026-05-06): cost-net sharpe cubin. Per spec §6.2 +/// per-side cost semantics: commission_per_rt charged at close (rt_ind=1); +/// half-spread × |pos| charged at entry AND exit on side_ind events (sums +/// to one full spread per round-trip); ofi_lambda × |pos| × |ofi| same +/// per-side. Reads ofi_lambda from ISV[OFI_IMPACT_LAMBDA_INDEX=407]; +/// emits mean cost-per-bar to ISV[COST_PER_BAR_AVG_INDEX=408]. Same +/// kernel reads LobBar fields from synthetic markets (Phase 2A) and real +/// fxcache LOB (prod) — dev/prod parity per Q3. +pub static SP15_COST_NET_SHARPE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/cost_net_sharpe_kernel.cubin")); + +/// SP15 Phase 1.2 (2026-05-06): launcher for the cost-net sharpe kernel. +/// Free function (matches `launch_sp15_sharpe_per_bar` precedent) so unit +/// / oracle tests can drive the kernel directly without the trainer struct; +/// production callers invoke the same launcher. +/// +/// Per-side cost semantics (spec §6.2): +/// `cost_t = commission_per_rt × rt_ind[t]` +/// `+ half_spread[t] × |position[t]| × side_ind[t]` +/// `+ ofi_lambda × |position[t]| × |ofi[t]| × side_ind[t]` +/// +/// `pnl`, `half_spread`, `ofi`, `position`, `isv`, `out` are device f32 +/// pointers. `side_ind` and `rt_ind` are device u32 pointers. `out` MUST +/// point to ≥3 f32 slots: `[mean_pnl_net, std, sharpe_net]`. `isv` MUST +/// be the ISV bus (≥443 f32 slots) — slot 407 is read for `ofi_lambda`, +/// slot 408 is written with `mean(cost_t)`. +#[allow(clippy::too_many_arguments)] +pub fn launch_sp15_cost_net_sharpe( + stream: &Arc, + pnl: cudarc::driver::sys::CUdeviceptr, + half_spread: cudarc::driver::sys::CUdeviceptr, + ofi: cudarc::driver::sys::CUdeviceptr, + side_ind: cudarc::driver::sys::CUdeviceptr, + rt_ind: cudarc::driver::sys::CUdeviceptr, + position: cudarc::driver::sys::CUdeviceptr, + n: i32, + commission_per_rt: f32, + isv: cudarc::driver::sys::CUdeviceptr, + out: cudarc::driver::sys::CUdeviceptr, +) -> Result<(), MLError> { + let module = stream + .context() + .load_cubin(SP15_COST_NET_SHARPE_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!( + "load sp15_cost_net_sharpe cubin: {e}" + )))?; + let kernel = module + .load_function("cost_net_sharpe_kernel") + .map_err(|e| MLError::ModelError(format!( + "load cost_net_sharpe_kernel function: {e}" + )))?; + unsafe { + stream + .launch_builder(&kernel) + .arg(&pnl) + .arg(&half_spread) + .arg(&ofi) + .arg(&side_ind) + .arg(&rt_ind) + .arg(&position) + .arg(&n) + .arg(&commission_per_rt) + .arg(&isv) + .arg(&out) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!( + "launch cost_net_sharpe_kernel: {e}" + )))?; + } + Ok(()) +} + /// SP11 Fix 39 (2026-05-04, Task A2): SimHash novelty signal — lookup + /// update kernels sharing one cubin. Lookup reads /// `1/sqrt(1+count)` ∈ [0, 1] for each (state, action) bucket; update @@ -19727,6 +19803,21 @@ impl GpuDqnTrainer { *sig_ptr.add(HOLD_COST_INDEX) = HOLD_COST_BASE; *sig_ptr.add(HOLD_RATE_TARGET_INDEX) = HOLD_RATE_TARGET_DEFAULT; + // SP15 Phase 1.2 (2026-05-06): OFI-impact lambda + // (per spec §6.2). Constant Invariant-1 anchor — NOT + // a stateful EMA — so constructor-write per + // `feedback_isv_for_adaptive_bounds.md`. Per-fold ISV + // refit may overwrite this in later phases; the + // FoldReset arm rewrites the same default to keep the + // cost kernel functional during cold-start of each + // new fold (must never reach sentinel 0, which would + // silently disable the OFI-impact term). The matching + // COST_PER_BAR_AVG_INDEX=408 is a stateful kernel + // output (FoldReset sentinel 0; first kernel launch + // writes the true mean). + use crate::cuda_pipeline::sp15_isv_slots::OFI_IMPACT_LAMBDA_INDEX; + *sig_ptr.add(OFI_IMPACT_LAMBDA_INDEX) = 2.0e-4_f32; + // Layout fingerprint (ISV[58..60)). Compile-time structural hash of // the slot layout; checkpoint load fails-fast on mismatch. // Stored as a u64 split across two f32 lanes using raw bit-cast so diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 0789a5a87..429ad9021 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -1005,6 +1005,31 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "ISV[AUX_DIR_ACC_POST_OPEN_MIN_INDEX=394] — SP14 anti-gradient-hacking circuit breaker: running minimum of aux_dir_acc since Gate 1 last opened. Initialised to 1.0 (no minimum observed) when Gate 1 opens; updated each step post-open as `min = fminf(min, current_aux_dir_acc)`. If `open_min − current_aux_dir_acc > LOCKOUT_TRIGGER_DROP=0.05` while q_disagreement also rises by >LOCKOUT_TRIGGER_DIS_RISE=0.10, the lockout circuit fires (B.5) and writes LOCKOUT_EPOCHS=2.0 to GRADIENT_HACK_LOCKOUT_REMAINING_INDEX=395. FoldReset sentinel SENTINEL_POST_OPEN_MIN=1.0 (no minimum observed at fold start; the circuit breaker must accumulate post-open observations before it can fire). Produced by `alpha_grad_compute_kernel` (B.4). Spec: docs/superpowers/specs/2026-05-05-sp14-aux-q-wire-earned-gradient-flow.md.", }, + // ── SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots ──────────── + // Two ISV slots [407, 408]: + // - OFI_IMPACT_LAMBDA_INDEX=407: Invariant-1 anchor (NOT a + // stateful EMA). Constructor-seeded to 2.0e-4 per spec §6.2; + // FoldReset rewrites the same anchor (NOT sentinel 0) so the + // cost kernel functions correctly on the new fold's first + // launch. Per-fold ISV refit may overwrite this in later + // phases; the FoldReset arm restores the default to break + // cold-start of each new fold (must never reach sentinel 0, + // which would silently disable the OFI-impact term). + // - COST_PER_BAR_AVG_INDEX=408: stateful kernel output + // (mean(cost_t)). FoldReset sentinel 0; Pearl A bootstraps + // from the new fold's first kernel launch per + // `pearl_first_observation_bootstrap.md`. + // Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md. + RegistryEntry { + name: "sp15_ofi_impact_lambda", + category: ResetCategory::FoldReset, + description: "ISV[OFI_IMPACT_LAMBDA_INDEX=407] — SP15 Phase 1.2 (2026-05-06) OFI-impact lambda for the cost-net sharpe kernel (per spec §6.2). Per-side cost: `λ × |position| × |ofi| × side_ind`. Constant Invariant-1 anchor (2.0e-4) per `feedback_isv_for_adaptive_bounds.md` — NOT a stateful EMA. Constructor-seeded; FoldReset rewrites the same default so the cost kernel functions correctly on the new fold's first launch (must never reach sentinel 0, which would silently disable the OFI-impact term). Per-fold ISV refit may overwrite this in later phases; the reset path restores the default to break cold-start of each new fold. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §6.2.", + }, + RegistryEntry { + name: "sp15_cost_per_bar_avg", + category: ResetCategory::FoldReset, + description: "ISV[COST_PER_BAR_AVG_INDEX=408] — SP15 Phase 1.2 (2026-05-06) mean cost-per-bar emitted by the cost-net sharpe kernel (`sum(cost_t) / n` over the sharpe window). Stateful kernel output (NOT cross-fold persistent); FoldReset sentinel 0 so the new fold's first `launch_sp15_cost_net_sharpe` writes the new fold's true mean. Per `pearl_first_observation_bootstrap.md`. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §6.2.", + }, // SP5 Task A1: Wiener-state companion reset. The wiener_state_buf // covers SP4 (SP4_PRODUCER_COUNT=71 producers × 3 = 213 floats) + // SP5 (SP5_PRODUCER_COUNT × 3 floats) = (71 + SP5_PRODUCER_COUNT) × 3 diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 65d5a365e..ce0b1b8f4 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -8066,6 +8066,33 @@ impl DQNTrainer { ); } } + // SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots. + // OFI_IMPACT_LAMBDA_INDEX=407 is an Invariant-1 anchor (NOT + // a stateful EMA) — rewrite the constructor's value at fold + // boundary so the cost kernel functions correctly on the new + // fold's first launch (must never reach sentinel 0, which + // would silently disable the OFI-impact term). Per + // `feedback_isv_for_adaptive_bounds.md`. Per-fold ISV refit + // may overwrite this in later phases. + "sp15_ofi_impact_lambda" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp15_isv_slots::OFI_IMPACT_LAMBDA_INDEX; + fused.trainer().write_isv_signal_at( + OFI_IMPACT_LAMBDA_INDEX, + 2.0e-4, + ); + } + } + // COST_PER_BAR_AVG_INDEX=408 is a stateful kernel output — + // FoldReset sentinel 0 so Pearl A bootstraps from the new + // fold's first `launch_sp15_cost_net_sharpe` per + // `pearl_first_observation_bootstrap.md`. + "sp15_cost_per_bar_avg" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp15_isv_slots::COST_PER_BAR_AVG_INDEX; + fused.trainer().write_isv_signal_at(COST_PER_BAR_AVG_INDEX, 0.0); + } + } // SP11 Task A2 (2026-05-04): novelty visit-count hash table // reset arm — closes the A0 deferral per the registry entry's // docstring. 1M f32 slots, zero-filled via mapped-pinned host diff --git a/crates/ml/tests/sp15_phase1_oracle_tests.rs b/crates/ml/tests/sp15_phase1_oracle_tests.rs index f0dcf62c4..f5323a5a9 100644 --- a/crates/ml/tests/sp15_phase1_oracle_tests.rs +++ b/crates/ml/tests/sp15_phase1_oracle_tests.rs @@ -10,8 +10,22 @@ mod gpu { use std::sync::Arc; use cudarc::driver::{CudaContext, CudaStream}; - use ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_sharpe_per_bar; - use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer; + use ml::cuda_pipeline::gpu_dqn_trainer::{ + launch_sp15_cost_net_sharpe, launch_sp15_sharpe_per_bar, + }; + use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedU32Buffer}; + use ml::cuda_pipeline::sp15_isv_slots::{ + COST_PER_BAR_AVG_INDEX, OFI_IMPACT_LAMBDA_INDEX, SP15_SLOT_END, + }; + + /// Upper bound for any ISV index touched by SP15 oracle tests. Matches + /// `gpu_dqn_trainer::ISV_TOTAL_DIM` (which is `pub(crate)` and not + /// reachable from integration tests). The `sp15_isv_slots` regression + /// test `all_sp15_slots_fit_within_isv_total_dim` enforces that + /// `SP15_SLOT_END <= ISV_TOTAL_DIM`, so allocating an ISV buffer of + /// length `SP15_SLOT_END` is large enough for every SP15 slot the + /// kernel reads or writes. + const ISV_LEN: usize = SP15_SLOT_END; /// Build a CUDA stream against device 0. Mirrors the helper in /// `sp4_producer_unit_tests.rs` and `distributional_q_tests.rs`. If no @@ -88,4 +102,114 @@ mod gpu { let out = out_buf.read_all(); assert!((out[2] - 0.5).abs() < 0.05, "sharpe={}, expected ~0.5", out[2]); } + + /// Test 1.2 — single round-trip oracle for the cost-net sharpe kernel. + /// + /// Setup: n=10 bars, position +1 contract held throughout, half_spread + /// 0.25, commission $1.50/RT, ofi=0 (zero out the OFI-impact term so + /// the spread + commission contributions are isolated). + /// Side events: bar 0 (entry, side_ind=1) and bar 5 (exit, + /// side_ind=1, rt_ind=1). Gross PnL: 10.0 booked at bar 5 (the + /// close-out bar). + /// + /// Expected per spec §6.2: + /// total cost = 1.50 commission + 0.25 (entry half-spread) + /// + 0.25 (exit half-spread) = 2.00 + /// cost_per_bar_avg = 2.00 / 10 = 0.20 + /// mean(pnl) = 10.0 / 10 = 1.0 + /// mean_pnl_net = 1.0 − 0.20 = 0.80 + /// + /// Validates: + /// • Per-side cost semantics (entry + exit = full spread per RT). + /// • Commission charged at close (rt_ind=1 only on bar 5). + /// • OFI-impact term scales with `|position|·|ofi|·side_ind` + /// (here zero because `ofi=0`). + /// • COST_PER_BAR_AVG_INDEX (slot 408) emitted to ISV. + #[test] + #[ignore = "requires GPU"] + fn cost_net_sharpe_round_trip_charges_full_spread() { + let stream = make_test_stream(); + let n = 10usize; + + // Inputs. + let mut pnl = vec![0.0f32; n]; + pnl[5] = 10.0; + let half_spread = vec![0.25f32; n]; + let ofi = vec![0.0f32; n]; + let position = vec![1.0f32; n]; + let mut side_ind = vec![0u32; n]; + side_ind[0] = 1; + side_ind[5] = 1; + let mut rt_ind = vec![0u32; n]; + rt_ind[5] = 1; + + // Safety: CUDA context active via `make_test_stream` above. + let pnl_buf = unsafe { MappedF32Buffer::new(n) } + .expect("alloc MappedF32Buffer for pnl"); + pnl_buf.write_from_slice(&pnl); + let half_spread_buf = unsafe { MappedF32Buffer::new(n) } + .expect("alloc MappedF32Buffer for half_spread"); + half_spread_buf.write_from_slice(&half_spread); + let ofi_buf = unsafe { MappedF32Buffer::new(n) } + .expect("alloc MappedF32Buffer for ofi"); + ofi_buf.write_from_slice(&ofi); + let position_buf = unsafe { MappedF32Buffer::new(n) } + .expect("alloc MappedF32Buffer for position"); + position_buf.write_from_slice(&position); + let side_ind_buf = unsafe { MappedU32Buffer::new(n) } + .expect("alloc MappedU32Buffer for side_ind"); + side_ind_buf.write_from_slice(&side_ind); + let rt_ind_buf = unsafe { MappedU32Buffer::new(n) } + .expect("alloc MappedU32Buffer for rt_ind"); + rt_ind_buf.write_from_slice(&rt_ind); + + // ISV bus seeded with OFI_IMPACT_LAMBDA=0 (mirrors the spec's + // zero-OFI test variable). The constructor seeds slot 407 to 2e-4 + // in production; here we override to 0 so the spread term is the + // sole non-commission cost contributor. + let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) } + .expect("alloc MappedF32Buffer for isv bus"); + let mut isv_init = vec![0.0f32; ISV_LEN]; + isv_init[OFI_IMPACT_LAMBDA_INDEX] = 0.0; + isv_buf.write_from_slice(&isv_init); + + let out_buf = unsafe { MappedF32Buffer::new(3) } + .expect("alloc MappedF32Buffer for cost-net sharpe output"); + + launch_sp15_cost_net_sharpe( + &stream, + pnl_buf.dev_ptr, + half_spread_buf.dev_ptr, + ofi_buf.dev_ptr, + side_ind_buf.dev_ptr, + rt_ind_buf.dev_ptr, + position_buf.dev_ptr, + n as i32, + /* commission_per_rt = */ 1.50, + isv_buf.dev_ptr, + out_buf.dev_ptr, + ) + .expect("launch cost_net_sharpe_kernel"); + stream + .synchronize() + .expect("synchronize after cost_net_sharpe_kernel launch"); + + let isv_host = isv_buf.read_all(); + let cost_per_bar_avg = isv_host[COST_PER_BAR_AVG_INDEX]; + // Total cost = 1.50 commission + 0.25 + 0.25 = 2.00 over 10 bars + // → cost_per_bar_avg = 0.20. + assert!( + (cost_per_bar_avg - 0.20).abs() < 1e-3, + "cost_per_bar_avg = {}, expected 0.20", + cost_per_bar_avg + ); + + let out = out_buf.read_all(); + // mean_pnl_net = mean(pnl) − mean(cost) = 1.0 − 0.20 = 0.80. + assert!( + (out[0] - 0.80).abs() < 1e-2, + "mean_pnl_net = {}, expected 0.80", + out[0] + ); + } } diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index b049ce360..5d988541f 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -4,6 +4,8 @@ SP15 Phase 1.1 — unified per-bar sharpe kernel (2026-05-06): single GPU kernel `sharpe_per_bar_kernel.cu` computes mean/std/sharpe via 2-pass shared-memory tree-reduce (single-block, BLOCK=256, no atomicAdd). Replaces the SP14-era split between `sharpe_ema` (per-batch EMA, train) and `sharpe_annualised` (val × sqrt(525600)). Same formula and same window definition for train and val; annualisation is the host-side caller's responsibility (`sharpe_annualised = out[2] × sqrt(N_bars_per_year)`). **Phase 1.1 lands kernel + free-function launcher only**; consumer migration in `metrics.rs` / `training_loop.rs` is deferred to a follow-up commit per `feedback_no_partial_refactor.md` (atomic consumer migration is a load-bearing change too risky to bundle with kernel introduction; the kernel is verified-working in isolation first via the two GPU oracle tests below). Touched: `crates/ml/src/cuda_pipeline/sharpe_per_bar_kernel.cu` (new), `crates/ml/build.rs` (+1 cubin manifest entry), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (+1 `pub static SP15_SHARPE_PER_BAR_CUBIN` cubin slot + 1 `pub fn launch_sp15_sharpe_per_bar` free-function launcher matching the `sp4_histogram_p99_test_kernel` cold-path standalone-launcher precedent), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (+2 GPU oracle tests inside the existing `mod gpu` block: `unified_sharpe_kernel_zero_mean` validates mean=0/std=1/sharpe=0 on alternating ±1 PnL; `unified_sharpe_kernel_positive_drift` validates sharpe≈0.5 on `0.5 + alternating ±1` PnL). Both tests pass on local RTX 3050 Ti (sm_86) in 1.78s. cargo test -p ml --lib --features cuda: 946 passed / 13 failed — all 13 failures pre-existing on the parent commit (verified via `git stash` baseline run); zero introduced by this commit. Hard rules: `feedback_no_atomicadd` (block-tree-reduce only), `feedback_no_partial_refactor` (kernel + launcher land atomically; consumer migration follows as its own atomic commit), `feedback_no_stubs` (launcher returns `Result<(), MLError>` and is fully functional, not a placeholder), `feedback_no_htod_htoh_only_mapped_pinned` (oracle tests use `MappedF32Buffer` for both input PnL and output `[mean, std, sharpe]`). +SP15 Phase 1.2 — cost-net sharpe kernel (2026-05-06): single GPU kernel `cost_net_sharpe_kernel.cu` computes mean(pnl − cost) / std(pnl − cost) sharpe with per-side cost semantics per spec §6.2 (`cost_t = commission_per_rt × rt_ind[t] + half_spread[t] × |position[t]| × side_ind[t] + ofi_lambda × |position[t]| × |ofi[t]| × side_ind[t]`). Commission charged at close (rt_ind=1); half-spread × |pos| charged at entry AND exit (sums to one full spread per round-trip); OFI impact same per-side. Single-block BLOCK=256 2-pass shared-memory tree-reduce — no atomicAdd. Reads `ofi_lambda` from `ISV[OFI_IMPACT_LAMBDA_INDEX=407]`; emits mean cost-per-bar to `ISV[COST_PER_BAR_AVG_INDEX=408]`. Same kernel reads LobBar fields from synthetic markets (Phase 2A) and real fxcache LOB (prod) — dev/prod parity per Q3. **Phase 1.2 lands kernel + free-function launcher + ISV anchor seed + state-reset registry entries only**; consumer migration in metrics / val sites is deferred to a follow-up commit per `feedback_no_partial_refactor.md` (consumer migration is a load-bearing change; the kernel is verified-working in isolation first via the GPU oracle test below, mirroring the Phase 1.1 atomic pattern). Touched: `crates/ml/src/cuda_pipeline/cost_net_sharpe_kernel.cu` (new), `crates/ml/build.rs` (+1 cubin manifest entry), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (+1 `pub static SP15_COST_NET_SHARPE_CUBIN` cubin slot + 1 `pub fn launch_sp15_cost_net_sharpe` free-function launcher mirroring `launch_sp15_sharpe_per_bar`'s `stream.context().load_cubin` precedent + 1 constructor-write `*sig_ptr.add(OFI_IMPACT_LAMBDA_INDEX) = 2.0e-4` Invariant-1 anchor seed at the SP13/SP14 anchor block per `feedback_isv_for_adaptive_bounds.md`), `crates/ml/src/trainers/dqn/state_reset_registry.rs` (+2 `RegistryEntry` records: `sp15_ofi_impact_lambda` rewrites the constructor's 2.0e-4 default at fold boundary since this is an Invariant-1 anchor not a stateful EMA — must never reach sentinel 0 which would silently disable the OFI-impact term; `sp15_cost_per_bar_avg` is a stateful kernel output, FoldReset sentinel 0 so Pearl A bootstraps from the new fold's first launch per `pearl_first_observation_bootstrap.md`), `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (+2 dispatch arms for the 2 new registry entries — registry-arm contract enforced by `every_fold_and_soft_reset_entry_has_dispatch_arm` regression test), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (+1 GPU oracle test `cost_net_sharpe_round_trip_charges_full_spread` validates a single round-trip on n=10 bars: 1 entry @ bar 0 + 1 exit @ bar 5, position +1 contract, half_spread 0.25, commission 1.50, ofi_lambda overridden to 0 in the test ISV bus to isolate spread+commission. Expected: `cost_per_bar_avg = (1.50 + 0.25 + 0.25) / 10 = 0.20` and `mean_pnl_net = 10/10 − 0.20 = 0.80`). Test passes on local RTX 3050 Ti (sm_86) in 1.69s. cargo test -p ml --lib --features cuda: 946 passed / 13 failed — same 13 failures pre-existing on the parent commit `3667cd1b0` (Task 1.1 baseline); zero introduced by this commit. State-reset registry tests: 4/4 pass including `every_fold_and_soft_reset_entry_has_dispatch_arm` validating both new entries have matching dispatch arms. Hard rules: `feedback_no_atomicadd` (block-tree-reduce only), `feedback_no_partial_refactor` (kernel + launcher + anchor-seed + registry entries + dispatch arms land atomically; consumer migration follows as its own atomic commit), `feedback_no_stubs` (launcher returns `Result<(), MLError>` and is fully functional), `feedback_no_htod_htoh_only_mapped_pinned` (oracle test uses `MappedF32Buffer` for f32 inputs/output and `MappedU32Buffer` for u32 side/rt indicators), `feedback_isv_for_adaptive_bounds` (OFI lambda is a constant Invariant-1 anchor written at constructor + FoldReset, NOT a runtime EMA — per-fold ISV refit may overwrite this in later phases). + SP12 v3 — per-trade event-driven reward composition (2026-05-04): three architectural changes in one atomic commit per spec `2026-05-04-sp12-quality-over-quantity-reward-shaping.md` (committed at `ecab584c3`). Empirical motivation: 50-epoch validation `train-multi-seed-pmbwn` on commit `6a259942e` showed sharpe-gaming (sharpe=80 held while PnL declined 30% over 8 epochs; 62% trade rate, 46% win rate stuck). **Three causes** addressed in one commit per `feedback_no_partial_refactor`: (1) **Asymmetric bounded cap** at `experience_kernels.cu:2758` — replaces SP11's symmetric ±10 cap (commit `35db31089`) with `fmaxf(REWARD_NEG_CAP=-10, fminf(base_reward, REWARD_POS_CAP=+5))`. The 2:1 ratio matches Kahneman/Tversky meta-analysis pegging human loss aversion at ~2.0-2.25 — calibration to a known psychological constant (Invariant-1 numerical anchor in `state_layout.cuh`), not a tuned multiplier. SP11's stability work preserved (cap still bilateral). Per `pearl_audit_unboundedness_for_implicit_asymmetry`: SP11 erased the pre-fix implicit downside-asymmetry by symmetrizing the bound; the model became risk-neutral and locked into HFT noise extraction. SP12 widens only the negative cap back toward the pre-SP11 unbounded regime.