From 2e84fd35d48c2cd8b91d924b37e1717a5f4abbde Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sat, 2 May 2026 00:19:36 +0200 Subject: [PATCH] =?UTF-8?q?feat(sp5):=20Task=20A8=20=E2=80=94=20Pearl=201-?= =?UTF-8?q?ext=20per-branch=20num=5Fatoms=20=E2=80=94=20Layer=20A=20comple?= =?UTF-8?q?te?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Final SP5 Layer A producer. Per-branch C51 num_atoms derived from per-branch ATOM_V_HALF (Pearl 1, Task A1). Threshold cascade: v_half < 0.1 → 64 atoms (narrow Q, high resolution) v_half < 1.0 → 32 atoms (moderate) v_half ≥ 1.0 → 16 atoms (wide Q, modest resolution) 4 ISV slots [ATOM_NUM_ATOMS_BASE=274..278). Pearls A+D smooths the discrete output during transitions; Layer B's atoms_update consumer rounds to nearest valid count. producer_step_scratch_buf grew 203 → 207. wiener_state_buf already at 543 (sized at A1 for entire SP5 block). StateResetRegistry: 1 new FoldReset entry (sp5_atom_num_atoms). atoms_update consumer migration deferred to Layer B. LAYER A COMPLETE. 8 producers + 3 auxiliary kernels (q_branch_stats, grad_cosine_sim, q_skew_kurtosis) populating 110 ISV slots [174..286) with 4 cross-fold-persistent slots carved out (Kelly). Refs: SP5 spec 6e6e0fa11 line 1020 Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/build.rs | 7 + .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 112 +++++++++++- .../pearl_1_ext_num_atoms_kernel.cu | 56 ++++++ .../src/trainers/dqn/state_reset_registry.rs | 9 + .../src/trainers/dqn/trainer/training_loop.rs | 7 + crates/ml/tests/sp5_producer_unit_tests.rs | 159 ++++++++++++++++++ docs/dqn-wire-up-audit.md | 2 + 7 files changed, 350 insertions(+), 2 deletions(-) create mode 100644 crates/ml/src/cuda_pipeline/pearl_1_ext_num_atoms_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 9332b415c..131c53077 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -473,6 +473,13 @@ fn main() { // ATR tracking deferred to Layer C). Writes 4 floats to scratch_buf[199..203). // Followed by 4 apply_pearls_ad_kernel calls → ISV[TRAIL_DIST_PER_DIR_BASE=270..274). "pearl_8_trail_kernel.cu", + // SP5 Task A8 (2026-05-01): Pearl 1-ext per-branch C51 num_atoms from per-branch v_half. + // 4-thread single-block kernel (one per branch). Reads ISV[ATOM_V_HALF_BASE=178..182) + // (populated by Pearl 1 in Task A1). Threshold cascade: v_half < 0.1 → 64 atoms, + // v_half < 1.0 → 32 atoms, v_half >= 1.0 → 16 atoms. + // Writes 4 floats to scratch_buf[203..207) = scratch[SCRATCH_PEARL_1_EXT_NUM_ATOMS]. + // Followed by 4 apply_pearls_ad_kernel calls → ISV[ATOM_NUM_ATOMS_BASE=274..278). + "pearl_1_ext_num_atoms_kernel.cu", // SP4 GPU-only Pearls A+D applicator (2026-05-01). Single-thread // device-side loop applies Pearls A+D to N consecutive ISV slots // (each with its own Wiener-state triple) on the producer's diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index a45016ca7..8d7bfd9ee 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -383,6 +383,18 @@ static SP5_PEARL_6_KELLY_CUBIN: &[u8] = static SP5_PEARL_8_TRAIL_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/pearl_8_trail_kernel.cubin")); +/// SP5 Task A8 (2026-05-01): Pearl 1-ext per-branch C51 num_atoms from per-branch v_half. +/// 4-thread single-block kernel (one per branch). Reads ISV[ATOM_V_HALF_BASE=178..182) +/// (populated by Pearl 1 in Task A1). Threshold cascade: v_half < 0.1 → 64 atoms (narrow +/// Q range, high resolution), v_half < 1.0 → 32 atoms (moderate), v_half ≥ 1.0 → 16 atoms +/// (wide Q range, modest resolution). Atom counts are powers of 2; total cap 4×64=256 atoms +/// is within the HW shared-memory budget the C51 backward kernel uses (Invariant 1 anchors). +/// Output smoothed by Pearls A+D; Layer B's atoms_update consumer rounds at consume time. +/// Writes 4 floats to scratch_buf[SCRATCH_PEARL_1_EXT_NUM_ATOMS=203..207). +/// Followed by 4 apply_pearls_ad_kernel calls → ISV[ATOM_NUM_ATOMS_BASE=274..278). +static SP5_PEARL_1_EXT_NUM_ATOMS_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/pearl_1_ext_num_atoms_kernel.cubin")); + /// Plan 4 Task 3 (E.3): IQN multi-quantile diagnostic EMAs into ISV[99..103). /// 4-block (256 threads/block, shmem-reduce, no atomicAdd) kernel launched /// alongside `h_s2_rms_ema_update` from `training_loop.rs`. Reads the IQN @@ -992,8 +1004,10 @@ pub const SP5_WIENER_TOTAL_FLOATS: usize = /// 20 floats for pearl_5_iqn_tau τ output (SCRATCH_PEARL_5_TAU=179, [4 branches × 5 quantiles] at [179..199)) /// Task A7 adds: /// 4 floats for pearl_8_trail output (SCRATCH_PEARL_8_TRAIL=199, trail_dist × 4 directions at [199..203)) -/// Combined: 71 + 16 + 16 + 4 + 4 + 20 + 8 + 8 + 8 + 8 + 8 + 4 + 4 + 20 + 4 = 203 scratch slots [0..203). -pub const SP5_SCRATCH_TOTAL: usize = 203; +/// Task A8 adds: +/// 4 floats for pearl_1_ext_num_atoms output (SCRATCH_PEARL_1_EXT_NUM_ATOMS=203, num_atoms × 4 branches at [203..207)) +/// Combined: 71 + 16 + 16 + 4 + 4 + 20 + 8 + 8 + 8 + 8 + 8 + 4 + 4 + 20 + 4 + 4 = 207 scratch slots [0..207). +pub const SP5_SCRATCH_TOTAL: usize = 207; /// SP5 Task A7: scratch index base for pearl_8_trail_update trail_dist[4] output block. /// Slots [199..203): per-direction trail-stop distance (Short=0, Hold=1, Long=2, Flat=3). @@ -1001,6 +1015,13 @@ pub const SP5_SCRATCH_TOTAL: usize = 203; /// ISV[TRAIL_DIST_PER_DIR_BASE=270..274). pub const SCRATCH_PEARL_8_TRAIL: usize = 199; +/// SP5 Task A8: scratch index base for pearl_1_ext_num_atoms_update num_atoms[4] output block. +/// Slots [203..207): per-branch C51 num_atoms discrete value {64, 32, 16} (as f32). +/// Written by `pearl_1_ext_num_atoms_update`; consumed by apply_pearls_ad_kernel → +/// ISV[ATOM_NUM_ATOMS_BASE=274..278). Pearls A+D smooths the discrete output during +/// transitions; Layer B's atoms_update consumer rounds to nearest valid count at consume time. +pub const SCRATCH_PEARL_1_EXT_NUM_ATOMS: usize = 203; + /// SP5 Task A5: scratch index base for q_skew_kurtosis skew[4] output block. /// Slots [171..175): per-branch Q-distribution standardized skewness. /// Written by `q_skew_kurtosis_update`; consumed by `pearl_5_iqn_tau_update`. @@ -4279,6 +4300,14 @@ pub struct GpuDqnTrainer { /// ISV[TRAIL_DIST_PER_DIR_BASE=270..274). In fold-reset registry (FoldReset). /// Loaded from `pearl_8_trail_kernel.cubin`. pearl_8_trail_kernel: CudaFunction, + /// SP5 Task A8 (2026-05-01): Pearl 1-ext per-branch C51 num_atoms kernel. + /// 4-thread single-block kernel. Reads ISV[ATOM_V_HALF_BASE=178..182) (Pearl 1, Task A1); + /// maps to discrete {64, 32, 16} via threshold cascade (thresholds 0.1, 1.0 are Invariant 1 + /// anchors). Writes 4 floats to scratch_buf[SCRATCH_PEARL_1_EXT_NUM_ATOMS=203..207). + /// Chained apply_pearls_ad_kernel (4 launches, ALPHA_META=1e-3) → + /// ISV[ATOM_NUM_ATOMS_BASE=274..278). In fold-reset registry (FoldReset). + /// Loaded from `pearl_1_ext_num_atoms_kernel.cubin`. + pearl_1_ext_num_atoms_kernel: CudaFunction, /// SP5 Task A4 (2026-05-01): previous-step gradient buffer for cosine similarity. /// [total_params f32] mapped-pinned (feedback_no_htod_htoh_only_mapped_pinned). /// Zeroed at construction and at each fold boundary (Pearl A sentinel: cosine_sim=0 @@ -11373,6 +11402,73 @@ impl GpuDqnTrainer { Ok(()) } + /// SP5 Pearl 1-ext: per-branch C51 num_atoms from per-branch v_half (Pearl 1). + /// + /// Reads ATOM_V_HALF[178..182) (populated by Pearl 1 in A1), maps to discrete + /// {64, 32, 16} via threshold cascade: + /// v_half < 0.1 → 64 atoms (narrow Q range, high resolution) + /// v_half < 1.0 → 32 atoms (moderate Q range) + /// v_half ≥ 1.0 → 16 atoms (wide Q range, modest resolution) + /// Pearls A+D smooths the discrete output during transitions; Layer B's + /// atoms_update consumer rounds to the nearest valid count at consume time. + /// + /// Single producer kernel + 4 apply_pearls calls (one per branch). + /// Wiener offset formula: `SP4_PRODUCER_COUNT * 3 + (isv_slot - SP5_SLOT_BASE) * 3`. + pub(crate) fn launch_sp5_pearl_1_ext_num_atoms(&self) -> Result<(), MLError> { + use crate::cuda_pipeline::sp5_isv_slots::{ATOM_V_HALF_BASE, ATOM_NUM_ATOMS_BASE, SP5_SLOT_BASE}; + use crate::cuda_pipeline::sp4_wiener_ema::{launch_apply_pearls, ALPHA_META}; + + debug_assert!(self.isv_signals_dev_ptr != 0, + "launch_sp5_pearl_1_ext_num_atoms: isv_signals_dev_ptr must be allocated"); + + let scratch_dev = self.producer_step_scratch_buf.dev_ptr; + let isv_dev = self.isv_signals_dev_ptr; + let wiener_dev = self.wiener_state_buf.dev_ptr; + let v_half_base_i32 = ATOM_V_HALF_BASE as i32; + let num_atoms_base_i32 = SCRATCH_PEARL_1_EXT_NUM_ATOMS as i32; + + // Wiener base offset for SP5 slots: SP4_PRODUCER_COUNT * 3 = 213. + let base_wiener_offset: i32 = (SP4_PRODUCER_COUNT * 3) as i32; + + // Step 1: pearl_1_ext_num_atoms_update — reads ISV[ATOM_V_HALF_BASE..+4]; + // writes num_atoms[4 branches] → scratch[203..207). + unsafe { + self.stream + .launch_builder(&self.pearl_1_ext_num_atoms_kernel) + .arg(&isv_dev) + .arg(&v_half_base_i32) + .arg(&scratch_dev) + .arg(&num_atoms_base_i32) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (4, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("pearl_1_ext_num_atoms_update: {e}")))?; + } + + // Step 2: 4 apply_pearls calls (one per branch). + // branch b: scratch[203 + b] → ISV[ATOM_NUM_ATOMS_BASE + b]. + for b in 0..4_usize { + let isv_slot = (ATOM_NUM_ATOMS_BASE + b) as i32; + let scratch_idx = (SCRATCH_PEARL_1_EXT_NUM_ATOMS + b) as i32; + let wiener_off = base_wiener_offset + (isv_slot - SP5_SLOT_BASE as i32) * 3; + unsafe { + launch_apply_pearls( + &self.stream, + &self.apply_pearls_ad_kernel, + scratch_dev, scratch_idx, + isv_dev, isv_slot, + wiener_dev, wiener_off, + 1, + ALPHA_META, + )?; + } + } + + Ok(()) + } + /// Plan C Phase 2 follow-up A.2 (2026-04-29): launch `q_drift_rate_ema_update` /// — single-thread single-block ISV producer for ISV[Q_DRIFT_RATE_INDEX=129]. /// @@ -13478,6 +13574,17 @@ impl GpuDqnTrainer { .map_err(|e| MLError::ModelError(format!("pearl_8_trail_update load: {e}")))? }; + // SP5 Task A8 (2026-05-01): load pearl_1_ext_num_atoms_kernel. + // 4-thread single-block kernel. Reads ISV[ATOM_V_HALF_BASE=178..182) (Pearl 1, Task A1); + // writes num_atoms[4 branches] to scratch[SCRATCH_PEARL_1_EXT_NUM_ATOMS=203..207). + // FoldReset: ISV[274..278) zeroed at fold boundary (Pearl A sentinel). + let pearl_1_ext_num_atoms_kernel = { + let module = stream.context().load_cubin(SP5_PEARL_1_EXT_NUM_ATOMS_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("sp5 pearl_1_ext_num_atoms cubin load: {e}")))?; + module.load_function("pearl_1_ext_num_atoms_update") + .map_err(|e| MLError::ModelError(format!("pearl_1_ext_num_atoms_update load: {e}")))? + }; + // SP5 Task A4 (2026-05-01): allocate grad_prev_buf_per_group [total_params f32]. // Zeroed at construction; fold-boundary reset zeroes it again (Pearl A sentinel: // zero grad_prev → cosine_sim=0 on first step → β1/β2 start at envelope midpoints). @@ -16629,6 +16736,7 @@ impl GpuDqnTrainer { pearl_5_iqn_tau_kernel, pearl_6_kelly_kernel, pearl_8_trail_kernel, + pearl_1_ext_num_atoms_kernel, grad_prev_buf_per_group, group_param_offsets_buf, atoms_clip_count_buf, diff --git a/crates/ml/src/cuda_pipeline/pearl_1_ext_num_atoms_kernel.cu b/crates/ml/src/cuda_pipeline/pearl_1_ext_num_atoms_kernel.cu new file mode 100644 index 000000000..7dc5590bc --- /dev/null +++ b/crates/ml/src/cuda_pipeline/pearl_1_ext_num_atoms_kernel.cu @@ -0,0 +1,56 @@ +// crates/ml/src/cuda_pipeline/pearl_1_ext_num_atoms_kernel.cu +// +// SP5 Task A8: Pearl 1-ext — per-branch C51 num_atoms from per-branch v_half. +// +// Reads ATOM_V_HALF[b] (ISV[178..182), populated by Pearl 1 in Task A1). +// Maps to discrete num_atoms ∈ {16, 32, 64} based on v_half thresholds: +// v_half < 0.1 → 64 atoms (narrow Q range, high resolution) +// v_half < 1.0 → 32 atoms (moderate Q range) +// v_half ≥ 1.0 → 16 atoms (wide Q range, modest resolution) +// +// The discrete output is smoothed by Pearls A+D in the launcher's apply_pearls +// step. During transitions (e.g. v_half drifts across 1.0), the smoothed ISV +// value will be intermediate between {16, 32}; the C51 atoms_update consumer +// in Layer B rounds to the nearest valid count at consume time. +// +// 4-thread single-block (one per branch). No atomicAdd. +// __threadfence_system() after writes. +// +// Invariant 1 anchors: +// THRESHOLD_NARROW = 0.1f (high-resolution boundary) +// THRESHOLD_MODERATE = 1.0f (moderate-resolution boundary) +// ATOMS_HIGH = 64.0f, ATOMS_MID = 32.0f, ATOMS_LOW = 16.0f +// (atom counts are powers of 2; total cap 4 branches × 64 = 256 atoms, +// within the HW shared-memory budget the C51 backward kernel uses.) + +#include + +extern "C" __global__ void pearl_1_ext_num_atoms_update( + const float* __restrict__ isv_signals, + int v_half_isv_base, // ATOM_V_HALF_BASE = 178 + float* __restrict__ scratch_out, + int num_atoms_idx_base // 4 outputs, scratch_out[base..base+4) +) { + int b = threadIdx.x; + if (blockIdx.x != 0 || b >= 4) return; + + const float THRESHOLD_NARROW = 0.1f; // Invariant 1 anchor: high-res boundary + const float THRESHOLD_MODERATE = 1.0f; // Invariant 1 anchor: moderate-res boundary + const float ATOMS_HIGH = 64.0f; // Invariant 1 anchor: narrow-range atom count + const float ATOMS_MID = 32.0f; // Invariant 1 anchor: moderate-range atom count + const float ATOMS_LOW = 16.0f; // Invariant 1 anchor: wide-range atom count + + float v_half = isv_signals[v_half_isv_base + b]; + + float num_atoms; + if (v_half < THRESHOLD_NARROW) { + num_atoms = ATOMS_HIGH; + } else if (v_half < THRESHOLD_MODERATE) { + num_atoms = ATOMS_MID; + } else { + num_atoms = ATOMS_LOW; + } + + scratch_out[num_atoms_idx_base + b] = num_atoms; + __threadfence_system(); +} diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 0fca8dd49..a9620f4e1 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -684,6 +684,15 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "ISV[TRAIL_DIST_PER_DIR_BASE=270..274) — per-direction trail-stop distance. Short[270] and Long[272]: atr_abs × 2.0; Hold[271] and Flat[273]: 1.0. SP5 Pearl A sentinel 0 at fold boundary so the first pearl_8_trail_update launch fires first-observation replacement with the new fold's ATR (Task A7).", }, + // SP5 Task A8: per-branch C51 num_atoms (Pearl 1-ext). + // ISV[ATOM_NUM_ATOMS_BASE=274..278) = 4 floats (one per branch). + // Pearl A sentinel 0 at fold boundary → first pearl_1_ext_num_atoms_update launch + // fires the first-observation replacement with the new fold's v_half-derived count. + RegistryEntry { + name: "sp5_atom_num_atoms", + category: ResetCategory::FoldReset, + description: "ISV[ATOM_NUM_ATOMS_BASE=274..278) — per-branch C51 num_atoms (discrete {16, 32, 64} mapped from v_half thresholds, smoothed via Pearls A+D so transition-region values may be fractional; Layer B's atoms_update consumer rounds at consume time). SP5 Pearl A sentinel 0 at fold boundary triggers first-observation replacement (Task A8).", + }, ]; 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 700d50c9a..1ad926aee 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -3653,6 +3653,13 @@ impl DQNTrainer { } } } + // SP5 Task A8: pearl_1_ext_num_atoms → apply_pearls_ad ×4 + // Reads ATOM_V_HALF (Pearl 1 from A1, populated earlier in this step's launch chain). + // Maps v_half to discrete {64, 32, 16} via threshold cascade; Pearls A+D smooths. + // Writes ISV[ATOM_NUM_ATOMS_BASE=274..278) = num_atoms[4 branches] (ALPHA_META=1e-3). + if let Err(e) = fused.trainer().launch_sp5_pearl_1_ext_num_atoms() { + tracing::warn!(error = %e, "SP5 Pearl 1-ext producer launch failed"); + } } } diff --git a/crates/ml/tests/sp5_producer_unit_tests.rs b/crates/ml/tests/sp5_producer_unit_tests.rs index 872b87821..f0e6bdc0e 100644 --- a/crates/ml/tests/sp5_producer_unit_tests.rs +++ b/crates/ml/tests/sp5_producer_unit_tests.rs @@ -29,6 +29,10 @@ //! 14. `pearl_8_trail_per_direction_factors` //! 15. `pearl_8_trail_atr_floor_at_quiet_market` //! +//! Task A8 (Pearl 1-ext per-branch num_atoms): +//! 16. `pearl_1_ext_num_atoms_threshold_cascade` +//! 17. `pearl_1_ext_num_atoms_exact_threshold_falls_to_next_branch` +//! //! All tests use analytically-known synthetic inputs with no CPU reference //! implementation, per `feedback_no_cpu_test_fallbacks.md`. All mapped-pinned //! memory (no DtoH), all `#[ignore]`-gated for GPU. @@ -74,6 +78,9 @@ const SP5_PEARL_6_KELLY_CUBIN: &[u8] = const SP5_PEARL_8_TRAIL_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/pearl_8_trail_kernel.cubin")); +const SP5_PEARL_1_EXT_NUM_ATOMS_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/pearl_1_ext_num_atoms_kernel.cubin")); + // ── Helpers ─────────────────────────────────────────────────────────────────── fn load_pearl_3_sigma_kernel(stream: &Arc) -> CudaFunction { @@ -1907,3 +1914,155 @@ fn pearl_8_trail_atr_floor_at_quiet_market() { assert!((out[3] - trail_hold_flat).abs() < tol, "Flat: expected {trail_hold_flat:.5}, got {:.5}", out[3]); } + +// ── Task A8 (Pearl 1-ext num_atoms) ────────────────────────────────────────── + +fn load_pearl_1_ext_num_atoms_kernel(stream: &Arc) -> CudaFunction { + let module = stream + .context() + .load_cubin(SP5_PEARL_1_EXT_NUM_ATOMS_CUBIN.to_vec()) + .expect("load pearl_1_ext_num_atoms_kernel cubin"); + module + .load_function("pearl_1_ext_num_atoms_update") + .expect("load pearl_1_ext_num_atoms_update function") +} + +/// SP5 Task A8 Test 16: pearl_1_ext_num_atoms threshold cascade. +/// +/// Synthetic ISV[ATOM_V_HALF_BASE + b] for 4 branches: +/// b=0: v_half=0.05 → strictly less than 0.1 → 64 atoms (narrow) +/// b=1: v_half=0.5 → not < 0.1, but < 1.0 → 32 atoms (moderate) +/// b=2: v_half=2.0 → not < 1.0 → 16 atoms (wide) +/// b=3: v_half=0.099 → just below 0.1 boundary → 64 atoms (narrow) +/// +/// Verifies the threshold cascade boundaries. +/// No CPU reference implementation (feedback_no_cpu_test_fallbacks.md). +/// Analytically-known: all expected values are structural constants +/// ATOMS_HIGH=64, ATOMS_MID=32, ATOMS_LOW=16 (Invariant 1 anchors). +#[test] +#[ignore = "requires GPU"] +fn pearl_1_ext_num_atoms_threshold_cascade() { + let ctx = CudaContext::new(0).expect("CUDA context"); + let stream = ctx.new_stream().expect("stream"); + let kernel = load_pearl_1_ext_num_atoms_kernel(&stream); + + // ISV buffer: place v_half values at ATOM_V_HALF_BASE=178..182. + // Allocate enough slots to cover ISV[182) (4 slots from base 178 = indices 178..182). + let isv_size: usize = 182; + let mut isv_buf = unsafe { MappedF32Buffer::new(isv_size) }.expect("isv_buf alloc"); + isv_buf.host_slice_mut().fill(0.0_f32); + // ATOM_V_HALF_BASE = 178 + isv_buf.host_slice_mut()[178] = 0.05_f32; // b=0: < 0.1 → 64 atoms + isv_buf.host_slice_mut()[179] = 0.5_f32; // b=1: < 1.0 → 32 atoms + isv_buf.host_slice_mut()[180] = 2.0_f32; // b=2: >= 1.0 → 16 atoms + isv_buf.host_slice_mut()[181] = 0.099_f32; // b=3: < 0.1 → 64 atoms + + let scratch_size: usize = 4; + let base: i32 = 0; + let mut scratch = unsafe { MappedF32Buffer::new(scratch_size) }.expect("scratch alloc"); + + let isv_dev = isv_buf.dev_ptr; + let scr_dev = scratch.dev_ptr; + let v_half_base_i32: i32 = 178; + + unsafe { + stream + .launch_builder(&kernel) + .arg(&isv_dev) + .arg(&v_half_base_i32) + .arg(&scr_dev) + .arg(&base) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (4, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch pearl_1_ext_num_atoms_update"); + } + stream.synchronize().expect("sync after pearl_1_ext_num_atoms_update"); + + let out = scratch.read_all(); + + println!("b=0 (v_half=0.05): got {:.0} (expected 64 — narrow)", out[0]); + println!("b=1 (v_half=0.5): got {:.0} (expected 32 — moderate)", out[1]); + println!("b=2 (v_half=2.0): got {:.0} (expected 16 — wide)", out[2]); + println!("b=3 (v_half=0.099): got {:.0} (expected 64 — just-narrow)", out[3]); + + let tol = 0.5_f32; // integer comparison via f32 — tolerate no rounding + assert!((out[0] - 64.0_f32).abs() < tol, + "b=0: expected 64 atoms (v_half=0.05 < 0.1), got {:.0}", out[0]); + assert!((out[1] - 32.0_f32).abs() < tol, + "b=1: expected 32 atoms (0.1 <= v_half=0.5 < 1.0), got {:.0}", out[1]); + assert!((out[2] - 16.0_f32).abs() < tol, + "b=2: expected 16 atoms (v_half=2.0 >= 1.0), got {:.0}", out[2]); + assert!((out[3] - 64.0_f32).abs() < tol, + "b=3: expected 64 atoms (v_half=0.099 < 0.1, just below narrow boundary), got {:.0}", out[3]); +} + +/// SP5 Task A8 Test 17: pearl_1_ext_num_atoms exact-threshold boundary semantics. +/// +/// At v_half = 0.1 exactly: `0.1 < 0.1` is false → falls to next branch → 32 atoms. +/// At v_half = 1.0 exactly: `1.0 < 1.0` is false → falls to else branch → 16 atoms. +/// +/// Verifies the strictly-less-than threshold semantics (NOT <=). +/// No CPU reference implementation (feedback_no_cpu_test_fallbacks.md). +/// The two branches exercise: THRESHOLD_NARROW boundary (b=0) and +/// THRESHOLD_MODERATE boundary (b=1); the other two branches (b=2, b=3) are +/// set to unambiguous interior values to confirm they are not affected. +#[test] +#[ignore = "requires GPU"] +fn pearl_1_ext_num_atoms_exact_threshold_falls_to_next_branch() { + let ctx = CudaContext::new(0).expect("CUDA context"); + let stream = ctx.new_stream().expect("stream"); + let kernel = load_pearl_1_ext_num_atoms_kernel(&stream); + + let isv_size: usize = 182; + let mut isv_buf = unsafe { MappedF32Buffer::new(isv_size) }.expect("isv_buf alloc"); + isv_buf.host_slice_mut().fill(0.0_f32); + // ATOM_V_HALF_BASE = 178 + isv_buf.host_slice_mut()[178] = 0.1_f32; // b=0: exactly 0.1 — NOT < 0.1 → 32 atoms + isv_buf.host_slice_mut()[179] = 1.0_f32; // b=1: exactly 1.0 — NOT < 1.0 → 16 atoms + isv_buf.host_slice_mut()[180] = 0.05_f32; // b=2: clearly narrow → 64 atoms (sanity) + isv_buf.host_slice_mut()[181] = 5.0_f32; // b=3: clearly wide → 16 atoms (sanity) + + let scratch_size: usize = 4; + let base: i32 = 0; + let mut scratch = unsafe { MappedF32Buffer::new(scratch_size) }.expect("scratch alloc"); + + let isv_dev = isv_buf.dev_ptr; + let scr_dev = scratch.dev_ptr; + let v_half_base_i32: i32 = 178; + + unsafe { + stream + .launch_builder(&kernel) + .arg(&isv_dev) + .arg(&v_half_base_i32) + .arg(&scr_dev) + .arg(&base) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (4, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch pearl_1_ext_num_atoms_update (exact threshold)"); + } + stream.synchronize().expect("sync after pearl_1_ext_num_atoms_update (exact threshold)"); + + let out = scratch.read_all(); + + println!("b=0 (v_half=0.1 exact): got {:.0} (expected 32 — NOT < 0.1)", out[0]); + println!("b=1 (v_half=1.0 exact): got {:.0} (expected 16 — NOT < 1.0)", out[1]); + println!("b=2 (v_half=0.05): got {:.0} (expected 64 — sanity)", out[2]); + println!("b=3 (v_half=5.0): got {:.0} (expected 16 — sanity)", out[3]); + + let tol = 0.5_f32; + assert!((out[0] - 32.0_f32).abs() < tol, + "b=0 (v_half=0.1): expected 32 atoms (threshold is strictly <, not <=), got {:.0}", out[0]); + assert!((out[1] - 16.0_f32).abs() < tol, + "b=1 (v_half=1.0): expected 16 atoms (threshold is strictly <, not <=), got {:.0}", out[1]); + assert!((out[2] - 64.0_f32).abs() < tol, + "b=2 (v_half=0.05): expected 64 atoms (sanity check), got {:.0}", out[2]); + assert!((out[3] - 16.0_f32).abs() < tol, + "b=3 (v_half=5.0): expected 16 atoms (sanity check), got {:.0}", out[3]); +} diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index b8f440757..e992210a5 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,8 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +SP5 Task A8 — Pearl 1-ext per-branch C51 num_atoms GPU producer — LAYER A COMPLETE (2026-05-01): final SP5 Layer A producer. One new CUDA kernel (`pearl_1_ext_num_atoms_kernel.cu`) lands as a Layer A additive producer feeding ISV slots [274..278) with per-branch C51 num_atoms. No `atoms_update_kernel` consumer migration in this commit — Layer B wires the consumer when the per-branch atom count is read at C51 backward time. `pearl_1_ext_num_atoms_update`: single-block 4-thread kernel (one thread per branch: Direction/Magnitude/Order/Urgency). Reads ISV[ATOM_V_HALF_BASE=178..182) (populated by Pearl 1 in Task A1); applies threshold cascade: `v_half < 0.1 → 64 atoms` (narrow Q range, high resolution), `v_half < 1.0 → 32 atoms` (moderate), `v_half ≥ 1.0 → 16 atoms` (wide Q range, modest resolution). THRESHOLD_NARROW=0.1, THRESHOLD_MODERATE=1.0, ATOMS_HIGH=64, ATOMS_MID=32, ATOMS_LOW=16 are Invariant 1 anchors (powers of 2; total cap 4×64=256 atoms within HW shared-memory budget). Discrete output smoothed by Pearls A+D — during transitions (e.g. v_half drifts across 1.0), the smoothed ISV value is intermediate between {16, 32}; Layer B's atoms_update consumer rounds to nearest valid count at consume time. `launch_sp5_pearl_1_ext_num_atoms` fires the producer kernel writing scratch[203..207), then 4 `launch_apply_pearls` calls (one per branch → ISV[ATOM_NUM_ATOMS_BASE=274..278)) with ALPHA_META=1e-3 (default rate). Wiener offset formula: `SP4_PRODUCER_COUNT × 3 + (isv_slot − SP5_SLOT_BASE) × 3` where `SP4_PRODUCER_COUNT × 3 = 213`; for ISV slot 274 (first ATOM_NUM_ATOMS slot): wiener_off = 213 + (274-174)×3 = 513 — well within the 543-float wiener_state_buf. Buffer growth: `producer_step_scratch_buf` 203→207 slots (SP5_SCRATCH_TOTAL); 1 new scratch constant `SCRATCH_PEARL_1_EXT_NUM_ATOMS=203`. wiener_state_buf stays at 543 (sized at A1 for entire SP5 block). StateResetRegistry gains 1 new FoldReset entry: `sp5_atom_num_atoms` (ISV[274..278)) — Pearl A sentinel 0 at fold boundary fires first-observation replacement with new fold's v_half-derived count. Wire-up in training_loop.rs: `launch_sp5_pearl_1_ext_num_atoms()` called immediately after `launch_sp5_pearl_8_trail` block (both depend on Pearl 1 outputs; Pearl 8 runs first for ordering parity with A7). Two GPU-only `#[ignore = "requires GPU"]` unit tests: (16) `pearl_1_ext_num_atoms_threshold_cascade` — 4 synthetic v_half values spanning all branches and just-below boundary; verifies 64/32/16/64 cascade; (17) `pearl_1_ext_num_atoms_exact_threshold_falls_to_next_branch` — v_half=0.1 exactly → 32 (NOT 64), v_half=1.0 exactly → 16 (NOT 32), verifies strictly-less-than semantics. No CPU compute, no HtoD/HtoH, no atomicAdd, no stubs, no consumer migration. **LAYER A COMPLETE**: 8 SP5 producers (Pearls 1, 1-ext, 2, 3, 4, 5, 6, 8) + 3 auxiliary kernels (q_branch_stats, grad_cosine_sim, q_skew_kurtosis) populating 110 ISV slots [174..286) with 4 cross-fold-persistent slots carved out (Kelly [280..286)). Touched: `cuda_pipeline/pearl_1_ext_num_atoms_kernel.cu` (new), `build.rs` (+1 cubin entry), `cuda_pipeline/gpu_dqn_trainer.rs` (SP5_SCRATCH_TOTAL 203→207 + 1 new scratch constant `SCRATCH_PEARL_1_EXT_NUM_ATOMS=203` + 1 static cubin + 1 struct field + cubin loader + struct initializer + launch method), `trainers/dqn/state_reset_registry.rs` (+1 FoldReset entry), `trainers/dqn/trainer/training_loop.rs` (+5 LOC after Pearl 8 launch block), `tests/sp5_producer_unit_tests.rs` (+2 GPU-only tests + 1 cubin constant + 1 loader helper + module docstring). Hard rules: feedback_no_cpu_compute_strict, feedback_no_atomicadd, feedback_no_htod_htoh_only_mapped_pinned, feedback_no_stubs, feedback_no_partial_refactor (no consumer migration — Layer A additive only), feedback_isv_for_adaptive_bounds (thresholds/counts are Invariant 1 structural anchors, not tuned constants). + SP5 Task A7 fix-up — close two minor review findings (2026-05-01): combined spec-quality review caught two minor issues in the Pearl 8 commit. (1) The audit doc's description of test 14 incorrectly claimed `atr_norm` was chosen to give `atr_abs=10.0` and `Short/Long=20.0`. The actual test uses `atr_norm=0.5` → `log_atr=1.0` → `atr_abs=e^1≈2.7183` → `Short/Long≈5.4366`. Updated the audit doc test-14 description to match the actual values. (2) `training_loop.rs:3640` used `let _ = std::mem::ManuallyDrop::new(guard);` to keep the cudarc StreamGuard alive past the inner block — the pattern leaks the guard's borrow bookkeeping for the lifetime of the function rather than dropping it after the kernel launch. Replaced with the idiomatic `let (feat_dev_ptr, _guard) = features_buf.device_ptr(stream_ref);` pattern at the same scope as the launch (matches existing conventions for SP4 producer wire-ups), with a comment explaining the bookkeeping intent. Comment-only / pattern-cleanup changes; no behavior change. cargo check + cargo test --no-run both clean. SP5 Task A7 — Pearl 8 per-direction trail-stop distance GPU producer (Layer A, 2026-05-01): one new CUDA kernel (`pearl_8_trail_kernel.cu`) lands as a Layer A additive producer feeding ISV slots [270..274) with per-direction trail-stop distances. No `check_trailing_stop` consumer migration in this commit — Layer B wires the consumer when the per-direction trail distance is read at trade-event time. `pearl_8_trail_update`: single-block 4-thread kernel (one thread per direction: Short/Hold/Long/Flat). Reads features[bar_idx × market_dim + 9] (ATR_NORM column), denormalizes via the canonical fxcache scheme `atr_abs = max(0.01, exp(atr_norm × 16 − 7))` (constants 16, 7, 0.01 are Invariant 1 anchors from the existing fxcache normalization in `experience_kernels.cu:2701`). Per-direction logic: Short[270] = atr_abs × ATR_TRAIL_FACTOR=2.0 (industry-standard 2× ATR trail), Hold[271] = EPS_CLAMP_FLOOR=1.0 (Hold doesn't trail-stop; floor prevents zero), Long[272] = atr_abs × 2.0 (same as Short for Layer A), Flat[273] = EPS_CLAMP_FLOOR=1.0. **Layer A simplification**: Short and Long share the same current-bar ATR value. The spec sketch implied direction-conditional ATR (`atr_ema_per_dir[4]` from per-trade-event aggregation conditioned on direction context), but that requires infrastructure not yet in place. The 4-slot ISV layout preserves the consumer contract so a future Layer C extension can populate truly direction-specific values without breaking Layer B's `check_trailing_stop` reads. `launch_sp5_pearl_8_trail(current_bar_idx)` reads the trainer's `latest_features_dev_ptr` + `config().market_dim`, fires the producer kernel writing scratch[199..203), then 4 `launch_apply_pearls` calls (one per direction) chained on the same stream with ALPHA_META=1e-3 (default rate). Wiener offsets: `213 + (270-174)×3 = 501` through `213 + (273-174)×3 = 510` — well within the 543-float wiener_state_buf. Buffer growth: `producer_step_scratch_buf` 199→203 slots (SP5_SCRATCH_TOTAL); 1 new scratch constant `SCRATCH_PEARL_8_TRAIL=199`. wiener_state_buf stays at 543 (sized at A1 for entire SP5 block). StateResetRegistry gains 1 new FoldReset entry: `sp5_trail_dist_per_dir` (ISV[270..274)) — Pearl A sentinel 0 at fold boundary fires first-observation replacement. Wire-up in training_loop.rs: `launch_sp5_pearl_8_trail(current_bar_idx)` called at the per-step boundary alongside other SP5 producer launches; `current_bar_idx` is read from the existing per-step bar-iteration variable. Two GPU-only `#[ignore = "requires GPU"]` unit tests: (14) `pearl_8_trail_per_direction_factors` — synthetic atr_norm chosen so atr_abs=10.0; verifies Short/Long = 20.0, Hold/Flat = 1.0; (15) `pearl_8_trail_atr_floor_at_quiet_market` — atr_norm=-10 → log_atr=-167 → exp(-167)≈0 → max(_,0.01)=0.01; verifies Short/Long = 0.02 (floor protection). No CPU compute, no HtoD/HtoH, no atomicAdd, no stubs, no consumer migration. Touched: `cuda_pipeline/pearl_8_trail_kernel.cu` (new), `build.rs` (+1 cubin entry), `cuda_pipeline/gpu_dqn_trainer.rs` (SP5_SCRATCH_TOTAL 199→203 + 1 new scratch constant + 1 static cubin + 1 struct field + cubin loader + struct initializer + launch method), `trainers/dqn/state_reset_registry.rs` (+1 FoldReset entry), `trainers/dqn/trainer/training_loop.rs` (+5 LOC after Pearl 5 launch + ManuallyDrop guard + DevicePtr trait import), `tests/sp5_producer_unit_tests.rs` (+2 GPU-only tests + 1 cubin constant + 1 loader helper). Hard rules: feedback_no_cpu_compute_strict, feedback_no_atomicadd, feedback_no_htod_htoh_only_mapped_pinned, feedback_no_stubs (Layer A simplification is documented design choice, not a stub), feedback_no_partial_refactor (no consumer migration — Layer A additive only).