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

feat(sp20): Phase 1.4 (Path C) — kernel-arg refactor + aggregation kernel + wire-up

Browse Source 
Atomic Phase 1.4 commit per `feedback_no_partial_refactor`. Wires the
SP20 fused-producer chain (Stats → Aggregate → EMAs → Controllers) into
GpuExperienceCollector's per-rollout-step path with one new aggregation
kernel + Phase 1.2 EMA kernel-signature refactor + production wire-up
+ tests + audit/spec/plan amendments, all in one commit.

## Path C decision rationale

The Phase 1.2 EMA kernel originally took 9 scalar value-args. The Phase
1.4 wire-up site (per-rollout-step in GpuExperienceCollector) needs to
feed per-env GPU-resident trade signals (trade_close_per_sample,
step_ret_per_sample, hold_at_exit_per_sample, packed factored
actions_out) into the kernel — forbidden via host sync
(`feedback_cpu_is_read_only`) and via memcpy_dtoh/htod
(`feedback_no_htod_htoh_only_mapped_pinned`). Path C refactors the EMA
kernel to take a device struct pointer (`SP20EmaInputs* ema_inputs`) +
adds an aggregation kernel that writes the struct on the GPU. Phase 1.2
had zero production consumers yet, so the sig change + Phase 1.4 wire-up
land atomically.

## What this commit contains

1. **EMA kernel signature refactor** (Phase 1.2 → Path C):
   `sp20_emas_compute_kernel.cu` swaps 9 scalar args for a single
   `const SP20EmaInputs* __restrict__ ema_inputs` device pointer.
   Math semantics bit-identical. Rust `EmaInputs` mirrored as
   `#[repr(C)]` byte-for-byte; new `pack_inputs_into_f32_view` helper
   for tests + collectors that fill the struct via a mapped-pinned
   f32-aliased buffer.

2. **New aggregation kernel** (`sp20_aggregate_inputs_kernel.cu`):
   Per-env arrays (sliced to current rollout step) + sp20_stats outputs
   (p50/std) + aux_dir_acc_reduce_kernel output [0] → SP20EmaInputs
   struct. Block tree-reduce (4 stripes × bdim × 4 bytes shmem); no
   atomicAdd. Aggregation rules per spec §4.5:
   - is_close = OR over envs
   - is_win = (≥ 0.5 of closed envs were wins)
   - trade_duration = round(mean(hold_at_exit) over closed envs)
   - action_is_hold = strict majority (count*2 > n_envs) HOLD
   - alpha = 0.0  (Phase 2 forward ref — reward kernel)
   - per_bar_hold_reward = 0.0  (Phase 3.2 forward ref — Hold-cost dual)
   - aux_logits_p50/std/dir_acc forwarded from upstream

3. **Production wire-up in GpuExperienceCollector**:
   4 kernel handles + 4 mapped-pinned buffers (struct fields, alloc,
   init); per-rollout-step launch sequence after env_step (step 5c):
   `Stats → Aggregate → EMAs → Controllers`. All on the same stream;
   stream-implicit producer→consumer ordering. Gated on
   `isv_signals_dev_ptr != 0 && trainer_params_ptr != 0` (matches the
   existing SP14-β EGF chain pattern).

4. **Fold-boundary reset**:
   3 new RegistryEntry records (sp20_ema_inputs_buf,
   sp20_emas_internal_buf, sp20_emas_obs_count_buf) + 13 new dispatch
   arms in training_loop.rs (10 SP20 ISV slot resets + 3 buffer resets).
   Closes the pre-existing `every_fold_and_soft_reset_entry_has_dispatch_arm`
   regression (was failing on fresh check after the SP20 ISV slot
   registrations landed in commit 4249ebc96).

5. **HEALTH_DIAG emit**:
   Per-epoch `HEALTH_DIAG[N]: sp20_isv [loss_cap=… alpha_ema=…
   wr_ema=… hold_cost_scale=… target_hold_pct=… hold_pct_ema=…
   hold_reward_ema=… n_step=… aux_conf_threshold=… aux_gate_temp=…]`
   right after the existing q_disagreement_diag emit.

6. **Tests** (5 test files, all green on RTX 3050 Ti sm_86):
   - sp20_emas_compute_test.rs (4 GPU oracle): updated to use the
     post-Path-C buffer-arg API (mapped-pinned f32-aliased struct).
   - sp20_aggregate_inputs_test.rs (NEW, 6 GPU oracle): aggregation
     rule coverage including the Phase 2 / Phase 3.2 placeholder
     contract.
   - sp20_phase1_4_wireup_test.rs (NEW, 2 GPU oracle): end-to-end
     4-kernel chain integration test.
   - sp20_stats_compute_test.rs (4 GPU oracle): unchanged, regression.
   - sp20_controllers_compute_test.rs (7 GPU oracle): unchanged,
     regression.
   - 18 lib unit tests across the SP20 launchers.

7. **Phase 2 / Phase 3.2 forward references**:
   The `alpha` (Phase 2) and `per_bar_hold_reward` (Phase 3.2) fields
   are emitted as 0.0 placeholders and documented at:
   - `sp20_aggregate_inputs_kernel.cu:46-52` (in-kernel docstring)
   - `sp20_aggregate_inputs.rs:46-49` (launcher docstring)
   - `dqn-wire-up-audit.md` "Phase 2 / Phase 3.2 forward references"
   These are NOT stubs — Phase 2 / 3.2 will replace the kernel's 0.0
   writes with real signals atomically with their respective producers.
   The original Task 2.3 (host-side aggregation) is **subsumed** by
   the GPU-side aggregation kernel.

## Hard rules

- `feedback_no_partial_refactor` — kernel sig change + aggregation
  kernel + production wire-up + tests + audit/spec/plan amendments
  in one atomic commit
- `feedback_no_atomicadd` — aggregation kernel uses block tree-reduce
- `feedback_no_cpu_compute_strict` — every aggregation lives on GPU
- `feedback_no_htod_htoh_only_mapped_pinned` — every Phase 1.4 buffer
  is mapped-pinned with `cuMemHostAlloc(DEVICEMAP)` reachable via
  `dev_ptr`; no memcpy_dtoh/htod
- `pearl_first_observation_bootstrap` — counter-based bootstrap
  preserved; placeholder writes (0.0) keep the ALPHA / HOLD_REWARD
  EMAs at 0.0 sentinel until Phase 2 / 3.2 wire real producers
- `pearl_no_host_branches_in_captured_graph` — every kernel is single-
  block; no host branches; safe to capture in the per-step CUDA Graph
- `pearl_tests_must_prove_not_lock_observations` — integration test
  asserts invariants (slots populate, bounds respected) rather than
  locked observed values

## Verification

- `cargo check -p ml --features cuda` clean
- 18 lib unit tests for SP20 launchers pass
- 23 GPU oracle tests across 5 test files pass on RTX 3050 Ti (sm_86)
- `every_fold_and_soft_reset_entry_has_dispatch_arm` regression test
  now passes (was failing pre-change)
- 14 baseline lib test failures unchanged (none introduced by this
  commit)

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-09 20:49:07 +02:00
parent e96f7fcdd1
commit 235e838422
17 changed files with 2170 additions and 143 deletions
Download Patch File Download Diff File Expand all files Collapse all files

34
crates/ml/build.rs
View File

@@ -1392,9 +1392,17 @@ fn main() {
// re-bootstrap. Wiener-α floor 0.4 hardcoded as a structural // re-bootstrap. Wiener-α floor 0.4 hardcoded as a structural
// property per `pearl_wiener_alpha_floor_for_nonstationary`. // property per `pearl_wiener_alpha_floor_for_nonstationary`.
// No atomicAdd, no host branches — captureable in the per- // No atomicAdd, no host branches — captureable in the per-
// step CUDA Graph. Phase 1.4 wires the production launch // step CUDA Graph.
// site atomically with Kernel 2 (controllers) per // AMENDED 2026-05-09 (Path C, Phase 1.4): kernel signature
// `feedback_no_partial_refactor`. // refactored from 9 scalar value-args to a single device
// `SP20EmaInputs* ema_inputs` struct pointer so the upstream
// `sp20_aggregate_inputs_kernel` can populate the 9 fields
// from per-env GPU arrays without any host sync (forbidden
// by `feedback_cpu_is_read_only` + `feedback_no_htod_htoh_only_
// mapped_pinned`). Math semantics bit-identical. The kernel +
// launcher signature change + tests + production wire-up +
// aggregation kernel + audit/spec/plan amendments land
// atomically per `feedback_no_partial_refactor`.
"sp20_emas_compute_kernel.cu", "sp20_emas_compute_kernel.cu",
// SP20 Phase 1.3 (2026-05-09): 6 derived ISV controller // SP20 Phase 1.3 (2026-05-09): 6 derived ISV controller
// outputs producer. Component 5 / Kernel 2 of the SP20 fused- // outputs producer. Component 5 / Kernel 2 of the SP20 fused-
@@ -1415,6 +1423,26 @@ fn main() {
// 1.4 wires the production launch site atomically with // 1.4 wires the production launch site atomically with
// Kernel 1 + Kernel 3 per `feedback_no_partial_refactor`. // Kernel 1 + Kernel 3 per `feedback_no_partial_refactor`.
"sp20_controllers_compute_kernel.cu", "sp20_controllers_compute_kernel.cu",
// SP20 Phase 1.4 (2026-05-09, Path C): per-env → SP20EmaInputs
// aggregation kernel. Reads `trade_close_per_env`,
// `step_ret_per_env`, `hold_at_exit_per_env` (all f32/i32
// [N] arrays sliced to current_t from the experience
// collector's per-sample `[N, L]` buffers) + packed factored
// `actions_per_env` + upstream `sp20_stats_compute` outputs
// (p50/std) + `aux_dir_acc_reduce_kernel` output [0]; tree-
// reduces over envs and writes a single `SP20EmaInputs`
// 36-byte struct that the post-Path-C `sp20_emas_compute_kernel`
// reads via a device-pointer arg. Aggregation rules:
// is_close = OR over envs; is_win = (≥0.5 of closed envs were
// wins); trade_duration = round(mean over closed envs);
// action_is_hold = (>n_envs/2 envs decoded to DIR_HOLD);
// alpha / per_bar_hold_reward = 0.0 (Phase 2 / 3.2 forward
// references). Single-block, BLOCK=64 tree-reduce; no
// atomicAdd per `feedback_no_atomicadd`. Mandatory companion
// to the Phase 1.2 buffer-arg kernel-signature refactor —
// the two land atomically with the production wire-up + audit
// / spec / plan amendments per `feedback_no_partial_refactor`.
"sp20_aggregate_inputs_kernel.cu",
]; ];
// ALL kernels get common header (BF16 types + wrappers) // ALL kernels get common header (BF16 types + wrappers)

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

@@ -727,6 +727,43 @@ pub(crate) static SP17_ADVANTAGE_CLIP_BOUND_CUBIN: &[u8] =
// `feedback_no_legacy_aliases`. Coupling A (forward feature wire // `feedback_no_legacy_aliases`. Coupling A (forward feature wire
// `dir_concat_qaux_kernel`) is preserved — its cubin static stays. // `dir_concat_qaux_kernel`) is preserved — its cubin static stays.
/// SP20 Phase 1.1 (2026-05-09): aux_logits p50 + std producer cubin.
/// Component 5 / Kernel 3 of the SP20 fused-producer chain. Single-
/// block, BLOCK=256 kernel reading `aux_logits [B, K=2]` and emitting
/// `[p50, std]` of the per-row peak-confidence signal `aux_conf =
/// max_c softmax(logits) - 1/2`. Loaded by the experience collector
/// for the per-rollout-step SP20 chain (Phase 1.4).
pub(crate) static SP20_STATS_COMPUTE_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/sp20_stats_compute_kernel.cubin"));
/// SP20 Phase 1.2 (2026-05-09): 8 Wiener-α EMA producer cubin.
/// Component 5 / Kernel 1 of the SP20 fused-producer chain. Single-
/// block, single-thread kernel updating 4 ISV slots + 4 internal
/// scratch slots from a device `SP20EmaInputs` struct (post Path C
/// kernel-arg refactor, 2026-05-09). Loaded by the experience
/// collector for the per-rollout-step SP20 chain (Phase 1.4).
pub(crate) static SP20_EMAS_COMPUTE_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/sp20_emas_compute_kernel.cubin"));
/// SP20 Phase 1.3 (2026-05-09): 6 derived ISV controller outputs
/// producer cubin. Component 5 / Kernel 2 of the SP20 fused-producer
/// chain. Single-block, single-thread kernel reading the EMAs Kernel
/// 1 wrote and deriving LOSS_CAP / N_STEP / AUX_CONF_THRESHOLD /
/// AUX_GATE_TEMP / TARGET_HOLD_PCT / HOLD_COST_SCALE per spec §4.5.
/// Loaded by the experience collector for the per-rollout-step SP20
/// chain (Phase 1.4).
pub(crate) static SP20_CONTROLLERS_COMPUTE_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/sp20_controllers_compute_kernel.cubin"));
/// SP20 Phase 1.4 (2026-05-09, Path C): per-env → SP20EmaInputs
/// aggregation kernel cubin. Bridges the experience collector's
/// per-sample `[N, L]` GPU buffers (`trade_close_per_sample`,
/// `step_ret_per_sample`, `hold_at_exit_per_sample`, `actions_out`)
/// to the EMA kernel's struct-input contract. Single-block, BLOCK=64
/// tree-reduce per `feedback_no_atomicadd`.
pub(crate) static SP20_AGGREGATE_INPUTS_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/sp20_aggregate_inputs_kernel.cubin"));
/// SP14 B.6 (2026-05-05): pre-SGEMM concat kernel for direction Q-head input. /// SP14 B.6 (2026-05-05): pre-SGEMM concat kernel for direction Q-head input.
/// Reads `h_s2 [B, SH2]` + `aux_softmax_diff [1]` and writes /// Reads `h_s2 [B, SH2]` + `aux_softmax_diff [1]` and writes
/// `[h_s2 | softmax_diff] [B, SH2 + 1]` into `dir_qaux_concat_scratch`. /// `[h_s2 | softmax_diff] [B, SH2 + 1]` into `dir_qaux_concat_scratch`.

308
crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
View File

@@ -1415,6 +1415,69 @@ pub struct GpuExperienceCollector {
/// per-epoch boundary is the trigger; the hot-path action computation /// per-epoch boundary is the trigger; the hot-path action computation
/// itself is 100% GPU. /// itself is 100% GPU.
seed_phase_active_cache: bool, seed_phase_active_cache: bool,
// ─── SP20 Phase 1.4 (2026-05-09) — fused-producer chain ────────────
// Per-rollout-step Stats → Aggregate → EMAs → Controllers chain
// wired directly into the timestep loop (after `experience_env_step`
// and the existing aux-head forward + dir-acc reduce). Reads
// per-env GPU arrays (`trade_close_per_sample`, `step_ret_per_sample`,
// `hold_at_exit_per_sample`, packed factored `actions_out`) at the
// current step slice + the upstream aux-head outputs; writes 10 ISV
// slots [510..520) (4 EMAs + 6 derived controllers) every rollout
// step. Fold-boundary reset zeroes the 4-internal scratch + the
// 8-counter obs_count + the 9-f32 ema_inputs scratch via the
// `sp20_emas_*` registry-arm dispatches in `training_loop.rs`.
/// SP20 Phase 1.1 Kernel 3 (stats producer): aux confidence p50 +
/// std producer. Reads `aux_logits [B, K=2]` from
/// `exp_aux_nb_logits_buf` + writes `[p50, std]` into
/// `sp20_stats_output_buf`. Loaded from `SP20_STATS_COMPUTE_CUBIN`.
sp20_stats_compute_kernel: CudaFunction,
/// SP20 Phase 1.4 Path C aggregation kernel: per-env →
/// `SP20EmaInputs` reducer. Loaded from
/// `SP20_AGGREGATE_INPUTS_CUBIN`.
sp20_aggregate_inputs_kernel: CudaFunction,
/// SP20 Phase 1.2 Kernel 1 (EMA producer): updates 4 ISV EMAs +
/// 4 internal scratch EMAs from a `SP20EmaInputs` device struct.
/// Loaded from `SP20_EMAS_COMPUTE_CUBIN`.
sp20_emas_compute_kernel: CudaFunction,
/// SP20 Phase 1.3 Kernel 2 (controller producer): derives 6 ISV
/// controller slots from the EMAs Kernel 1 wrote. Loaded from
/// `SP20_CONTROLLERS_COMPUTE_CUBIN`.
sp20_controllers_compute_kernel: CudaFunction,
/// Mapped-pinned [2] f32 output of `sp20_stats_compute_kernel`:
/// `[p50, std]`. The aggregation kernel reads `[0]` for p50 and
/// `[1]` for std (passed as separate dev_ptrs offset by 4 bytes).
pub(crate) sp20_stats_output_buf: super::mapped_pinned::MappedF32Buffer,
/// Mapped-pinned [9] f32 byte-aliased to the `SP20EmaInputs` struct.
/// Output of the aggregation kernel; input to the EMA kernel. The
/// 9-f32 buffer aliases the device-side 36-byte struct one-for-one
/// (4 ints in slots [0..4) via bit-cast; 5 floats in slots [4..9)).
/// FoldReset writes 9 zeros (clears any inherited state — though
/// the kernel rewrites every byte each step, so this is purely
/// defensive).
pub(crate) sp20_ema_inputs_buf: super::mapped_pinned::MappedF32Buffer,
/// Mapped-pinned [4] f32 internal scratch for the EMA kernel.
/// Layout per `EMA_INTERNAL_*` constants: [TRADE_DURATION (0),
/// AUX_P50 (1), AUX_STD (2), AUX_DIR_ACC (3)]. Read by
/// `sp20_controllers_compute_kernel` to derive N_STEP /
/// TARGET_HOLD_PCT / AUX_GATE_TEMP / AUX_CONF_THRESHOLD.
/// FoldReset zeros all 4 slots (Pearl-A bootstrap on next launch).
pub(crate) sp20_emas_internal_buf: super::mapped_pinned::MappedF32Buffer,
/// Mapped-pinned [8] i32 per-EMA observation counter for the EMA
/// kernel. Layout per `OBS_COUNT_*` constants. Counter == 0 ⇒
/// next firing observation REPLACES sentinel directly; counter >
/// 0 ⇒ Wiener-blend with α=0.4 floor. FoldReset zeros all 8
/// counters so the new fold's first per-EMA firing observation
/// re-bootstraps the EMA from sentinel.
pub(crate) sp20_emas_obs_count_buf: super::mapped_pinned::MappedI32Buffer,
} }
impl Drop for GpuExperienceCollector { impl Drop for GpuExperienceCollector {
@@ -2098,6 +2161,105 @@ impl GpuExperienceCollector {
// deleted atomically with the kernel file. The q_disagreement_update // deleted atomically with the kernel file. The q_disagreement_update
// handle loaded above remains for the HEALTH_DIAG diagnostic chain. // handle loaded above remains for the HEALTH_DIAG diagnostic chain.
// SP20 Phase 1.4 (2026-05-09, Path C): pre-load the 4 SP20
// fused-producer kernel handles on the collector's stream.
// Same loading pattern as the SP13/SP14 EGF chain above —
// each handle is resolved once at construct time so the
// per-rollout-step launches add no per-step host work.
// Per-step ordering on `self.stream` (CUDA stream-implicit
// producer→consumer):
// 1. sp20_stats_compute — aux_logits → [p50, std]
// 2. sp20_aggregate_inputs — per-env arrays + stats + dir_acc
// → SP20EmaInputs struct
// 3. sp20_emas_compute — struct → 4 ISV EMAs + 4 internal
// 4. sp20_controllers_compute — EMAs → 6 ISV controllers
let sp20_stats_compute_kernel = {
use super::gpu_dqn_trainer::SP20_STATS_COMPUTE_CUBIN;
let m = stream.context()
.load_cubin(SP20_STATS_COMPUTE_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!(
"sp20: stats_compute cubin (collector): {e}"
)))?;
m.load_function("sp20_stats_compute_kernel")
.map_err(|e| MLError::ModelError(format!(
"sp20: stats_compute_kernel load (collector): {e}"
)))?
};
let sp20_aggregate_inputs_kernel = {
use super::gpu_dqn_trainer::SP20_AGGREGATE_INPUTS_CUBIN;
let m = stream.context()
.load_cubin(SP20_AGGREGATE_INPUTS_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!(
"sp20: aggregate_inputs cubin (collector): {e}"
)))?;
m.load_function("sp20_aggregate_inputs_kernel")
.map_err(|e| MLError::ModelError(format!(
"sp20: aggregate_inputs_kernel load (collector): {e}"
)))?
};
let sp20_emas_compute_kernel = {
use super::gpu_dqn_trainer::SP20_EMAS_COMPUTE_CUBIN;
let m = stream.context()
.load_cubin(SP20_EMAS_COMPUTE_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!(
"sp20: emas_compute cubin (collector): {e}"
)))?;
m.load_function("sp20_emas_compute_kernel")
.map_err(|e| MLError::ModelError(format!(
"sp20: emas_compute_kernel load (collector): {e}"
)))?
};
let sp20_controllers_compute_kernel = {
use super::gpu_dqn_trainer::SP20_CONTROLLERS_COMPUTE_CUBIN;
let m = stream.context()
.load_cubin(SP20_CONTROLLERS_COMPUTE_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!(
"sp20: controllers_compute cubin (collector): {e}"
)))?;
m.load_function("sp20_controllers_compute_kernel")
.map_err(|e| MLError::ModelError(format!(
"sp20: controllers_compute_kernel load (collector): {e}"
)))?
};
// SP20 Phase 1.4 mapped-pinned scratch buffers. Allocate up
// front so the per-step launches read/write fixed dev pointers
// (no per-step alloc, no per-step HtoD copy).
let sp20_stats_output_buf = unsafe {
super::mapped_pinned::MappedF32Buffer::new(2)
}.map_err(|e| MLError::ModelError(format!(
"sp20: alloc sp20_stats_output_buf (2 f32): {e}"
)))?;
// Initialise to a sentinel that the EMA kernel reads cleanly
// on the cold-start path (per-step EMAs always fire, so the
// first observation REPLACES sentinel directly per
// `pearl_first_observation_bootstrap`).
sp20_stats_output_buf.write_from_slice(&[0.0_f32, 0.0_f32]);
use super::sp20_emas_compute::{
EMA_INTERNAL_SLOT_COUNT, OBS_COUNT_SLOTS, SP20_EMA_INPUTS_F32_LEN,
};
let sp20_ema_inputs_buf = unsafe {
super::mapped_pinned::MappedF32Buffer::new(SP20_EMA_INPUTS_F32_LEN)
}.map_err(|e| MLError::ModelError(format!(
"sp20: alloc sp20_ema_inputs_buf ({SP20_EMA_INPUTS_F32_LEN} f32): {e}"
)))?;
sp20_ema_inputs_buf.write_from_slice(&vec![0.0_f32; SP20_EMA_INPUTS_F32_LEN]);
let sp20_emas_internal_buf = unsafe {
super::mapped_pinned::MappedF32Buffer::new(EMA_INTERNAL_SLOT_COUNT)
}.map_err(|e| MLError::ModelError(format!(
"sp20: alloc sp20_emas_internal_buf ({EMA_INTERNAL_SLOT_COUNT} f32): {e}"
)))?;
sp20_emas_internal_buf.write_from_slice(&vec![0.0_f32; EMA_INTERNAL_SLOT_COUNT]);
let sp20_emas_obs_count_buf = unsafe {
super::mapped_pinned::MappedI32Buffer::new(OBS_COUNT_SLOTS)
}.map_err(|e| MLError::ModelError(format!(
"sp20: alloc sp20_emas_obs_count_buf ({OBS_COUNT_SLOTS} i32): {e}"
)))?;
sp20_emas_obs_count_buf.write_from_slice(&vec![0_i32; OBS_COUNT_SLOTS]);
// SP14 β-migration step 2 (2026-05-07): collector-side aux-head // SP14 β-migration step 2 (2026-05-07): collector-side aux-head
// forward infrastructure. The aux next-bar head's forward computes // forward infrastructure. The aux next-bar head's forward computes
// a `[N, K=2]` softmax tile from `h_s2` + the trainer's params at // a `[N, K=2]` softmax tile from `h_s2` + the trainer's params at
@@ -2721,6 +2883,16 @@ impl GpuExperienceCollector {
last_cf_ratio_eff: 0.5, // D4/N4: standard cf_ratio at healthy state last_cf_ratio_eff: 0.5, // D4/N4: standard cf_ratio at healthy state
contrarian_active_cache: 0, // D7/N7: off by default contrarian_active_cache: 0, // D7/N7: off by default
seed_phase_active_cache: false, // B.3 Plan 3 Task 8: cold-start in network mode; trainer flips on at epoch boundary if SEED_STEPS_DONE < TARGET seed_phase_active_cache: false, // B.3 Plan 3 Task 8: cold-start in network mode; trainer flips on at epoch boundary if SEED_STEPS_DONE < TARGET
// ── SP20 Phase 1.4 (2026-05-09) — fused-producer chain ──
sp20_stats_compute_kernel,
sp20_aggregate_inputs_kernel,
sp20_emas_compute_kernel,
sp20_controllers_compute_kernel,
sp20_stats_output_buf,
sp20_ema_inputs_buf,
sp20_emas_internal_buf,
sp20_emas_obs_count_buf,
reward_rank_kernel, reward_rank_kernel,
reward_compute_abs_sharpe_kernel, reward_compute_abs_sharpe_kernel,
bitonic_sort_step_kernel, bitonic_sort_step_kernel,
@@ -6114,6 +6286,142 @@ impl GpuExperienceCollector {
// `sp15_dd_trajectory_prev_dd_dev_ptr` setter / pointer // `sp15_dd_trajectory_prev_dd_dev_ptr` setter / pointer
// gate were deleted as part of this commit. // gate were deleted as part of this commit.
// ── 5c. SP20 Phase 1.4 (2026-05-09, Path C) — fused-producer chain
//
// Stats → Aggregate → EMAs → Controllers, all on
// `self.stream` after `experience_env_step` so the per-env
// arrays we read (trade_close, step_ret, hold_at_exit,
// actions_out) are already populated for the current `t`.
// The aux-head outputs (exp_aux_nb_logits_buf for stats,
// exp_aux_dir_acc_buf[0] for dir_acc) are populated by
// step 3a-β above. Stream-implicit producer→consumer
// ordering: every kernel here reads the prior kernel's
// outputs without explicit event sync.
//
// Gated on `isv_signals_dev_ptr != 0` (the EMAs + controller
// launches both write/read ISV; without it wired the EMAs
// would have no destination). Also requires
// `trainer_params_ptr != 0` so the aux head forward in
// step 3 wrote real logits (not the alloc_zeros sentinel).
if self.isv_signals_dev_ptr != 0 && self.trainer_params_ptr != 0 {
use crate::cuda_pipeline::sp20_aggregate_inputs::launch_sp20_aggregate_inputs;
use crate::cuda_pipeline::sp20_emas_compute::launch_sp20_emas_compute;
use crate::cuda_pipeline::sp20_controllers_compute::launch_sp20_controllers_compute;
use crate::cuda_pipeline::sp20_stats_compute::launch_sp20_stats_compute;
let n_envs_i32_sp20 = n_episodes as i32;
let aux_logits_dev = self.exp_aux_nb_logits_buf.raw_ptr();
let stats_out_dev = self.sp20_stats_output_buf.dev_ptr;
// The stats kernel writes [p50, std] as 2 f32s; the
// aggregate kernel reads p50 from offset 0 and std
// from offset 1*sizeof(f32) = 4 bytes.
let stats_p50_dev: u64 = stats_out_dev;
let stats_std_dev: u64 = stats_out_dev
.checked_add(std::mem::size_of::<f32>() as u64)
.expect("sp20: stats_std_dev offset overflow");
// 1. sp20_stats_compute — aux_logits → [p50, std].
unsafe {
launch_sp20_stats_compute(
&self.stream,
&self.sp20_stats_compute_kernel,
aux_logits_dev,
stats_out_dev,
n_envs_i32_sp20,
)?;
}
// 2. sp20_aggregate_inputs — per-env arrays + stats +
// dir_acc → SP20EmaInputs struct.
//
// Pointer arithmetic: per-sample buffers are flat
// `[N, L]` with `out_off = i*L + t`. The kernel strides
// by `env_stride = L = alloc_timesteps` between envs,
// so we pass `base_ptr + t * sizeof(elem)` for each
// array (placing env-0's slice at element 0).
let l = self.alloc_timesteps;
let env_stride: i32 = l as i32;
let t_offset_i32: u64 =
(t as u64).checked_mul(std::mem::size_of::<i32>() as u64)
.expect("sp20: t_offset_i32 overflow");
let t_offset_f32: u64 =
(t as u64).checked_mul(std::mem::size_of::<f32>() as u64)
.expect("sp20: t_offset_f32 overflow");
let trade_close_t_dev = self.trade_close_per_sample.raw_ptr()
.checked_add(t_offset_i32)
.expect("sp20: trade_close_t_dev overflow");
let step_ret_t_dev = self.step_ret_per_sample.raw_ptr()
.checked_add(t_offset_f32)
.expect("sp20: step_ret_t_dev overflow");
let hold_at_exit_t_dev = self.hold_at_exit_per_sample.raw_ptr()
.checked_add(t_offset_f32)
.expect("sp20: hold_at_exit_t_dev overflow");
let actions_t_dev = self.actions_out.raw_ptr()
.checked_add(t_offset_i32)
.expect("sp20: actions_t_dev overflow");
// dir_divisor = MAGNITUDES * ORD * URG (branch_sizes[1..4]).
// hold_action_id = DIR_HOLD = 1 (per state_layout.cuh —
// SHORT=0, HOLD=1, LONG=2, FLAT=3). Same decoding as
// `hold_rate_observer_kernel.cu`.
let dir_divisor: i32 = (self.branch_sizes[1]
* self.branch_sizes[2]
* self.branch_sizes[3]) as i32;
let hold_action_id: i32 = 1; // DIR_HOLD
// aux_dir_acc lives at the existing exp_aux_dir_acc_buf
// [0] (populated by aux_dir_acc_reduce_kernel up at
// step 3a-β). Reused as Phase 1.4's dir_acc input.
let aux_dir_acc_dev = self.exp_aux_dir_acc_buf.dev_ptr;
let ema_inputs_dev = self.sp20_ema_inputs_buf.dev_ptr;
unsafe {
launch_sp20_aggregate_inputs(
&self.stream,
&self.sp20_aggregate_inputs_kernel,
trade_close_t_dev,
step_ret_t_dev,
hold_at_exit_t_dev,
actions_t_dev,
env_stride,
stats_p50_dev,
stats_std_dev,
aux_dir_acc_dev,
n_envs_i32_sp20,
dir_divisor,
hold_action_id,
ema_inputs_dev,
)?;
}
// 3. sp20_emas_compute — struct → 4 ISV EMAs +
// 4 internal scratch.
let emas_internal_dev = self.sp20_emas_internal_buf.dev_ptr;
let emas_obs_count_dev = self.sp20_emas_obs_count_buf.dev_ptr;
unsafe {
launch_sp20_emas_compute(
&self.stream,
&self.sp20_emas_compute_kernel,
self.isv_signals_dev_ptr,
ema_inputs_dev,
emas_internal_dev,
emas_obs_count_dev,
)?;
}
// 4. sp20_controllers_compute — EMAs → 6 ISV
// controllers.
unsafe {
launch_sp20_controllers_compute(
&self.stream,
&self.sp20_controllers_compute_kernel,
self.isv_signals_dev_ptr,
emas_internal_dev,
)?;
}
}
// ── 6. Advance GPU-resident step counter ──────────────────── // ── 6. Advance GPU-resident step counter ────────────────────
// Single-thread kernel: step_counter_gpu[0] += 1. // Single-thread kernel: step_counter_gpu[0] += 1.
// Runs on the same stream — serialized after env_step, before // Runs on the same stream — serialized after env_step, before

17
crates/ml/src/cuda_pipeline/mapped_pinned.rs
View File

@@ -95,6 +95,23 @@ impl MappedI32Buffer {
} }
} }
} }
/// Mutable host slice over the mapped pages — `&mut [i32]` aliasing
/// `host_ptr..host_ptr+len`. Lets callers structure writes (e.g.
/// fold-boundary `host_slice_mut().fill(0)`) without per-element
/// `write_volatile`. The kernel sees the writes after a
/// stream-sync barrier (mapped pinned coherence — same memory,
/// different alias).
///
/// # Safety
/// `host_ptr` is valid for `self.len` i32s by construction; the
/// caller must not invalidate the allocation while the slice is
/// live and must respect the unique-mutable-borrow rule for
/// `&mut [T]`.
pub fn host_slice_mut(&mut self) -> &mut [i32] {
// Safety: see invariants above.
unsafe { std::slice::from_raw_parts_mut(self.host_ptr, self.len) }
}
} }
impl Drop for MappedI32Buffer { impl Drop for MappedI32Buffer {

11
crates/ml/src/cuda_pipeline/mod.rs
View File

@@ -109,6 +109,17 @@ pub mod sp20_emas_compute;
// 1.4 wires the production launch site atomically with the rest of // 1.4 wires the production launch site atomically with the rest of
// the SP20 chain per `feedback_no_partial_refactor`. // the SP20 chain per `feedback_no_partial_refactor`.
pub mod sp20_controllers_compute; pub mod sp20_controllers_compute;
// SP20 Phase 1.4 (2026-05-09, Path C): per-env → SP20EmaInputs
// aggregation kernel launcher. Reads `trade_close_per_sample`,
// `step_ret_per_sample`, `hold_at_exit_per_sample`, packed factored
// `actions_out` (all `[N, L]` per-sample arrays sliced to the current
// step) plus the upstream `sp20_stats_compute` outputs (`p50`/`std`)
// and `aux_dir_acc_reduce` output [0]; aggregates them into a single
// 9-field `SP20EmaInputs` struct buffer that
// `sp20_emas_compute_kernel` (post Path C kernel-arg refactor) reads.
// Lands atomically with the kernel-arg refactor + production wire-up
// + audit/spec/plan amendments per `feedback_no_partial_refactor`.
pub mod sp20_aggregate_inputs;
// `launch_apply_pearls` is `pub(crate)` and consumed only inside this // `launch_apply_pearls` is `pub(crate)` and consumed only inside this
// crate via `use crate::cuda_pipeline::sp4_wiener_ema::launch_apply_pearls`. // crate via `use crate::cuda_pipeline::sp4_wiener_ema::launch_apply_pearls`.
// `pearls_ad_update` stays `pub` for the test-oracle path in // `pearls_ad_update` stays `pub` for the test-oracle path in

204
crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs Normal file
View File

@@ -0,0 +1,204 @@
#![allow(unsafe_code)] // CUDA kernel launch via cudarc requires unsafe.
//! SP20 Phase 1.4 (2026-05-09) — `sp20_aggregate_inputs_kernel` launcher.
//!
//! Path C aggregation step. Aggregates per-env GPU arrays produced by
//! `experience_env_step` (`trade_close`, `step_ret`, `hold_at_exit`,
//! packed factored `actions`) plus the upstream `sp20_stats_compute`
//! outputs (`aux_logits_p50` / `aux_logits_std`) plus the existing
//! `aux_dir_acc_reduce_kernel` output [0] (`aux_dir_acc`) into the
//! 9-field [`super::sp20_emas_compute::EmaInputs`] struct that
//! `sp20_emas_compute_kernel` reads.
//!
//! Aggregation rules (one struct, one rollout-step):
//!
//! | Field | Rule |
//! |------------------------|---------------------------------------------------------------------------------------|
//! | `is_close` | `1` if any env's `trade_close[env] != 0`, else `0` |
//! | `is_win` | `1` if `is_close == 1` AND fraction of closed envs with `step_ret > 0` ≥ 0.5, else `0` |
//! | `trade_duration` | `round(mean(hold_at_exit) over closed envs)` if `is_close == 1`, else `0` |
//! | `action_is_hold` | `1` if `count(decoded_dir == HOLD) > n_envs/2`, else `0` |
//! | `alpha` | `0.0` — **Phase 2 forward reference** (R_event - hold_baseline) |
//! | `per_bar_hold_reward` | `0.0` — **Phase 3.2 forward reference** (-aux_conf * cost_scale) |
//! | `aux_logits_p50` | `aux_logits_p50_dev[0]` (sp20_stats output) |
//! | `aux_logits_std` | `aux_logits_std_dev[0]` (sp20_stats output) |
//! | `aux_dir_acc` | `aux_dir_acc_dev[0]` (existing aux_dir_acc_reduce_kernel output [0]) |
//!
//! ## Hard rules covered
//!
//! - `feedback_no_atomicadd` — block tree-reduce only. The 4 reductions
//! (closed-count, wins-count, hold-count, hold_at_exit-sum) share one
//! shmem tile and reduce in a single loop pass.
//! - `feedback_no_cpu_compute_strict` — every aggregation lives in the
//! kernel; the launcher only passes pointers + `n_envs` + constants.
//! - `feedback_no_htod_htoh_only_mapped_pinned` — output is a device
//! pointer (mapped-pinned scratch in the production wire-up); kernel
//! emits `__threadfence_system()` after the writes.
//! - `pearl_no_host_branches_in_captured_graph` — single-block kernel,
//! no host branches; safe to capture in the per-step CUDA Graph.
//! - `feedback_no_partial_refactor` — kernel + launcher + test + build
//! entry land atomically with the Phase 1.2 buffer-arg refactor +
//! the production wire-up + audit/spec/plan amendments.
//!
//! ## Phase 2 / Phase 3.2 forward references
//!
//! The `alpha` and `per_bar_hold_reward` fields are intentionally
//! emitted as `0.0` here:
//! - `alpha = R_event - hold_baseline` is computed by the SP20 reward
//! kernel (Phase 2). Phase 1.4 writes `0.0` so the ALPHA_EMA's
//! bootstrap counter advances on every closed trade and the EMA
//! stays at its 0.0 sentinel until Phase 2 wires the real producer.
//! - `per_bar_hold_reward = -aux_conf * cost_scale` is computed by
//! the SP20 Hold-cost dual emission (Phase 3.2). Phase 1.4 writes
//! `0.0` so the HOLD_REWARD_EMA's bootstrap counter advances on
//! Hold-action bars and the EMA stays at its 0.0 sentinel.
//!
//! These are NOT stubs — they are documented forward references where
//! the producer is being shipped in a later phase per
//! `feedback_no_partial_refactor` (the Phase 1.4 commit's atomic
//! scope is the SP20 ISV state-tracker chain, not the reward kernel).
/// Block size for the aggregation reduction. 32 threads is enough for
/// `n_envs` up to a few thousand (one warp tree-reduce per reduction).
/// Bumped to 64 for slightly better occupancy at modest n_envs.
pub const SP20_AGGREGATE_BLOCK: u32 = 64;
/// Compute the dynamic shared-memory bytes the aggregation kernel
/// needs for the configured block size: 4 stripes × `bdim` × 4 bytes
/// (3 i32 stripes for closed/wins/hold counts + 1 f32 stripe for the
/// hold_at_exit sum).
#[must_use]
pub fn dynamic_shmem_bytes() -> u32 {
SP20_AGGREGATE_BLOCK
.checked_mul(4)
.and_then(|x| x.checked_mul(std::mem::size_of::<i32>() as u32))
.expect("SP20 aggregate: shmem bytes overflow")
}
/// Number of `f32` slots an upstream `MappedF32Buffer` for the EMA
/// inputs struct needs to hold (= `SP20_EMA_INPUTS_F32_LEN` = 9). Kept
/// here as a re-export for collector wire-up sites that prefer to
/// import from the aggregation module to avoid pulling in the full
/// EMA launcher module.
pub const SP20_EMA_INPUTS_F32_LEN: usize =
super::sp20_emas_compute::SP20_EMA_INPUTS_F32_LEN;
/// Launch `sp20_aggregate_inputs_kernel` on the producer's stream.
///
/// Single-block, [`SP20_AGGREGATE_BLOCK`]-thread launch. `kernel` MUST
/// be the loaded `sp20_aggregate_inputs_kernel` `CudaFunction`.
///
/// # Pointer & stride contract
///
/// The 4 per-env arrays (`trade_close_dev`, `step_ret_dev`,
/// `hold_at_exit_dev`, `actions_dev`) MUST already be offset to point
/// at the slice for the **current rollout step**. The flat layout for
/// the per-sample buffers in `experience_env_step` is `out_off = env*L + t`,
/// so the production wire-up passes `base_dev_ptr + (current_t * sizeof(T))`
/// for each, and `env_stride = L` (alloc_timesteps).
///
/// ### Decoding contract
///
/// `dir_divisor` MUST be the production `NUM_MAGNITUDES * NUM_ORD * NUM_URG`
/// (typically `27 = 3*3*3`). `hold_action_id` MUST be `DIR_HOLD = 1`
/// (per `state_layout.cuh`). The kernel decodes:
///
/// `dir = action_idx / dir_divisor`
///
/// matching `experience_action_select` and `hold_rate_observer_kernel`.
///
/// # Safety
///
/// All device pointers MUST be valid for the current stream and remain
/// valid until the kernel completes (or until `stream.synchronize()`).
/// Stream ordering with `sp20_stats_compute_kernel` (which writes the
/// p50/std outputs) and `aux_dir_acc_reduce_kernel` (which writes the
/// dir_acc output) is the caller's responsibility — same-stream is
/// sufficient (CUDA stream-ordering invariant).
pub unsafe fn launch_sp20_aggregate_inputs(
stream: &cudarc::driver::CudaStream,
kernel: &cudarc::driver::CudaFunction,
trade_close_dev: u64,
step_ret_dev: u64,
hold_at_exit_dev: u64,
actions_dev: u64,
env_stride: i32,
aux_logits_p50_dev: u64,
aux_logits_std_dev: u64,
aux_dir_acc_dev: u64,
n_envs: i32,
dir_divisor: i32,
hold_action_id: i32,
out_inputs_dev: u64,
) -> Result<(), crate::MLError> {
use cudarc::driver::{LaunchConfig, PushKernelArg};
debug_assert!(trade_close_dev != 0);
debug_assert!(step_ret_dev != 0);
debug_assert!(hold_at_exit_dev != 0);
debug_assert!(actions_dev != 0);
debug_assert!(aux_logits_p50_dev != 0);
debug_assert!(aux_logits_std_dev != 0);
debug_assert!(aux_dir_acc_dev != 0);
debug_assert!(out_inputs_dev != 0);
debug_assert!(n_envs > 0, "n_envs must be > 0, got {n_envs}");
debug_assert!(env_stride > 0, "env_stride must be > 0, got {env_stride}");
debug_assert!(dir_divisor > 0, "dir_divisor must be > 0, got {dir_divisor}");
debug_assert!(hold_action_id >= 0,
"hold_action_id must be >= 0, got {hold_action_id}");
let cfg = LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (SP20_AGGREGATE_BLOCK, 1, 1),
shared_mem_bytes: dynamic_shmem_bytes(),
};
stream
.launch_builder(kernel)
.arg(&trade_close_dev)
.arg(&step_ret_dev)
.arg(&hold_at_exit_dev)
.arg(&actions_dev)
.arg(&env_stride)
.arg(&aux_logits_p50_dev)
.arg(&aux_logits_std_dev)
.arg(&aux_dir_acc_dev)
.arg(&n_envs)
.arg(&dir_divisor)
.arg(&hold_action_id)
.arg(&out_inputs_dev)
.launch(cfg)
.map_err(|e| {
crate::MLError::ModelError(format!(
"sp20_aggregate_inputs_kernel launch: {e}"
))
})?;
Ok(())
}
#[cfg(test)]
mod tests {
use super::*;
#[test]
fn block_size_is_warp_multiple() {
assert!(SP20_AGGREGATE_BLOCK >= 32);
assert_eq!(SP20_AGGREGATE_BLOCK % 32, 0);
}
#[test]
fn shmem_bytes_match_layout() {
// 4 stripes × bdim × 4 bytes.
let expected = SP20_AGGREGATE_BLOCK as usize * 4 * 4;
assert_eq!(dynamic_shmem_bytes() as usize, expected);
}
#[test]
fn ema_inputs_f32_len_re_export_locked() {
// Re-export agreement with the EMA launcher.
assert_eq!(
SP20_EMA_INPUTS_F32_LEN,
super::super::sp20_emas_compute::SP20_EMA_INPUTS_F32_LEN
);
assert_eq!(SP20_EMA_INPUTS_F32_LEN, 9);
}
}

259
crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu Normal file
View File

@@ -0,0 +1,259 @@
/* ══════════════════════════════════════════════════════════════════════════
* SP20 Phase 1.4 (2026-05-09) — sp20_aggregate_inputs kernel.
*
* Path C aggregation step. Reads per-env GPU arrays produced by
* `experience_env_step` (trade_close, step_ret, hold_at_exit, actions)
* + the upstream `sp20_stats_compute` outputs (aux_logits_p50/std) +
* the existing `aux_dir_acc_reduce_kernel` output [0] (aux_dir_acc),
* aggregates them into the 9-field `SP20EmaInputs` struct that
* `sp20_emas_compute_kernel` reads.
*
* Aggregation rules (one struct, one rollout-step):
*
* is_close = 1 if any env's trade_close_per_env[env] != 0, else 0.
* is_win = 1 if (is_close == 1) AND fraction of closed envs with
* step_ret_per_env[env] > 0 is >= 0.5, else 0.
* trade_duration = round(mean(hold_at_exit_per_env[env]) for env
* where trade_close_per_env[env] != 0)
* if is_close == 1, else 0.
* alpha = 0.0 (Phase 2 forward reference — `R_event - hold_baseline`
* is computed by the SP20 reward kernel that lands in
* Phase 2 alongside the hold-baseline circular buffer.
* Phase 1.4 writes 0.0 here as a documented placeholder
* so the EMA's bootstrap counter advances and the WR/HOLD
* EMAs can fire on real signals while ALPHA_EMA stays at
* its 0.0 sentinel.)
* per_bar_hold_reward = 0.0 (Phase 3.2 forward reference — the
* `r_per_bar = -aux_conf * cost_scale` formula needs the
* per-bar aux_conf signal that Phase 3.2's per-bar aux
* forward emits. Phase 1.4 writes 0.0 so the HOLD_REWARD
* EMA's bootstrap counter advances on Hold-action bars
* while staying at its 0.0 sentinel until Phase 3.2.)
* aux_logits_p50 = aux_logits_p50_dev[0] (sp20_stats output)
* aux_logits_std = aux_logits_std_dev[0] (sp20_stats output)
* aux_dir_acc = aux_dir_acc_dev[0] (existing aux_dir_acc_buf[0]
* from aux_dir_acc_reduce_kernel)
* action_is_hold = 1 if count(actions_per_env[env] decoded direction
* == HOLD) > n_envs / 2, else 0.
*
* The action decoding uses the production packed factored
* action_idx encoding from `experience_action_select`:
*
* action_idx = dir * (NUM_MAGNITUDES * NUM_ORD * NUM_URG)
* + mag * (NUM_ORD * NUM_URG)
* + ord * NUM_URG
* + urg
*
* so `dir = action_idx / (NUM_MAGNITUDES * NUM_ORD * NUM_URG)`. The
* caller passes the divisor as `dir_divisor` (typically 27 = 3*3*3
* for the production branch sizes) and the Hold direction id as
* `hold_action_id` (= 1 = DIR_HOLD per `state_layout.cuh`).
*
* ── Pearls + invariants ────────────────────────────────────────────────
* - `feedback_no_atomicadd` — block tree-reduce only. The 3 reductions
* (closed-count, wins-count, hold-count) share one shmem tile
* sequentially with explicit `__syncthreads()` between phases. The
* `hold_at_exit` mean is computed as a separate sum reduction
* immediately after; total of 4 tree-reductions, all bdim×i32 or
* bdim×f32 in shmem.
* - `feedback_no_cpu_compute_strict` — every aggregation lives in
* this kernel; the launcher only passes pointers + n_envs +
* constants.
* - `feedback_no_htod_htoh_only_mapped_pinned` — the output struct
* buffer is mapped-pinned (or device-only — both work since the
* consumer kernel reads via dev_ptr). Kernel emits
* `__threadfence_system()` after the writes so the downstream
* `sp20_emas_compute_kernel` sees the updates after the in-stream
* ordering invariant.
* - `pearl_no_host_branches_in_captured_graph` — single-block kernel,
* no host branches; safe to capture in the per-step CUDA Graph.
* - `feedback_no_partial_refactor` — kernel + launcher + test +
* build entry land atomically with the Phase 1.2 buffer-arg
* refactor + the production wire-up + audit/spec/plan
* amendments in Phase 1.4.
* - `pearl_first_observation_bootstrap` — the placeholder fields
* (`alpha = 0.0`, `per_bar_hold_reward = 0.0`) emit the bootstrap
* sentinel so Phase 1.2's EMAs see "no firing" semantics for
* those EMAs until Phase 2 / Phase 3.2 wire the real producers.
*
* Launch contract:
* grid_dim = (1, 1, 1)
* block_dim = (BLOCK_SIZE, 1, 1) — typically 32 or 64 for n_envs
* up to a few hundred.
* shared mem = 4 × bdim × max(sizeof(int), sizeof(float)) = 4 × bdim × 4
* ══════════════════════════════════════════════════════════════════════════ */
#include <cuda_runtime.h>
/* Same struct layout as `sp20_emas_compute_kernel.cu`. The two .cu
* files are compiled separately into different cubins; redeclaring the
* struct here keeps each cubin self-contained and avoids a shared
* .cuh header that would couple the build manifest. The Rust side's
* single source of truth (`EmaInputs` in `sp20_emas_compute.rs` with
* `#[repr(C)]`) ensures both kernels see the same byte layout. */
struct SP20EmaInputs {
int is_close;
int is_win;
int trade_duration;
int action_is_hold;
float alpha;
float per_bar_hold_reward;
float aux_logits_p50;
float aux_logits_std;
float aux_dir_acc;
};
#ifndef SP20_AGG_BLOCK
#define SP20_AGG_BLOCK 64
#endif
extern "C" __global__ void sp20_aggregate_inputs_kernel(
/* Per-env input arrays (read at offset env_id for env_id in [0, n_envs)).
* Each is a slice of the [N, L] per-sample buffers at the current
* timestep — the Rust launcher passes a pointer offset by
* `env_offset_for_step = current_t` (since the per-sample layout is
* `[i, t]` flat with stride L between envs, we pass the full base
* pointer + a stride-L iteration through env_id below). */
const int* __restrict__ trade_close_per_env, /* [n_envs * stride] */
const float* __restrict__ step_ret_per_env, /* [n_envs * stride] */
const float* __restrict__ hold_at_exit_per_env, /* [n_envs * stride] */
const int* __restrict__ actions_per_env, /* [n_envs * stride] */
/* Stride between consecutive envs in the four per-env arrays. The
* `experience_env_step` flat layout is `out_off = env*L + t`, so
* the launcher passes `stride = L` (alloc_timesteps) and adds
* `current_t` to each base pointer before launch — i.e., the
* arrays here are already pointing at the slice for the current
* timestep, and we stride by L between envs. */
int env_stride,
/* Upstream sp20_stats_compute outputs — single-element device
* pointers for the median / std of aux_logits at this step. */
const float* __restrict__ aux_logits_p50_dev, /* [1] */
const float* __restrict__ aux_logits_std_dev, /* [1] */
/* Existing aux_dir_acc_reduce_kernel output [0] — the directional
* accuracy of the aux head this step. */
const float* __restrict__ aux_dir_acc_dev, /* [1] */
int n_envs,
/* Action decoding parameters (production packed factored action
* encoding). `dir_divisor = NUM_MAGNITUDES * NUM_ORD * NUM_URG`,
* `hold_action_id = DIR_HOLD = 1`. */
int dir_divisor,
int hold_action_id,
/* Output struct (single 36-byte struct on device; the EMA kernel
* reads it as `SP20EmaInputs* ema_inputs`). */
SP20EmaInputs* __restrict__ out_inputs)
{
/* Single-block contract — guard against accidental multi-block
* launches that would corrupt the shmem reductions. */
if (blockIdx.x != 0) return;
const int tid = threadIdx.x;
const int bdim = blockDim.x;
/* Dynamic shared memory: 4 stripes of bdim × 4 bytes each. We use
* `int` storage for the 3 count reductions and `float` storage for
* the hold-at-exit sum (alias-compatible because both are 4 bytes
* wide; we cast between views as needed). */
extern __shared__ int sh_int[];
/* Layout: [closed_count | wins_count | hold_count | hold_at_exit_sum_f32] */
int* sh_closed_count = &sh_int[0];
int* sh_wins_count = &sh_int[bdim];
int* sh_hold_count = &sh_int[2 * bdim];
float* sh_hold_at_exit_sum = (float*)&sh_int[3 * bdim];
/* Per-thread strided accumulation over [0, n_envs). */
int local_closed_count = 0;
int local_wins_count = 0;
int local_hold_count = 0;
float local_hold_at_exit_sum = 0.0f;
for (int env = tid; env < n_envs; env += bdim) {
const int tc = trade_close_per_env[env * env_stride];
const float sr = step_ret_per_env[env * env_stride];
const float he = hold_at_exit_per_env[env * env_stride];
const int action_idx = actions_per_env[env * env_stride];
/* dir = action_idx / dir_divisor — same as
* `hold_rate_observer_kernel.cu`'s decoding. */
const int dir = action_idx / dir_divisor;
if (tc != 0) {
local_closed_count += 1;
if (sr > 0.0f) local_wins_count += 1;
local_hold_at_exit_sum += he;
}
if (dir == hold_action_id) local_hold_count += 1;
}
sh_closed_count[tid] = local_closed_count;
sh_wins_count[tid] = local_wins_count;
sh_hold_count[tid] = local_hold_count;
sh_hold_at_exit_sum[tid] = local_hold_at_exit_sum;
__syncthreads();
/* Tree reduction across `bdim`. Standard log2(bdim) pattern. We
* reduce all four stripes in one loop body — single barrier per
* pass keeps the cost equal to a single reduction. */
for (int s = bdim / 2; s > 0; s >>= 1) {
if (tid < s) {
sh_closed_count[tid] += sh_closed_count[tid + s];
sh_wins_count[tid] += sh_wins_count[tid + s];
sh_hold_count[tid] += sh_hold_count[tid + s];
sh_hold_at_exit_sum[tid] += sh_hold_at_exit_sum[tid + s];
}
__syncthreads();
}
/* Thread 0 writes the output struct. */
if (tid == 0) {
const int closed_count = sh_closed_count[0];
const int wins_count = sh_wins_count[0];
const int hold_count = sh_hold_count[0];
const float hold_at_exit_sum_f = sh_hold_at_exit_sum[0];
const int is_close_out = (closed_count > 0) ? 1 : 0;
/* is_win: fraction of closed envs that won >= 0.5. */
int is_win_out = 0;
if (is_close_out == 1) {
/* (wins_count / closed_count) >= 0.5
* ⇔ 2 * wins_count >= closed_count (avoids fp). */
is_win_out = (2 * wins_count >= closed_count) ? 1 : 0;
}
/* trade_duration: round(mean of hold_at_exit over closed envs). */
int trade_duration_out = 0;
if (is_close_out == 1) {
const float mean_he = hold_at_exit_sum_f / (float)closed_count;
trade_duration_out = (int)__float2int_rn(mean_he);
}
/* action_is_hold: majority-vote over envs. > n/2 threshold
* matches the spec (strict majority, ties ⇒ 0). */
const int action_is_hold_out = (hold_count * 2 > n_envs) ? 1 : 0;
/* Read the upstream stats / dir-acc outputs. Stream-ordering
* invariant guarantees these are visible to this kernel's
* read because the launcher places sp20_stats_compute and
* aux_dir_acc_reduce BEFORE this aggregation kernel on the
* same stream (see Phase 1.4 collector wire-up). */
const float p50_in = aux_logits_p50_dev[0];
const float std_in = aux_logits_std_dev[0];
const float dir_acc_in = aux_dir_acc_dev[0];
/* Write the 9 fields. The float fields `alpha` and
* `per_bar_hold_reward` are intentionally zero — Phase 2 /
* Phase 3.2 wire the real producers; until then the EMAs that
* blend those fields stay at sentinel 0.0 (per
* `pearl_first_observation_bootstrap`). */
out_inputs->is_close = is_close_out;
out_inputs->is_win = is_win_out;
out_inputs->trade_duration = trade_duration_out;
out_inputs->action_is_hold = action_is_hold_out;
out_inputs->alpha = 0.0f; /* Phase 2 forward ref */
out_inputs->per_bar_hold_reward = 0.0f; /* Phase 3.2 fwd ref */
out_inputs->aux_logits_p50 = p50_in;
out_inputs->aux_logits_std = std_in;
out_inputs->aux_dir_acc = dir_acc_in;
__threadfence_system();
}
}

246
crates/ml/src/cuda_pipeline/sp20_emas_compute.rs
View File

@@ -19,7 +19,7 @@
//! | EMA | Storage | Gate | Observation source | //! | EMA | Storage | Gate | Observation source |
//! |-----|---------|------|--------------------| //! |-----|---------|------|--------------------|
//! | `ALPHA_EMA` | ISV[511] | `is_close == 1` | `alpha = R_event - hold_baseline` | //! | `ALPHA_EMA` | ISV[511] | `is_close == 1` | `alpha = R_event - hold_baseline` |
//! | `WR_EMA` | ISV[512] | `is_close == 1` | `is_win` ∈ {0.0, 1.0} | //! | `WR_EMA` | ISV[512] | `is_close == 1` | `is_win` ∈ {0, 1} |
//! | `TRADE_DUR` | internal[0]| `is_close == 1` | `trade_duration` (bars) | //! | `TRADE_DUR` | internal[0]| `is_close == 1` | `trade_duration` (bars) |
//! | `HOLD_PCT_EMA` | ISV[515] | per-step | `(action_is_hold == 1) ? 1 : 0` | //! | `HOLD_PCT_EMA` | ISV[515] | per-step | `(action_is_hold == 1) ? 1 : 0` |
//! | `HOLD_REWARD_EMA`| ISV[516] | `action_is_hold == 1` | `r_per_bar = -aux_conf * cost_scale` | //! | `HOLD_REWARD_EMA`| ISV[516] | `action_is_hold == 1` | `r_per_bar = -aux_conf * cost_scale` |
@@ -36,6 +36,21 @@
//! [`WIENER_ALPHA_FLOOR`] = 0.4 floor per //! [`WIENER_ALPHA_FLOOR`] = 0.4 floor per
//! `pearl_wiener_alpha_floor_for_nonstationary`. //! `pearl_wiener_alpha_floor_for_nonstationary`.
//! //!
//! ## Path C kernel-arg refactor (AMENDED 2026-05-09, Phase 1.4)
//!
//! Phase 1.4's wire-up needs to feed per-env GPU-resident trade signals
//! (`trade_close_per_sample`, `step_ret_per_sample`, etc.) into the EMA
//! kernel without round-tripping through the CPU
//! (`feedback_cpu_is_read_only`, `feedback_no_htod_htoh_only_mapped_pinned`).
//! Path C: the kernel now reads its 9 fields from a device-side
//! [`EmaInputs`] struct pointer; the upstream `sp20_aggregate_inputs_kernel`
//! aggregates per-env arrays into the struct on the GPU. Math
//! semantics are bit-identical with the original 9-scalar-arg signature
//! (every blend formula reads the same fields by name, only the
//! kernel-arg signature changed). The Phase 1.2 tests + the production
//! wire-up + the aggregation kernel + audit/spec/plan amendments land
//! atomically per `feedback_no_partial_refactor`.
//!
//! ## Hard rules covered //! ## Hard rules covered
//! //!
//! - `pearl_first_observation_bootstrap` — counter-based bootstrap. //! - `pearl_first_observation_bootstrap` — counter-based bootstrap.
@@ -45,20 +60,21 @@
//! - `feedback_no_atomicadd` — single-thread kernel, no concurrent //! - `feedback_no_atomicadd` — single-thread kernel, no concurrent
//! writes anywhere. //! writes anywhere.
//! - `feedback_no_cpu_compute_strict` — every blend lives in the //! - `feedback_no_cpu_compute_strict` — every blend lives in the
//! kernel; the launcher only packs scalar inputs into kernel args. //! kernel; the launcher only passes pointers. The `EmaInputs` struct
//! is filled by the upstream aggregation kernel reading per-env
//! GPU buffers.
//! - `feedback_no_htod_htoh_only_mapped_pinned` — both ISV and the //! - `feedback_no_htod_htoh_only_mapped_pinned` — both ISV and the
//! `internal` / `obs_count` buffers are mapped-pinned; kernel emits //! `internal` / `obs_count` buffers are mapped-pinned; tests fill
//! `__threadfence_system()` for PCIe-visible coherence. //! the input struct via mapped-pinned [`MappedF32Buffer`]
//! aliasing. Kernel emits `__threadfence_system()` for PCIe-visible
//! coherence.
//! - `pearl_no_host_branches_in_captured_graph` — single-block, //! - `pearl_no_host_branches_in_captured_graph` — single-block,
//! single-thread kernel; safe to capture in the per-step CUDA //! single-thread kernel; safe to capture in the per-step CUDA
//! Graph. //! Graph.
//! //! - `feedback_no_partial_refactor` — Phase 1.2 lands kernel +
//! ## Phase 1.2 wiring scope //! launcher + tests + build entry; Phase 1.4 lands Path C kernel-
//! //! arg refactor + aggregation kernel + production wire-up
//! Kernel + launcher + tests + build entry **only**. The training-loop //! atomically.
//! call site lands in Phase 1.4 atomically with the rest of the SP20
//! reward chain (Kernel 2 = controllers, Kernel 3 = stats already
//! landed in Phase 1.1) per `feedback_no_partial_refactor`.
/// Wiener-α floor per `pearl_wiener_alpha_floor_for_nonstationary`. NOT /// Wiener-α floor per `pearl_wiener_alpha_floor_for_nonstationary`. NOT
/// a tuned constant — structural floor that preserves catch-up /// a tuned constant — structural floor that preserves catch-up
@@ -120,43 +136,99 @@ pub const OBS_COUNT_AUX_STD: usize = 6;
/// Observation-counter slot: aux_dir_acc_ema (internal). /// Observation-counter slot: aux_dir_acc_ema (internal).
pub const OBS_COUNT_AUX_DIR_ACC: usize = 7; pub const OBS_COUNT_AUX_DIR_ACC: usize = 7;
/// Per-step inputs to `launch_sp20_emas_compute`. Packs the 9 scalar /// Per-step EMA inputs as the kernel reads them.
/// fields the kernel needs into a single struct so the launcher ///
/// signature stays manageable and the call sites are /// **Path C 2026-05-09**: the kernel now reads these 9 fields from a
/// readable. /// device buffer pointer (`SP20EmaInputs*` in the .cu file), so this
/// Rust struct is the host-side mirror of the device-side layout.
/// Mark `#[repr(C)]` so the field order/alignment matches the kernel's
/// `struct SP20EmaInputs` byte-for-byte.
/// ///
/// Field semantics: /// Field semantics:
/// - `is_close`: 1 if a trade closed this step, 0 otherwise. Gates /// - `is_close`: 1 if a trade closed this step, 0 otherwise. Gates
/// the per-trade EMAs (WR, ALPHA, TRADE_DURATION). /// the per-trade EMAs (WR, ALPHA, TRADE_DURATION).
/// - `is_win`: 1.0 if the closed trade was a win, 0.0 if loss. /// - `is_win`: 1 if the closed trade was a win, 0 otherwise.
/// Only meaningful when `is_close == 1`. /// Only meaningful when `is_close == 1`. Stored as i32 to match
/// - `alpha`: `R_event - hold_baseline` from the SP20 spec /// the device-struct field width.
/// Component 1 reward kernel. Only meaningful when
/// `is_close == 1`.
/// - `trade_duration`: bars in the closed trade. Only meaningful /// - `trade_duration`: bars in the closed trade. Only meaningful
/// when `is_close == 1`. /// when `is_close == 1`. Stored as i32 (rounded mean across envs
/// - `action_is_hold`: 1 if the action this step was Hold, 0 /// in the production aggregation path).
/// otherwise. Gates the HOLD_REWARD_EMA; also drives the /// - `action_is_hold`: 1 if the (majority) action this step was
/// Hold, 0 otherwise. Gates the HOLD_REWARD_EMA; also drives the
/// HOLD_PCT_EMA observation. /// HOLD_PCT_EMA observation.
/// - `r_per_bar`: per-bar Hold reward = `-aux_conf * cost_scale`. /// - `alpha`: `R_event - hold_baseline` from the SP20 spec
/// Only meaningful when `action_is_hold == 1`. /// Component 1 reward kernel. Phase 2 wires the real producer;
/// - `aux_p50_in`: per-step `aux_logits_p50` from Phase 1.1's /// Phase 1.4 wire-up populates this with sentinel 0.0 (forward
/// reference, see audit doc + `sp20_aggregate_inputs_kernel.cu`).
/// - `per_bar_hold_reward`: per-bar Hold reward = `-aux_conf *
/// cost_scale`. Only meaningful when `action_is_hold == 1`.
/// Phase 3.2 wires the real producer; Phase 1.4 populates this
/// with sentinel 0.0 (forward reference).
/// - `aux_logits_p50`: per-step `aux_logits_p50` from Phase 1.1's
/// `sp20_stats_compute` (mapped-pinned `[2]` output, slot 0). /// `sp20_stats_compute` (mapped-pinned `[2]` output, slot 0).
/// - `aux_std_in`: per-step `aux_logits_std` from Phase 1.1 /// - `aux_logits_std`: per-step `aux_logits_std` from Phase 1.1
/// (slot 1 of the same output). /// (slot 1 of the same output).
/// - `aux_dir_acc`: 1.0 if the aux head predicted the right /// - `aux_dir_acc`: 1.0 if the aux head predicted the right
/// direction at this step, 0.0 otherwise. /// direction at this step, 0.0 otherwise. Phase 1.4 reads this
#[derive(Copy, Clone, Debug)] /// from the existing `exp_aux_dir_acc_buf[0]` infra populated
/// by `aux_dir_acc_reduce_kernel`.
#[repr(C)]
#[derive(Copy, Clone, Debug, Default)]
pub struct EmaInputs { pub struct EmaInputs {
pub is_close: i32, pub is_close: i32,
pub is_win: f32, pub is_win: i32,
pub alpha: f32, pub trade_duration: i32,
pub trade_duration: f32, pub action_is_hold: i32,
pub action_is_hold: i32, pub alpha: f32,
pub r_per_bar: f32, pub per_bar_hold_reward: f32,
pub aux_p50_in: f32, pub aux_logits_p50: f32,
pub aux_std_in: f32, pub aux_logits_std: f32,
pub aux_dir_acc: f32, pub aux_dir_acc: f32,
}
/// Number of bytes the device-side `SP20EmaInputs` struct occupies.
/// Used by both the launcher (when sizing the input buffer) and tests.
/// `9 × 4 = 36` bytes (4 ints + 5 floats, all 4-byte aligned).
pub const SP20_EMA_INPUTS_BYTES: usize = std::mem::size_of::<EmaInputs>();
/// Number of `f32`s the EMA-inputs buffer should be allocated as if
/// allocated through `MappedF32Buffer` (since the buffer is byte-aliased
/// by the kernel as a struct, the buffer length in `f32` units is
/// `SP20_EMA_INPUTS_BYTES / 4 = 9`).
pub const SP20_EMA_INPUTS_F32_LEN: usize = SP20_EMA_INPUTS_BYTES / 4;
/// Pack `EmaInputs` into a 9-`f32` mapped-pinned buffer view. The
/// device-side kernel reads the buffer as a `SP20EmaInputs*` struct;
/// the buffer must be sized at least [`SP20_EMA_INPUTS_F32_LEN`] f32s
/// (9), and writes are bit-pattern transmutes (i32/f32 share f32
/// alignment).
///
/// # Safety
///
/// `out` must be sized at least [`SP20_EMA_INPUTS_F32_LEN`]. The caller
/// owns the mapped-pinned buffer; the kernel sees the writes after a
/// stream-sync barrier (mapped-pinned coherence — same memory, different
/// alias).
pub fn pack_inputs_into_f32_view(out: &mut [f32], inputs: &EmaInputs) {
debug_assert!(
out.len() >= SP20_EMA_INPUTS_F32_LEN,
"pack_inputs_into_f32_view: out len ({}) < SP20_EMA_INPUTS_F32_LEN ({})",
out.len(), SP20_EMA_INPUTS_F32_LEN,
);
// Match the device-side `SP20EmaInputs` layout: 4 i32 + 5 f32.
// Bit-pattern transmute: the device kernel reads the f32-aliased
// buffer as i32 fields by struct offset (the .cu file's
// `int is_close` is a 4-byte read off the same offset our
// `f32::from_bits(...)` write lands at).
out[0] = f32::from_bits(inputs.is_close as u32);
out[1] = f32::from_bits(inputs.is_win as u32);
out[2] = f32::from_bits(inputs.trade_duration as u32);
out[3] = f32::from_bits(inputs.action_is_hold as u32);
out[4] = inputs.alpha;
out[5] = inputs.per_bar_hold_reward;
out[6] = inputs.aux_logits_p50;
out[7] = inputs.aux_logits_std;
out[8] = inputs.aux_dir_acc;
} }
/// Launch `sp20_emas_compute_kernel` on the producer's stream. /// Launch `sp20_emas_compute_kernel` on the producer's stream.
@@ -169,44 +241,43 @@ pub struct EmaInputs {
/// allocations: /// allocations:
/// - `isv_dev`: ISV bus, at least /// - `isv_dev`: ISV bus, at least
/// [`crate::cuda_pipeline::sp14_isv_slots::SP20_ISV_SLOT_END`] f32s. /// [`crate::cuda_pipeline::sp14_isv_slots::SP20_ISV_SLOT_END`] f32s.
/// - `ema_inputs_dev`: device pointer to a single
/// [`SP20_EMA_INPUTS_BYTES`]-byte struct (`SP20EmaInputs`). Either
/// a mapped-pinned scratch buffer the caller wrote via
/// [`pack_inputs_into_f32_view`], or the output of the upstream
/// `sp20_aggregate_inputs_kernel` (Phase 1.4 production wire).
/// - `internal_dev`: [`EMA_INTERNAL_SLOT_COUNT`]-long f32 scratch. /// - `internal_dev`: [`EMA_INTERNAL_SLOT_COUNT`]-long f32 scratch.
/// - `obs_count_dev`: [`OBS_COUNT_SLOTS`]-long i32 counters. /// - `obs_count_dev`: [`OBS_COUNT_SLOTS`]-long i32 counters.
/// ///
/// `inputs` is the per-step gather of all 9 scalar EMA observations;
/// see [`EmaInputs`].
///
/// # Safety /// # Safety
/// ///
/// All three buffer pointers MUST be valid mapped-pinned device /// All four buffer pointers MUST be valid mapped-pinned device
/// pointers obtained via `MappedF32Buffer::dev_ptr` / /// pointers obtained via `MappedF32Buffer::dev_ptr` /
/// `MappedI32Buffer::dev_ptr`. The caller MUST issue this launch on /// `MappedI32Buffer::dev_ptr` (or the upstream aggregation kernel's
/// the SAME stream as the upstream `sp20_stats_compute` (so the /// output buffer). The caller MUST issue this launch on the SAME
/// stream-ordering invariant guarantees `aux_p50_in` / `aux_std_in` /// stream as the upstream `sp20_stats_compute` AND
/// reads see the latest stats output). Stream ordering with the /// `sp20_aggregate_inputs_kernel` (so the stream-ordering invariant
/// trade-physics kernels that produce `is_close` / `is_win` / etc. is /// guarantees `aux_logits_p50` / `aux_logits_std` / aggregated
/// the caller's responsibility — Phase 1.4 lands the production wire /// per-env signals are visible to this kernel's reads). Phase 1.4
/// site with that ordering enforced. /// lands the production wire site with that ordering enforced.
pub unsafe fn launch_sp20_emas_compute( pub unsafe fn launch_sp20_emas_compute(
stream: &cudarc::driver::CudaStream, stream: &cudarc::driver::CudaStream,
kernel: &cudarc::driver::CudaFunction, kernel: &cudarc::driver::CudaFunction,
isv_dev: u64, isv_dev: u64,
ema_inputs_dev: u64,
internal_dev: u64, internal_dev: u64,
obs_count_dev: u64, obs_count_dev: u64,
inputs: EmaInputs,
) -> Result<(), crate::MLError> { ) -> Result<(), crate::MLError> {
use cudarc::driver::{LaunchConfig, PushKernelArg}; use cudarc::driver::{LaunchConfig, PushKernelArg};
debug_assert!(isv_dev != 0, debug_assert!(isv_dev != 0,
"launch_sp20_emas_compute: isv_dev must be a valid device pointer"); "launch_sp20_emas_compute: isv_dev must be a valid device pointer");
debug_assert!(ema_inputs_dev != 0,
"launch_sp20_emas_compute: ema_inputs_dev must be a valid device pointer");
debug_assert!(internal_dev != 0, debug_assert!(internal_dev != 0,
"launch_sp20_emas_compute: internal_dev must be a valid device pointer"); "launch_sp20_emas_compute: internal_dev must be a valid device pointer");
debug_assert!(obs_count_dev != 0, debug_assert!(obs_count_dev != 0,
"launch_sp20_emas_compute: obs_count_dev must be a valid device pointer"); "launch_sp20_emas_compute: obs_count_dev must be a valid device pointer");
debug_assert!(inputs.is_close == 0 || inputs.is_close == 1,
"launch_sp20_emas_compute: is_close must be 0 or 1, got {}", inputs.is_close);
debug_assert!(inputs.action_is_hold == 0 || inputs.action_is_hold == 1,
"launch_sp20_emas_compute: action_is_hold must be 0 or 1, got {}",
inputs.action_is_hold);
// Slot indices passed in as kernel args so the kernel stays // Slot indices passed in as kernel args so the kernel stays
// decoupled from SP14 slot-number drift (the contract is enforced // decoupled from SP14 slot-number drift (the contract is enforced
@@ -225,21 +296,13 @@ pub unsafe fn launch_sp20_emas_compute(
stream stream
.launch_builder(kernel) .launch_builder(kernel)
.arg(&isv_dev) .arg(&isv_dev)
.arg(&ema_inputs_dev)
.arg(&internal_dev) .arg(&internal_dev)
.arg(&obs_count_dev) .arg(&obs_count_dev)
.arg(&isv_idx_alpha) .arg(&isv_idx_alpha)
.arg(&isv_idx_wr) .arg(&isv_idx_wr)
.arg(&isv_idx_hold_pct) .arg(&isv_idx_hold_pct)
.arg(&isv_idx_hold_reward) .arg(&isv_idx_hold_reward)
.arg(&inputs.is_close)
.arg(&inputs.is_win)
.arg(&inputs.alpha)
.arg(&inputs.trade_duration)
.arg(&inputs.action_is_hold)
.arg(&inputs.r_per_bar)
.arg(&inputs.aux_p50_in)
.arg(&inputs.aux_std_in)
.arg(&inputs.aux_dir_acc)
.launch(cfg) .launch(cfg)
.map_err(|e| { .map_err(|e| {
crate::MLError::ModelError(format!("sp20_emas_compute_kernel launch: {e}")) crate::MLError::ModelError(format!("sp20_emas_compute_kernel launch: {e}"))
@@ -290,21 +353,11 @@ mod tests {
#[test] #[test]
fn ema_inputs_default_is_safe_sentinel() { fn ema_inputs_default_is_safe_sentinel() {
// A zero-initialised `EmaInputs` has all gates closed // A `Default::default()` `EmaInputs` has all gates closed
// (is_close = 0, action_is_hold = 0) and all observations at // (is_close = 0, action_is_hold = 0) and all observations at
// 0 — equivalent to "no firing this step" for all 8 EMAs. // 0 — equivalent to "no firing this step" for all 8 EMAs.
// This is the canonical pre-warmup / cold-start input shape. // This is the canonical pre-warmup / cold-start input shape.
let inp = EmaInputs { let inp = EmaInputs::default();
is_close: 0,
is_win: 0.0,
alpha: 0.0,
trade_duration: 0.0,
action_is_hold: 0,
r_per_bar: 0.0,
aux_p50_in: 0.0,
aux_std_in: 0.0,
aux_dir_acc: 0.0,
};
// Per-step EMAs still fire on a zero-input call, but the // Per-step EMAs still fire on a zero-input call, but the
// observation is 0 so the EMA Wiener-blends toward 0 — that // observation is 0 so the EMA Wiener-blends toward 0 — that
// is correct behavior for cold-start, and is verified by the // is correct behavior for cold-start, and is verified by the
@@ -312,4 +365,45 @@ mod tests {
assert_eq!(inp.is_close, 0); assert_eq!(inp.is_close, 0);
assert_eq!(inp.action_is_hold, 0); assert_eq!(inp.action_is_hold, 0);
} }
#[test]
fn ema_inputs_bytes_locked() {
// Path C kernel-arg refactor: the device-side struct is
// 4 ints + 5 floats = 36 bytes. The Rust mirror MUST match.
// If the device struct ever grows a field, this assertion
// catches the drift early — the launcher's
// `SP20_EMA_INPUTS_F32_LEN` derives from this.
assert_eq!(SP20_EMA_INPUTS_BYTES, 36);
assert_eq!(SP20_EMA_INPUTS_F32_LEN, 9);
}
#[test]
fn pack_inputs_round_trips() {
// Sanity: pack into a 9-f32 buffer and read back the bit
// patterns; verify i32 fields land in the right slots.
let inp = EmaInputs {
is_close: 1,
is_win: 1,
trade_duration: 7,
action_is_hold: 0,
alpha: 0.5,
per_bar_hold_reward: -0.05,
aux_logits_p50: 0.2,
aux_logits_std: 0.1,
aux_dir_acc: 1.0,
};
let mut buf = [0.0_f32; SP20_EMA_INPUTS_F32_LEN];
pack_inputs_into_f32_view(&mut buf, &inp);
// Verify i32 slots come back via to_bits round trip.
assert_eq!(buf[0].to_bits() as i32, 1);
assert_eq!(buf[1].to_bits() as i32, 1);
assert_eq!(buf[2].to_bits() as i32, 7);
assert_eq!(buf[3].to_bits() as i32, 0);
// Float slots are direct.
assert_eq!(buf[4], 0.5);
assert_eq!(buf[5], -0.05);
assert_eq!(buf[6], 0.2);
assert_eq!(buf[7], 0.1);
assert_eq!(buf[8], 1.0);
}
} }

104
crates/ml/src/cuda_pipeline/sp20_emas_compute_kernel.cu
View File

@@ -9,6 +9,21 @@
* consumes to derive `LOSS_CAP`, `HOLD_COST_SCALE`, `TARGET_HOLD_PCT`, * consumes to derive `LOSS_CAP`, `HOLD_COST_SCALE`, `TARGET_HOLD_PCT`,
* `N_STEP`, `AUX_CONF_THRESHOLD`, `AUX_GATE_TEMP`. * `N_STEP`, `AUX_CONF_THRESHOLD`, `AUX_GATE_TEMP`.
* *
* AMENDED 2026-05-09 (Path C): the kernel originally took 9 scalar
* value-args (is_close, is_win, …, aux_dir_acc). The Phase 1.4 wire-up
* needs to feed per-env GPU-resident trade signals into these inputs;
* since CPU compute is forbidden (`feedback_cpu_is_read_only`,
* `feedback_no_htod_htoh_only_mapped_pinned`), the only zero-CPU way
* to populate the inputs is via an upstream aggregation kernel writing
* to a device buffer. Path C: the kernel now reads the 9 fields from a
* device buffer pointer (`SP20EmaInputs* ema_inputs`); the upstream
* `sp20_aggregate_inputs_kernel` aggregates per-env arrays into the
* struct on the GPU. Math semantics are bit-identical: every blend
* formula reads the same fields by name, only the kernel-arg signature
* changed. Phase 1.2's tests + the production wire-up + the aggregation
* kernel + audit/spec/plan amendments land in one atomic commit per
* `feedback_no_partial_refactor`.
*
* EMAs maintained (8 total — see launcher header for slot indices): * EMAs maintained (8 total — see launcher header for slot indices):
* *
* Per-trade (gated on is_close == 1): * Per-trade (gated on is_close == 1):
@@ -47,17 +62,20 @@
* - `feedback_no_atomicadd` — single-thread kernel; no concurrent * - `feedback_no_atomicadd` — single-thread kernel; no concurrent
* writes anywhere, atomicAdd never used. * writes anywhere, atomicAdd never used.
* - `feedback_no_htod_htoh_only_mapped_pinned` — both ISV (mapped- * - `feedback_no_htod_htoh_only_mapped_pinned` — both ISV (mapped-
* pinned) and `internal` / `obs_count` (this commit's two new * pinned), `internal` / `obs_count` (mapped-pinned), and
* mapped-pinned buffers) are device-visible via DEVICEMAP; kernel * `ema_inputs` (mapped-pinned struct buffer or upstream-kernel-
* emits `__threadfence_system()` after the writes for PCIe-visible * written device buffer) are device-visible; kernel emits
* `__threadfence_system()` after the writes for PCIe-visible
* coherence. * coherence.
* - `feedback_no_cpu_compute_strict` — every blend lives in this * - `feedback_no_cpu_compute_strict` — every blend lives in this
* kernel; the launcher's only host work is to package the per-step * kernel; the launcher's only host work is to pass pointers. The
* scalar inputs into kernel arguments. * `ema_inputs` struct is filled by the upstream `sp20_aggregate_
* inputs_kernel` (Phase 1.4) reading per-env GPU buffers.
* - `feedback_no_partial_refactor` — Phase 1.2 lands kernel + launcher * - `feedback_no_partial_refactor` — Phase 1.2 lands kernel + launcher
* + tests + build entry atomically. The production wire-up (Phase * + tests + build entry atomically. Phase 1.4 wires the production
* 1.4) consumes the launcher and the audit doc updates per * consumer (Path C kernel-arg refactor + aggregation kernel +
* Invariant 7. * collector launch site + audit/spec/plan amendments) atomically
* in one commit.
* *
* Algorithm (single block, single thread): * Algorithm (single block, single thread):
* *
@@ -83,6 +101,13 @@
* Args: * Args:
* isv — [ISV_TOTAL_DIM] f32 mapped-pinned dev ptr (slots * isv — [ISV_TOTAL_DIM] f32 mapped-pinned dev ptr (slots
* 511, 512, 515, 516 written). * 511, 512, 515, 516 written).
* ema_inputs — pointer to a single `SP20EmaInputs` struct on the
* device. The 9 fields (`is_close`, `is_win`, etc.)
* are read directly. Production caller (Phase 1.4)
* fills this via the upstream `sp20_aggregate_inputs_
* kernel`; tests fill it via mapped-pinned
* `MappedF32Buffer` byte-aliased over the struct
* layout.
* internal — [4] f32 mapped-pinned dev ptr (the 4 internal * internal — [4] f32 mapped-pinned dev ptr (the 4 internal
* scratch slots: TRADE_DUR, AUX_P50, AUX_STD, * scratch slots: TRADE_DUR, AUX_P50, AUX_STD,
* AUX_DIR_ACC). Layout matches launcher constants. * AUX_DIR_ACC). Layout matches launcher constants.
@@ -93,14 +118,32 @@
* — ISV slot indices for the 4 ISV EMAs (caller passes * — ISV slot indices for the 4 ISV EMAs (caller passes
* the SP14 constants so the kernel stays decoupled * the SP14 constants so the kernel stays decoupled
* from slot-number drift). * from slot-number drift).
* is_close, is_win, alpha, trade_duration, action_is_hold,
* r_per_bar, aux_p50_in, aux_std_in, aux_dir_acc_in — per-step
* inputs (see EMA list above for which fields gate
* which EMA).
* ══════════════════════════════════════════════════════════════════════════ */ * ══════════════════════════════════════════════════════════════════════════ */
#include <cuda_runtime.h> #include <cuda_runtime.h>
/* SP20 Phase 1.2 + 1.4 (Path C, 2026-05-09): Device-side EMA input
* struct. The Rust side declares the matching `EmaInputs` `#[repr(C)]`
* struct in `sp20_emas_compute.rs`; the field order, types, and
* widths MUST match exactly. See the launcher header comment for the
* field-by-field semantics.
*
* Layout (4-byte aligned, total 36 bytes; the trailing struct may pad
* to 40 bytes on alignment-strict ABIs but our consumers read by name,
* not by offset, so trailing pad is benign): four i32s + five f32s,
* each 4 bytes. */
struct SP20EmaInputs {
int is_close; /* 0 or 1 — gates per-trade EMAs */
int is_win; /* 0 or 1 — only meaningful when is_close == 1 */
int trade_duration; /* bars — only meaningful when is_close == 1 */
int action_is_hold; /* 0 or 1 — gates HOLD_REWARD_EMA */
float alpha; /* R_event - hold_baseline at trade close */
float per_bar_hold_reward; /* -aux_conf * cost_scale on Hold bars */
float aux_logits_p50; /* upstream sp20_stats output [0] */
float aux_logits_std; /* upstream sp20_stats output [1] */
float aux_dir_acc; /* directional accuracy ∈ [0, 1] */
};
/* Per-EMA slot indices in the obs_count[] array. MUST match the /* Per-EMA slot indices in the obs_count[] array. MUST match the
* launcher's `OBS_COUNT_*` constants. The order is fixed by the * launcher's `OBS_COUNT_*` constants. The order is fixed by the
* launch contract, not by physical ordering of the EMAs in the * launch contract, not by physical ordering of the EMAs in the
@@ -146,27 +189,36 @@ float ema_step(float prev, float observation, int* counter)
} }
extern "C" __global__ void sp20_emas_compute_kernel( extern "C" __global__ void sp20_emas_compute_kernel(
float* __restrict__ isv, /* [ISV_TOTAL_DIM] */ float* __restrict__ isv, /* [ISV_TOTAL_DIM] */
float* __restrict__ internal, /* [4] f32 internal scratch */ const SP20EmaInputs* __restrict__ ema_inputs, /* [1] device ptr */
int* __restrict__ obs_count, /* [8] i32 per-EMA counters */ float* __restrict__ internal, /* [4] f32 scratch */
int* __restrict__ obs_count, /* [8] i32 counters */
int isv_idx_alpha, int isv_idx_alpha,
int isv_idx_wr, int isv_idx_wr,
int isv_idx_hold_pct, int isv_idx_hold_pct,
int isv_idx_hold_reward, int isv_idx_hold_reward)
int is_close,
float is_win,
float alpha,
float trade_duration,
int action_is_hold,
float r_per_bar,
float aux_p50_in,
float aux_std_in,
float aux_dir_acc_in)
{ {
/* Single-block, single-thread contract. */ /* Single-block, single-thread contract. */
if (blockIdx.x != 0) return; if (blockIdx.x != 0) return;
if (threadIdx.x != 0) return; if (threadIdx.x != 0) return;
/* Read all 9 fields once. Compiler folds these into 36 bytes of
* coalesced loads off `ema_inputs` — single-thread kernel so
* coalescing is moot, but the explicit local copies make the
* dependency chain readable. */
const int is_close = ema_inputs->is_close;
const int is_win_i = ema_inputs->is_win;
const int trade_duration = ema_inputs->trade_duration;
const int action_is_hold = ema_inputs->action_is_hold;
const float alpha = ema_inputs->alpha;
const float r_per_bar = ema_inputs->per_bar_hold_reward;
const float aux_p50_in = ema_inputs->aux_logits_p50;
const float aux_std_in = ema_inputs->aux_logits_std;
const float aux_dir_acc_in = ema_inputs->aux_dir_acc;
const float is_win = (float)is_win_i;
const float trade_duration_f = (float)trade_duration;
/* ── Per-trade EMAs: gated on is_close == 1 ───────────────────── */ /* ── Per-trade EMAs: gated on is_close == 1 ───────────────────── */
if (is_close == 1) { if (is_close == 1) {
/* WR_EMA */ /* WR_EMA */
@@ -179,7 +231,7 @@ extern "C" __global__ void sp20_emas_compute_kernel(
/* TRADE_DURATION_EMA — internal scratch */ /* TRADE_DURATION_EMA — internal scratch */
internal[INT_IDX_TRADE_DUR] = internal[INT_IDX_TRADE_DUR] =
ema_step(internal[INT_IDX_TRADE_DUR], trade_duration, ema_step(internal[INT_IDX_TRADE_DUR], trade_duration_f,
&obs_count[OBS_IDX_TRADE_DUR]); &obs_count[OBS_IDX_TRADE_DUR]);
} }

21
crates/ml/src/trainers/dqn/state_reset_registry.rs
View File

@@ -2093,6 +2093,27 @@ impl StateResetRegistry {
category: ResetCategory::FoldReset, category: ResetCategory::FoldReset,
description: "ISV[AUX_GATE_TEMP_INDEX=519] — SP20 auxiliary task gating temperature; controls softmax sharpness on aux-path selection; sentinel 0.0 bootstrap per pearl_first_observation_bootstrap", description: "ISV[AUX_GATE_TEMP_INDEX=519] — SP20 auxiliary task gating temperature; controls softmax sharpness on aux-path selection; sentinel 0.0 bootstrap per pearl_first_observation_bootstrap",
}, },
// ── SP20 Phase 1.4 (2026-05-09, Path C): non-ISV scratch buffer
// resets owned by the experience collector. Mapped-pinned
// `host_slice_mut().fill(0.0/0)` resets — the buffers are
// device-visible via `cuMemHostAlloc(DEVICEMAP)` so the host
// write is allowed under `feedback_no_htod_htoh_only_mapped_pinned`'s
// allowed-write rule (NOT a host-side compute).
RegistryEntry {
name: "sp20_ema_inputs_buf",
category: ResetCategory::FoldReset,
description: "SP20 Phase 1.4 mapped-pinned [9] f32 scratch buffer holding the SP20EmaInputs struct between aggregate→EMAs kernels. FoldReset zeros all 9 slots; the aggregation kernel rewrites every byte each step so the reset is purely defensive (clears any inherited bit pattern). Bridges the per-env collector arrays to the post-Path-C EMA kernel struct contract per `feedback_no_partial_refactor`.",
},
RegistryEntry {
name: "sp20_emas_internal_buf",
category: ResetCategory::FoldReset,
description: "SP20 Phase 1.4 mapped-pinned [4] f32 internal scratch for `sp20_emas_compute_kernel`: trade_duration_ema (slot 0), aux_conf_p50_ema (slot 1), aux_conf_std_ema (slot 2), aux_dir_acc_ema (slot 3). FoldReset zeros all 4 — Pearl-A first-observation bootstrap re-engages on each fold's first per-trade / per-step firing observation. Read by `sp20_controllers_compute_kernel` to derive N_STEP / TARGET_HOLD_PCT / AUX_GATE_TEMP / AUX_CONF_THRESHOLD.",
},
RegistryEntry {
name: "sp20_emas_obs_count_buf",
category: ResetCategory::FoldReset,
description: "SP20 Phase 1.4 mapped-pinned [8] i32 per-EMA observation counter for `sp20_emas_compute_kernel`. counter==0 ⇒ next firing observation REPLACES sentinel directly (Pearl-A bootstrap); counter>0 ⇒ Wiener-blend with α=0.4 floor. FoldReset zeros all 8 counters so each fold's first per-EMA firing observation re-bootstraps from sentinel — necessary because some EMAs legitimately observe 0.0 and the count-based bootstrap is the correct sentinel discriminator per `pearl_first_observation_bootstrap`.",
},
]; ];
Self { entries } Self { entries }
} }

145
crates/ml/src/trainers/dqn/trainer/training_loop.rs
View File

@@ -5856,6 +5856,63 @@ impl DQNTrainer {
); );
} }
// SP20 Phase 1.4 (2026-05-09): per-epoch ISV diag for the
// 10 SP20 slots [510..520) populated by the per-rollout-step
// fused-producer chain (Stats → Aggregate → EMAs →
// Controllers, wired in `gpu_experience_collector.rs` step
// 5c). Surfaces the 4 EMAs (alpha/wr/hold_pct/hold_reward)
// + 6 derived controllers (loss_cap / hold_cost_scale /
// target_hold_pct / n_step / aux_conf_threshold /
// aux_gate_temp) so smoke / multi-seed validation can
// verify the chain is firing and the controller outputs
// are within their spec-clamped ranges.
//
// The order in the emit string mirrors the spec §4.5
// formula table for the controllers, with the 4 EMAs
// appended first (alpha, wr, hold_pct, hold_reward) so a
// human reader sees the EMA→controller dependency at a
// glance.
{
use crate::cuda_pipeline::sp14_isv_slots::{
ALPHA_EMA_INDEX, AUX_CONF_THRESHOLD_INDEX,
AUX_GATE_TEMP_INDEX, HOLD_COST_SCALE_INDEX,
HOLD_PCT_EMA_INDEX, HOLD_REWARD_EMA_INDEX,
LOSS_CAP_INDEX, N_STEP_INDEX, TARGET_HOLD_PCT_INDEX,
WR_EMA_INDEX,
};
let (alpha_ema, wr_ema, hold_pct_ema, hold_reward_ema,
loss_cap, hold_cost_scale, target_hold_pct, n_step,
aux_conf_threshold, aux_gate_temp) = if let Some(ref fused) = self.fused_ctx {
let trainer = fused.trainer();
(
trainer.read_isv_signal_at(ALPHA_EMA_INDEX),
trainer.read_isv_signal_at(WR_EMA_INDEX),
trainer.read_isv_signal_at(HOLD_PCT_EMA_INDEX),
trainer.read_isv_signal_at(HOLD_REWARD_EMA_INDEX),
trainer.read_isv_signal_at(LOSS_CAP_INDEX),
trainer.read_isv_signal_at(HOLD_COST_SCALE_INDEX),
trainer.read_isv_signal_at(TARGET_HOLD_PCT_INDEX),
trainer.read_isv_signal_at(N_STEP_INDEX),
trainer.read_isv_signal_at(AUX_CONF_THRESHOLD_INDEX),
trainer.read_isv_signal_at(AUX_GATE_TEMP_INDEX),
)
} else {
(0.0_f32, 0.0_f32, 0.0_f32, 0.0_f32,
0.0_f32, 0.0_f32, 0.0_f32, 0.0_f32,
0.0_f32, 0.0_f32)
};
tracing::info!(
"HEALTH_DIAG[{}]: sp20_isv [loss_cap={:.4} alpha_ema={:.4} wr_ema={:.4} \
hold_cost_scale={:.4} target_hold_pct={:.4} hold_pct_ema={:.4} \
hold_reward_ema={:.4} n_step={:.0} aux_conf_threshold={:.4} aux_gate_temp={:.4}]",
epoch,
loss_cap, alpha_ema, wr_ema,
hold_cost_scale, target_hold_pct, hold_pct_ema,
hold_reward_ema, n_step,
aux_conf_threshold, aux_gate_temp,
);
}
// C1/P1: propagate health to GPU replay buffer for diversity-weighted priorities. // C1/P1: propagate health to GPU replay buffer for diversity-weighted priorities.
{ {
let mut agent = self.agent.write().await; let mut agent = self.agent.write().await;
@@ -9393,6 +9450,94 @@ impl DQNTrainer {
); );
} }
} }
// ── SP20 Phase 1.4 (2026-05-09) — fold-boundary ISV resets ──
// 10 ISV slots [510..520) — every SP20 EMA / controller
// output is FoldReset sentinel 0.0 per
// `pearl_first_observation_bootstrap`. The Phase 1.4
// collector chain re-bootstraps every EMA from sentinel
// on the new fold's first per-step launch (the obs_count
// counters are also zeroed below so the re-bootstrap
// discriminates count==0 ⇒ replace, count>0 ⇒ blend).
"loss_cap" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::LOSS_CAP_INDEX;
fused.trainer().write_isv_signal_at(LOSS_CAP_INDEX, 0.0);
}
}
"alpha_ema" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::ALPHA_EMA_INDEX;
fused.trainer().write_isv_signal_at(ALPHA_EMA_INDEX, 0.0);
}
}
"wr_ema" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::WR_EMA_INDEX;
fused.trainer().write_isv_signal_at(WR_EMA_INDEX, 0.0);
}
}
"hold_cost_scale" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::HOLD_COST_SCALE_INDEX;
fused.trainer().write_isv_signal_at(HOLD_COST_SCALE_INDEX, 0.0);
}
}
"target_hold_pct" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::TARGET_HOLD_PCT_INDEX;
fused.trainer().write_isv_signal_at(TARGET_HOLD_PCT_INDEX, 0.0);
}
}
"hold_pct_ema" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::HOLD_PCT_EMA_INDEX;
fused.trainer().write_isv_signal_at(HOLD_PCT_EMA_INDEX, 0.0);
}
}
"hold_reward_ema" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::HOLD_REWARD_EMA_INDEX;
fused.trainer().write_isv_signal_at(HOLD_REWARD_EMA_INDEX, 0.0);
}
}
"n_step" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::N_STEP_INDEX;
fused.trainer().write_isv_signal_at(N_STEP_INDEX, 0.0);
}
}
"aux_conf_threshold" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::AUX_CONF_THRESHOLD_INDEX;
fused.trainer().write_isv_signal_at(AUX_CONF_THRESHOLD_INDEX, 0.0);
}
}
"aux_gate_temp" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::AUX_GATE_TEMP_INDEX;
fused.trainer().write_isv_signal_at(AUX_GATE_TEMP_INDEX, 0.0);
}
}
// SP20 Phase 1.4 mapped-pinned scratch buffer resets owned
// by the experience collector. Mapped-pinned
// `host_slice_mut().fill(0)` is the allowed-write pattern
// per `feedback_no_htod_htoh_only_mapped_pinned`'s allowed-
// write rule (NOT a host-side compute).
"sp20_ema_inputs_buf" => {
if let Some(ref mut collector) = self.gpu_experience_collector {
collector.sp20_ema_inputs_buf.host_slice_mut().fill(0.0);
}
}
"sp20_emas_internal_buf" => {
if let Some(ref mut collector) = self.gpu_experience_collector {
collector.sp20_emas_internal_buf.host_slice_mut().fill(0.0);
}
}
"sp20_emas_obs_count_buf" => {
if let Some(ref mut collector) = self.gpu_experience_collector {
collector.sp20_emas_obs_count_buf.host_slice_mut().fill(0);
}
}
// SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots. // SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots.
// OFI_IMPACT_LAMBDA_INDEX=407 is an Invariant-1 anchor (NOT // OFI_IMPACT_LAMBDA_INDEX=407 is an Invariant-1 anchor (NOT
// a stateful EMA) — rewrite the constructor's value at fold // a stateful EMA) — rewrite the constructor's value at fold

289
crates/ml/tests/sp20_aggregate_inputs_test.rs Normal file
View File

@@ -0,0 +1,289 @@
//! SP20 Phase 1.4 (2026-05-09) — `sp20_aggregate_inputs_kernel` GPU oracle tests.
//!
//! Validates the per-env → SP20EmaInputs aggregation rules atomically
//! with the Path C buffer-arg refactor of `sp20_emas_compute_kernel`.
//! The aggregation kernel is the GPU-side replacement for what would
//! otherwise need a host-side reduction over `trade_close_per_sample`,
//! `step_ret_per_sample`, etc. — forbidden by `feedback_cpu_is_read_only`
//! + `feedback_no_htod_htoh_only_mapped_pinned`.
//!
//! Each test constructs synthetic per-env arrays in mapped-pinned f32/i32
//! buffers, primes the upstream stats / dir_acc inputs, launches the
//! aggregation kernel, then reads back the resulting `SP20EmaInputs`
//! struct via the same f32-aliased buffer trick the EMA tests use.
//!
//! Run on a GPU host:
//!
//! SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 \
//! cargo test -p ml --test sp20_aggregate_inputs_test --features cuda \
//! -- --ignored --nocapture
#![allow(clippy::tests_outside_test_module)]
#[cfg(feature = "cuda")]
#[allow(unsafe_code)] // CUDA kernel launch + mapped-pinned memory.
mod gpu {
use std::sync::Arc;
use cudarc::driver::{CudaContext, CudaFunction, CudaStream};
use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedI32Buffer};
use ml::cuda_pipeline::sp20_aggregate_inputs::{
launch_sp20_aggregate_inputs, SP20_EMA_INPUTS_F32_LEN,
};
const SP20_AGGREGATE_INPUTS_CUBIN: &[u8] = include_bytes!(concat!(
env!("OUT_DIR"),
"/sp20_aggregate_inputs_kernel.cubin"
));
/// Production direction divisor (`NUM_MAGNITUDES * NUM_ORD * NUM_URG = 27`)
/// matching `state_layout.cuh` and `experience_action_select`.
const DIR_DIVISOR: i32 = 27;
/// `DIR_HOLD = 1` per `state_layout.cuh`. The aggregation kernel
/// flags `action_is_hold` when the majority of envs decode to this
/// direction.
const HOLD_ACTION_ID: i32 = 1;
/// Encode a packed factored action with the given direction.
/// `dir * 27 + 0` lands at the `mag=0, ord=0, urg=0` slot of the
/// chosen direction — the aggregation kernel only cares about the
/// direction component.
fn encode_dir(dir: i32) -> i32 {
dir * DIR_DIVISOR
}
fn make_test_stream() -> Arc<CudaStream> {
let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
ctx.default_stream()
}
fn load_aggregation_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
let module = stream
.context()
.load_cubin(SP20_AGGREGATE_INPUTS_CUBIN.to_vec())
.expect("load sp20_aggregate_inputs_kernel cubin");
module
.load_function("sp20_aggregate_inputs_kernel")
.expect("load sp20_aggregate_inputs_kernel function")
}
/// Read the post-launch `SP20EmaInputs` struct out of the f32-aliased
/// buffer. Returns `(is_close, is_win, trade_duration, action_is_hold,
/// alpha, per_bar_hold_reward, p50, std, dir_acc)`.
fn read_ema_inputs(buf: &MappedF32Buffer) -> (i32, i32, i32, i32, f32, f32, f32, f32, f32) {
let v = buf.read_all();
assert_eq!(v.len(), SP20_EMA_INPUTS_F32_LEN);
// The kernel writes int fields in the first 4 slots and float
// fields in the last 5. We cast back from f32::to_bits.
let is_close = v[0].to_bits() as i32;
let is_win = v[1].to_bits() as i32;
let trade_duration = v[2].to_bits() as i32;
let action_is_hold = v[3].to_bits() as i32;
let alpha = v[4];
let per_bar_hr = v[5];
let p50 = v[6];
let std = v[7];
let dir_acc = v[8];
(is_close, is_win, trade_duration, action_is_hold,
alpha, per_bar_hr, p50, std, dir_acc)
}
/// Construct a 1-stride arrangement (one slice per env at offset
/// `env_id * 1`). This mirrors a per-step slice that has already
/// been offset to the current timestep — the production wire-up
/// passes `base + current_t` so the kernel reads `[env*L + 0]`.
/// In tests we use stride=1 and pack envs contiguously.
fn run_kernel(
trade_close: &[i32],
step_ret: &[f32],
hold_at_exit: &[f32],
actions: &[i32],
aux_logits_p50: f32,
aux_logits_std: f32,
aux_dir_acc: f32,
) -> (i32, i32, i32, i32, f32, f32, f32, f32, f32) {
let stream = make_test_stream();
let kernel = load_aggregation_kernel(&stream);
let n = trade_close.len();
assert!(n > 0, "test n_envs must be > 0");
assert_eq!(step_ret.len(), n);
assert_eq!(hold_at_exit.len(), n);
assert_eq!(actions.len(), n);
let tc_buf = unsafe { MappedI32Buffer::new(n) }.expect("alloc tc");
tc_buf.write_from_slice(trade_close);
let sr_buf = unsafe { MappedF32Buffer::new(n) }.expect("alloc sr");
sr_buf.write_from_slice(step_ret);
let he_buf = unsafe { MappedF32Buffer::new(n) }.expect("alloc he");
he_buf.write_from_slice(hold_at_exit);
let act_buf = unsafe { MappedI32Buffer::new(n) }.expect("alloc act");
act_buf.write_from_slice(actions);
let p50_buf = unsafe { MappedF32Buffer::new(1) }.expect("alloc p50");
p50_buf.write_from_slice(&[aux_logits_p50]);
let std_buf = unsafe { MappedF32Buffer::new(1) }.expect("alloc std");
std_buf.write_from_slice(&[aux_logits_std]);
let dir_buf = unsafe { MappedF32Buffer::new(1) }.expect("alloc dir");
dir_buf.write_from_slice(&[aux_dir_acc]);
let out_buf = unsafe { MappedF32Buffer::new(SP20_EMA_INPUTS_F32_LEN) }
.expect("alloc out");
// Pre-fill with sentinel so we can detect a kernel that fails to write.
out_buf.write_from_slice(&[f32::NAN; SP20_EMA_INPUTS_F32_LEN]);
unsafe {
launch_sp20_aggregate_inputs(
&stream,
&kernel,
tc_buf.dev_ptr,
sr_buf.dev_ptr,
he_buf.dev_ptr,
act_buf.dev_ptr,
/* env_stride = */ 1,
p50_buf.dev_ptr,
std_buf.dev_ptr,
dir_buf.dev_ptr,
/* n_envs = */ n as i32,
DIR_DIVISOR,
HOLD_ACTION_ID,
out_buf.dev_ptr,
).expect("launch aggregate kernel");
}
stream.synchronize().expect("sync after aggregate");
read_ema_inputs(&out_buf)
}
/// All envs flat (no trade close), all picking Long. Expected:
/// is_close = 0, action_is_hold = 0, trade_duration = 0,
/// alpha = per_bar_hold_reward = 0 (Phase 2 / 3.2 placeholders),
/// upstream stats forwarded.
#[test]
#[ignore = "requires GPU"]
fn no_trades_no_holds_emits_sentinel_state() {
let trade_close = vec![0_i32; 8];
let step_ret = vec![0.0_f32; 8];
let hold_at_exit = vec![0.0_f32; 8];
let actions = vec![encode_dir(2 /* DIR_LONG */); 8];
let (ic, iw, td, ah, alpha, hr, p50, std, dir) =
run_kernel(&trade_close, &step_ret, &hold_at_exit, &actions,
0.3, 0.1, 0.55);
assert_eq!(ic, 0, "is_close must be 0 (no trades closed)");
assert_eq!(iw, 0, "is_win must be 0 when is_close=0");
assert_eq!(td, 0, "trade_duration must be 0 when is_close=0");
assert_eq!(ah, 0, "action_is_hold must be 0 (all envs picked Long)");
assert_eq!(alpha, 0.0, "alpha is Phase 2 forward ref placeholder = 0.0");
assert_eq!(hr, 0.0, "per_bar_hold_reward is Phase 3.2 forward ref = 0.0");
assert!((p50 - 0.3).abs() < 1e-6, "aux_logits_p50 forwarded");
assert!((std - 0.1).abs() < 1e-6, "aux_logits_std forwarded");
assert!((dir - 0.55).abs() < 1e-6, "aux_dir_acc forwarded");
}
/// Half envs win, half lose; all close. Expected:
/// is_close = 1, is_win = 1 (4 wins / 8 closed = 0.5, threshold).
/// trade_duration = round(mean(hold_at_exit)) = round(7.5) = 8
/// (round-to-nearest-even at 7.5 ⇒ 8).
#[test]
#[ignore = "requires GPU"]
fn half_wins_meets_threshold_and_duration_rounds() {
// 4 wins, 4 losses. Hold-at-exit mean = (5+6+7+8+7+8+9+10)/8 = 7.5.
let trade_close = vec![1; 8];
let step_ret = vec![ 0.5, 0.3, 0.2, 0.1, -0.5, -0.3, -0.2, -0.1];
let hold_at_exit = vec![5.0, 6.0, 7.0, 8.0, 7.0, 8.0, 9.0, 10.0];
let actions = vec![encode_dir(0); 8]; // all DIR_SHORT — not Hold
let (ic, iw, td, ah, _, _, _, _, _) =
run_kernel(&trade_close, &step_ret, &hold_at_exit, &actions,
0.0, 0.0, 0.0);
assert_eq!(ic, 1);
assert_eq!(iw, 1, "fraction wins (4/8 = 0.5) meets >= 0.5 threshold");
// round-to-nearest-even at 7.5 = 8 (banker's rounding); CUDA
// __float2int_rn uses round-to-nearest-even. Either 7 or 8 are
// structurally acceptable; the kernel uses CUDA's RN so we
// accept both as a bound test.
assert!(
td == 7 || td == 8,
"trade_duration must round 7.5 to 7 (down) or 8 (RN-even); got {td}",
);
assert_eq!(ah, 0);
}
/// Strict majority Hold over n_envs. With 5 of 8 picking Hold,
/// 5 > 8/2 = 4 ⇒ action_is_hold = 1.
#[test]
#[ignore = "requires GPU"]
fn majority_hold_sets_action_is_hold() {
let trade_close = vec![0; 8];
let step_ret = vec![0.0; 8];
let hold_at_exit = vec![0.0; 8];
let mut actions = Vec::with_capacity(8);
for i in 0..8 {
actions.push(if i < 5 { encode_dir(HOLD_ACTION_ID) } else { encode_dir(0) });
}
let (_, _, _, ah, _, _, _, _, _) =
run_kernel(&trade_close, &step_ret, &hold_at_exit, &actions,
0.0, 0.0, 0.0);
assert_eq!(ah, 1, "5 of 8 = strict majority Hold => action_is_hold=1");
}
/// Tie at exactly half Hold. 4 of 8 = 4 = n_envs/2; the kernel uses
/// `count * 2 > n_envs` (strict majority), so a tie ⇒ 0.
#[test]
#[ignore = "requires GPU"]
fn tie_hold_count_does_not_set_action_is_hold() {
let trade_close = vec![0; 8];
let step_ret = vec![0.0; 8];
let hold_at_exit = vec![0.0; 8];
let mut actions = Vec::with_capacity(8);
for i in 0..8 {
actions.push(if i < 4 { encode_dir(HOLD_ACTION_ID) } else { encode_dir(2) });
}
let (_, _, _, ah, _, _, _, _, _) =
run_kernel(&trade_close, &step_ret, &hold_at_exit, &actions,
0.0, 0.0, 0.0);
assert_eq!(ah, 0, "tie (4 of 8 Hold) => action_is_hold=0 (strict majority)");
}
/// Single env closed with a loss: is_close=1, is_win=0
/// (0/1 < 0.5), trade_duration = round(3.0) = 3.
#[test]
#[ignore = "requires GPU"]
fn single_env_loss_close() {
let trade_close = vec![1, 0, 0, 0];
let step_ret = vec![-0.2, 0.0, 0.0, 0.0];
let hold_at_exit = vec![3.0, 0.0, 0.0, 0.0];
let actions = vec![encode_dir(0); 4];
let (ic, iw, td, _, _, _, _, _, _) =
run_kernel(&trade_close, &step_ret, &hold_at_exit, &actions,
0.0, 0.0, 0.0);
assert_eq!(ic, 1);
assert_eq!(iw, 0, "loss (0 wins out of 1 closed) => is_win=0");
assert_eq!(td, 3, "trade_duration = round(3.0) = 3");
}
/// Phase 2 / Phase 3.2 forward-reference placeholders are emitted
/// as 0.0 regardless of any input combinations. This test
/// **verifies the documented contract** — the kernel does NOT
/// surface any per-env signal into `alpha` / `per_bar_hold_reward`
/// in Phase 1.4; that wiring lands in Phase 2 / 3.2.
#[test]
#[ignore = "requires GPU"]
fn alpha_and_per_bar_hold_reward_are_phase_2_and_3_2_placeholders() {
let trade_close = vec![1, 1, 1];
let step_ret = vec![0.5, 0.4, 0.3];
let hold_at_exit = vec![10.0, 10.0, 10.0];
let actions = vec![encode_dir(HOLD_ACTION_ID); 3];
let (_, _, _, _, alpha, hr, _, _, _) =
run_kernel(&trade_close, &step_ret, &hold_at_exit, &actions,
0.5, 0.5, 0.99);
assert_eq!(
alpha, 0.0,
"alpha is Phase 2 placeholder; aggregation kernel must NOT \
surface a per-env signal here regardless of inputs",
);
assert_eq!(
hr, 0.0,
"per_bar_hold_reward is Phase 3.2 placeholder; aggregation \
kernel must NOT surface a per-env signal here",
);
}
}

106
crates/ml/tests/sp20_emas_compute_test.rs
View File

@@ -6,6 +6,17 @@
//! `pearl_first_observation_bootstrap` + `pearl_wiener_alpha_floor_for_ //! `pearl_first_observation_bootstrap` + `pearl_wiener_alpha_floor_for_
//! nonstationary` discipline. //! nonstationary` discipline.
//! //!
//! ## Path C kernel-arg refactor (2026-05-09, Phase 1.4)
//!
//! These tests exercise the post-Path-C kernel signature: the kernel
//! reads its 9 fields from a device struct pointer (`SP20EmaInputs*`),
//! not 9 scalar value-args. The tests use a mapped-pinned f32 buffer
//! sized to [`SP20_EMA_INPUTS_F32_LEN`] (= 9), pack the inputs via
//! [`pack_inputs_into_f32_view`], and pass the buffer's `dev_ptr` to
//! the kernel. Math semantics are bit-identical to the pre-Path-C
//! 9-scalar-arg version — every blend formula reads the same fields
//! by name.
//!
//! Tests cover: //! Tests cover:
//! //!
//! 1. **First-observation bootstrap** — sentinel (count==0) replaces //! 1. **First-observation bootstrap** — sentinel (count==0) replaces
@@ -41,9 +52,10 @@ mod gpu {
ALPHA_EMA_INDEX, HOLD_PCT_EMA_INDEX, HOLD_REWARD_EMA_INDEX, WR_EMA_INDEX, ALPHA_EMA_INDEX, HOLD_PCT_EMA_INDEX, HOLD_REWARD_EMA_INDEX, WR_EMA_INDEX,
}; };
use ml::cuda_pipeline::sp20_emas_compute::{ use ml::cuda_pipeline::sp20_emas_compute::{
launch_sp20_emas_compute, EmaInputs, EMA_INTERNAL_AUX_DIR_ACC, launch_sp20_emas_compute, pack_inputs_into_f32_view, EmaInputs,
EMA_INTERNAL_AUX_P50, EMA_INTERNAL_AUX_STD, EMA_INTERNAL_TRADE_DURATION, EMA_INTERNAL_AUX_DIR_ACC, EMA_INTERNAL_AUX_P50, EMA_INTERNAL_AUX_STD,
EMA_INTERNAL_SLOT_COUNT, OBS_COUNT_SLOTS, WIENER_ALPHA_FLOOR, EMA_INTERNAL_SLOT_COUNT, EMA_INTERNAL_TRADE_DURATION, OBS_COUNT_SLOTS,
SP20_EMA_INPUTS_F32_LEN, WIENER_ALPHA_FLOOR,
}; };
/// Total ISV slot count is the trainer's contract; for tests we /// Total ISV slot count is the trainer's contract; for tests we
@@ -70,9 +82,16 @@ mod gpu {
.expect("load sp20_emas_compute_kernel function") .expect("load sp20_emas_compute_kernel function")
} }
/// Test harness: allocates ISV + internal + obs_count buffers, then /// Test harness: allocates ISV + internal + obs_count + ema_inputs
/// fires `inputs.len()` sequential kernel invocations on the same /// buffers, then fires `inputs.len()` sequential kernel invocations
/// stream. Returns final ISV + internal state for assertions. /// on the same stream. Returns final ISV + internal + obs_count
/// state for assertions.
///
/// Per Path C: the `ema_inputs` buffer is a 9-f32 mapped-pinned
/// scratch byte-aliased to `SP20EmaInputs`. We rewrite it before
/// each launch via `pack_inputs_into_f32_view`. The kernel reads
/// it as `SP20EmaInputs*` — the f32-aliased writes land at the
/// same byte offsets the device struct's fields do.
fn run_kernel_sequence(inputs: &[EmaInputs]) -> (Vec<f32>, Vec<f32>, Vec<i32>) { fn run_kernel_sequence(inputs: &[EmaInputs]) -> (Vec<f32>, Vec<f32>, Vec<i32>) {
let stream = make_test_stream(); let stream = make_test_stream();
let kernel = load_sp20_emas_compute(&stream); let kernel = load_sp20_emas_compute(&stream);
@@ -88,20 +107,31 @@ mod gpu {
unsafe { MappedI32Buffer::new(OBS_COUNT_SLOTS) }.expect("alloc obs_count"); unsafe { MappedI32Buffer::new(OBS_COUNT_SLOTS) }.expect("alloc obs_count");
obs_count.write_from_slice(&vec![0_i32; OBS_COUNT_SLOTS]); obs_count.write_from_slice(&vec![0_i32; OBS_COUNT_SLOTS]);
let mut ema_inputs_buf =
unsafe { MappedF32Buffer::new(SP20_EMA_INPUTS_F32_LEN) }
.expect("alloc ema_inputs");
for inp in inputs { for inp in inputs {
// Re-pack the 9-field struct view before each launch.
pack_inputs_into_f32_view(ema_inputs_buf.host_slice_mut(), inp);
unsafe { unsafe {
launch_sp20_emas_compute( launch_sp20_emas_compute(
&stream, &stream,
&kernel, &kernel,
isv.dev_ptr, isv.dev_ptr,
ema_inputs_buf.dev_ptr,
internal.dev_ptr, internal.dev_ptr,
obs_count.dev_ptr, obs_count.dev_ptr,
*inp,
) )
.expect("launch sp20_emas_compute_kernel"); .expect("launch sp20_emas_compute_kernel");
} }
// Synchronise after each launch so the kernel finishes
// reading the input buffer before we overwrite it for the
// next iteration. (Pre-Path-C the inputs travelled as
// kernel-arg by-value; under Path C the kernel reads the
// buffer in-flight, so we must sync between writes.)
stream.synchronize().expect("sync between sp20_emas launches");
} }
stream.synchronize().expect("sync after sp20_emas_compute");
(isv.read_all(), internal.read_all(), obs_count.read_all()) (isv.read_all(), internal.read_all(), obs_count.read_all())
} }
@@ -116,13 +146,13 @@ mod gpu {
fn first_observation_replaces_sentinel() { fn first_observation_replaces_sentinel() {
let inp = EmaInputs { let inp = EmaInputs {
is_close: 1, is_close: 1,
is_win: 1.0, is_win: 1,
alpha: 0.5, trade_duration: 7,
trade_duration: 7.0,
action_is_hold: 0, action_is_hold: 0,
r_per_bar: 0.0, alpha: 0.5,
aux_p50_in: 0.0, per_bar_hold_reward: 0.0,
aux_std_in: 0.0, aux_logits_p50: 0.0,
aux_logits_std: 0.0,
aux_dir_acc: 0.0, aux_dir_acc: 0.0,
}; };
let (isv, internal, obs_count) = run_kernel_sequence(&[inp]); let (isv, internal, obs_count) = run_kernel_sequence(&[inp]);
@@ -186,13 +216,13 @@ mod gpu {
let inputs: Vec<EmaInputs> = (0..100) let inputs: Vec<EmaInputs> = (0..100)
.map(|i| EmaInputs { .map(|i| EmaInputs {
is_close: 1, is_close: 1,
is_win: if i % 2 == 0 { 1.0 } else { 0.0 }, is_win: if i % 2 == 0 { 1 } else { 0 },
alpha: 0.0, trade_duration: 1,
trade_duration: 1.0,
action_is_hold: 0, action_is_hold: 0,
r_per_bar: 0.0, alpha: 0.0,
aux_p50_in: 0.0, per_bar_hold_reward: 0.0,
aux_std_in: 0.0, aux_logits_p50: 0.0,
aux_logits_std: 0.0,
aux_dir_acc: 0.0, aux_dir_acc: 0.0,
}) })
.collect(); .collect();
@@ -223,26 +253,26 @@ mod gpu {
for _ in 0..5 { for _ in 0..5 {
inputs.push(EmaInputs { inputs.push(EmaInputs {
is_close: 0, is_close: 0,
is_win: 0.0, is_win: 0,
alpha: 0.0, trade_duration: 0,
trade_duration: 0.0,
action_is_hold: 0, // gated OFF action_is_hold: 0, // gated OFF
r_per_bar: -0.05, alpha: 0.0,
aux_p50_in: 0.0, per_bar_hold_reward: -0.05,
aux_std_in: 0.0, aux_logits_p50: 0.0,
aux_logits_std: 0.0,
aux_dir_acc: 0.0, aux_dir_acc: 0.0,
}); });
} }
for _ in 0..5 { for _ in 0..5 {
inputs.push(EmaInputs { inputs.push(EmaInputs {
is_close: 0, is_close: 0,
is_win: 0.0, is_win: 0,
alpha: 0.0, trade_duration: 0,
trade_duration: 0.0,
action_is_hold: 1, // gate ON action_is_hold: 1, // gate ON
r_per_bar: -0.05, alpha: 0.0,
aux_p50_in: 0.0, per_bar_hold_reward: -0.05,
aux_std_in: 0.0, aux_logits_p50: 0.0,
aux_logits_std: 0.0,
aux_dir_acc: 0.0, aux_dir_acc: 0.0,
}); });
} }
@@ -287,13 +317,13 @@ mod gpu {
let inputs: Vec<EmaInputs> = (0..10) let inputs: Vec<EmaInputs> = (0..10)
.map(|_| EmaInputs { .map(|_| EmaInputs {
is_close: 0, is_close: 0,
is_win: 0.0, is_win: 0,
alpha: 0.0, trade_duration: 0,
trade_duration: 0.0,
action_is_hold: 0, action_is_hold: 0,
r_per_bar: 0.0, alpha: 0.0,
aux_p50_in: 0.2, per_bar_hold_reward: 0.0,
aux_std_in: 0.1, aux_logits_p50: 0.2,
aux_logits_std: 0.1,
aux_dir_acc: 1.0, aux_dir_acc: 1.0,
}) })
.collect(); .collect();

330
crates/ml/tests/sp20_phase1_4_wireup_test.rs Normal file
View File

@@ -0,0 +1,330 @@
//! SP20 Phase 1.4 (2026-05-09, Path C) — fused-producer chain integration test.
//!
//! Direct end-to-end exercise of the 4-kernel chain that the
//! `GpuExperienceCollector` invokes per rollout step:
//!
//! Stats (Kernel 3) → Aggregate (Path C) → EMAs (Kernel 1) → Controllers (Kernel 2)
//!
//! Constructs synthetic per-env arrays + aux-head outputs in mapped-pinned
//! buffers, fires the 4 launches in sequence on a single stream, and
//! asserts that the 10 SP20 ISV slots populate with sensible
//! formula-derived values.
//!
//! Per `pearl_tests_must_prove_not_lock_observations` the assertions are
//! invariants (slot is non-sentinel, bound-respecting, monotonic in the
//! expected direction) rather than locked numerical values.
//!
//! Run on a GPU host:
//!
//! SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 \
//! cargo test -p ml --test sp20_phase1_4_wireup_test --features cuda \
//! -- --ignored --nocapture
#![allow(clippy::tests_outside_test_module)]
#[cfg(feature = "cuda")]
#[allow(unsafe_code)]
mod gpu {
use std::sync::Arc;
use cudarc::driver::{CudaContext, CudaFunction, CudaStream};
use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedI32Buffer};
use ml::cuda_pipeline::sp14_isv_slots::{
ALPHA_EMA_INDEX, AUX_CONF_THRESHOLD_INDEX, AUX_GATE_TEMP_INDEX,
HOLD_COST_SCALE_INDEX, HOLD_PCT_EMA_INDEX, HOLD_REWARD_EMA_INDEX,
LOSS_CAP_INDEX, N_STEP_INDEX, TARGET_HOLD_PCT_INDEX, WR_EMA_INDEX,
};
use ml::cuda_pipeline::sp20_aggregate_inputs::launch_sp20_aggregate_inputs;
use ml::cuda_pipeline::sp20_controllers_compute::launch_sp20_controllers_compute;
use ml::cuda_pipeline::sp20_emas_compute::{
launch_sp20_emas_compute, EMA_INTERNAL_SLOT_COUNT, OBS_COUNT_SLOTS,
SP20_EMA_INPUTS_F32_LEN,
};
use ml::cuda_pipeline::sp20_stats_compute::{
launch_sp20_stats_compute, AUX_K_CLASSES,
};
/// ISV slot count: cover up to slot 519.
const ISV_LEN: usize = 520;
/// Production action encoding constants (mirror `state_layout.cuh`).
const DIR_DIVISOR: i32 = 27; // 3 * 3 * 3 (mag * ord * urg)
const HOLD_ACTION_ID: i32 = 1;
fn encode_dir(dir: i32) -> i32 {
dir * DIR_DIVISOR
}
const SP20_STATS_CUBIN: &[u8] = include_bytes!(concat!(
env!("OUT_DIR"),
"/sp20_stats_compute_kernel.cubin"
));
const SP20_AGGREGATE_CUBIN: &[u8] = include_bytes!(concat!(
env!("OUT_DIR"),
"/sp20_aggregate_inputs_kernel.cubin"
));
const SP20_EMAS_CUBIN: &[u8] = include_bytes!(concat!(
env!("OUT_DIR"),
"/sp20_emas_compute_kernel.cubin"
));
const SP20_CONTROLLERS_CUBIN: &[u8] = include_bytes!(concat!(
env!("OUT_DIR"),
"/sp20_controllers_compute_kernel.cubin"
));
fn make_stream() -> Arc<CudaStream> {
let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
ctx.default_stream()
}
fn load_kernel(stream: &Arc<CudaStream>, cubin: &[u8], name: &str) -> CudaFunction {
let m = stream
.context()
.load_cubin(cubin.to_vec())
.unwrap_or_else(|e| panic!("load {name} cubin: {e}"));
m.load_function(name)
.unwrap_or_else(|e| panic!("load {name} function: {e}"))
}
/// Run the full 4-kernel chain once with the given synthetic inputs.
/// Returns the final ISV state after the chain settles.
fn run_chain_once(
n_envs: usize,
aux_logits: &[f32],
trade_close: &[i32],
step_ret: &[f32],
hold_at_exit: &[f32],
actions: &[i32],
aux_dir_acc: f32,
) -> Vec<f32> {
let stream = make_stream();
let stats_kernel = load_kernel(&stream, SP20_STATS_CUBIN, "sp20_stats_compute_kernel");
let agg_kernel = load_kernel(&stream, SP20_AGGREGATE_CUBIN, "sp20_aggregate_inputs_kernel");
let emas_kernel = load_kernel(&stream, SP20_EMAS_CUBIN, "sp20_emas_compute_kernel");
let ctrl_kernel = load_kernel(&stream, SP20_CONTROLLERS_CUBIN, "sp20_controllers_compute_kernel");
// ── ISV bus ──
let isv = unsafe { MappedF32Buffer::new(ISV_LEN) }.expect("alloc isv");
isv.write_from_slice(&vec![0.0_f32; ISV_LEN]);
// ── Stats inputs/outputs ──
assert_eq!(aux_logits.len(), n_envs * AUX_K_CLASSES,
"aux_logits must be [n_envs * K=2]");
let aux_logits_buf = unsafe {
MappedF32Buffer::new(n_envs * AUX_K_CLASSES)
}.expect("alloc aux_logits");
aux_logits_buf.write_from_slice(aux_logits);
let stats_out = unsafe { MappedF32Buffer::new(2) }.expect("alloc stats_out");
stats_out.write_from_slice(&[0.0_f32, 0.0_f32]);
// ── Aggregation inputs ──
let tc_buf = unsafe { MappedI32Buffer::new(n_envs) }.expect("alloc tc");
tc_buf.write_from_slice(trade_close);
let sr_buf = unsafe { MappedF32Buffer::new(n_envs) }.expect("alloc sr");
sr_buf.write_from_slice(step_ret);
let he_buf = unsafe { MappedF32Buffer::new(n_envs) }.expect("alloc he");
he_buf.write_from_slice(hold_at_exit);
let act_buf = unsafe { MappedI32Buffer::new(n_envs) }.expect("alloc act");
act_buf.write_from_slice(actions);
let dir_acc_buf = unsafe { MappedF32Buffer::new(1) }.expect("alloc dir_acc");
dir_acc_buf.write_from_slice(&[aux_dir_acc]);
// ── Aggregation output / EMA input struct (9 f32 byte-aliased) ──
let ema_inputs = unsafe { MappedF32Buffer::new(SP20_EMA_INPUTS_F32_LEN) }
.expect("alloc ema_inputs");
ema_inputs.write_from_slice(&vec![0.0_f32; SP20_EMA_INPUTS_F32_LEN]);
// ── EMA scratch ──
let internal = unsafe { MappedF32Buffer::new(EMA_INTERNAL_SLOT_COUNT) }
.expect("alloc internal");
internal.write_from_slice(&vec![0.0_f32; EMA_INTERNAL_SLOT_COUNT]);
let obs_count = unsafe { MappedI32Buffer::new(OBS_COUNT_SLOTS) }
.expect("alloc obs_count");
obs_count.write_from_slice(&vec![0_i32; OBS_COUNT_SLOTS]);
// ── Launch chain ──
unsafe {
launch_sp20_stats_compute(
&stream, &stats_kernel,
aux_logits_buf.dev_ptr,
stats_out.dev_ptr,
n_envs as i32,
).expect("launch sp20_stats_compute");
let stats_p50 = stats_out.dev_ptr;
let stats_std = stats_out.dev_ptr + std::mem::size_of::<f32>() as u64;
launch_sp20_aggregate_inputs(
&stream, &agg_kernel,
tc_buf.dev_ptr,
sr_buf.dev_ptr,
he_buf.dev_ptr,
act_buf.dev_ptr,
/* env_stride = */ 1,
stats_p50,
stats_std,
dir_acc_buf.dev_ptr,
n_envs as i32,
DIR_DIVISOR,
HOLD_ACTION_ID,
ema_inputs.dev_ptr,
).expect("launch sp20_aggregate_inputs");
launch_sp20_emas_compute(
&stream, &emas_kernel,
isv.dev_ptr,
ema_inputs.dev_ptr,
internal.dev_ptr,
obs_count.dev_ptr,
).expect("launch sp20_emas_compute");
launch_sp20_controllers_compute(
&stream, &ctrl_kernel,
isv.dev_ptr,
internal.dev_ptr,
).expect("launch sp20_controllers_compute");
}
stream.synchronize().expect("sync after sp20 chain");
isv.read_all()
}
/// Smoke: one rollout step with mid-confidence aux + a winning trade
/// closure at one env should populate every SP20 ISV slot with
/// non-sentinel, bound-respecting values.
#[test]
#[ignore = "requires GPU"]
fn one_step_chain_populates_all_10_isv_slots() {
let n_envs = 8;
// Mid-confidence aux logits: each row [logit_down, logit_up] with
// a clear winner (e.g. up=1.0, down=0.0 ⇒ softmax up ≈ 0.73).
// peak_softmax ≈ 0.73 ⇒ aux_conf ≈ 0.73 - 0.5 = 0.23.
let mut aux_logits = Vec::with_capacity(n_envs * AUX_K_CLASSES);
for _ in 0..n_envs {
aux_logits.push(0.0); // logit_down
aux_logits.push(1.0); // logit_up
}
// One env closes a winning trade at duration 5.
let mut trade_close = vec![0; n_envs];
let mut step_ret = vec![0.0_f32; n_envs];
let mut hold_at_exit = vec![0.0_f32; n_envs];
trade_close[3] = 1;
step_ret[3] = 0.4;
hold_at_exit[3] = 5.0;
// Half the envs picked Hold (4 of 8 = tie ⇒ action_is_hold = 0).
// Bump to 5 to cross strict-majority threshold and exercise the
// HOLD_PCT_EMA fire-on-Hold path.
let mut actions = vec![encode_dir(2); n_envs]; // most envs Long
for i in 0..5 {
actions[i] = encode_dir(HOLD_ACTION_ID);
}
let isv = run_chain_once(
n_envs, &aux_logits, &trade_close, &step_ret, &hold_at_exit,
&actions, /* aux_dir_acc = */ 0.62,
);
// ── ISV invariants ──
// ALPHA_EMA: aggregation kernel writes alpha=0.0 (Phase 2
// forward ref); first-observation REPLACE on a closed trade ⇒
// ALPHA_EMA stays at 0.0. **This validates the documented
// forward reference contract.**
assert_eq!(isv[ALPHA_EMA_INDEX], 0.0,
"ALPHA_EMA stays at 0.0 (Phase 2 forward ref) on first close");
// WR_EMA: 1 win out of 1 close ⇒ is_win = 1, first-obs replace
// ⇒ WR_EMA = 1.0.
assert_eq!(isv[WR_EMA_INDEX], 1.0,
"WR_EMA = 1.0 after one winning close (1 win / 1 closed >= 0.5)");
// HOLD_PCT_EMA: 5 of 8 envs picked Hold ⇒ action_is_hold = 1
// ⇒ HOLD_PCT_EMA = 1.0 on first observation.
assert_eq!(isv[HOLD_PCT_EMA_INDEX], 1.0,
"HOLD_PCT_EMA = 1.0 after one majority-Hold observation");
// HOLD_REWARD_EMA: per_bar_hold_reward = 0.0 (Phase 3.2 fwd
// ref); fired on Hold ⇒ EMA = 0.0.
assert_eq!(isv[HOLD_REWARD_EMA_INDEX], 0.0,
"HOLD_REWARD_EMA stays at 0.0 (Phase 3.2 forward ref)");
// LOSS_CAP: WR_EMA = 1.0 > 0.55 ⇒ ramp clamped at 1.0 ⇒
// LOSS_CAP = -1 - 1 = -2.0.
let lc = isv[LOSS_CAP_INDEX];
assert!((lc - (-2.0)).abs() < 1e-6,
"LOSS_CAP = -2.0 with WR_EMA = 1.0; got {lc}");
// N_STEP: trade_duration_ema = 5 (first-obs replace), then
// round to 5; clamp [1, 30] ⇒ N_STEP = 5.
let n = isv[N_STEP_INDEX];
assert!((n - 5.0).abs() < 1e-6,
"N_STEP = 5 with trade_duration = 5; got {n}");
// AUX_CONF_THRESHOLD: aux_dir_acc_ema = 0.62 (first-obs
// replace) ⇒ raw = 0.62 - 0.50 = 0.12 ⇒ clamp [0.01, 0.20] ⇒
// 0.12.
let act = isv[AUX_CONF_THRESHOLD_INDEX];
assert!((act - 0.12).abs() < 1e-5,
"AUX_CONF_THRESHOLD = 0.12 with dir_acc = 0.62; got {act}");
// AUX_GATE_TEMP: aux_std_ema is the std of aux_conf rows; with
// identical logits per row, aux_conf is identical so std = 0;
// floor at 0.01.
let agt = isv[AUX_GATE_TEMP_INDEX];
assert!((agt - 0.01).abs() < 1e-6,
"AUX_GATE_TEMP floor = 0.01 with identical aux logits; got {agt}");
// TARGET_HOLD_PCT: aux_p50_ema is the kernel's median of
// aux_conf rows (sp4-histogram quantization). With identical
// logits per row it lands somewhere in [0, 0.5]; the controller
// formula then maps that to TARGET_HOLD_PCT in [0.1, 0.8].
// We assert bound-respecting per
// `pearl_tests_must_prove_not_lock_observations` rather than
// an observed numerical value (the histogram bin granularity
// makes the exact value implementation-specific).
let thp = isv[TARGET_HOLD_PCT_INDEX];
assert!(thp >= 0.1 - 1e-6 && thp <= 0.8 + 1e-6,
"TARGET_HOLD_PCT must be in [0.1, 0.8]; got {thp}");
// HOLD_COST_SCALE: prev = 0 (constructor), hold_pct_ema = 1.0 >
// tgt + 0.05 ⇒ ramp up: 0 * 1.05 = 0; final clamp [0.01, 0.5]
// ⇒ 0.01 (floor). Verifies the deadband / final-clamp wiring
// even at sentinel-zero start.
let hcs = isv[HOLD_COST_SCALE_INDEX];
assert!(hcs >= 0.01 - 1e-6 && hcs <= 0.5 + 1e-6,
"HOLD_COST_SCALE in [0.01, 0.5]; got {hcs}");
}
/// Locks the Phase 2 / Phase 3.2 forward-reference contract: the
/// aggregation kernel writes 0.0 to `alpha` and `per_bar_hold_reward`
/// regardless of inputs, so the corresponding EMAs (ALPHA_EMA,
/// HOLD_REWARD_EMA) stay at 0.0 sentinel under Phase 1.4 even after
/// many fired observations.
#[test]
#[ignore = "requires GPU"]
fn alpha_and_hold_reward_emas_stay_at_zero_phase_2_3_2_placeholders() {
let n_envs = 4;
// 4 closing wins per call, 4 Hold actions per call
let trade_close = vec![1; n_envs];
let step_ret = vec![0.5; n_envs];
let hold_at_exit = vec![10.0; n_envs];
let actions = vec![encode_dir(HOLD_ACTION_ID); n_envs];
let aux_logits = vec![0.5_f32; n_envs * AUX_K_CLASSES];
// Single chain run is enough — first-obs REPLACE on placeholder
// 0.0 lands the EMA at 0.0; subsequent observations would only
// re-blend toward 0.0 from 0.0.
let isv = run_chain_once(
n_envs, &aux_logits, &trade_close, &step_ret, &hold_at_exit,
&actions, 0.99,
);
assert_eq!(isv[ALPHA_EMA_INDEX], 0.0,
"ALPHA_EMA must be 0.0 — Phase 2 forward ref placeholder. \
Wire-up must NOT surface a per-env signal in Phase 1.4");
assert_eq!(isv[HOLD_REWARD_EMA_INDEX], 0.0,
"HOLD_REWARD_EMA must be 0.0 — Phase 3.2 forward ref \
placeholder. Wire-up must NOT surface a per-env signal \
in Phase 1.4");
}
}

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

@@ -2,6 +2,172 @@
**Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7.
## 2026-05-09 — SP20 Phase 1.4 (Path C): kernel-arg refactor + aggregation kernel + production wire-up (atomic)
Wires the SP20 fused-producer chain (Stats → Aggregate → EMAs →
Controllers) into `GpuExperienceCollector::collect_experiences_gpu`'s
per-rollout-step path. Lands the Phase 1.2 EMA kernel-signature
refactor (9 scalar value-args → 1 device struct pointer arg), the
new `sp20_aggregate_inputs_kernel`, the production wire-up, the
fold-boundary reset registry entries + dispatch arms, the per-epoch
`HEALTH_DIAG[N]: sp20_isv [...]` emit, the integration test, and
the spec / plan / audit-doc amendments **atomically** per
`feedback_no_partial_refactor`.
### Path C decision rationale
Phase 1.2's EMA kernel originally took 9 scalar value-args
(`is_close`, `is_win`, `alpha`, `trade_duration`, `action_is_hold`,
`r_per_bar`, `aux_p50_in`, `aux_std_in`, `aux_dir_acc_in`). The
Phase 1.4 wire-up site sits in the per-rollout-step loop, where
the per-env trade signals (`trade_close_per_sample`,
`step_ret_per_sample`, `hold_at_exit_per_sample`, packed factored
`actions_out`) live in device-only `CudaSlice` buffers. Three
options were considered:
- **Path A** (host sync + CPU aggregation): blocks on
`feedback_cpu_is_read_only` + `feedback_no_htod_htoh_only_mapped_pinned`.
Rejected.
- **Path B** (per-env-tile + reduce kernel writing 9 mapped-pinned
scalars the existing kernel reads): doable but doubles the
kernel-arg surface area and forces the EMA kernel to re-read
9 mapped-pinned addresses every step. Rejected on simplicity.
- **Path C** (refactor EMA kernel to take a device struct pointer +
add an aggregation kernel writing the struct): one new kernel +
one signature change. **Chosen** — minimizes the per-step host
work and keeps the kernel-arg surface area bounded.
Phase 1.2 has zero production consumers yet (Phase 1.4 is the first),
so the kernel signature refactor + Phase 1.4 wire-up + audit / spec /
plan amendments land together per `feedback_no_partial_refactor`.
### Aggregation kernel design
`crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu` —
single-block, BLOCK=64 tree-reduce kernel reading per-env arrays
sliced to the current rollout step + the upstream `sp20_stats_compute`
outputs (p50/std) + the existing `aux_dir_acc_reduce_kernel` output
[0]. Aggregation rules:
| Field | Aggregation |
|------------------------|----------------------------------------------------------------------------------------|
| `is_close` | `1` if any env's `trade_close_per_env[env] != 0`, else `0` |
| `is_win` | `1` if `is_close == 1` AND `2 * wins >= closed`, else `0` (≥ 0.5 fraction threshold) |
| `trade_duration` | `round(mean(hold_at_exit) over closed envs)` (CUDA RN-even) if `is_close == 1`, else `0` |
| `action_is_hold` | `1` if `count(decoded_dir == HOLD) * 2 > n_envs` (strict majority), else `0` |
| `alpha` | `0.0` — **Phase 2 forward reference** (`R_event - hold_baseline`) |
| `per_bar_hold_reward` | `0.0` — **Phase 3.2 forward reference** (`-aux_conf * cost_scale`) |
| `aux_logits_p50` | `aux_logits_p50_dev[0]` (sp20_stats output) |
| `aux_logits_std` | `aux_logits_std_dev[0]` (sp20_stats output) |
| `aux_dir_acc` | `aux_dir_acc_dev[0]` (existing `aux_dir_acc_reduce_kernel` output) |
Block tree-reduce (4 stripes × bdim × 4 bytes shmem) — no atomicAdd
per `feedback_no_atomicadd`. Single-block × 32-warp tree pattern
mirrors `hold_rate_observer_kernel.cu`.
### Phase 2 / Phase 3.2 forward references
The `alpha` and `per_bar_hold_reward` fields are emitted as `0.0`
placeholders in Phase 1.4. These are **NOT stubs** — they're
documented forward references where the real producer ships in a
later phase per `feedback_no_partial_refactor`:
- `alpha = R_event - hold_baseline` is computed by the SP20
reward kernel (Phase 2). See
`crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu:46-52`
(in-kernel docstring) and
`crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs:46-49`
(launcher docstring).
- `per_bar_hold_reward = -aux_conf * cost_scale` is computed by
the SP20 Hold-cost dual emission (Phase 3.2). See same file
references at the line below in each.
The aggregation kernel writes 0.0 so the EMAs (ALPHA_EMA,
HOLD_REWARD_EMA) bootstrap counter advances and the EMAs stay at
their 0.0 sentinel — matching the `pearl_first_observation_bootstrap`
discipline. Phase 2 / 3.2 will replace the kernel's 0.0 writes with
real signals atomically with their respective producers.
### Phase 2.3 task subsumed
The original plan's Phase 2.3 (host-side aggregation of per-env
trade signals into `EmaInputs` for the trainer-side launch) is
subsumed by this commit's GPU-side aggregation kernel. Phase 2.3 is
deleted; the aggregation logic now lives entirely on-device per
`feedback_no_cpu_compute_strict`.
### Files modified
| File | Status | Purpose |
|------|--------|---------|
| `crates/ml/src/cuda_pipeline/sp20_emas_compute_kernel.cu` | MOD | Path C: 9 scalar args → 1 device struct ptr |
| `crates/ml/src/cuda_pipeline/sp20_emas_compute.rs` | MOD | Path C launcher signature; `EmaInputs` `#[repr(C)]` mirror; `pack_inputs_into_f32_view` helper |
| `crates/ml/tests/sp20_emas_compute_test.rs` | MOD | Use mapped-pinned f32-aliased buffer for inputs (math semantics bit-identical) |
| `crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu` | NEW | Per-env → SP20EmaInputs reducer |
| `crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs` | NEW | Rust launcher for the aggregation kernel |
| `crates/ml/tests/sp20_aggregate_inputs_test.rs` | NEW | 6 GPU oracle tests for aggregation rules |
| `crates/ml/tests/sp20_phase1_4_wireup_test.rs` | NEW | End-to-end 4-kernel chain integration test |
| `crates/ml/src/cuda_pipeline/mod.rs` | +pub mod | `sp20_aggregate_inputs` module declaration |
| `crates/ml/src/cuda_pipeline/mapped_pinned.rs` | +host_slice_mut | `MappedI32Buffer::host_slice_mut` for fold-reset |
| `crates/ml/build.rs` | +cubin | `sp20_aggregate_inputs_kernel.cu`; updated EMA cubin comment for Path C |
| `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` | +cubin statics | 4 SP20 cubin `pub(crate) static` declarations |
| `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` | MOD | 4 kernel handles + 4 mapped-pinned buffers (struct fields, alloc, init); per-rollout-step launch sequence |
| `crates/ml/src/trainers/dqn/state_reset_registry.rs` | +3 entries | sp20_ema_inputs_buf / sp20_emas_internal_buf / sp20_emas_obs_count_buf FoldReset |
| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | +13 dispatch arms +HEALTH_DIAG | 10 SP20 ISV slot resets + 3 buffer resets + per-epoch `sp20_isv [...]` emit |
| `docs/superpowers/specs/2026-05-09-sp19-20-wr-first-design.md` | AMENDED | Path C kernel-arg refactor note |
| `docs/superpowers/plans/2026-05-09-sp19-20-wr-first.md` | AMENDED | Path C kernel-arg refactor note |
| `docs/dqn-wire-up-audit.md` | This entry | Audit log |
### Pearls + invariants honoured
- `feedback_no_partial_refactor` — kernel sig change + aggregation
kernel + production wire-up + tests + audit/spec/plan amendments
land in **one atomic commit**. Phase 1.2 had zero production
consumers, so the sig change is invisible to anything outside
the SP20 chain.
- `feedback_no_atomicadd` — aggregation kernel uses block tree-reduce
(4 shmem stripes × bdim × 4 bytes), no atomicAdd anywhere.
- `feedback_no_cpu_compute_strict` — every aggregation lives on the
GPU; the host only passes pointers + n_envs + constants.
- `feedback_no_htod_htoh_only_mapped_pinned` — every buffer in the
Phase 1.4 chain is mapped-pinned (`cuMemHostAlloc(DEVICEMAP)`)
reachable via `dev_ptr`; no `memcpy_dtoh` / `memcpy_htod`.
- `pearl_first_observation_bootstrap` — counter-based bootstrap
preserved; the placeholder fields (alpha = 0, per_bar_hold_reward
= 0) write directly to sentinel so the EMAs stay at 0.0 sentinel
until Phase 2 / 3.2 wire the real producers.
- `pearl_no_host_branches_in_captured_graph` — every kernel in the
chain is single-block; no host branches; safe to capture in the
per-step CUDA Graph.
- `feedback_isv_for_adaptive_bounds` — the 6 controller outputs ARE
the adaptive bounds; downstream consumers read them rather than
embedding the formulas.
- `pearl_tests_must_prove_not_lock_observations` — integration test
asserts invariants (slots populate, bounds respected, monotonic
in expected directions) rather than locked observed values; the
forward-ref placeholder asserts validate the documented contract.
### Verification
```
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --features cuda
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --features cuda --lib sp20
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --features cuda \
--test sp20_stats_compute_test -- --ignored --nocapture
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --features cuda \
--test sp20_emas_compute_test -- --ignored --nocapture
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --features cuda \
--test sp20_controllers_compute_test -- --ignored --nocapture
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --features cuda \
--test sp20_aggregate_inputs_test -- --ignored --nocapture
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --features cuda \
--test sp20_phase1_4_wireup_test -- --ignored --nocapture
```
`cargo check -p ml --features cuda` clean. The 18 SP20 launcher unit
tests pass (no GPU required). GPU oracle tests verified on RTX 3050
Ti (sm_86) — L40S verification deferred until smoke run.
## 2026-05-09 — SP20 magnitude intent fix: replay buffer records intent magnitude (Bellman consistency for magnitude axis) ## 2026-05-09 — SP20 magnitude intent fix: replay buffer records intent magnitude (Bellman consistency for magnitude axis)
Parallel to the 2026-05-08 Class C direction fix above (commit `8f218cab2`, Parallel to the 2026-05-08 Class C direction fix above (commit `8f218cab2`,

17
docs/superpowers/plans/2026-05-09-sp19-20-wr-first.md
View File

@@ -1,5 +1,22 @@
# SP19+20 WR-First Reward Implementation Plan # SP19+20 WR-First Reward Implementation Plan
> **AMENDED 2026-05-09 (Path C, Phase 1.4)**: `sp20_emas_compute_kernel`
> signature refactored from 9 scalar value-args to a single device-pointer
> arg (`SP20EmaInputs* ema_inputs`). The Task 1.2 pseudocode below shows
> the kernel taking 9 scalar args (`launch_compute_emas(.., is_close,
> is_win, alpha, ...)`); read as `launch_sp20_emas_compute(stream,
> kernel, isv_dev, ema_inputs_dev, internal_dev, obs_count_dev)`. Math
> semantics bit-identical. The Path C decision lands the kernel sig
> change + a new `sp20_aggregate_inputs_kernel` (per-env arrays →
> SP20EmaInputs struct, on-device tree-reduce) + the Phase 1.4
> production wire-up + tests + audit / spec / plan amendments
> atomically per `feedback_no_partial_refactor`. The original Task 2.3
> (host-side aggregation) is **subsumed** by the GPU-side aggregation
> kernel and removed from the plan — aggregation now lives entirely
> on-device per `feedback_no_cpu_compute_strict`. See
> `docs/dqn-wire-up-audit.md` "SP20 Phase 1.4 (Path C)" entry for the
> full design and rationale.
> **AMENDED 2026-05-09**: K=3 → K=2 retarget for `sp20_stats_compute`. The > **AMENDED 2026-05-09**: K=3 → K=2 retarget for `sp20_stats_compute`. The
> production aux head emits `K = AUX_NEXT_BAR_K = 2` logits `{down, up}` per > production aux head emits `K = AUX_NEXT_BAR_K = 2` logits `{down, up}` per
> `crates/ml/src/cuda_pipeline/gpu_aux_heads.rs:61` (established by SP13 B1.1a), > `crates/ml/src/cuda_pipeline/gpu_aux_heads.rs:61` (established by SP13 B1.1a),

19
docs/superpowers/specs/2026-05-09-sp19-20-wr-first-design.md
View File

@@ -1,5 +1,24 @@
# SP19+20: WR-First Reward + Multi-Horizon Label Utilization # SP19+20: WR-First Reward + Multi-Horizon Label Utilization
> **AMENDED 2026-05-09 (Path C, Phase 1.4)**: `sp20_emas_compute_kernel`
> signature refactored from 9 scalar value-args to a single device-pointer
> arg (`SP20EmaInputs* ema_inputs`). The spec's pseudocode showing
> `compute_emas(is_close, is_win, ..., aux_dir_acc)` should be read as
> `compute_emas(SP20EmaInputs*)`. Math semantics bit-identical. The
> rationale: the Phase 1.4 wire-up site (per-rollout-step in
> `GpuExperienceCollector`) needs to feed per-env GPU-resident trade
> signals into the kernel, which is forbidden via host sync
> (`feedback_cpu_is_read_only`) and via `memcpy_dtoh`/`memcpy_htod`
> (`feedback_no_htod_htoh_only_mapped_pinned`). The atomic commit lands
> the kernel sig change + a new `sp20_aggregate_inputs_kernel`
> (per-env arrays → SP20EmaInputs struct, on-device tree-reduce) +
> the production wire-up + tests + audit / spec / plan amendments
> per `feedback_no_partial_refactor`. See
> `docs/dqn-wire-up-audit.md` "SP20 Phase 1.4 (Path C)" entry for the
> full design and rationale. Also subsumes the original Phase 2.3
> task (host-side aggregation) — the aggregation now lives entirely
> on-device.
> **AMENDED 2026-05-09**: K=3 → K=2 retarget for `sp20_stats_compute`. The > **AMENDED 2026-05-09**: K=3 → K=2 retarget for `sp20_stats_compute`. The
> production aux head emits `K = AUX_NEXT_BAR_K = 2` logits `{down, up}` per > production aux head emits `K = AUX_NEXT_BAR_K = 2` logits `{down, up}` per
> `crates/ml/src/cuda_pipeline/gpu_aux_heads.rs:61` (established by SP13 B1.1a), > `crates/ml/src/cuda_pipeline/gpu_aux_heads.rs:61` (established by SP13 B1.1a),