From 235e838422495aef6178487dc53eed56d90c31e8 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sat, 9 May 2026 20:49:07 +0200 Subject: [PATCH] =?UTF-8?q?feat(sp20):=20Phase=201.4=20(Path=20C)=20?= =?UTF-8?q?=E2=80=94=20kernel-arg=20refactor=20+=20aggregation=20kernel=20?= =?UTF-8?q?+=20wire-up?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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) --- crates/ml/build.rs | 34 +- .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 37 ++ .../cuda_pipeline/gpu_experience_collector.rs | 308 ++++++++++++++++ crates/ml/src/cuda_pipeline/mapped_pinned.rs | 17 + crates/ml/src/cuda_pipeline/mod.rs | 11 + .../cuda_pipeline/sp20_aggregate_inputs.rs | 204 +++++++++++ .../sp20_aggregate_inputs_kernel.cu | 259 ++++++++++++++ .../ml/src/cuda_pipeline/sp20_emas_compute.rs | 246 +++++++++---- .../cuda_pipeline/sp20_emas_compute_kernel.cu | 104 ++++-- .../src/trainers/dqn/state_reset_registry.rs | 21 ++ .../src/trainers/dqn/trainer/training_loop.rs | 145 ++++++++ crates/ml/tests/sp20_aggregate_inputs_test.rs | 289 +++++++++++++++ crates/ml/tests/sp20_emas_compute_test.rs | 106 ++++-- crates/ml/tests/sp20_phase1_4_wireup_test.rs | 330 ++++++++++++++++++ docs/dqn-wire-up-audit.md | 166 +++++++++ .../plans/2026-05-09-sp19-20-wr-first.md | 17 + .../2026-05-09-sp19-20-wr-first-design.md | 19 + 17 files changed, 2170 insertions(+), 143 deletions(-) create mode 100644 crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs create mode 100644 crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu create mode 100644 crates/ml/tests/sp20_aggregate_inputs_test.rs create mode 100644 crates/ml/tests/sp20_phase1_4_wireup_test.rs diff --git a/crates/ml/build.rs b/crates/ml/build.rs index ddec9fa6d..014d60284 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -1392,9 +1392,17 @@ fn main() { // re-bootstrap. Wiener-α floor 0.4 hardcoded as a structural // property per `pearl_wiener_alpha_floor_for_nonstationary`. // No atomicAdd, no host branches — captureable in the per- - // step CUDA Graph. Phase 1.4 wires the production launch - // site atomically with Kernel 2 (controllers) per - // `feedback_no_partial_refactor`. + // step CUDA Graph. + // AMENDED 2026-05-09 (Path C, Phase 1.4): kernel signature + // 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 Phase 1.3 (2026-05-09): 6 derived ISV controller // 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 // Kernel 1 + Kernel 3 per `feedback_no_partial_refactor`. "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) diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 486d9a5f5..93f5e65de 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -727,6 +727,43 @@ pub(crate) static SP17_ADVANTAGE_CLIP_BOUND_CUBIN: &[u8] = // `feedback_no_legacy_aliases`. Coupling A (forward feature wire // `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. /// Reads `h_s2 [B, SH2]` + `aux_softmax_diff [1]` and writes /// `[h_s2 | softmax_diff] [B, SH2 + 1]` into `dir_qaux_concat_scratch`. diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 184773da1..ca854f8c4 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -1415,6 +1415,69 @@ pub struct GpuExperienceCollector { /// per-epoch boundary is the trigger; the hot-path action computation /// itself is 100% GPU. 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 { @@ -2098,6 +2161,105 @@ impl GpuExperienceCollector { // deleted atomically with the kernel file. The q_disagreement_update // 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 // forward infrastructure. The aux next-bar head's forward computes // 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 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 + + // ── 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_compute_abs_sharpe_kernel, bitonic_sort_step_kernel, @@ -6114,6 +6286,142 @@ impl GpuExperienceCollector { // `sp15_dd_trajectory_prev_dd_dev_ptr` setter / pointer // 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::() 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::() as u64) + .expect("sp20: t_offset_i32 overflow"); + let t_offset_f32: u64 = + (t as u64).checked_mul(std::mem::size_of::() 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 ──────────────────── // Single-thread kernel: step_counter_gpu[0] += 1. // Runs on the same stream — serialized after env_step, before diff --git a/crates/ml/src/cuda_pipeline/mapped_pinned.rs b/crates/ml/src/cuda_pipeline/mapped_pinned.rs index adb4e9012..1346f63b7 100644 --- a/crates/ml/src/cuda_pipeline/mapped_pinned.rs +++ b/crates/ml/src/cuda_pipeline/mapped_pinned.rs @@ -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 { diff --git a/crates/ml/src/cuda_pipeline/mod.rs b/crates/ml/src/cuda_pipeline/mod.rs index ca0939589..67968a674 100644 --- a/crates/ml/src/cuda_pipeline/mod.rs +++ b/crates/ml/src/cuda_pipeline/mod.rs @@ -109,6 +109,17 @@ pub mod sp20_emas_compute; // 1.4 wires the production launch site atomically with the rest of // the SP20 chain per `feedback_no_partial_refactor`. 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 // crate via `use crate::cuda_pipeline::sp4_wiener_ema::launch_apply_pearls`. // `pearls_ad_update` stays `pub` for the test-oracle path in diff --git a/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs b/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs new file mode 100644 index 000000000..30ee1041e --- /dev/null +++ b/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs @@ -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::() 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); + } +} diff --git a/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu b/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu new file mode 100644 index 000000000..b9caa5e1e --- /dev/null +++ b/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu @@ -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 + +/* 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(); + } +} diff --git a/crates/ml/src/cuda_pipeline/sp20_emas_compute.rs b/crates/ml/src/cuda_pipeline/sp20_emas_compute.rs index 5d133b5d4..719d00aa9 100644 --- a/crates/ml/src/cuda_pipeline/sp20_emas_compute.rs +++ b/crates/ml/src/cuda_pipeline/sp20_emas_compute.rs @@ -19,7 +19,7 @@ //! | EMA | Storage | Gate | Observation source | //! |-----|---------|------|--------------------| //! | `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) | //! | `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` | @@ -36,6 +36,21 @@ //! [`WIENER_ALPHA_FLOOR`] = 0.4 floor per //! `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 //! //! - `pearl_first_observation_bootstrap` — counter-based bootstrap. @@ -45,20 +60,21 @@ //! - `feedback_no_atomicadd` — single-thread kernel, no concurrent //! writes anywhere. //! - `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 -//! `internal` / `obs_count` buffers are mapped-pinned; kernel emits -//! `__threadfence_system()` for PCIe-visible coherence. +//! `internal` / `obs_count` buffers are mapped-pinned; tests fill +//! 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, //! single-thread kernel; safe to capture in the per-step CUDA //! Graph. -//! -//! ## Phase 1.2 wiring scope -//! -//! Kernel + launcher + tests + build entry **only**. The training-loop -//! 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`. +//! - `feedback_no_partial_refactor` — Phase 1.2 lands kernel + +//! launcher + tests + build entry; Phase 1.4 lands Path C kernel- +//! arg refactor + aggregation kernel + production wire-up +//! atomically. /// Wiener-α floor per `pearl_wiener_alpha_floor_for_nonstationary`. NOT /// 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). pub const OBS_COUNT_AUX_DIR_ACC: usize = 7; -/// Per-step inputs to `launch_sp20_emas_compute`. Packs the 9 scalar -/// fields the kernel needs into a single struct so the launcher -/// signature stays manageable and the call sites are -/// readable. +/// Per-step EMA inputs as the kernel reads them. +/// +/// **Path C 2026-05-09**: the kernel now reads these 9 fields from a +/// 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: /// - `is_close`: 1 if a trade closed this step, 0 otherwise. Gates /// the per-trade EMAs (WR, ALPHA, TRADE_DURATION). -/// - `is_win`: 1.0 if the closed trade was a win, 0.0 if loss. -/// Only meaningful when `is_close == 1`. -/// - `alpha`: `R_event - hold_baseline` from the SP20 spec -/// Component 1 reward kernel. Only meaningful when -/// `is_close == 1`. +/// - `is_win`: 1 if the closed trade was a win, 0 otherwise. +/// Only meaningful when `is_close == 1`. Stored as i32 to match +/// the device-struct field width. /// - `trade_duration`: bars in the closed trade. Only meaningful -/// when `is_close == 1`. -/// - `action_is_hold`: 1 if the action this step was Hold, 0 -/// otherwise. Gates the HOLD_REWARD_EMA; also drives the +/// when `is_close == 1`. Stored as i32 (rounded mean across envs +/// in the production aggregation path). +/// - `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. -/// - `r_per_bar`: per-bar Hold reward = `-aux_conf * cost_scale`. -/// Only meaningful when `action_is_hold == 1`. -/// - `aux_p50_in`: per-step `aux_logits_p50` from Phase 1.1's +/// - `alpha`: `R_event - hold_baseline` from the SP20 spec +/// Component 1 reward kernel. Phase 2 wires the real producer; +/// 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). -/// - `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). /// - `aux_dir_acc`: 1.0 if the aux head predicted the right -/// direction at this step, 0.0 otherwise. -#[derive(Copy, Clone, Debug)] +/// direction at this step, 0.0 otherwise. Phase 1.4 reads this +/// 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 is_close: i32, - pub is_win: f32, - pub alpha: f32, - pub trade_duration: f32, - pub action_is_hold: i32, - pub r_per_bar: f32, - pub aux_p50_in: f32, - pub aux_std_in: f32, - pub aux_dir_acc: f32, + pub is_close: i32, + pub is_win: i32, + pub trade_duration: i32, + pub action_is_hold: i32, + pub alpha: f32, + pub per_bar_hold_reward: f32, + pub aux_logits_p50: f32, + pub aux_logits_std: 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::(); + +/// 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. @@ -169,44 +241,43 @@ pub struct EmaInputs { /// allocations: /// - `isv_dev`: ISV bus, at least /// [`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. /// - `obs_count_dev`: [`OBS_COUNT_SLOTS`]-long i32 counters. /// -/// `inputs` is the per-step gather of all 9 scalar EMA observations; -/// see [`EmaInputs`]. -/// /// # 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` / -/// `MappedI32Buffer::dev_ptr`. The caller MUST issue this launch on -/// the SAME stream as the upstream `sp20_stats_compute` (so the -/// stream-ordering invariant guarantees `aux_p50_in` / `aux_std_in` -/// reads see the latest stats output). Stream ordering with the -/// trade-physics kernels that produce `is_close` / `is_win` / etc. is -/// the caller's responsibility — Phase 1.4 lands the production wire -/// site with that ordering enforced. +/// `MappedI32Buffer::dev_ptr` (or the upstream aggregation kernel's +/// output buffer). The caller MUST issue this launch on the SAME +/// stream as the upstream `sp20_stats_compute` AND +/// `sp20_aggregate_inputs_kernel` (so the stream-ordering invariant +/// guarantees `aux_logits_p50` / `aux_logits_std` / aggregated +/// per-env signals are visible to this kernel's reads). Phase 1.4 +/// lands the production wire site with that ordering enforced. pub unsafe fn launch_sp20_emas_compute( stream: &cudarc::driver::CudaStream, kernel: &cudarc::driver::CudaFunction, isv_dev: u64, + ema_inputs_dev: u64, internal_dev: u64, obs_count_dev: u64, - inputs: EmaInputs, ) -> Result<(), crate::MLError> { use cudarc::driver::{LaunchConfig, PushKernelArg}; debug_assert!(isv_dev != 0, "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, "launch_sp20_emas_compute: internal_dev must be a valid device pointer"); debug_assert!(obs_count_dev != 0, "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 // decoupled from SP14 slot-number drift (the contract is enforced @@ -225,21 +296,13 @@ pub unsafe fn launch_sp20_emas_compute( stream .launch_builder(kernel) .arg(&isv_dev) + .arg(&ema_inputs_dev) .arg(&internal_dev) .arg(&obs_count_dev) .arg(&isv_idx_alpha) .arg(&isv_idx_wr) .arg(&isv_idx_hold_pct) .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) .map_err(|e| { crate::MLError::ModelError(format!("sp20_emas_compute_kernel launch: {e}")) @@ -290,21 +353,11 @@ mod tests { #[test] 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 // 0 — equivalent to "no firing this step" for all 8 EMAs. // This is the canonical pre-warmup / cold-start input shape. - let inp = EmaInputs { - 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, - }; + let inp = EmaInputs::default(); // Per-step EMAs still fire on a zero-input call, but the // observation is 0 so the EMA Wiener-blends toward 0 — that // 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.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); + } } diff --git a/crates/ml/src/cuda_pipeline/sp20_emas_compute_kernel.cu b/crates/ml/src/cuda_pipeline/sp20_emas_compute_kernel.cu index 1ef1369b1..4ef597d7a 100644 --- a/crates/ml/src/cuda_pipeline/sp20_emas_compute_kernel.cu +++ b/crates/ml/src/cuda_pipeline/sp20_emas_compute_kernel.cu @@ -9,6 +9,21 @@ * consumes to derive `LOSS_CAP`, `HOLD_COST_SCALE`, `TARGET_HOLD_PCT`, * `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): * * Per-trade (gated on is_close == 1): @@ -47,17 +62,20 @@ * - `feedback_no_atomicadd` — single-thread kernel; no concurrent * writes anywhere, atomicAdd never used. * - `feedback_no_htod_htoh_only_mapped_pinned` — both ISV (mapped- - * pinned) and `internal` / `obs_count` (this commit's two new - * mapped-pinned buffers) are device-visible via DEVICEMAP; kernel - * emits `__threadfence_system()` after the writes for PCIe-visible + * pinned), `internal` / `obs_count` (mapped-pinned), and + * `ema_inputs` (mapped-pinned struct buffer or upstream-kernel- + * written device buffer) are device-visible; kernel emits + * `__threadfence_system()` after the writes for PCIe-visible * coherence. * - `feedback_no_cpu_compute_strict` — every blend lives in this - * kernel; the launcher's only host work is to package the per-step - * scalar inputs into kernel arguments. + * kernel; the launcher's only host work is to pass pointers. The + * `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 - * + tests + build entry atomically. The production wire-up (Phase - * 1.4) consumes the launcher and the audit doc updates per - * Invariant 7. + * + tests + build entry atomically. Phase 1.4 wires the production + * consumer (Path C kernel-arg refactor + aggregation kernel + + * collector launch site + audit/spec/plan amendments) atomically + * in one commit. * * Algorithm (single block, single thread): * @@ -83,6 +101,13 @@ * Args: * isv — [ISV_TOTAL_DIM] f32 mapped-pinned dev ptr (slots * 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 * scratch slots: TRADE_DUR, AUX_P50, AUX_STD, * AUX_DIR_ACC). Layout matches launcher constants. @@ -93,14 +118,32 @@ * — ISV slot indices for the 4 ISV EMAs (caller passes * the SP14 constants so the kernel stays decoupled * 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 +/* 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 * launcher's `OBS_COUNT_*` constants. The order is fixed by 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( - float* __restrict__ isv, /* [ISV_TOTAL_DIM] */ - float* __restrict__ internal, /* [4] f32 internal scratch */ - int* __restrict__ obs_count, /* [8] i32 per-EMA counters */ + float* __restrict__ isv, /* [ISV_TOTAL_DIM] */ + const SP20EmaInputs* __restrict__ ema_inputs, /* [1] device ptr */ + float* __restrict__ internal, /* [4] f32 scratch */ + int* __restrict__ obs_count, /* [8] i32 counters */ int isv_idx_alpha, int isv_idx_wr, int isv_idx_hold_pct, - 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) + int isv_idx_hold_reward) { /* Single-block, single-thread contract. */ if (blockIdx.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 ───────────────────── */ if (is_close == 1) { /* WR_EMA */ @@ -179,7 +231,7 @@ extern "C" __global__ void sp20_emas_compute_kernel( /* TRADE_DURATION_EMA — internal scratch */ 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]); } diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index ef398e5f2..44187bfdb 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -2093,6 +2093,27 @@ impl StateResetRegistry { 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", }, + // ── 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 } } diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index dac23eee4..309747a73 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -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. { 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. // OFI_IMPACT_LAMBDA_INDEX=407 is an Invariant-1 anchor (NOT // a stateful EMA) — rewrite the constructor's value at fold diff --git a/crates/ml/tests/sp20_aggregate_inputs_test.rs b/crates/ml/tests/sp20_aggregate_inputs_test.rs new file mode 100644 index 000000000..d384cbb14 --- /dev/null +++ b/crates/ml/tests/sp20_aggregate_inputs_test.rs @@ -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 { + let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?"); + ctx.default_stream() + } + + fn load_aggregation_kernel(stream: &Arc) -> 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", + ); + } +} diff --git a/crates/ml/tests/sp20_emas_compute_test.rs b/crates/ml/tests/sp20_emas_compute_test.rs index 88fded482..79b5cd180 100644 --- a/crates/ml/tests/sp20_emas_compute_test.rs +++ b/crates/ml/tests/sp20_emas_compute_test.rs @@ -6,6 +6,17 @@ //! `pearl_first_observation_bootstrap` + `pearl_wiener_alpha_floor_for_ //! 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: //! //! 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, }; use ml::cuda_pipeline::sp20_emas_compute::{ - launch_sp20_emas_compute, EmaInputs, EMA_INTERNAL_AUX_DIR_ACC, - EMA_INTERNAL_AUX_P50, EMA_INTERNAL_AUX_STD, EMA_INTERNAL_TRADE_DURATION, - EMA_INTERNAL_SLOT_COUNT, OBS_COUNT_SLOTS, WIENER_ALPHA_FLOOR, + launch_sp20_emas_compute, pack_inputs_into_f32_view, EmaInputs, + EMA_INTERNAL_AUX_DIR_ACC, EMA_INTERNAL_AUX_P50, EMA_INTERNAL_AUX_STD, + 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 @@ -70,9 +82,16 @@ mod gpu { .expect("load sp20_emas_compute_kernel function") } - /// Test harness: allocates ISV + internal + obs_count buffers, then - /// fires `inputs.len()` sequential kernel invocations on the same - /// stream. Returns final ISV + internal state for assertions. + /// Test harness: allocates ISV + internal + obs_count + ema_inputs + /// buffers, then fires `inputs.len()` sequential kernel invocations + /// 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, Vec, Vec) { let stream = make_test_stream(); let kernel = load_sp20_emas_compute(&stream); @@ -88,20 +107,31 @@ mod gpu { unsafe { MappedI32Buffer::new(OBS_COUNT_SLOTS) }.expect("alloc obs_count"); 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 { + // Re-pack the 9-field struct view before each launch. + pack_inputs_into_f32_view(ema_inputs_buf.host_slice_mut(), inp); unsafe { launch_sp20_emas_compute( &stream, &kernel, isv.dev_ptr, + ema_inputs_buf.dev_ptr, internal.dev_ptr, obs_count.dev_ptr, - *inp, ) .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()) } @@ -116,13 +146,13 @@ mod gpu { fn first_observation_replaces_sentinel() { let inp = EmaInputs { is_close: 1, - is_win: 1.0, - alpha: 0.5, - trade_duration: 7.0, + is_win: 1, + trade_duration: 7, action_is_hold: 0, - r_per_bar: 0.0, - aux_p50_in: 0.0, - aux_std_in: 0.0, + alpha: 0.5, + per_bar_hold_reward: 0.0, + aux_logits_p50: 0.0, + aux_logits_std: 0.0, aux_dir_acc: 0.0, }; let (isv, internal, obs_count) = run_kernel_sequence(&[inp]); @@ -186,13 +216,13 @@ mod gpu { let inputs: Vec = (0..100) .map(|i| EmaInputs { is_close: 1, - is_win: if i % 2 == 0 { 1.0 } else { 0.0 }, - alpha: 0.0, - trade_duration: 1.0, + is_win: if i % 2 == 0 { 1 } else { 0 }, + trade_duration: 1, action_is_hold: 0, - r_per_bar: 0.0, - aux_p50_in: 0.0, - aux_std_in: 0.0, + alpha: 0.0, + per_bar_hold_reward: 0.0, + aux_logits_p50: 0.0, + aux_logits_std: 0.0, aux_dir_acc: 0.0, }) .collect(); @@ -223,26 +253,26 @@ mod gpu { for _ in 0..5 { inputs.push(EmaInputs { is_close: 0, - is_win: 0.0, - alpha: 0.0, - trade_duration: 0.0, + is_win: 0, + trade_duration: 0, action_is_hold: 0, // gated OFF - r_per_bar: -0.05, - aux_p50_in: 0.0, - aux_std_in: 0.0, + alpha: 0.0, + per_bar_hold_reward: -0.05, + aux_logits_p50: 0.0, + aux_logits_std: 0.0, aux_dir_acc: 0.0, }); } for _ in 0..5 { inputs.push(EmaInputs { is_close: 0, - is_win: 0.0, - alpha: 0.0, - trade_duration: 0.0, + is_win: 0, + trade_duration: 0, action_is_hold: 1, // gate ON - r_per_bar: -0.05, - aux_p50_in: 0.0, - aux_std_in: 0.0, + alpha: 0.0, + per_bar_hold_reward: -0.05, + aux_logits_p50: 0.0, + aux_logits_std: 0.0, aux_dir_acc: 0.0, }); } @@ -287,13 +317,13 @@ mod gpu { let inputs: Vec = (0..10) .map(|_| EmaInputs { is_close: 0, - is_win: 0.0, - alpha: 0.0, - trade_duration: 0.0, + is_win: 0, + trade_duration: 0, action_is_hold: 0, - r_per_bar: 0.0, - aux_p50_in: 0.2, - aux_std_in: 0.1, + alpha: 0.0, + per_bar_hold_reward: 0.0, + aux_logits_p50: 0.2, + aux_logits_std: 0.1, aux_dir_acc: 1.0, }) .collect(); diff --git a/crates/ml/tests/sp20_phase1_4_wireup_test.rs b/crates/ml/tests/sp20_phase1_4_wireup_test.rs new file mode 100644 index 000000000..d8574116c --- /dev/null +++ b/crates/ml/tests/sp20_phase1_4_wireup_test.rs @@ -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 { + let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?"); + ctx.default_stream() + } + + fn load_kernel(stream: &Arc, 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 { + 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::() 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"); + } +} diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 7a0ad7ca5..7118c53ce 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,172 @@ **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) Parallel to the 2026-05-08 Class C direction fix above (commit `8f218cab2`, diff --git a/docs/superpowers/plans/2026-05-09-sp19-20-wr-first.md b/docs/superpowers/plans/2026-05-09-sp19-20-wr-first.md index 434248bfd..314f485ca 100644 --- a/docs/superpowers/plans/2026-05-09-sp19-20-wr-first.md +++ b/docs/superpowers/plans/2026-05-09-sp19-20-wr-first.md @@ -1,5 +1,22 @@ # 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 > 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), diff --git a/docs/superpowers/specs/2026-05-09-sp19-20-wr-first-design.md b/docs/superpowers/specs/2026-05-09-sp19-20-wr-first-design.md index e01e356f8..456d4e1ba 100644 --- a/docs/superpowers/specs/2026-05-09-sp19-20-wr-first-design.md +++ b/docs/superpowers/specs/2026-05-09-sp19-20-wr-first-design.md @@ -1,5 +1,24 @@ # 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 > 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),