feat(sp5): Task A8 — Pearl 1-ext per-branch num_atoms — Layer A complete
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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
|
||||
|
||||
@@ -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,
|
||||
|
||||
56
crates/ml/src/cuda_pipeline/pearl_1_ext_num_atoms_kernel.cu
Normal file
56
crates/ml/src/cuda_pipeline/pearl_1_ext_num_atoms_kernel.cu
Normal file
@@ -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 <cuda_runtime.h>
|
||||
|
||||
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();
|
||||
}
|
||||
@@ -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 }
|
||||
}
|
||||
|
||||
@@ -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");
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@@ -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<CudaStream>) -> 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<CudaStream>) -> 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]);
|
||||
}
|
||||
|
||||
@@ -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).
|
||||
|
||||
Reference in New Issue
Block a user