gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

feat(sp15-wave3a): kernel-side foundation — baseline output buffers + position_history derivation

Browse Source 
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-06 23:03:30 +02:00
parent 334b496647
commit e968f4ded9
8 changed files with 500 additions and 137 deletions
Download Patch File Download Diff File Expand all files Collapse all files

28
crates/ml/build.rs
View File

@@ -854,7 +854,7 @@ fn main() {
// DD_PCT (406). The 1e-4 calmar floor eliminates the saturation- // DD_PCT (406). The 1e-4 calmar floor eliminates the saturation-
// at-100 artifact previously seen in the train-dd4xl HEALTH_DIAG. // at-100 artifact previously seen in the train-dd4xl HEALTH_DIAG.
"dd_state_kernel.cu", "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 // counterfactual baselines per spec §6.4. Single source file with
// 4 `extern "C" __global__` symbols (baseline_buyhold_kernel, // 4 `extern "C" __global__` symbols (baseline_buyhold_kernel,
// baseline_hold_only_kernel, baseline_naive_momentum_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 // BLOCK=256 with a templated 2-pass shared-memory tree-reduce
// (no atomicAdd). Reads price/half_spread/ofi arrays + commission // (no atomicAdd). Reads price/half_spread/ofi arrays + commission
// scalar; computes constant-policy or last-bar-return-based // scalar; computes constant-policy or last-bar-return-based
// actions; writes sharpe_net to its own ISV slot (409 buyhold, // actions; writes per-window `[mean, std, raw_sharpe]` to a
// 410 hold_only, 412 naive_momentum, 416 naive_reversion). // caller-provided output buffer (Wave 3a contract change — was
// Trunk-shared baselines (random_dir_kelly slot 411, aux_only // ISV-scalar slots 409/410/412/416 pre-Wave-3a). Trunk-shared
// slot 413, mag_quarter_fixed slot 414, trail_only slot 415) are // baselines (random_dir_kelly slot 411, aux_only slot 413,
// out of scope for this commit — they need partial-policy- // mag_quarter_fixed slot 414, trail_only slot 415) are out of
// forward access from the main eval pass and are deferred to // scope for this commit — they need partial-policy-forward
// Task 1.4.b. // access from the main eval pass and are deferred to Task 1.4.b.
"baseline_kernels.cu", "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 // SP15 Phase 1.5 (2026-05-06): dd_pct foundational state input
// concat. LAYOUT FINGERPRINT BREAK — pre-SP15 checkpoints WILL // concat. LAYOUT FINGERPRINT BREAK — pre-SP15 checkpoints WILL
// NOT LOAD after this commit. Greenfield OK per spec Q1. Per // NOT LOAD after this commit. Greenfield OK per spec Q1. Per

63
crates/ml/src/cuda_pipeline/action_decoding_helpers.cuh Normal file
View File

@@ -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

84
crates/ml/src/cuda_pipeline/baseline_kernels.cu
View File

@@ -1,12 +1,23 @@
// crates/ml/src/cuda_pipeline/baseline_kernels.cu // 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 // Per spec §6.4. Pure-CUDA kernels (no trunk-share). Each kernel reads
// the price/spread/ofi arrays + a commission scalar, computes the // the price/spread/ofi arrays + a commission scalar, computes the
// baseline's per-bar PnL stream from a constant or last-bar-return-based // 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 // action policy, then runs a 2-pass shared-memory tree-reduce mean/std
// (no atomicAdd per `feedback_no_atomicadd`) and writes sharpe_net to // (no atomicAdd per `feedback_no_atomicadd`) and writes per-window
// its allocated ISV slot. // `[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 // Single source file — multiple `extern "C" __global__` symbols share
// one cubin per the build.rs 1:1-source-to-cubin convention. The four // 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- // different `get_function()` symbols, mirroring the multi-kernel-per-
// file pattern from SP11 (`novelty_simhash_kernel.cu`). // 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, // 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 // mag_quarter_fixed slot 414, trail_only slot 415) are out of scope
// for this commit — they need partial-policy-forward access from the // 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 // the model-under-test (both pay full RT commission + entry/exit
// spread, neither pays OFI impact for this baseline class). // 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 // 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 // the per-bar PnL stream produced by `ActionFn`. Writes `[mean, std,
// 0 only — caller stores it to the appropriate ISV slot. // raw_sharpe]` to `out[0..3]` from thread 0 only.
// //
// Template parameter `ActionFn` is a CUDA functor: // Template parameter `ActionFn` is a CUDA functor:
// __device__ int operator()(const float* prices, int i, int n) const; // __device__ int operator()(const float* prices, int i, int n) const;
// returns +1 (long), 0 (flat/hold), -1 (short) for bar `i`. // returns +1 (long), 0 (flat/hold), -1 (short) for bar `i`.
template <typename ActionFn> template <typename ActionFn>
__device__ float compute_baseline_sharpe( __device__ void compute_baseline_sharpe(
const float* __restrict__ prices, const float* __restrict__ prices,
const float* __restrict__ half_spread, const float* __restrict__ half_spread,
const float* __restrict__ /* ofi */, // unused for this baseline class (lambda=0) const float* __restrict__ /* ofi */, // unused for this baseline class (lambda=0)
int n, int n,
float commission_per_rt, float commission_per_rt,
ActionFn action_fn ActionFn action_fn,
float* out /* [mean, std, raw_sharpe] */
) { ) {
const int tid = (int)threadIdx.x; const int tid = (int)threadIdx.x;
const int BLOCK = (int)blockDim.x; const int BLOCK = (int)blockDim.x;
@@ -117,9 +118,12 @@ __device__ float compute_baseline_sharpe(
if (tid == 0) { if (tid == 0) {
const float var = (n > 1) ? reduce_buf[0] / (float)(n - 1) : 0.0f; const float var = (n > 1) ? reduce_buf[0] / (float)(n - 1) : 0.0f;
const float std = sqrtf(fmaxf(var, 1e-12f)); 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. // CUDA functor structs — one per baseline policy.
@@ -158,16 +162,12 @@ extern "C" __global__ void baseline_buyhold_kernel(
const float* __restrict__ ofi, const float* __restrict__ ofi,
int n, int n,
float commission_per_rt, float commission_per_rt,
float* __restrict__ isv float* __restrict__ out /* [mean, std, raw_sharpe] */
) { ) {
if (blockIdx.x != 0) return; if (blockIdx.x != 0) return;
const float sharpe = compute_baseline_sharpe( compute_baseline_sharpe(
prices, half_spread, ofi, n, commission_per_rt, BuyholdAction() 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( extern "C" __global__ void baseline_hold_only_kernel(
@@ -176,16 +176,12 @@ extern "C" __global__ void baseline_hold_only_kernel(
const float* __restrict__ ofi, const float* __restrict__ ofi,
int n, int n,
float commission_per_rt, float commission_per_rt,
float* __restrict__ isv float* __restrict__ out /* [mean, std, raw_sharpe] */
) { ) {
if (blockIdx.x != 0) return; if (blockIdx.x != 0) return;
const float sharpe = compute_baseline_sharpe( compute_baseline_sharpe(
prices, half_spread, ofi, n, commission_per_rt, HoldOnlyAction() 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( extern "C" __global__ void baseline_naive_momentum_kernel(
@@ -194,16 +190,12 @@ extern "C" __global__ void baseline_naive_momentum_kernel(
const float* __restrict__ ofi, const float* __restrict__ ofi,
int n, int n,
float commission_per_rt, float commission_per_rt,
float* __restrict__ isv float* __restrict__ out /* [mean, std, raw_sharpe] */
) { ) {
if (blockIdx.x != 0) return; if (blockIdx.x != 0) return;
const float sharpe = compute_baseline_sharpe( compute_baseline_sharpe(
prices, half_spread, ofi, n, commission_per_rt, MomentumAction() 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( extern "C" __global__ void baseline_naive_reversion_kernel(
@@ -212,14 +204,10 @@ extern "C" __global__ void baseline_naive_reversion_kernel(
const float* __restrict__ ofi, const float* __restrict__ ofi,
int n, int n,
float commission_per_rt, float commission_per_rt,
float* __restrict__ isv float* __restrict__ out /* [mean, std, raw_sharpe] */
) { ) {
if (blockIdx.x != 0) return; if (blockIdx.x != 0) return;
const float sharpe = compute_baseline_sharpe( compute_baseline_sharpe(
prices, half_spread, ofi, n, commission_per_rt, ReversionAction() prices, half_spread, ofi, n, commission_per_rt, ReversionAction(), out
); );
if (threadIdx.x == 0) {
isv[REVERSION_SHARPE_SLOT] = sharpe;
__threadfence_system();
}
} }

165
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
View File

@@ -993,35 +993,48 @@ pub fn launch_sp15_dd_pct_concat(
Ok(()) 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`, /// policy counterfactual baselines (`baseline_buyhold_kernel`,
/// `baseline_hold_only_kernel`, `baseline_naive_momentum_kernel`, /// `baseline_hold_only_kernel`, `baseline_naive_momentum_kernel`,
/// `baseline_naive_reversion_kernel`). Per spec §6.4. Each baseline /// `baseline_naive_reversion_kernel`). Per spec §6.4. Each baseline
/// computes its sharpe_net by running a constant-policy or last-bar- /// computes its per-window `[mean, std, raw_sharpe]` by running a
/// return-based action stream through a 2-pass shared-memory tree-reduce /// constant-policy or last-bar-return-based action stream through a
/// (single-block BLOCK=256, no atomicAdd) and writing to its own ISV /// 2-pass shared-memory tree-reduce (single-block BLOCK=256, no atomicAdd)
/// slot. The four launchers below all load this single cubin and resolve /// and writing to a caller-provided output buffer. Wave 3a contract
/// different `get_function()` symbols — multi-kernel-per-file pattern /// change: ISV-scalar writes to slots 409/410/412/416 are removed entirely
/// per `novelty_simhash_kernel.cu` precedent. /// — 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, /// 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 /// 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 /// 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 /// eval pass and are deferred to Task 1.4.b. Their slots stay at sentinel
/// 0.0 in the meantime. /// 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] = pub static SP15_BASELINE_KERNELS_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/baseline_kernels.cubin")); 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` /// counterfactual baseline. Free function (matches `launch_sp15_dd_state`
/// precedent) so unit/oracle tests can drive the kernel directly without /// precedent) so unit/oracle tests can drive the kernel directly without
/// the trainer struct; production callers (eval-pass baseline launches) /// the trainer struct; production callers (eval-pass baseline launches
/// invoke the same launcher. /// in Wave 3b's `GpuBacktestEvaluator::new`) invoke the same launcher.
/// ///
/// Policy: position=+1 every bar (no exits). Writes sharpe_net to /// Policy: position=+1 every bar (no exits). Writes per-window
/// `ISV[BASELINE_BUYHOLD_SHARPE_INDEX=409]`. `prices`, `half_spread`, /// `[mean, std, raw_sharpe]` to `out[0..3]`. `prices`, `half_spread`,
/// `ofi`, `isv` are device f32 pointers; `n` is bar count. `isv` MUST /// `ofi`, `out` are device f32 pointers; `n` is bar count. `out` MUST
/// be the ISV bus (≥443 f32 slots). /// 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( pub fn launch_sp15_baseline_buyhold(
stream: &Arc<CudaStream>, stream: &Arc<CudaStream>,
prices: cudarc::driver::sys::CUdeviceptr, prices: cudarc::driver::sys::CUdeviceptr,
@@ -1029,7 +1042,7 @@ pub fn launch_sp15_baseline_buyhold(
ofi: cudarc::driver::sys::CUdeviceptr, ofi: cudarc::driver::sys::CUdeviceptr,
n: i32, n: i32,
commission_per_rt: f32, commission_per_rt: f32,
isv: cudarc::driver::sys::CUdeviceptr, out: cudarc::driver::sys::CUdeviceptr,
) -> Result<(), MLError> { ) -> Result<(), MLError> {
let module = stream let module = stream
.context() .context()
@@ -1050,7 +1063,7 @@ pub fn launch_sp15_baseline_buyhold(
.arg(&ofi) .arg(&ofi)
.arg(&n) .arg(&n)
.arg(&commission_per_rt) .arg(&commission_per_rt)
.arg(&isv) .arg(&out)
.launch(LaunchConfig { .launch(LaunchConfig {
grid_dim: (1, 1, 1), grid_dim: (1, 1, 1),
block_dim: (256, 1, 1), block_dim: (256, 1, 1),
@@ -1063,12 +1076,12 @@ pub fn launch_sp15_baseline_buyhold(
Ok(()) 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 /// 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 /// 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 kernel's `(std > 0) ? mean / std : 0` guard emits sharpe=0, which
/// the spec-mandated sentinel for an undefined sharpe. /// is the spec-mandated sentinel for an undefined sharpe.
pub fn launch_sp15_baseline_hold_only( pub fn launch_sp15_baseline_hold_only(
stream: &Arc<CudaStream>, stream: &Arc<CudaStream>,
prices: cudarc::driver::sys::CUdeviceptr, prices: cudarc::driver::sys::CUdeviceptr,
@@ -1076,7 +1089,7 @@ pub fn launch_sp15_baseline_hold_only(
ofi: cudarc::driver::sys::CUdeviceptr, ofi: cudarc::driver::sys::CUdeviceptr,
n: i32, n: i32,
commission_per_rt: f32, commission_per_rt: f32,
isv: cudarc::driver::sys::CUdeviceptr, out: cudarc::driver::sys::CUdeviceptr,
) -> Result<(), MLError> { ) -> Result<(), MLError> {
let module = stream let module = stream
.context() .context()
@@ -1097,7 +1110,7 @@ pub fn launch_sp15_baseline_hold_only(
.arg(&ofi) .arg(&ofi)
.arg(&n) .arg(&n)
.arg(&commission_per_rt) .arg(&commission_per_rt)
.arg(&isv) .arg(&out)
.launch(LaunchConfig { .launch(LaunchConfig {
grid_dim: (1, 1, 1), grid_dim: (1, 1, 1),
block_dim: (256, 1, 1), block_dim: (256, 1, 1),
@@ -1110,10 +1123,10 @@ pub fn launch_sp15_baseline_hold_only(
Ok(()) 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] − /// counterfactual baseline. Policy: `direction = sign(prices[i] −
/// prices[i-1])` with |position|=1. Writes sharpe_net to /// prices[i-1])` with |position|=1. Writes per-window `[mean, std,
/// `ISV[BASELINE_NAIVE_MOMENTUM_SHARPE_INDEX=412]`. /// raw_sharpe]` to `out[0..3]`.
pub fn launch_sp15_baseline_naive_momentum( pub fn launch_sp15_baseline_naive_momentum(
stream: &Arc<CudaStream>, stream: &Arc<CudaStream>,
prices: cudarc::driver::sys::CUdeviceptr, prices: cudarc::driver::sys::CUdeviceptr,
@@ -1121,7 +1134,7 @@ pub fn launch_sp15_baseline_naive_momentum(
ofi: cudarc::driver::sys::CUdeviceptr, ofi: cudarc::driver::sys::CUdeviceptr,
n: i32, n: i32,
commission_per_rt: f32, commission_per_rt: f32,
isv: cudarc::driver::sys::CUdeviceptr, out: cudarc::driver::sys::CUdeviceptr,
) -> Result<(), MLError> { ) -> Result<(), MLError> {
let module = stream let module = stream
.context() .context()
@@ -1142,7 +1155,7 @@ pub fn launch_sp15_baseline_naive_momentum(
.arg(&ofi) .arg(&ofi)
.arg(&n) .arg(&n)
.arg(&commission_per_rt) .arg(&commission_per_rt)
.arg(&isv) .arg(&out)
.launch(LaunchConfig { .launch(LaunchConfig {
grid_dim: (1, 1, 1), grid_dim: (1, 1, 1),
block_dim: (256, 1, 1), block_dim: (256, 1, 1),
@@ -1155,10 +1168,10 @@ pub fn launch_sp15_baseline_naive_momentum(
Ok(()) 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] − /// counterfactual baseline. Policy: `direction = -sign(prices[i] −
/// prices[i-1])` with |position|=1 (mirror of naive_momentum). Writes /// 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( pub fn launch_sp15_baseline_naive_reversion(
stream: &Arc<CudaStream>, stream: &Arc<CudaStream>,
prices: cudarc::driver::sys::CUdeviceptr, prices: cudarc::driver::sys::CUdeviceptr,
@@ -1166,7 +1179,7 @@ pub fn launch_sp15_baseline_naive_reversion(
ofi: cudarc::driver::sys::CUdeviceptr, ofi: cudarc::driver::sys::CUdeviceptr,
n: i32, n: i32,
commission_per_rt: f32, commission_per_rt: f32,
isv: cudarc::driver::sys::CUdeviceptr, out: cudarc::driver::sys::CUdeviceptr,
) -> Result<(), MLError> { ) -> Result<(), MLError> {
let module = stream let module = stream
.context() .context()
@@ -1187,7 +1200,7 @@ pub fn launch_sp15_baseline_naive_reversion(
.arg(&ofi) .arg(&ofi)
.arg(&n) .arg(&n)
.arg(&commission_per_rt) .arg(&commission_per_rt)
.arg(&isv) .arg(&out)
.launch(LaunchConfig { .launch(LaunchConfig {
grid_dim: (1, 1, 1), grid_dim: (1, 1, 1),
block_dim: (256, 1, 1), block_dim: (256, 1, 1),
@@ -1200,6 +1213,91 @@ pub fn launch_sp15_baseline_naive_reversion(
Ok(()) 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<CudaStream>,
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 α /// SP15 Wave 2 (2026-05-06): cubin for the per-step ISV-driven α
/// producer (`alpha_split_producer_kernel`). Replaces the pre-Wave-2 /// producer (`alpha_split_producer_kernel`). Replaces the pre-Wave-2
/// `SP15_R_QUALITY_DISCIPLINE_SPLIT_CUBIN` which housed both the /// `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;\ 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;\ 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;\ 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_RANDOM_DIR_KELLY_SHARPE=411;BASELINE_NAIVE_MOMENTUM_SHARPE=412;\
BASELINE_AUX_ONLY_SHARPE=413;BASELINE_MAG_QUARTER_FIXED_SHARPE=414;\ 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;\ ALPHA_SPLIT=417;GRAD_NORM_QUALITY=418;GRAD_NORM_DISCIPLINE=419;\
LAMBDA_DD=420;DD_THRESHOLD=421;DD_PENALTY_GRAD_NORM=422;\ LAMBDA_DD=420;DD_THRESHOLD=421;DD_PENALTY_GRAD_NORM=422;\
REGRET_EMA=423;LAMBDA_REGRET=424;REGRET_GRAD_NORM=425;\ REGRET_EMA=423;LAMBDA_REGRET=424;REGRET_GRAD_NORM=425;\

76
crates/ml/src/cuda_pipeline/position_history_derivation_kernel.cu Normal file
View File

@@ -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();
}

35
crates/ml/src/cuda_pipeline/sp15_isv_slots.rs
View File

@@ -24,15 +24,29 @@ pub const DD_PCT_INDEX: usize = 406;
pub const OFI_IMPACT_LAMBDA_INDEX: usize = 407; pub const OFI_IMPACT_LAMBDA_INDEX: usize = 407;
pub const COST_PER_BAR_AVG_INDEX: usize = 408; pub const COST_PER_BAR_AVG_INDEX: usize = 408;
// === Phase 1.4: 8 counterfactual baselines (8 slots) === // === Phase 1.4 + Wave 3a: 8 counterfactual baseline slots (4 retired, 4 reserved) ===
pub const BASELINE_BUYHOLD_SHARPE_INDEX: usize = 409; //
pub const BASELINE_HOLD_ONLY_SHARPE_INDEX: usize = 410; // 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_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_AUX_ONLY_SHARPE_INDEX: usize = 413;
pub const BASELINE_MAG_QUARTER_FIXED_SHARPE_INDEX: usize = 414; pub const BASELINE_MAG_QUARTER_FIXED_SHARPE_INDEX: usize = 414;
pub const BASELINE_TRAIL_ONLY_SHARPE_INDEX: usize = 415; 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) === // === Phase 3.1: r_quality + r_discipline split (3 slots) ===
pub const ALPHA_SPLIT_INDEX: usize = 417; pub const ALPHA_SPLIT_INDEX: usize = 417;
@@ -108,8 +122,15 @@ mod tests {
assert_eq!(DD_CURRENT_INDEX, 401); assert_eq!(DD_CURRENT_INDEX, 401);
assert_eq!(DD_PCT_INDEX, 406); assert_eq!(DD_PCT_INDEX, 406);
assert_eq!(OFI_IMPACT_LAMBDA_INDEX, 407); assert_eq!(OFI_IMPACT_LAMBDA_INDEX, 407);
assert_eq!(BASELINE_BUYHOLD_SHARPE_INDEX, 409); // Wave 3a (2026-05-06): BASELINE_BUYHOLD_SHARPE_INDEX (409) +
assert_eq!(BASELINE_NAIVE_REVERSION_SHARPE_INDEX, 416); // 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!(ALPHA_SPLIT_INDEX, 417);
assert_eq!(LAMBDA_DD_INDEX, 420); assert_eq!(LAMBDA_DD_INDEX, 420);
assert_eq!(REGRET_EMA_INDEX, 423); assert_eq!(REGRET_EMA_INDEX, 423);

184
crates/ml/tests/sp15_phase1_oracle_tests.rs
View File

@@ -13,15 +13,13 @@ mod gpu {
use ml::cuda_pipeline::gpu_dqn_trainer::{ use ml::cuda_pipeline::gpu_dqn_trainer::{
launch_sp15_baseline_buyhold, launch_sp15_baseline_hold_only, launch_sp15_baseline_buyhold, launch_sp15_baseline_hold_only,
launch_sp15_baseline_naive_momentum, launch_sp15_baseline_naive_reversion, 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::{ use ml::cuda_pipeline::sp15_isv_slots::{
ALPHA_SPLIT_INDEX, BASELINE_BUYHOLD_SHARPE_INDEX, BASELINE_HOLD_ONLY_SHARPE_INDEX, ALPHA_SPLIT_INDEX, COST_PER_BAR_AVG_INDEX, DD_CURRENT_INDEX, DD_MAX_INDEX, DD_PCT_INDEX,
BASELINE_NAIVE_MOMENTUM_SHARPE_INDEX, BASELINE_NAIVE_REVERSION_SHARPE_INDEX, DD_RECOVERY_BARS_INDEX, OFI_IMPACT_LAMBDA_INDEX, SP15_SLOT_END,
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,
}; };
/// Upper bound for any ISV index touched by SP15 oracle tests. Matches /// 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 /// rules out a sharpe-near-zero degenerate baseline; upper 1.0
/// catches accidental zero-cost or reduction bugs that would /// catches accidental zero-cost or reduction bugs that would
/// inflate the sharpe. /// 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] #[test]
#[ignore = "requires GPU"] #[ignore = "requires GPU"]
fn baseline_buyhold_positive_on_drift() { fn baseline_buyhold_positive_on_drift() {
@@ -365,11 +367,9 @@ mod gpu {
.expect("alloc MappedF32Buffer for ofi"); .expect("alloc MappedF32Buffer for ofi");
ofi_buf.write_from_slice(&ofi); ofi_buf.write_from_slice(&ofi);
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) } let out_buf = unsafe { MappedF32Buffer::new(3) }
.expect("alloc MappedF32Buffer for isv bus"); .expect("alloc MappedF32Buffer for baseline output [mean, std, raw_sharpe]");
let mut isv_init = vec![0.0f32; ISV_LEN]; out_buf.write_from_slice(&[0.0f32; 3]);
isv_init[OFI_IMPACT_LAMBDA_INDEX] = 0.0;
isv_buf.write_from_slice(&isv_init);
launch_sp15_baseline_buyhold( launch_sp15_baseline_buyhold(
&stream, &stream,
@@ -378,24 +378,24 @@ mod gpu {
ofi_buf.dev_ptr, ofi_buf.dev_ptr,
n as i32, n as i32,
/* commission_per_rt = */ 0.0, /* commission_per_rt = */ 0.0,
isv_buf.dev_ptr, out_buf.dev_ptr,
) )
.expect("launch baseline_buyhold_kernel"); .expect("launch baseline_buyhold_kernel");
stream stream
.synchronize() .synchronize()
.expect("synchronize after baseline_buyhold_kernel launch"); .expect("synchronize after baseline_buyhold_kernel launch");
let isv = isv_buf.read_all(); let out = out_buf.read_all();
let sharpe = isv[BASELINE_BUYHOLD_SHARPE_INDEX]; let sharpe = out[2];
assert!( assert!(
sharpe > 0.2, sharpe > 0.2,
"buyhold sharpe = {}, expected > 0.2 on +drift series", "buyhold sharpe = {}, expected > 0.2 on +drift series (mean={}, std={})",
sharpe sharpe, out[0], out[1]
); );
assert!( assert!(
sharpe < 1.0, sharpe < 1.0,
"buyhold sharpe = {}, expected < 1.0 (sanity bound)", "buyhold sharpe = {}, expected < 1.0 (sanity bound; mean={}, std={})",
sharpe sharpe, out[0], out[1]
); );
} }
@@ -405,6 +405,9 @@ mod gpu {
/// zero on every bar. The kernel's reduction yields mean=0 std=0 /// zero on every bar. The kernel's reduction yields mean=0 std=0
/// and the `(std > 0) ? mean / std : 0` guard returns 0, the /// and the `(std > 0) ? mean / std : 0` guard returns 0, the
/// spec-mandated sentinel for an undefined sharpe (no trades). /// 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] #[test]
#[ignore = "requires GPU"] #[ignore = "requires GPU"]
fn baseline_hold_only_emits_zero() { fn baseline_hold_only_emits_zero() {
@@ -424,9 +427,9 @@ mod gpu {
let ofi_buf = unsafe { MappedF32Buffer::new(n) } let ofi_buf = unsafe { MappedF32Buffer::new(n) }
.expect("alloc MappedF32Buffer for ofi"); .expect("alloc MappedF32Buffer for ofi");
ofi_buf.write_from_slice(&ofi); ofi_buf.write_from_slice(&ofi);
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) } let out_buf = unsafe { MappedF32Buffer::new(3) }
.expect("alloc MappedF32Buffer for isv bus"); .expect("alloc MappedF32Buffer for baseline output");
isv_buf.write_from_slice(&vec![0.0f32; ISV_LEN]); out_buf.write_from_slice(&[0.0f32; 3]);
launch_sp15_baseline_hold_only( launch_sp15_baseline_hold_only(
&stream, &stream,
@@ -435,19 +438,26 @@ mod gpu {
ofi_buf.dev_ptr, ofi_buf.dev_ptr,
n as i32, n as i32,
/* commission_per_rt = */ 0.0, /* commission_per_rt = */ 0.0,
isv_buf.dev_ptr, out_buf.dev_ptr,
) )
.expect("launch baseline_hold_only_kernel"); .expect("launch baseline_hold_only_kernel");
stream stream
.synchronize() .synchronize()
.expect("synchronize after baseline_hold_only_kernel launch"); .expect("synchronize after baseline_hold_only_kernel launch");
let isv = isv_buf.read_all(); let out = out_buf.read_all();
let sharpe = isv[BASELINE_HOLD_ONLY_SHARPE_INDEX]; let sharpe = out[2];
assert!( assert!(
sharpe.abs() < 1e-5, sharpe.abs() < 1e-5,
"hold_only sharpe = {}, expected ~0 (no trades)", "hold_only sharpe = {}, expected ~0 (no trades; mean={}, std={})",
sharpe 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 /// sharpe values should sum to ~0. The series is deliberately
/// mean-reverting (reversion outperforms) but the symmetry test is /// mean-reverting (reversion outperforms) but the symmetry test is
/// invariant to which side wins. /// 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] #[test]
#[ignore = "requires GPU"] #[ignore = "requires GPU"]
fn baseline_momentum_reversion_symmetry() { fn baseline_momentum_reversion_symmetry() {
@@ -486,9 +499,12 @@ mod gpu {
.expect("alloc MappedF32Buffer for ofi"); .expect("alloc MappedF32Buffer for ofi");
ofi_buf.write_from_slice(&ofi); ofi_buf.write_from_slice(&ofi);
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) } let momentum_out = unsafe { MappedF32Buffer::new(3) }
.expect("alloc MappedF32Buffer for isv bus"); .expect("alloc MappedF32Buffer for momentum output");
isv_buf.write_from_slice(&vec![0.0f32; ISV_LEN]); 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( launch_sp15_baseline_naive_momentum(
&stream, &stream,
@@ -497,13 +513,9 @@ mod gpu {
ofi_buf.dev_ptr, ofi_buf.dev_ptr,
n as i32, n as i32,
0.0, 0.0,
isv_buf.dev_ptr, momentum_out.dev_ptr,
) )
.expect("launch baseline_naive_momentum_kernel"); .expect("launch baseline_naive_momentum_kernel");
stream
.synchronize()
.expect("synchronize after baseline_naive_momentum_kernel launch");
launch_sp15_baseline_naive_reversion( launch_sp15_baseline_naive_reversion(
&stream, &stream,
prices_buf.dev_ptr, prices_buf.dev_ptr,
@@ -511,16 +523,15 @@ mod gpu {
ofi_buf.dev_ptr, ofi_buf.dev_ptr,
n as i32, n as i32,
0.0, 0.0,
isv_buf.dev_ptr, reversion_out.dev_ptr,
) )
.expect("launch baseline_naive_reversion_kernel"); .expect("launch baseline_naive_reversion_kernel");
stream stream
.synchronize() .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 = momentum_out.read_all()[2];
let momentum_sharpe = isv[BASELINE_NAIVE_MOMENTUM_SHARPE_INDEX]; let reversion_sharpe = reversion_out.read_all()[2];
let reversion_sharpe = isv[BASELINE_NAIVE_REVERSION_SHARPE_INDEX];
// Sign-flipped policies: their sum should be near zero (anti- // Sign-flipped policies: their sum should be near zero (anti-
// correlated picks on the same per-bar return stream). Bound 0.5 // 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 /// Test 1.5 — dd_pct foundational state input concat (LAYOUT FINGERPRINT
/// BREAK). /// BREAK).
/// ///

2
docs/dqn-wire-up-audit.md
View File

File diff suppressed because one or more lines are too long