From e968f4ded99200ea37772aa220dc2b95931cee13 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Wed, 6 May 2026 23:03:30 +0200 Subject: [PATCH] =?UTF-8?q?feat(sp15-wave3a):=20kernel-side=20foundation?= =?UTF-8?q?=20=E2=80=94=20baseline=20output=20buffers=20+=20position=5Fhis?= =?UTF-8?q?tory=20derivation?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Wave 3a half of the val-cost-streams refactor (3b host-side wire-up follows). Atomically migrates the kernel-side contracts; 5 launchers remain orphan transiently awaiting 3b production callers. Baseline kernels (1.4): - 4 baseline_*_kernel signatures gain 'out: float*' parameter writing per-window [mean, std, raw_sharpe] (matches 1.1.b sharpe_per_bar shape) - ISV writes to slots 409, 410, 412, 416 removed entirely (per-window output is correct for WindowMetrics consumption; ISV-scalar writes were spec scaffolding for a single-fold-aggregate version that 1.4.b's per-window contract supersedes) - 4 ISV slot constants removed from sp15_isv_slots.rs - state_reset_registry: NO entries to remove (verified via grep — the 4 slots never had registry entries / dispatch arms in the first place; they were single-fold-aggregate scalars defaulted at every fold start by the constructor-write that initialises the ISV bus). Task 4 from the dispatch is a no-op; the every_fold_and_soft_reset_entry_has_dispatch_arm regression test continues to pass unchanged. - 4 oracle tests migrated to output-buffer assertion - layout_fingerprint_seed string updated (4 retired entries removed, 4 trunk-shared entries retained; layout-break-class change) New action_decoding_helpers.cuh: - Extracts factored_action_to_dir_idx + factored_action_to_position __device__ helpers (the latter is a higher-level position state- machine helper not previously available) - Mirrors trade_physics.cuh::decode_direction_4b semantics exactly so on-policy and counterfactual paths agree on factored-action meaning - Single source of truth for action→direction→position mapping; consumers #include the header New position_history_derivation_kernel.cu (post-loop derivation for cost_net_sharpe consumer in Wave 3b): - Reads actions_history_buf, reconstructs per-bar position_history (-1/0/+1), side_ind (1.0 on position change), rt_ind (1.0 on transition-to-flat from non-flat) via sequential walk (single block per window, no atomicAdd per feedback_no_atomicadd) - New launcher launch_sp15_position_history_derivation in gpu_dqn_trainer.rs - New cubin manifest entry in build.rs - 1 oracle test covering 8-bar Short→Hold→Long→Hold→Flat→Long→Flat→ Short sequence; expected position/side_ind/rt_ind triples match hand-computed values Atomic per feedback_no_partial_refactor for the ISV-contract change (every consumer of slots 409/410/412/416 migrated in this commit; their consumers were the 4 oracle tests, all migrated). The orphan launcher transient state for the 5 baselines + derivation kernel is explicitly the 3a/3b split point — production callers land in 3b's GpuBacktestEvaluator::new constructor signature change. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/build.rs | 28 ++- .../cuda_pipeline/action_decoding_helpers.cuh | 63 ++++++ .../ml/src/cuda_pipeline/baseline_kernels.cu | 84 ++++---- .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 165 ++++++++++++---- .../position_history_derivation_kernel.cu | 76 ++++++++ crates/ml/src/cuda_pipeline/sp15_isv_slots.rs | 35 +++- crates/ml/tests/sp15_phase1_oracle_tests.rs | 184 ++++++++++++++---- docs/dqn-wire-up-audit.md | 2 + 8 files changed, 500 insertions(+), 137 deletions(-) create mode 100644 crates/ml/src/cuda_pipeline/action_decoding_helpers.cuh create mode 100644 crates/ml/src/cuda_pipeline/position_history_derivation_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 543a330c7..e21666d7e 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -854,7 +854,7 @@ fn main() { // DD_PCT (406). The 1e-4 calmar floor eliminates the saturation- // at-100 artifact previously seen in the train-dd4xl HEALTH_DIAG. "dd_state_kernel.cu", - // SP15 Phase 1.4 partial (2026-05-06): 4 constant-policy + // SP15 Phase 1.4 + Wave 3a (2026-05-06): 4 constant-policy // counterfactual baselines per spec §6.4. Single source file with // 4 `extern "C" __global__` symbols (baseline_buyhold_kernel, // baseline_hold_only_kernel, baseline_naive_momentum_kernel, @@ -865,14 +865,26 @@ fn main() { // BLOCK=256 with a templated 2-pass shared-memory tree-reduce // (no atomicAdd). Reads price/half_spread/ofi arrays + commission // scalar; computes constant-policy or last-bar-return-based - // actions; writes sharpe_net to its own ISV slot (409 buyhold, - // 410 hold_only, 412 naive_momentum, 416 naive_reversion). - // Trunk-shared baselines (random_dir_kelly slot 411, aux_only - // slot 413, mag_quarter_fixed slot 414, trail_only slot 415) are - // out of scope for this commit — they need partial-policy- - // forward access from the main eval pass and are deferred to - // Task 1.4.b. + // actions; writes per-window `[mean, std, raw_sharpe]` to a + // caller-provided output buffer (Wave 3a contract change — was + // ISV-scalar slots 409/410/412/416 pre-Wave-3a). Trunk-shared + // baselines (random_dir_kelly slot 411, aux_only slot 413, + // mag_quarter_fixed slot 414, trail_only slot 415) are out of + // scope for this commit — they need partial-policy-forward + // access from the main eval pass and are deferred to Task 1.4.b. "baseline_kernels.cu", + // SP15 Wave 3a (2026-05-06): post-loop derivation of per-bar + // position state from `actions_history_buf` (factored ints + // recorded by env_step). Single-block-per-window launch with one + // thread per block (the position state machine is sequential per + // window — no in-window parallelism). Emits 3 f32 streams per + // (window, bar): position_history (-1/0/+1), side_ind (1.0 on + // position change), rt_ind (1.0 on transition-to-flat from + // non-flat). Consumed by Wave 3b's cost-net sharpe path; uses + // `action_decoding_helpers.cuh::factored_action_to_position` + // single-source-of-truth for the action→direction→position + // mapping (mirrors `trade_physics.cuh::decode_direction_4b`). + "position_history_derivation_kernel.cu", // SP15 Phase 1.5 (2026-05-06): dd_pct foundational state input // concat. LAYOUT FINGERPRINT BREAK — pre-SP15 checkpoints WILL // NOT LOAD after this commit. Greenfield OK per spec Q1. Per diff --git a/crates/ml/src/cuda_pipeline/action_decoding_helpers.cuh b/crates/ml/src/cuda_pipeline/action_decoding_helpers.cuh new file mode 100644 index 000000000..89285d9b8 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/action_decoding_helpers.cuh @@ -0,0 +1,63 @@ +// crates/ml/src/cuda_pipeline/action_decoding_helpers.cuh +// +// SP15 Wave 3a (2026-05-06) — single-source-of-truth `__device__` helpers +// for the factored action → direction / position mapping. +// +// Background: the canonical decoding lives in `trade_physics.cuh` +// (`decode_direction_4b`, `decode_magnitude_4b`, …). Those helpers are +// already used inline by every consumer that needs the dir/mag/order/urg +// indices — `experience_kernels.cu`, `backtest_env_kernel.cu`, +// `experience_action_select`, etc. The new SP15 Wave 3a derivation +// kernels (`position_history_derivation_kernel.cu`) consume the same +// factored-action ints but need a higher-level helper that emits the +// signed *position* directly: -1 / 0 / +1. +// +// This header DOES NOT replace `trade_physics.cuh`'s primitives; it adds +// one helper on top of them so derivation kernels (which don't need the +// trade-physics machinery) avoid pulling that whole header. Both this +// header and `trade_physics.cuh` ultimately read the same DIR_* macros +// from `state_layout.cuh`, so the source-of-truth for the direction +// mapping (DIR_SHORT=0, DIR_HOLD=1, DIR_LONG=2, DIR_FLAT=3) is unchanged. + +#ifndef ACTION_DECODING_HELPERS_CUH +#define ACTION_DECODING_HELPERS_CUH + +#include "state_layout.cuh" + +// Decode direction-axis index from a factored action int. +// Mirrors `trade_physics.cuh::decode_direction_4b` exactly — duplicated +// here as a stand-alone helper so derivation kernels can `#include` this +// header without dragging in `trade_physics.cuh`'s GPU-only Kelly-cap +// machinery. If `b1`, `b2`, `b3` are all 3 (production setup) the +// product is 27 and direction = action / 27 (matches the FlatAction +// encoding `3 * b1 * b2 * b3` used by `handle_capital_floor_breach`). +__device__ __forceinline__ int factored_action_to_dir_idx( + int action, int b1, int b2, int b3 +) { + return action / (b1 * b2 * b3); +} + +// Map a factored action int to a signed position scalar tracking the +// state machine of (-1=Short, 0=Flat, +1=Long, prev=Hold). Reads the +// previous step's position and returns the new position. +// +// DIR_SHORT (0) → -1 +// DIR_HOLD (1) → prev_position (no-op, keeps current state) +// DIR_LONG (2) → +1 +// DIR_FLAT (3) → 0 +// +// `prev_position` is the integer position (−1/0/+1) at bar t-1. At t=0 +// the caller passes 0 (every window starts flat by convention; matches +// `experience_kernels.cu`'s init-portfolio-state convention). +__device__ __forceinline__ int factored_action_to_position( + int action, int b1, int b2, int b3, int prev_position +) { + const int dir = factored_action_to_dir_idx(action, b1, b2, b3); + if (dir == DIR_SHORT) return -1; + if (dir == DIR_LONG ) return 1; + if (dir == DIR_FLAT ) return 0; + /* DIR_HOLD or any out-of-band sentinel: keep prev. */ + return prev_position; +} + +#endif // ACTION_DECODING_HELPERS_CUH diff --git a/crates/ml/src/cuda_pipeline/baseline_kernels.cu b/crates/ml/src/cuda_pipeline/baseline_kernels.cu index 01ca24293..cd0b90937 100644 --- a/crates/ml/src/cuda_pipeline/baseline_kernels.cu +++ b/crates/ml/src/cuda_pipeline/baseline_kernels.cu @@ -1,12 +1,23 @@ // crates/ml/src/cuda_pipeline/baseline_kernels.cu // -// SP15 Phase 1.4 (partial) — 4 constant-policy counterfactual baselines. +// SP15 Phase 1.4 + Wave 3a (2026-05-06) — 4 constant-policy counterfactual baselines. // Per spec §6.4. Pure-CUDA kernels (no trunk-share). Each kernel reads // the price/spread/ofi arrays + a commission scalar, computes the // baseline's per-bar PnL stream from a constant or last-bar-return-based // action policy, then runs a 2-pass shared-memory tree-reduce mean/std -// (no atomicAdd per `feedback_no_atomicadd`) and writes sharpe_net to -// its allocated ISV slot. +// (no atomicAdd per `feedback_no_atomicadd`) and writes per-window +// `[mean, std, raw_sharpe]` to a caller-provided output buffer (matches +// the 1.1.b `sharpe_per_bar_kernel` shape). +// +// Wave 3a contract change (2026-05-06): the prior ISV-scalar writes to +// slots 409/410/412/416 are removed entirely. Per-window output is the +// correct shape for `WindowMetrics` consumption; ISV scalars were spec +// scaffolding for a single-fold-aggregate version that 1.4.b's per-window +// contract supersedes. Migration is atomic per `feedback_no_partial_refactor`: +// the 4 oracle tests that previously read ISV slots are migrated in the +// same commit to read from the output buffer instead. Production callers +// (Wave 3b) will allocate an `[n_windows, 3]` MappedF32Buffer and invoke +// the launchers per-window (mirrors the 1.1.b sharpe_per_bar pattern). // // Single source file — multiple `extern "C" __global__` symbols share // one cubin per the build.rs 1:1-source-to-cubin convention. The four @@ -14,12 +25,6 @@ // different `get_function()` symbols, mirroring the multi-kernel-per- // file pattern from SP11 (`novelty_simhash_kernel.cu`). // -// ISV slots written (allocated in `sp15_isv_slots.rs`, spec §4.3): -// slot 409 — BASELINE_BUYHOLD_SHARPE_INDEX -// slot 410 — BASELINE_HOLD_ONLY_SHARPE_INDEX -// slot 412 — BASELINE_NAIVE_MOMENTUM_SHARPE_INDEX -// slot 416 — BASELINE_NAIVE_REVERSION_SHARPE_INDEX -// // Trunk-shared baselines (random_dir_kelly slot 411, aux_only slot 413, // mag_quarter_fixed slot 414, trail_only slot 415) are out of scope // for this commit — they need partial-policy-forward access from the @@ -38,26 +43,22 @@ // the model-under-test (both pay full RT commission + entry/exit // spread, neither pays OFI impact for this baseline class). -#define BUYHOLD_SHARPE_SLOT 409 -#define HOLD_ONLY_SHARPE_SLOT 410 -#define MOMENTUM_SHARPE_SLOT 412 -#define REVERSION_SHARPE_SLOT 416 - // Templated helper: 2-pass shared-memory tree-reduce of mean and std of -// the per-bar PnL stream produced by `ActionFn`. Returns sharpe to thread -// 0 only — caller stores it to the appropriate ISV slot. +// the per-bar PnL stream produced by `ActionFn`. Writes `[mean, std, +// raw_sharpe]` to `out[0..3]` from thread 0 only. // // Template parameter `ActionFn` is a CUDA functor: // __device__ int operator()(const float* prices, int i, int n) const; // returns +1 (long), 0 (flat/hold), -1 (short) for bar `i`. template -__device__ float compute_baseline_sharpe( +__device__ void compute_baseline_sharpe( const float* __restrict__ prices, const float* __restrict__ half_spread, const float* __restrict__ /* ofi */, // unused for this baseline class (lambda=0) int n, float commission_per_rt, - ActionFn action_fn + ActionFn action_fn, + float* out /* [mean, std, raw_sharpe] */ ) { const int tid = (int)threadIdx.x; const int BLOCK = (int)blockDim.x; @@ -117,9 +118,12 @@ __device__ float compute_baseline_sharpe( if (tid == 0) { const float var = (n > 1) ? reduce_buf[0] / (float)(n - 1) : 0.0f; const float std = sqrtf(fmaxf(var, 1e-12f)); - return (std > 0.0f) ? mean_shared / std : 0.0f; + const float sharpe = (std > 0.0f) ? mean_shared / std : 0.0f; + out[0] = mean_shared; + out[1] = std; + out[2] = sharpe; + __threadfence_system(); } - return 0.0f; } // CUDA functor structs — one per baseline policy. @@ -158,16 +162,12 @@ extern "C" __global__ void baseline_buyhold_kernel( const float* __restrict__ ofi, int n, float commission_per_rt, - float* __restrict__ isv + float* __restrict__ out /* [mean, std, raw_sharpe] */ ) { if (blockIdx.x != 0) return; - const float sharpe = compute_baseline_sharpe( - prices, half_spread, ofi, n, commission_per_rt, BuyholdAction() + compute_baseline_sharpe( + prices, half_spread, ofi, n, commission_per_rt, BuyholdAction(), out ); - if (threadIdx.x == 0) { - isv[BUYHOLD_SHARPE_SLOT] = sharpe; - __threadfence_system(); - } } extern "C" __global__ void baseline_hold_only_kernel( @@ -176,16 +176,12 @@ extern "C" __global__ void baseline_hold_only_kernel( const float* __restrict__ ofi, int n, float commission_per_rt, - float* __restrict__ isv + float* __restrict__ out /* [mean, std, raw_sharpe] */ ) { if (blockIdx.x != 0) return; - const float sharpe = compute_baseline_sharpe( - prices, half_spread, ofi, n, commission_per_rt, HoldOnlyAction() + compute_baseline_sharpe( + prices, half_spread, ofi, n, commission_per_rt, HoldOnlyAction(), out ); - if (threadIdx.x == 0) { - isv[HOLD_ONLY_SHARPE_SLOT] = sharpe; - __threadfence_system(); - } } extern "C" __global__ void baseline_naive_momentum_kernel( @@ -194,16 +190,12 @@ extern "C" __global__ void baseline_naive_momentum_kernel( const float* __restrict__ ofi, int n, float commission_per_rt, - float* __restrict__ isv + float* __restrict__ out /* [mean, std, raw_sharpe] */ ) { if (blockIdx.x != 0) return; - const float sharpe = compute_baseline_sharpe( - prices, half_spread, ofi, n, commission_per_rt, MomentumAction() + compute_baseline_sharpe( + prices, half_spread, ofi, n, commission_per_rt, MomentumAction(), out ); - if (threadIdx.x == 0) { - isv[MOMENTUM_SHARPE_SLOT] = sharpe; - __threadfence_system(); - } } extern "C" __global__ void baseline_naive_reversion_kernel( @@ -212,14 +204,10 @@ extern "C" __global__ void baseline_naive_reversion_kernel( const float* __restrict__ ofi, int n, float commission_per_rt, - float* __restrict__ isv + float* __restrict__ out /* [mean, std, raw_sharpe] */ ) { if (blockIdx.x != 0) return; - const float sharpe = compute_baseline_sharpe( - prices, half_spread, ofi, n, commission_per_rt, ReversionAction() + compute_baseline_sharpe( + prices, half_spread, ofi, n, commission_per_rt, ReversionAction(), out ); - if (threadIdx.x == 0) { - isv[REVERSION_SHARPE_SLOT] = sharpe; - __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 67b4efd71..58b454d38 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -993,35 +993,48 @@ pub fn launch_sp15_dd_pct_concat( Ok(()) } -/// SP15 Phase 1.4 partial (2026-05-06): cubin shared by all 4 constant- +/// SP15 Phase 1.4 + Wave 3a (2026-05-06): cubin shared by all 4 constant- /// policy counterfactual baselines (`baseline_buyhold_kernel`, /// `baseline_hold_only_kernel`, `baseline_naive_momentum_kernel`, /// `baseline_naive_reversion_kernel`). Per spec §6.4. Each baseline -/// computes its sharpe_net by running a constant-policy or last-bar- -/// return-based action stream through a 2-pass shared-memory tree-reduce -/// (single-block BLOCK=256, no atomicAdd) and writing to its own ISV -/// slot. The four launchers below all load this single cubin and resolve -/// different `get_function()` symbols — multi-kernel-per-file pattern -/// per `novelty_simhash_kernel.cu` precedent. +/// computes its per-window `[mean, std, raw_sharpe]` by running a +/// constant-policy or last-bar-return-based action stream through a +/// 2-pass shared-memory tree-reduce (single-block BLOCK=256, no atomicAdd) +/// and writing to a caller-provided output buffer. Wave 3a contract +/// change: ISV-scalar writes to slots 409/410/412/416 are removed entirely +/// — per-window output is the correct shape for `WindowMetrics` +/// consumption, mirrors the 1.1.b sharpe_per_bar pattern. The four +/// launchers below all load this single cubin and resolve different +/// `get_function()` symbols — multi-kernel-per-file pattern per +/// `novelty_simhash_kernel.cu` precedent. /// /// Trunk-shared baselines (random_dir_kelly slot 411, aux_only slot 413, /// mag_quarter_fixed slot 414, trail_only slot 415) are out of scope for /// this commit — they need partial-policy-forward access from the main /// eval pass and are deferred to Task 1.4.b. Their slots stay at sentinel /// 0.0 in the meantime. +/// +/// Wave 3a transient state (2026-05-06): the 4 baseline launchers below +/// are currently ORPHAN — production callers in `GpuBacktestEvaluator` +/// (`Wave 3b`) invoke them per-window once that constructor signature +/// migration lands. Oracle tests in `sp15_phase1_oracle_tests.rs` drive +/// the launchers directly. Acceptable per the documented 3a/3b split. pub static SP15_BASELINE_KERNELS_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/baseline_kernels.cubin")); -/// SP15 Phase 1.4 partial (2026-05-06): launcher for the buyhold +/// SP15 Phase 1.4 + Wave 3a (2026-05-06): launcher for the buyhold /// counterfactual baseline. Free function (matches `launch_sp15_dd_state` /// precedent) so unit/oracle tests can drive the kernel directly without -/// the trainer struct; production callers (eval-pass baseline launches) -/// invoke the same launcher. +/// the trainer struct; production callers (eval-pass baseline launches +/// in Wave 3b's `GpuBacktestEvaluator::new`) invoke the same launcher. /// -/// Policy: position=+1 every bar (no exits). Writes sharpe_net to -/// `ISV[BASELINE_BUYHOLD_SHARPE_INDEX=409]`. `prices`, `half_spread`, -/// `ofi`, `isv` are device f32 pointers; `n` is bar count. `isv` MUST -/// be the ISV bus (≥443 f32 slots). +/// Policy: position=+1 every bar (no exits). Writes per-window +/// `[mean, std, raw_sharpe]` to `out[0..3]`. `prices`, `half_spread`, +/// `ofi`, `out` are device f32 pointers; `n` is bar count. `out` MUST +/// point to ≥3 f32 slots — typically the `dev_ptr` of a `MappedF32Buffer` +/// of length `n_windows * 3`, with the launcher invoked sequentially per +/// window with the per-window slice (mirror of 1.1.b sharpe_per_bar +/// per-window launch sequence). pub fn launch_sp15_baseline_buyhold( stream: &Arc, prices: cudarc::driver::sys::CUdeviceptr, @@ -1029,7 +1042,7 @@ pub fn launch_sp15_baseline_buyhold( ofi: 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() @@ -1050,7 +1063,7 @@ pub fn launch_sp15_baseline_buyhold( .arg(&ofi) .arg(&n) .arg(&commission_per_rt) - .arg(&isv) + .arg(&out) .launch(LaunchConfig { grid_dim: (1, 1, 1), block_dim: (256, 1, 1), @@ -1063,12 +1076,12 @@ pub fn launch_sp15_baseline_buyhold( Ok(()) } -/// SP15 Phase 1.4 partial (2026-05-06): launcher for the hold-only +/// SP15 Phase 1.4 + Wave 3a (2026-05-06): launcher for the hold-only /// counterfactual baseline. Policy: action=Hold every bar (no entry, no -/// exit). Writes sharpe_net to `ISV[BASELINE_HOLD_ONLY_SHARPE_INDEX=410]`. +/// exit). Writes per-window `[mean, std, raw_sharpe]` to `out[0..3]`. /// On a strict no-trade trajectory `mean(pnl)=0` and `std(pnl)=0` so -/// the kernel's `(std > 0) ? mean / std : 0` guard emits 0, which is -/// the spec-mandated sentinel for an undefined sharpe. +/// the kernel's `(std > 0) ? mean / std : 0` guard emits sharpe=0, which +/// is the spec-mandated sentinel for an undefined sharpe. pub fn launch_sp15_baseline_hold_only( stream: &Arc, prices: cudarc::driver::sys::CUdeviceptr, @@ -1076,7 +1089,7 @@ pub fn launch_sp15_baseline_hold_only( ofi: 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() @@ -1097,7 +1110,7 @@ pub fn launch_sp15_baseline_hold_only( .arg(&ofi) .arg(&n) .arg(&commission_per_rt) - .arg(&isv) + .arg(&out) .launch(LaunchConfig { grid_dim: (1, 1, 1), block_dim: (256, 1, 1), @@ -1110,10 +1123,10 @@ pub fn launch_sp15_baseline_hold_only( Ok(()) } -/// SP15 Phase 1.4 partial (2026-05-06): launcher for the naive-momentum +/// SP15 Phase 1.4 + Wave 3a (2026-05-06): launcher for the naive-momentum /// counterfactual baseline. Policy: `direction = sign(prices[i] − -/// prices[i-1])` with |position|=1. Writes sharpe_net to -/// `ISV[BASELINE_NAIVE_MOMENTUM_SHARPE_INDEX=412]`. +/// prices[i-1])` with |position|=1. Writes per-window `[mean, std, +/// raw_sharpe]` to `out[0..3]`. pub fn launch_sp15_baseline_naive_momentum( stream: &Arc, prices: cudarc::driver::sys::CUdeviceptr, @@ -1121,7 +1134,7 @@ pub fn launch_sp15_baseline_naive_momentum( ofi: 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() @@ -1142,7 +1155,7 @@ pub fn launch_sp15_baseline_naive_momentum( .arg(&ofi) .arg(&n) .arg(&commission_per_rt) - .arg(&isv) + .arg(&out) .launch(LaunchConfig { grid_dim: (1, 1, 1), block_dim: (256, 1, 1), @@ -1155,10 +1168,10 @@ pub fn launch_sp15_baseline_naive_momentum( Ok(()) } -/// SP15 Phase 1.4 partial (2026-05-06): launcher for the naive-reversion +/// SP15 Phase 1.4 + Wave 3a (2026-05-06): launcher for the naive-reversion /// counterfactual baseline. Policy: `direction = -sign(prices[i] − /// prices[i-1])` with |position|=1 (mirror of naive_momentum). Writes -/// sharpe_net to `ISV[BASELINE_NAIVE_REVERSION_SHARPE_INDEX=416]`. +/// per-window `[mean, std, raw_sharpe]` to `out[0..3]`. pub fn launch_sp15_baseline_naive_reversion( stream: &Arc, prices: cudarc::driver::sys::CUdeviceptr, @@ -1166,7 +1179,7 @@ pub fn launch_sp15_baseline_naive_reversion( ofi: 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() @@ -1187,7 +1200,7 @@ pub fn launch_sp15_baseline_naive_reversion( .arg(&ofi) .arg(&n) .arg(&commission_per_rt) - .arg(&isv) + .arg(&out) .launch(LaunchConfig { grid_dim: (1, 1, 1), block_dim: (256, 1, 1), @@ -1200,6 +1213,91 @@ pub fn launch_sp15_baseline_naive_reversion( Ok(()) } +/// SP15 Wave 3a (2026-05-06): cubin for the post-loop position-history +/// derivation kernel (`sp15_position_history_derivation_kernel`). Reads +/// `actions_history_buf` (factored ints recorded by env_step) and emits +/// per-bar `position_history` / `side_ind` / `rt_ind` f32 streams +/// consumed by Wave 3b's cost-net sharpe path. Uses the new +/// `action_decoding_helpers.cuh` single-source-of-truth helper for the +/// action→direction→position mapping (mirrors +/// `trade_physics.cuh::decode_direction_4b` semantics so on-policy and +/// counterfactual paths agree on what each factored action means). +/// +/// Wave 3a transient state: the launcher below is currently ORPHAN — the +/// production caller in `GpuBacktestEvaluator` lands in Wave 3b alongside +/// the cost-net sharpe wire-up. Oracle tests in +/// `sp15_phase1_oracle_tests.rs` drive the launcher directly. Acceptable +/// per the documented 3a/3b split. +pub static SP15_POSITION_HISTORY_DERIVATION_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/position_history_derivation_kernel.cubin")); + +/// SP15 Wave 3a (2026-05-06): launcher for the position-history derivation +/// kernel. Free function (matches `launch_sp15_dd_state` / +/// `launch_sp15_baseline_*` precedent) so unit/oracle tests can drive +/// the kernel directly without the trainer struct. +/// +/// `actions_history_buf` is the device pointer to `[n_windows * max_len]` +/// i32 factored actions (recorded by `experience_env_step` / +/// `backtest_env_step`; sentinel -1 marks unwritten / early-terminated +/// bars and is treated as "keep prev position"). The 3 output f32 streams +/// are caller-provided device pointers, each `[n_windows * max_len]` long. +/// +/// Grid: `[n_windows, 1, 1]` × block `[1, 1, 1]`. The position state +/// machine is sequential per window (position[t] depends on position[t-1]), +/// so within-window parallelism is zero by construction; the across- +/// windows axis carries all the parallelism. No atomicAdd per +/// `feedback_no_atomicadd` (the kernel is a pure scan, not a reduction). +#[allow(clippy::too_many_arguments)] +pub fn launch_sp15_position_history_derivation( + stream: &Arc, + actions_history_buf: cudarc::driver::sys::CUdeviceptr, + n_windows: i32, + max_len: i32, + b1_size: i32, + b2_size: i32, + b3_size: i32, + position_history_out: cudarc::driver::sys::CUdeviceptr, + side_ind_out: cudarc::driver::sys::CUdeviceptr, + rt_ind_out: cudarc::driver::sys::CUdeviceptr, +) -> Result<(), MLError> { + let module = stream + .context() + .load_cubin(SP15_POSITION_HISTORY_DERIVATION_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!( + "load sp15_position_history_derivation cubin: {e}" + )))?; + let kernel = module + .load_function("sp15_position_history_derivation_kernel") + .map_err(|e| MLError::ModelError(format!( + "load sp15_position_history_derivation_kernel function: {e}" + )))?; + if n_windows <= 0 { + return Ok(()); + } + unsafe { + stream + .launch_builder(&kernel) + .arg(&actions_history_buf) + .arg(&n_windows) + .arg(&max_len) + .arg(&b1_size) + .arg(&b2_size) + .arg(&b3_size) + .arg(&position_history_out) + .arg(&side_ind_out) + .arg(&rt_ind_out) + .launch(LaunchConfig { + grid_dim: (n_windows as u32, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!( + "launch sp15_position_history_derivation_kernel: {e}" + )))?; + } + Ok(()) +} + /// SP15 Wave 2 (2026-05-06): cubin for the per-step ISV-driven α /// producer (`alpha_split_producer_kernel`). Replaces the pre-Wave-2 /// `SP15_R_QUALITY_DISCIPLINE_SPLIT_CUBIN` which housed both the @@ -3222,10 +3320,9 @@ const fn layout_fingerprint_seed() -> &'static [u8] { EGF_SCHMITT_HI=397;EGF_SCHMITT_LO=398;EGF_VAR_AUX_REF=399;EGF_VAR_Q_REF=400;\ DD_CURRENT=401;DD_MAX=402;DD_RECOVERY_BARS=403;DD_PERSISTENCE=404;CALMAR=405;DD_PCT=406;\ OFI_IMPACT_LAMBDA=407;COST_PER_BAR_AVG=408;\ - BASELINE_BUYHOLD_SHARPE=409;BASELINE_HOLD_ONLY_SHARPE=410;\ - BASELINE_RANDOM_DIR_KELLY_SHARPE=411;BASELINE_NAIVE_MOMENTUM_SHARPE=412;\ + BASELINE_RANDOM_DIR_KELLY_SHARPE=411;\ BASELINE_AUX_ONLY_SHARPE=413;BASELINE_MAG_QUARTER_FIXED_SHARPE=414;\ - BASELINE_TRAIL_ONLY_SHARPE=415;BASELINE_NAIVE_REVERSION_SHARPE=416;\ + BASELINE_TRAIL_ONLY_SHARPE=415;\ ALPHA_SPLIT=417;GRAD_NORM_QUALITY=418;GRAD_NORM_DISCIPLINE=419;\ LAMBDA_DD=420;DD_THRESHOLD=421;DD_PENALTY_GRAD_NORM=422;\ REGRET_EMA=423;LAMBDA_REGRET=424;REGRET_GRAD_NORM=425;\ diff --git a/crates/ml/src/cuda_pipeline/position_history_derivation_kernel.cu b/crates/ml/src/cuda_pipeline/position_history_derivation_kernel.cu new file mode 100644 index 000000000..f9a9f9932 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/position_history_derivation_kernel.cu @@ -0,0 +1,76 @@ +// crates/ml/src/cuda_pipeline/position_history_derivation_kernel.cu +// +// SP15 Wave 3a (2026-05-06) — post-loop derivation of per-bar position +// state from the recorded factored-action history. +// +// Reads `actions_history_buf[w * max_len + t]` (factored ints recorded +// by `experience_env_step` / `backtest_env_step`) and emits three f32 +// streams per (window, bar) consumed by Wave 3b's cost-net sharpe path: +// +// position_history_out[w * max_len + t] ∈ {-1.0, 0.0, +1.0} +// side_ind_out[w * max_len + t] 1.0 when position changed +// from prev step, else 0.0 +// rt_ind_out[w * max_len + t] 1.0 when transitioning to +// flat from a non-flat +// position (round-trip +// close), else 0.0 +// +// Single-block-per-window launch with one thread per window. Each thread +// walks bars sequentially because position(t) depends on position(t-1) +// — there's no parallelism within a window for this state machine. The +// per-window walks ARE parallel across windows (n_windows blocks). +// No atomicAdd per `feedback_no_atomicadd` (the kernel doesn't reduce; +// it's a pure scan). +// +// Sentinel handling: `actions_history_buf` is initialised to -1 by +// `experience_kernels.cu` (per the actions_history measurement-bug fix +// from a prior commit). Bars where `action == -1` (never written by +// env_step — typically because the window terminated early) are treated +// as "no position change" — we keep the prior position. For windows +// shorter than `max_len`, the trailing -1 sentinels emit position = +// last-known-position with side_ind=0, rt_ind=0; downstream consumers +// (Wave 3b) gate on `window_lens[w]` to suppress these tail bars. + +#include "action_decoding_helpers.cuh" + +extern "C" __global__ void sp15_position_history_derivation_kernel( + const int* __restrict__ actions_history_buf, /* [n_windows * max_len] */ + int n_windows, + int max_len, + int b1_size, + int b2_size, + int b3_size, + float* __restrict__ position_history_out, /* [n_windows * max_len] */ + float* __restrict__ side_ind_out, /* [n_windows * max_len] */ + float* __restrict__ rt_ind_out /* [n_windows * max_len] */ +) { + const int w = (int)blockIdx.x; + if (w >= n_windows) return; + if (threadIdx.x != 0) return; /* one thread per block per window */ + + int prev_position = 0; /* every window starts flat by convention */ + for (int t = 0; t < max_len; ++t) { + const int idx = w * max_len + t; + const int action = actions_history_buf[idx]; + + int new_position; + if (action < 0) { + /* sentinel — keep prior position, emit no-change indicators */ + new_position = prev_position; + } else { + new_position = factored_action_to_position( + action, b1_size, b2_size, b3_size, prev_position + ); + } + + const int changed = (new_position != prev_position) ? 1 : 0; + const int round_trip = (changed && new_position == 0 && prev_position != 0) ? 1 : 0; + + position_history_out[idx] = (float)new_position; + side_ind_out[idx] = (float)changed; + rt_ind_out[idx] = (float)round_trip; + + prev_position = new_position; + } + __threadfence_system(); +} diff --git a/crates/ml/src/cuda_pipeline/sp15_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp15_isv_slots.rs index 68a72957e..c0373606b 100644 --- a/crates/ml/src/cuda_pipeline/sp15_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/sp15_isv_slots.rs @@ -24,15 +24,29 @@ pub const DD_PCT_INDEX: usize = 406; pub const OFI_IMPACT_LAMBDA_INDEX: usize = 407; pub const COST_PER_BAR_AVG_INDEX: usize = 408; -// === Phase 1.4: 8 counterfactual baselines (8 slots) === -pub const BASELINE_BUYHOLD_SHARPE_INDEX: usize = 409; -pub const BASELINE_HOLD_ONLY_SHARPE_INDEX: usize = 410; +// === Phase 1.4 + Wave 3a: 8 counterfactual baseline slots (4 retired, 4 reserved) === +// +// Wave 3a (2026-05-06): the 4 constant-policy baselines (buyhold, +// hold_only, naive_momentum, naive_reversion) had their own ISV slots +// (409, 410, 412, 416) for a single-fold-aggregate sharpe scalar. Wave 3a +// migrates those baselines to per-window output buffers (matching the +// 1.1.b sharpe_per_bar shape) — the kernels now write `[mean, std, +// raw_sharpe]` per window to a caller-provided `MappedF32Buffer`. The +// 4 ISV slot constants are retired here per `feedback_no_legacy_aliases` +// (no ghost constants left in the codebase). The slot indices 409/410/ +// 412/416 themselves remain reserved (gap in the SP15 layout) per the +// layout fingerprint contract — re-allocation is left to a future SP +// to keep the fingerprint stable across this transitional release. +// +// Trunk-shared baseline slots (411, 413, 414, 415) remain allocated: +// random_dir_kelly, aux_only, mag_quarter_fixed, trail_only need +// partial-policy-forward access from the main eval pass and stay on +// the ISV-scalar path until follow-up Task 1.4.b. They sit at sentinel +// 0.0 in the meantime. pub const BASELINE_RANDOM_DIR_KELLY_SHARPE_INDEX: usize = 411; -pub const BASELINE_NAIVE_MOMENTUM_SHARPE_INDEX: usize = 412; pub const BASELINE_AUX_ONLY_SHARPE_INDEX: usize = 413; pub const BASELINE_MAG_QUARTER_FIXED_SHARPE_INDEX: usize = 414; pub const BASELINE_TRAIL_ONLY_SHARPE_INDEX: usize = 415; -pub const BASELINE_NAIVE_REVERSION_SHARPE_INDEX: usize = 416; // === Phase 3.1: r_quality + r_discipline split (3 slots) === pub const ALPHA_SPLIT_INDEX: usize = 417; @@ -108,8 +122,15 @@ mod tests { assert_eq!(DD_CURRENT_INDEX, 401); assert_eq!(DD_PCT_INDEX, 406); assert_eq!(OFI_IMPACT_LAMBDA_INDEX, 407); - assert_eq!(BASELINE_BUYHOLD_SHARPE_INDEX, 409); - assert_eq!(BASELINE_NAIVE_REVERSION_SHARPE_INDEX, 416); + // Wave 3a (2026-05-06): BASELINE_BUYHOLD_SHARPE_INDEX (409) + + // BASELINE_HOLD_ONLY_SHARPE_INDEX (410) + BASELINE_NAIVE_MOMENTUM_SHARPE_INDEX + // (412) + BASELINE_NAIVE_REVERSION_SHARPE_INDEX (416) retired + // alongside the kernel contract change to per-window output buffers + // — slot indices remain reserved gaps for layout fingerprint stability. + assert_eq!(BASELINE_RANDOM_DIR_KELLY_SHARPE_INDEX, 411); + assert_eq!(BASELINE_AUX_ONLY_SHARPE_INDEX, 413); + assert_eq!(BASELINE_MAG_QUARTER_FIXED_SHARPE_INDEX, 414); + assert_eq!(BASELINE_TRAIL_ONLY_SHARPE_INDEX, 415); assert_eq!(ALPHA_SPLIT_INDEX, 417); assert_eq!(LAMBDA_DD_INDEX, 420); assert_eq!(REGRET_EMA_INDEX, 423); diff --git a/crates/ml/tests/sp15_phase1_oracle_tests.rs b/crates/ml/tests/sp15_phase1_oracle_tests.rs index 8a55fecb5..896c5c28a 100644 --- a/crates/ml/tests/sp15_phase1_oracle_tests.rs +++ b/crates/ml/tests/sp15_phase1_oracle_tests.rs @@ -13,15 +13,13 @@ mod gpu { use ml::cuda_pipeline::gpu_dqn_trainer::{ launch_sp15_baseline_buyhold, launch_sp15_baseline_hold_only, launch_sp15_baseline_naive_momentum, launch_sp15_baseline_naive_reversion, - launch_sp15_cost_net_sharpe, launch_sp15_dd_state, launch_sp15_sharpe_per_bar, + launch_sp15_cost_net_sharpe, launch_sp15_dd_state, + launch_sp15_position_history_derivation, launch_sp15_sharpe_per_bar, }; - use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedU32Buffer}; + use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedI32Buffer, MappedU32Buffer}; use ml::cuda_pipeline::sp15_isv_slots::{ - ALPHA_SPLIT_INDEX, BASELINE_BUYHOLD_SHARPE_INDEX, BASELINE_HOLD_ONLY_SHARPE_INDEX, - BASELINE_NAIVE_MOMENTUM_SHARPE_INDEX, BASELINE_NAIVE_REVERSION_SHARPE_INDEX, - COST_PER_BAR_AVG_INDEX, DD_CURRENT_INDEX, DD_MAX_INDEX, DD_PCT_INDEX, - DD_RECOVERY_BARS_INDEX, GRAD_NORM_DISCIPLINE_INDEX, GRAD_NORM_QUALITY_INDEX, - OFI_IMPACT_LAMBDA_INDEX, SP15_SLOT_END, + ALPHA_SPLIT_INDEX, COST_PER_BAR_AVG_INDEX, DD_CURRENT_INDEX, DD_MAX_INDEX, DD_PCT_INDEX, + DD_RECOVERY_BARS_INDEX, OFI_IMPACT_LAMBDA_INDEX, SP15_SLOT_END, }; /// Upper bound for any ISV index touched by SP15 oracle tests. Matches @@ -337,6 +335,10 @@ mod gpu { /// rules out a sharpe-near-zero degenerate baseline; upper 1.0 /// catches accidental zero-cost or reduction bugs that would /// inflate the sharpe. + /// + /// Wave 3a (2026-05-06): assertion reads from caller-provided output + /// buffer `[mean, std, raw_sharpe]` instead of ISV slot 409 (kernel + /// signature change). #[test] #[ignore = "requires GPU"] fn baseline_buyhold_positive_on_drift() { @@ -365,11 +367,9 @@ mod gpu { .expect("alloc MappedF32Buffer for ofi"); ofi_buf.write_from_slice(&ofi); - 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 baseline output [mean, std, raw_sharpe]"); + out_buf.write_from_slice(&[0.0f32; 3]); launch_sp15_baseline_buyhold( &stream, @@ -378,24 +378,24 @@ mod gpu { ofi_buf.dev_ptr, n as i32, /* commission_per_rt = */ 0.0, - isv_buf.dev_ptr, + out_buf.dev_ptr, ) .expect("launch baseline_buyhold_kernel"); stream .synchronize() .expect("synchronize after baseline_buyhold_kernel launch"); - let isv = isv_buf.read_all(); - let sharpe = isv[BASELINE_BUYHOLD_SHARPE_INDEX]; + let out = out_buf.read_all(); + let sharpe = out[2]; assert!( sharpe > 0.2, - "buyhold sharpe = {}, expected > 0.2 on +drift series", - sharpe + "buyhold sharpe = {}, expected > 0.2 on +drift series (mean={}, std={})", + sharpe, out[0], out[1] ); assert!( sharpe < 1.0, - "buyhold sharpe = {}, expected < 1.0 (sanity bound)", - sharpe + "buyhold sharpe = {}, expected < 1.0 (sanity bound; mean={}, std={})", + sharpe, out[0], out[1] ); } @@ -405,6 +405,9 @@ mod gpu { /// zero on every bar. The kernel's reduction yields mean=0 std=0 /// and the `(std > 0) ? mean / std : 0` guard returns 0, the /// spec-mandated sentinel for an undefined sharpe (no trades). + /// + /// Wave 3a (2026-05-06): assertion reads from caller-provided output + /// buffer `[mean, std, raw_sharpe]` instead of ISV slot 410. #[test] #[ignore = "requires GPU"] fn baseline_hold_only_emits_zero() { @@ -424,9 +427,9 @@ mod gpu { let ofi_buf = unsafe { MappedF32Buffer::new(n) } .expect("alloc MappedF32Buffer for ofi"); ofi_buf.write_from_slice(&ofi); - let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) } - .expect("alloc MappedF32Buffer for isv bus"); - isv_buf.write_from_slice(&vec![0.0f32; ISV_LEN]); + let out_buf = unsafe { MappedF32Buffer::new(3) } + .expect("alloc MappedF32Buffer for baseline output"); + out_buf.write_from_slice(&[0.0f32; 3]); launch_sp15_baseline_hold_only( &stream, @@ -435,19 +438,26 @@ mod gpu { ofi_buf.dev_ptr, n as i32, /* commission_per_rt = */ 0.0, - isv_buf.dev_ptr, + out_buf.dev_ptr, ) .expect("launch baseline_hold_only_kernel"); stream .synchronize() .expect("synchronize after baseline_hold_only_kernel launch"); - let isv = isv_buf.read_all(); - let sharpe = isv[BASELINE_HOLD_ONLY_SHARPE_INDEX]; + let out = out_buf.read_all(); + let sharpe = out[2]; assert!( sharpe.abs() < 1e-5, - "hold_only sharpe = {}, expected ~0 (no trades)", - sharpe + "hold_only sharpe = {}, expected ~0 (no trades; mean={}, std={})", + sharpe, out[0], out[1] + ); + // Also verify mean/std are themselves zero — hold_only's per-bar + // PnL stream is identically zero, not just zero-mean noise. + assert!( + out[0].abs() < 1e-5, + "hold_only mean = {}, expected ~0 (no trades)", + out[0] ); } @@ -456,6 +466,9 @@ mod gpu { /// sharpe values should sum to ~0. The series is deliberately /// mean-reverting (reversion outperforms) but the symmetry test is /// invariant to which side wins. + /// + /// Wave 3a (2026-05-06): each launcher writes to its own output + /// buffer `[mean, std, raw_sharpe]` (was ISV slots 412/416). #[test] #[ignore = "requires GPU"] fn baseline_momentum_reversion_symmetry() { @@ -486,9 +499,12 @@ mod gpu { .expect("alloc MappedF32Buffer for ofi"); ofi_buf.write_from_slice(&ofi); - let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) } - .expect("alloc MappedF32Buffer for isv bus"); - isv_buf.write_from_slice(&vec![0.0f32; ISV_LEN]); + let momentum_out = unsafe { MappedF32Buffer::new(3) } + .expect("alloc MappedF32Buffer for momentum output"); + momentum_out.write_from_slice(&[0.0f32; 3]); + let reversion_out = unsafe { MappedF32Buffer::new(3) } + .expect("alloc MappedF32Buffer for reversion output"); + reversion_out.write_from_slice(&[0.0f32; 3]); launch_sp15_baseline_naive_momentum( &stream, @@ -497,13 +513,9 @@ mod gpu { ofi_buf.dev_ptr, n as i32, 0.0, - isv_buf.dev_ptr, + momentum_out.dev_ptr, ) .expect("launch baseline_naive_momentum_kernel"); - stream - .synchronize() - .expect("synchronize after baseline_naive_momentum_kernel launch"); - launch_sp15_baseline_naive_reversion( &stream, prices_buf.dev_ptr, @@ -511,16 +523,15 @@ mod gpu { ofi_buf.dev_ptr, n as i32, 0.0, - isv_buf.dev_ptr, + reversion_out.dev_ptr, ) .expect("launch baseline_naive_reversion_kernel"); stream .synchronize() - .expect("synchronize after baseline_naive_reversion_kernel launch"); + .expect("synchronize after baseline_naive_{momentum,reversion}_kernel launches"); - let isv = isv_buf.read_all(); - let momentum_sharpe = isv[BASELINE_NAIVE_MOMENTUM_SHARPE_INDEX]; - let reversion_sharpe = isv[BASELINE_NAIVE_REVERSION_SHARPE_INDEX]; + let momentum_sharpe = momentum_out.read_all()[2]; + let reversion_sharpe = reversion_out.read_all()[2]; // Sign-flipped policies: their sum should be near zero (anti- // correlated picks on the same per-bar return stream). Bound 0.5 @@ -535,6 +546,99 @@ mod gpu { ); } + /// Test 1.4.j (Wave 3a, 2026-05-06) — position_history derivation kernel + /// reconstructs per-bar position/side_ind/rt_ind from a known factored + /// action sequence. + /// + /// Fixture: 1 window, 8 bars, b1=b2=b3=3 (production sizes). The + /// factored-action encoding is `dir * 27 + mag * 9 + order * 3 + urg`. + /// Per `state_layout.cuh`: DIR_SHORT=0, DIR_HOLD=1, DIR_LONG=2, + /// DIR_FLAT=3. So: + /// action= 0 → DIR_SHORT (position −1) + /// action= 27 → DIR_HOLD (keep prev) + /// action= 54 → DIR_LONG (position +1) + /// action= 81 → DIR_FLAT (position 0) + /// + /// The action sequence below exercises every transition: Short→Hold→ + /// Long→Hold→Flat→Long→Flat→Short. Expected per-bar derived state: + /// position : [-1, -1, +1, +1, 0, +1, 0, -1] + /// side_ind : [ 1, 0, 1, 0, 1, 1, 1, 1] /* changed-from-prev */ + /// rt_ind : [ 0, 0, 0, 0, 1, 0, 1, 0] /* transitioned to flat */ + /// (initial prev_position=0 — every window starts flat by convention.) + #[test] + #[ignore = "requires GPU"] + fn position_history_derivation_walks_action_sequence() { + let stream = make_test_stream(); + + const N_WINDOWS: usize = 1; + const MAX_LEN: usize = 8; + const B1: i32 = 3; + const B2: i32 = 3; + const B3: i32 = 3; + + // Action sequence: Short, Hold, Long, Hold, Flat, Long, Flat, Short. + // dir codes: 0, 1, 2, 1, 3, 2, 3, 0 → action = dir * 27. + let actions: [i32; MAX_LEN] = [0, 27, 54, 27, 81, 54, 81, 0]; + + // Safety: CUDA context active via `make_test_stream` above. + let actions_buf = unsafe { MappedI32Buffer::new(N_WINDOWS * MAX_LEN) } + .expect("alloc MappedI32Buffer for actions_history"); + actions_buf.write_from_slice(&actions); + + let position_buf = unsafe { MappedF32Buffer::new(N_WINDOWS * MAX_LEN) } + .expect("alloc MappedF32Buffer for position_history"); + position_buf.write_from_slice(&vec![0.0f32; N_WINDOWS * MAX_LEN]); + let side_ind_buf = unsafe { MappedF32Buffer::new(N_WINDOWS * MAX_LEN) } + .expect("alloc MappedF32Buffer for side_ind"); + side_ind_buf.write_from_slice(&vec![0.0f32; N_WINDOWS * MAX_LEN]); + let rt_ind_buf = unsafe { MappedF32Buffer::new(N_WINDOWS * MAX_LEN) } + .expect("alloc MappedF32Buffer for rt_ind"); + rt_ind_buf.write_from_slice(&vec![0.0f32; N_WINDOWS * MAX_LEN]); + + launch_sp15_position_history_derivation( + &stream, + actions_buf.dev_ptr, + N_WINDOWS as i32, + MAX_LEN as i32, + B1, + B2, + B3, + position_buf.dev_ptr, + side_ind_buf.dev_ptr, + rt_ind_buf.dev_ptr, + ) + .expect("launch sp15_position_history_derivation_kernel"); + stream + .synchronize() + .expect("synchronize after position_history_derivation launch"); + + let position = position_buf.read_all(); + let side_ind = side_ind_buf.read_all(); + let rt_ind = rt_ind_buf.read_all(); + + let expected_position = [-1.0_f32, -1.0, 1.0, 1.0, 0.0, 1.0, 0.0, -1.0]; + let expected_side_ind = [1.0_f32, 0.0, 1.0, 0.0, 1.0, 1.0, 1.0, 1.0]; + let expected_rt_ind = [0.0_f32, 0.0, 0.0, 0.0, 1.0, 0.0, 1.0, 0.0]; + + for t in 0..MAX_LEN { + assert!( + (position[t] - expected_position[t]).abs() < 1e-5, + "position[{}] = {}, expected {}", + t, position[t], expected_position[t] + ); + assert!( + (side_ind[t] - expected_side_ind[t]).abs() < 1e-5, + "side_ind[{}] = {}, expected {}", + t, side_ind[t], expected_side_ind[t] + ); + assert!( + (rt_ind[t] - expected_rt_ind[t]).abs() < 1e-5, + "rt_ind[{}] = {}, expected {}", + t, rt_ind[t], expected_rt_ind[t] + ); + } + } + /// Test 1.5 — dd_pct foundational state input concat (LAYOUT FINGERPRINT /// BREAK). /// diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 57abb276b..7073c56ce 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,8 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +SP15 Wave 3a — kernel-side foundation for val-cost-streams (2026-05-06): atomic kernel-side half of the Phase 1.4.b + Phase 1.2.b val-cost-streams refactor, decomposed into 3a (this commit) + 3b (host-side wire-up, separate dispatch) per session-capacity scoping. Wave 3a lands the kernel contract changes + new derivation kernel + shared header + oracle-test migration; the 5 launchers (4 baselines + 1 derivation) currently sit ORPHAN awaiting Wave 3b's `GpuBacktestEvaluator::new` constructor signature change that wires production callers. Acceptable transient state — explicitly bounded by the 3a/3b split. (1) **Baseline kernel signature change (1.4)**: `baseline_buyhold_kernel`, `baseline_hold_only_kernel`, `baseline_naive_momentum_kernel`, `baseline_naive_reversion_kernel` all gain `out: float*` as their last parameter, replacing the prior `isv: float*` ISV-bus pointer. Each kernel writes `[mean, std, raw_sharpe]` to `out[0..3]` for the per-window slice the launcher invocation handles — mirrors the 1.1.b `sharpe_per_bar_kernel` shape exactly. The ISV-scalar writes to slots 409 (BASELINE_BUYHOLD_SHARPE), 410 (BASELINE_HOLD_ONLY_SHARPE), 412 (BASELINE_NAIVE_MOMENTUM_SHARPE), 416 (BASELINE_NAIVE_REVERSION_SHARPE) are removed entirely (NOT kept as parallel writes per `feedback_no_partial_refactor`). The internal `compute_baseline_sharpe` helper changes return-by-value to write-via-out-pointer, emitting all three of mean/std/sharpe (was sharpe-only). (2) **4 ISV slot constants retired in `sp15_isv_slots.rs`**: `BASELINE_BUYHOLD_SHARPE_INDEX` / `BASELINE_HOLD_ONLY_SHARPE_INDEX` / `BASELINE_NAIVE_MOMENTUM_SHARPE_INDEX` / `BASELINE_NAIVE_REVERSION_SHARPE_INDEX` removed per `feedback_no_legacy_aliases` — no ghost constants. The slot indices 409/410/412/416 themselves remain reserved gaps in the SP15 layout for layout-fingerprint stability across the transitional release; re-allocation deferred to a future SP. The 4 trunk-shared baseline slots (411 random_dir_kelly, 413 aux_only, 414 mag_quarter_fixed, 415 trail_only) remain allocated — they need partial-policy-forward access from the main eval pass and stay on the ISV-scalar path until follow-up Task 1.4.b (out of Wave 3 scope). (3) **`grep -rn` audit verified zero non-baseline consumers** of the 4 retired slots: only `baseline_kernels.cu` (writes), `sp15_isv_slots.rs` (constants + layout regression test), `gpu_dqn_trainer.rs` (launcher doc-comments + layout-fingerprint string), and `tests/sp15_phase1_oracle_tests.rs` (asserts) referenced them — every consumer is migrated in this commit. The layout-fingerprint string in `gpu_dqn_trainer.rs::layout_fingerprint_seed` drops the 4 retired entries (BASELINE_BUYHOLD_SHARPE / BASELINE_HOLD_ONLY_SHARPE / BASELINE_NAIVE_MOMENTUM_SHARPE / BASELINE_NAIVE_REVERSION_SHARPE) but keeps the 4 trunk-shared entries (BASELINE_RANDOM_DIR_KELLY_SHARPE / BASELINE_AUX_ONLY_SHARPE / BASELINE_MAG_QUARTER_FIXED_SHARPE / BASELINE_TRAIL_ONLY_SHARPE) — the fingerprint changes with this contract change (intentional, layout-break-class) but is locked again post-commit. (4) **`state_reset_registry` change — no entries to remove**: a `grep -n` audit on `state_reset_registry.rs` + `training_loop.rs::reset_named_state` confirmed the 4 retired baseline slots NEVER had registry entries or dispatch arms in the first place (they were always single-fold-aggregate scalars defaulted at every fold start by the constructor-write that initialises the ISV bus). Task #4 from the dispatch is therefore a no-op — there are no fold-reset arms to remove. The `every_fold_and_soft_reset_entry_has_dispatch_arm` regression test continues to pass unchanged. (5) **4 launcher signatures updated in `gpu_dqn_trainer.rs`**: `launch_sp15_baseline_buyhold` / `launch_sp15_baseline_hold_only` / `launch_sp15_baseline_naive_momentum` / `launch_sp15_baseline_naive_reversion` each replace their `isv: CUdeviceptr` parameter with `out: CUdeviceptr` (per-window output buffer of `[mean, std, raw_sharpe]`). Production callers (Wave 3b) will allocate an `[n_windows × 3]` MappedF32Buffer per baseline and invoke each launcher per-window with the per-window slice (mirror of 1.1.b sharpe_per_bar per-window launch sequence). Doc-comments updated to reflect the contract change and explicitly mark the launchers as ORPHAN (transient 3a-only state). (6) **`action_decoding_helpers.cuh` (NEW)** — single-source-of-truth for the factored-action → direction → position mapping. Two `__device__ __forceinline__` helpers: `factored_action_to_dir_idx(action, b1, b2, b3)` mirrors `trade_physics.cuh::decode_direction_4b` exactly (returns DIR_SHORT=0 / DIR_HOLD=1 / DIR_LONG=2 / DIR_FLAT=3), and `factored_action_to_position(action, b1, b2, b3, prev_position)` lifts the dir-idx to a signed position scalar with explicit Hold-keeps-prev semantics. The header `#include "state_layout.cuh"` for the canonical DIR_* macros, so this and `trade_physics.cuh` share the source-of-truth for direction encoding. Existing consumers (`backtest_env_kernel.cu`, `experience_kernels.cu`) already use `decode_direction_4b` inline; the new derivation kernel is the only consumer of the new helper today. The header does NOT replace `trade_physics.cuh`'s primitives — it adds a higher-level helper on top so derivation kernels avoid pulling in the full Kelly-cap machinery. (7) **`position_history_derivation_kernel.cu` (NEW)** — post-loop derivation of per-bar `position_history` (-1/0/+1), `side_ind` (1.0 on position change), and `rt_ind` (1.0 on transition-to-flat from non-flat) from the recorded `actions_history_buf`. Single-block-per-window launch (`grid=[n_windows,1,1]`, `block=[1,1,1]`): the position state machine is sequential per window because `position(t)` depends on `position(t-1)`, so within-window parallelism is zero by construction; the across-windows axis carries all parallelism. No atomicAdd per `feedback_no_atomicadd` (the kernel is a pure scan, not a reduction). Sentinel handling: `actions_history_buf` is initialised to -1 by `experience_kernels.cu` (per the actions_history measurement-bug fix from a prior commit). Bars where `action == -1` are treated as "no position change" — keep prior position, emit zero side_ind / rt_ind. Downstream Wave 3b consumers gate on `window_lens[w]` to suppress the trailing-sentinel tail. (8) **`build.rs` cubin manifest** gains `position_history_derivation_kernel.cu`; the existing `baseline_kernels.cu` entry's comment block is updated to reflect the per-window output-buffer contract (was ISV-scalar pre-3a). Both compile clean on local RTX 3050 Ti (sm_86) via `CUDA_COMPUTE_CAP=86`. (9) **`launch_sp15_position_history_derivation` (NEW launcher)** in `gpu_dqn_trainer.rs`: free function (matches `launch_sp15_dd_state` / `launch_sp15_baseline_*` precedent) so unit/oracle tests can drive the kernel without the trainer struct. Caches no `CudaFunction` — loads the cubin per-call, same precedent as the baseline launchers (this is a once-per-eval call, not a hot-path producer). New static `SP15_POSITION_HISTORY_DERIVATION_CUBIN` follows the existing static-pub include-bytes pattern. Marked ORPHAN in doc-comment with explicit Wave 3b reference. (10) **5 oracle tests migrated / added in `tests/sp15_phase1_oracle_tests.rs`**: the 4 existing baseline tests (`baseline_buyhold_positive_on_drift`, `baseline_hold_only_emits_zero`, `baseline_momentum_reversion_symmetry`) now allocate a 3-f32 `MappedF32Buffer` per baseline (or two for the symmetry test) and assert against `out[2]` (raw_sharpe) instead of `isv[BASELINE_*_SHARPE_INDEX]`; the hold_only test additionally asserts `out[0] == 0` (hold_only's per-bar PnL stream is identically zero, not just zero-mean noise — exercises the new mean/std outputs that pure-sharpe coverage missed). One new test `position_history_derivation_walks_action_sequence`: 1 window × 8 bars, action sequence Short→Hold→Long→Hold→Flat→Long→Flat→Short (factored ints 0, 27, 54, 27, 81, 54, 81, 0 with b1=b2=b3=3); expected `position = [-1,-1,+1,+1,0,+1,0,-1]`, `side_ind = [1,0,1,0,1,1,1,1]`, `rt_ind = [0,0,0,0,1,0,1,0]`. All 5 tests pass on RTX 3050 Ti via `CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored` (8 oracle tests total, 0 failures). (11) **Transient state — 5 ORPHAN launchers awaiting Wave 3b**: `launch_sp15_baseline_buyhold` / `launch_sp15_baseline_hold_only` / `launch_sp15_baseline_naive_momentum` / `launch_sp15_baseline_naive_reversion` / `launch_sp15_position_history_derivation` are all reachable only from oracle tests in this commit. Wave 3b lands the production callers atomically: `GpuBacktestEvaluator::new` constructor signature gains the per-baseline output buffers; eval-time per-window launch sequence becomes `… → launch_sp15_baseline_*` (×4 per window) `→ launch_sp15_position_history_derivation` (once over all windows) `→ launch_sp15_cost_net_sharpe` (×n_windows, consuming the position_history outputs). The 3a/3b decomposition is documented here so reviewers can connect the kernel-side foundation to the host-side wire-up across the two commits. Touched: `crates/ml/src/cuda_pipeline/baseline_kernels.cu` (4 kernel signatures + helper template + extern-C entry-point bodies), `crates/ml/src/cuda_pipeline/sp15_isv_slots.rs` (4 const removals + 4 layout-test assertions removed + 4 retained-slot assertions added), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (4 launcher param-rename + doc-comment updates + layout-fingerprint string update + new SP15_POSITION_HISTORY_DERIVATION_CUBIN static + new launch_sp15_position_history_derivation), `crates/ml/src/cuda_pipeline/action_decoding_helpers.cuh` (NEW — 2 device helpers, 60 LOC), `crates/ml/src/cuda_pipeline/position_history_derivation_kernel.cu` (NEW — single extern-C kernel, ~70 LOC), `crates/ml/build.rs` (1 new cubin entry + baseline_kernels comment refresh), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (4 baseline-test bodies migrated to output-buffer assertion + 1 new derivation oracle test + import block trimmed). Hard rules: `feedback_no_partial_refactor` (4 ISV slots removed + 4 launcher signatures + 4 oracle tests + new derivation kernel + new shared header + audit doc all in this commit; the only "partial" element is the 5 ORPHAN launchers, explicitly bounded by the documented 3a/3b split), `feedback_no_legacy_aliases` (4 ISV slot constants gone — no compatibility-shim aliases), `feedback_no_atomicadd` (baseline kernels still tree-reduce; derivation kernel is a pure scan with one thread per window), `feedback_no_cpu_compute_strict` (the new derivation kernel runs entirely on the GPU; output buffers are mapped-pinned for the host-readback path used by Wave 3b's HEALTH_DIAG emission), `feedback_no_stubs` (the 5 launchers are real GPU launches against real data; the ORPHAN status is the absence of a production caller, not a stub-return), `pearl_no_host_branches_in_captured_graph` (the derivation kernel is launched OUTSIDE the experience-collection CUDA Graph capture region — Wave 3b will launch it once at end-of-fold from a per-fold scratch path mirroring 1.1.b's sharpe-per-bar end-of-fold launch), `pearl_first_observation_bootstrap` (every window starts with `prev_position = 0` per derivation kernel convention; matches `experience_env_step`'s portfolio-state init that every window starts flat). Verified: `SQLX_OFFLINE=true cargo check -p ml --features cuda` clean (18 unrelated pre-existing warnings); `cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored baseline` 3 tests green; `cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored position_history` 1 test green; `cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored unified_sharpe` 3 tests green; `cargo test -p ml --features cuda --lib` holds the 946 pass / 13 fail baseline (no regression). + SP15 Wave 2 — fused post-SP11 reward-axis composer (Option β layered architecture, 2026-05-06): closes the deferred-consumer gap left by Phase 3.1 (`r_quality_discipline_split_kernel`, commit `5d36f3238`), Phase 3.3 (`dd_penalty_kernel`), and Phase 3.5.2 (`dd_asymmetric_reward_kernel`) — three SP15 reward-axis kernels that all landed only as standalone scalar producers awaiting deferred consumer wiring per `feedback_no_partial_refactor.md`. Wave 2 chooses Option β (layered, composable post-modifier) over Option α (replace SP11 entirely): SP11 B1b stays canonical "trader-quality" composer that writes `out_rewards`; the SP15 reward-axis composition becomes a fused PER-(i,t) post-modifier read-modify-writing the same buffer in place. This preserves SP11's z-score mag-ratio contract (canary tests untouched) while landing all three deferred SP15 consumers atomically with zero parallel paths. (1) **Architectural decision — fused per-(i,t) post-modifier, NOT three sequential scalar launches**: the three deferred kernels (split composer + dd_penalty + dd_asymmetric_reward) each produced one f32 from ISV reads + a couple of arithmetic ops. Launching three single-thread/single-block kernels per training step plus three ↔-buffer round-trips between them was wasted launch overhead and added three CPU-sync/scratch-pointer plumbing chains. The fused `compute_sp15_final_reward_kernel` does the four-stage composition (α-blend → DD asymmetric reward → quadratic DD penalty → SP12 v3 cap) as inline `__device__` helpers in `sp15_reward_axis_helpers.cuh`, parallel over `[N*2*L]` slots — 3 launches → 1 launch, ~50× more parallelism per step. (2) **Per-step launch sequence in `gpu_experience_collector.rs::collect_experiences_gpu`** (after step 5 dd_state launch from Phase 1.3.b): `launch_experience_env_step` → `launch_sp15_alpha_split_producer` (per-step scalar producer of α from grad-norm ratio, OWNS warm-count increment per Q4) → `launch_sp15_final_reward` (the new fused composer that reads `out_rewards` + `r_discipline_per_sample` + `slot_completed_normally_per_sample` + ISV and writes back). Both new SP15 launches gated on `isv_signals_dev_ptr != 0 && sp15_alpha_warm_count_dev_ptr != 0` — test scaffolds may run env_step without ISV wiring, in which case both buffers receive their default values and the fused composer would be a no-op anyway. (3) **`compute_sp15_final_reward_kernel.cu` (new)**: `extern "C" __global__ void compute_sp15_final_reward_kernel(float* out_rewards, const float* r_discipline_out, const int* slot_completed_normally, float* isv, int N, int L)`. Per-(i,t) parallel over `[N*2*L]`; thread `idx` reads `out_rewards[idx]`, applies `sp15_alpha_blend` → `sp15_dd_asymmetric_reward` → `sp15_dd_penalty` → `sp15_apply_sp12_cap` from `sp15_reward_axis_helpers.cuh`, writes back. NaN/Inf guard skips the write (preserves SP11's value). (4) **`sp15_reward_axis_helpers.cuh` (new)**: 4 `__device__` helpers — `sp15_alpha_blend(r_in, r_disc, isv)` reads ALPHA_SPLIT_INDEX=417, returns `α × r_in + (1 − α) × r_disc`; `sp15_dd_asymmetric_reward(r_in, isv, boost_diag_out)` reads DD_PCT_INDEX=406 + DD_ASYMMETRY_LAMBDA_INDEX=430, returns `r × (1 + λ × dd_pct)` for r>0 else r unchanged, optionally writes R_GAIN_DD_BOOST_INDEX=431; `sp15_dd_penalty(r_in, isv)` reads DD_CURRENT_INDEX=401 + DD_THRESHOLD_INDEX=421 + LAMBDA_DD_INDEX=420, returns `r − λ_dd × max(0, dd − thr)²`; `sp15_apply_sp12_cap(r_in)` returns `fmaxf(REWARD_NEG_CAP, fminf(REWARD_POS_CAP, r))` where the macros come from `state_layout.cuh`. The diagnostic write to slot 431 happens ONLY from a single representative thread (idx==0) to avoid concurrent writes. (5) **Q1 resolution — SP12 caps are `state_layout.cuh` macros, NOT lifted to ISV**: per the user resolution, `REWARD_NEG_CAP=-10.0f` / `REWARD_POS_CAP=+5.0f` are spec'd constants per the SP12 v3 design (asymmetric loss aversion); lifting them to ISV would be a separate refactor outside Wave 2 scope. The fused kernel `#include "state_layout.cuh"` and uses the macros directly via the `sp15_apply_sp12_cap` helper, mirroring the existing `compute_asymmetric_capped_pnl` pattern in `trade_physics.cuh`. (6) **Q2 resolution — uniform shaping over on-policy AND CF slots**: `out_rewards[cf_off]` contains `cf_reward_weighted = w_cf × r_cf` (SP11 controller weight already applied). The fused kernel's per-(i,t) parallelism iterates over the full `[N*2*L]` range; both on-policy and CF slots get the same DD-aware shaping. CF slots map back to the per-(i,t) `r_discipline_out` and `slot_completed_normally` via `idx % (N*L)` — the DD context is per-step, not per-action (`dd_pct` doesn't depend on which action you took, it depends on what state you're in). The model learns "in this DD context, this counterfactual action would have been better/worse" with consistent SP15 shaping. (7) **Q3 resolution — `slot_completed_normally_per_sample[N*L]` flag preserves early-return semantics**: SP11's reward composer takes early-return paths at `experience_kernels.cu:2142` (data-end → `out_rewards = 0.0`) and `:2203` (blown-account → `out_rewards = -10.0`). Those terminal-episode sentinel rewards MUST NOT be SP15-shaped (the bounded clamp would still pass them through unchanged at +0/-10, but applying α-blend to them with a regret-EMA r_discipline would corrupt the sentinel). The new flag buffer defaults to 0 at every kernel entry (alongside the other `*_per_sample` defaults) and is set to 1 ONLY at the end of the normal reward-composition path (right before `out_rewards[out_off] = reward`). The fused kernel `if (slot_completed_normally[idx % (N*L)] == 0) return;` skips early-return slots. The flag is shared between on-policy and CF: if on-policy at (i,t) was an early-return, the CF slot at (i,t) is also skipped. (8) **Q4 resolution — `alpha_split_producer_kernel` OWNS the warm-count increment**: pre-Wave-2 the deleted scalar composer owned the increment. Moving the composer into the fused per-(i,t) kernel would over-tick by N×2×L per step. Per Q4 the warm-count increment moves to the `alpha_split_producer_kernel` — the per-step scalar producer launched once per step from `gpu_experience_collector.rs`, preserving the original "exactly one increment per step" semantics. The producer kernel was renamed from `r_quality_discipline_split_kernel.cu` to `alpha_split_producer_kernel.cu` and dropped the deleted scalar composer; it now runs the warm-count increment unconditionally and gates the formula write on `warm >= N_WARM=100`. (9) **`experience_env_step` signature change — 2 new output params**: `r_discipline_out: float*` ([N*L] f32, NULL-tolerant) gets `isv_signals_ptr[REGRET_EMA_INDEX=423]` written at the end of normal reward composition; `slot_completed_normally_out: int*` ([N*L] i32, NULL-tolerant) gets 0 default at entry and 1 at the same finalisation point. Both writes happen RIGHT BEFORE `out_rewards[out_off] = reward` so the fused composer's reads see consistent values. The CudaSlice/i32 buffers are allocated on `GpuExperienceCollector` mirroring the existing `cf_flip_per_sample` / `slot_live_per_sample` pattern (per-step ephemeral; defaulted at every kernel entry; no fold-boundary registry reset needed). i32 (not u8) for `slot_completed_normally` because no `MappedU8Buffer` exists in `mapped_pinned.rs` and the existing `cf_flip_per_sample` / `trade_close_per_sample` int-flag pattern is the established choice. (10) **3 deleted kernel files**: `r_quality_discipline_split_kernel.cu` (composer + producer; replaced by renamed `alpha_split_producer_kernel.cu` keeping ONLY the producer), `dd_penalty_kernel.cu` (replaced by `sp15_dd_penalty` `__device__` helper), `dd_asymmetric_reward_kernel.cu` (replaced by `sp15_dd_asymmetric_reward` helper). 3 deleted launchers in `gpu_dqn_trainer.rs`: `launch_sp15_r_quality_discipline_split` (composer, scalar variant), `launch_sp15_dd_penalty`, `launch_sp15_dd_asymmetric_reward`. 3 deleted CUBIN statics: `SP15_R_QUALITY_DISCIPLINE_SPLIT_CUBIN`, `SP15_DD_PENALTY_CUBIN`, `SP15_DD_ASYMMETRIC_REWARD_CUBIN`. 2 new CUBIN statics: `SP15_ALPHA_SPLIT_PRODUCER_CUBIN`, `SP15_FINAL_REWARD_CUBIN`. 1 new launcher: `launch_sp15_final_reward`. The retained `launch_sp15_alpha_split_producer` now loads the new `SP15_ALPHA_SPLIT_PRODUCER_CUBIN`. (11) **State reset registry — no new entries**: `r_discipline_per_sample` and `slot_completed_normally_per_sample` are per-step ephemeral (the env_step kernel defaults them at every entry), so cross-fold leakage is impossible without an explicit registry reset. The existing Phase 3.1 ISV slot 417/418/419 + `sp15_alpha_warm_count` registry entries already cover the α path; the existing 420/421/422 entries cover the DD-penalty path; the existing 430/431/432 entries cover the DD-asymmetric path. None of those need new entries — only their CONSUMERS changed (from deleted scalar kernels to inline helpers). (12) **6 new oracle tests** in `sp15_phase1_oracle_tests.rs::mod gpu`: `final_reward_alpha_blend_at_cold_start` (Stage 1: α=0.5 sentinel produces -0.5 from r=1, r_disc=-2), `final_reward_dd_penalty_above_threshold` (Stage 3: dd=0.10, thr=0.05, λ=10 → r=1.0 - 0.025 = 0.975), `final_reward_dd_asymmetric_gain` (Stage 2: r=4 + dd_pct=0.5 + λ=0.5 → r=5.0 boosted; R_GAIN_DD_BOOST=1.25 written by idx==0 thread), `final_reward_dd_asymmetric_loss` (Stage 2: r=-3 unchanged through gain branch), `final_reward_sp12_cap_clamps_both_directions` (Stage 4: r=+20 → +5, r=-20 → -10 via macros), `final_reward_skips_early_return_slots` (Q3: slot_completed=0 + r=-10 sentinel UNCHANGED; slot_completed=1 + r=2.5 SP15-shaped). All 6 drive `compute_sp15_final_reward_kernel` directly with hand-crafted `out_rewards` + `r_discipline_out` + ISV bundles. The 9 OLD oracle tests for the deleted scalar kernels (`r_split_uses_sentinel_alpha_at_cold_start`, `r_quality_subtracts_explicit_cost`, `dd_penalty_quadratic_above_threshold`, `dd_penalty_zero_below_threshold`, 3 `dd_asymmetric_reward_*`) are deleted with the kernels — the new fused-kernel tests cover the same behavioral surface end-to-end. The 2 retained `regret_*` tests (regret_signal_kernel) and 4 `cooldown_*` / 3 `plasticity_*` / 2 `dd_trajectory_*` tests are unchanged (those kernels still have orphan-launcher status pending their own consumer wiring). (13) **Phase 3.2 cost-on-trade-close subtraction NOT replicated in Wave 2**: pre-Wave-2 the deleted composer subtracted `cost_t × trade_close_indicator` from `r_quality` BEFORE the α-blend. The Wave 2 layered architecture re-uses SP11's `out_rewards[out_off]` as `r_in` — and SP11 already charges the model at evaluation-time costs via `cost_anneal_ptr` + the existing inventory / churn / spread-cost shaping. Re-inserting a cost subtraction at the SP15 layer would double-count. The `r_quality_subtracts_explicit_cost` oracle test is deleted accordingly. Touched: `crates/ml/src/cuda_pipeline/r_quality_discipline_split_kernel.cu` (DELETED), `crates/ml/src/cuda_pipeline/dd_penalty_kernel.cu` (DELETED), `crates/ml/src/cuda_pipeline/dd_asymmetric_reward_kernel.cu` (DELETED), `crates/ml/src/cuda_pipeline/alpha_split_producer_kernel.cu` (NEW — producer only, with warm-count increment moved here per Q4), `crates/ml/src/cuda_pipeline/sp15_reward_axis_helpers.cuh` (NEW — 4 `__device__` helpers + slot index `#define`s), `crates/ml/src/cuda_pipeline/compute_sp15_final_reward_kernel.cu` (NEW — fused per-(i,t) composer), `crates/ml/src/cuda_pipeline/experience_kernels.cu` (signature: 2 new output params; defaults at entry; writes at end-of-normal-path), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (3 deleted CUBIN statics, 3 deleted launchers, 2 new CUBIN statics, 1 new launcher; producer launcher renamed/refactored), `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (2 new buffer fields, 2 new alloc lines, 2 new env_step args, 1 new setter `set_sp15_alpha_warm_count_ptr`, +30-line Step 5b block invoking both SP15 launches), `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (1 new wire-up block setting the warm-count dev_ptr on the collector), `crates/ml/build.rs` (3 deleted cubin manifest entries replaced with REMOVED comment blocks; 2 new entries for `alpha_split_producer_kernel.cu` + `compute_sp15_final_reward_kernel.cu`), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (9 deleted scalar-kernel tests; 6 new fused-kernel oracle tests). Hard rules: `feedback_no_partial_refactor` (3 phases' deferred consumers + 3 deletions + 2 new files + 6 new tests + audit doc all in this commit; no parallel paths, no feature flags), `feedback_wire_everything_up` (closes 3 SP15 phase orphan launchers in one commit), `feedback_no_legacy_aliases` (3 scalar-kernel files + their CUBIN+launcher scaffolding all deleted in same commit as their replacement; no compatibility shim left), `feedback_no_atomicadd` (fused kernel is per-(i,t) parallel, pure scalar arithmetic; no reductions, no atomics), `pearl_audit_unboundedness_for_implicit_asymmetry` (gain-only DD multiplier + asymmetric NEG/POS caps preserve loss aversion through the layered pipeline), `pearl_symmetric_clamp_audit` (SP12 cap is bilateral via `fmaxf(NEG, fminf(POS, x))` even though the bounds are intentionally asymmetric per spec), `pearl_no_host_branches_in_captured_graph` (the new SP15 launches happen inside `collect_experiences_gpu::launch_timestep_loop` per-step, OUTSIDE the experience-fwd CUDA Graph capture region — same precedent as Phase 1.3.b's dd_state launch). Verified: NEEDS_VERIFICATION pending L40S CUDA build + oracle-test pass (RTX 3050 Ti smoke verification on next push); SQLX_OFFLINE cargo check + ml-lib non-CUDA suite expected to hold the 946 pass / 13 fail baseline (SP15 buffers are CUDA-feature-gated paths). SP15 Phase 3.5.b + 3.5.3.b — wire hold_floor (inline) + cooldown mask into experience_action_select (2026-05-06): closes the deferred-consumer gap left by Phase 3.5 (`hold_floor_kernel.cu` + `launch_sp15_hold_floor`, commit `5d36f3238`) and Phase 3.5.3 (`cooldown_kernel.cu` + `launch_sp15_cooldown`, commit `649128e73`). Both phases landed only the producer + state-machinery scaffolding; the actual action-selection consumers were deferred per `feedback_no_partial_refactor.md` to keep the kernel-by-kernel landings small. Wave 1.B bundles both deferred consumer migrations atomically. (1) **Architectural decision — hold_floor is now an INLINE `__device__` computation, NOT a separate kernel call**: the standalone `hold_floor_kernel.cu` produced one f32 (`α × σ(k × (entropy − ε₀))`) per launch into a scratch buffer for `experience_action_select` to read back. Launching a kernel to write a single scalar that the next kernel immediately consumes is unnecessary indirection. The bounded-sigmoid math is trivial (3 ISV reads + one `sigmoidf` + one multiply); inlining it eliminates a launch, eliminates a scratch buffer, and removes a parallel-path liability per `feedback_wire_everything_up.md` + `feedback_no_legacy_aliases.md`. The 4 ISV slots (`HOLD_FLOOR_ALPHA=426`, `HOLD_FLOOR_K=427`, `HOLD_FLOOR_EPS0=428`, `ENTROPY_DIST_REF=429`) and their state_reset_registry entries + dispatch arms remain — only the launch path is removed. (2) **Entropy source — per-step direction-policy entropy from `softmax(e_dir)`**: the kernel already computes `e_dir[d] = E[Q(d)]` (joint C51 expected Q) for every direction during Pass 1 of the Thompson direction selector. Pass 1.5 (new) takes the stable softmax of those 4 e_dir floats and computes Shannon entropy `−Σ p[d] log p[d]` directly in registers — no atomic, no shared memory, no cross-thread reduction (one thread per sample). This is the natural per-thread "policy uncertainty" signal: high entropy = the posterior is near-uniform (model uncertain about WHERE to trade) → hold_floor lifts Hold's q_eff so Hold becomes the principled uncertainty expression rather than the distributional default; low entropy = the posterior concentrates on one direction (model confident) → hold_floor ≈ 0 → Hold gets no boost. (3) **q_eff_dir scratch — preserves e_dir for downstream consumers**: hold_floor is added to a local `q_eff_dir[4]` (= e_dir + hold_floor at DIR_HOLD), NOT to `e_dir[]` directly. The conviction + q_gap consumers downstream (`out_conviction`, `out_magnitude_conviction`, `out_q_gaps`) keep reading raw `e_dir` so the policy-confidence signals reflect the model's expressed preference — adding hold_floor there would corrupt the Kelly-cap warmup floor with a meta-confidence "be cautious" mask that's not what the consumer is asking for. (4) **Cooldown mask — hard short-circuit, NOT temperature-blended**: when `ISV[COOLDOWN_BARS_REMAINING=435] > 0`, Pass 2 is skipped entirely and `dir_idx = DIR_HOLD` is set directly. The alternative (q_eff_dir[non-Hold] = -INFINITY pre-blend) breaks under the Thompson temperature blend `q_eff = q_eff_dir + temp·(q_sample - q_eff_dir)`: at -INFINITY the algebra produces NaN; at -1e30 the formula reduces to `q_eff = q_sample` at temp=1 (the default, pure Thompson), defeating the mask. Hard short-circuit is the only correct semantics — and matches the spec's "force Hold for M bars" language exactly. The Philox draws Pass 2 would have made for direction sampling are not advanced on the cooldown path; this shifts the mag/ord/urg streams' state by a constant amount per cooldown step but reproducibility-at-fixed-seed is preserved because the cooldown gate is a deterministic function of the ISV state at this thread's bar. (5) **Deletion of `hold_floor_kernel.cu` + `launch_sp15_hold_floor` + `SP15_HOLD_FLOOR_CUBIN` + cubin manifest entry** per `feedback_no_legacy_aliases`: scaffolding for a deferred consumer becomes dead code once the consumer lands inline. The build.rs entry is replaced with a comment block documenting the move; the launcher is replaced with a stub doc comment in gpu_dqn_trainer.rs explaining the inline migration; the standalone `hold_floor_bounded_sigmoid` oracle test is removed (the math is now exercised end-to-end through the new `action_select_applies_hold_floor_inline` test). (6) **3 new oracle tests** in `crates/ml/tests/sp15_phase1_oracle_tests.rs::mod gpu`: `action_select_applies_hold_floor_inline` (near-uniform e_dir → high entropy ≈ ln(4) → hold_floor ≈ 0.49 lifts Hold's q_eff past Long's; argmax = Hold), `action_select_forces_hold_during_cooldown` (Long-preferred e_dir + cooldown=5.0 → mask short-circuit forces Hold regardless of e_dir), `action_select_no_force_hold_when_cooldown_zero` (Long-preferred e_dir + cooldown=0 + low-entropy → standard argmax wins, picked dir = Long). All 3 tests load the production `experience_kernels.cubin`, set up minimal C51 logits with peaked atoms (peak_logit=50 → softmax ~delta on closest atom, `q_sample` ≈ `e_dir`), and call `experience_action_select` end-to-end via cudarc. Tests pass on RTX 3050 Ti. (7) **Phase 3.5 + Phase 3.5.3 deferred consumers eliminated** per `feedback_wire_everything_up.md`. The cooldown_kernel itself (which maintains the consecutive_losses streak counter and decrements COOLDOWN_BARS_REMAINING per bar) is unchanged and still wired in production — only the consumer that reads its output is now connected to action-selection logic. Touched: `crates/ml/src/cuda_pipeline/experience_kernels.cu` (inline hold_floor + cooldown short-circuit added at the start of Pass 2 in `experience_action_select`; `e_dir` preserved for downstream consumers; ~80 LOC added), `crates/ml/src/cuda_pipeline/hold_floor_kernel.cu` (DELETED), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (`launch_sp15_hold_floor` + `SP15_HOLD_FLOOR_CUBIN` removed; replaced by a non-doc comment block referencing the inline migration), `crates/ml/build.rs` (`"hold_floor_kernel.cu"` cubin manifest entry removed; replaced with a comment block explaining the move), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (removed `hold_floor_bounded_sigmoid` test; added 3 new `action_select_*` tests + `ProdActionSelectFixture` re-implementation in the integration-test module since the production fixture is `pub(crate)` and not reachable). Hard rules: `feedback_no_partial_refactor` (3.5.b + 3.5.3.b consumers + hold_floor_kernel deletion + 3 oracle tests + audit doc all in this commit; no parallel paths, no feature flags), `feedback_no_cpu_compute_strict` (hold_floor + cooldown mask are GPU-resident inside `experience_action_select` — no CPU forwards, no CPU reductions), `feedback_no_stubs` (the inline computation is the actual policy bias signal end-to-end, not a return-zero placeholder), `feedback_wire_everything_up` + `feedback_no_legacy_aliases` (hold_floor_kernel + launcher deleted because the consumer wiring landed inline; ISV slots + state_reset_registry entries + dispatch arms remain because they're still consumed by the inline path), `pearl_no_host_branches_in_captured_graph` (changes are inside `experience_action_select` which is already inside the existing CUDA Graph capture; no new launches added or removed from the captured stream — only one fewer kernel launch upstream of action_select on the cooldown side because the standalone hold_floor_kernel is gone, but it was always launched OUTSIDE the experience-collection capture region). Verified: `cargo check -p ml --features cuda` clean (18 unrelated pre-existing warnings); `cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored` runs 6 tests green (3 cooldown + 3 new action_select); `cargo test -p ml --features cuda --lib` holds 946 pass / 13 fail = baseline (no regression).