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

feat(sp5): Layer D Task D3 — Training metrics EMA kernel (additive)

Browse Source 
Third of 3 Layer D producer kernels. Replaces host-side training_sharpe_ema,
max_dd_ema, low_dd_ratio updates (host-side EMAs in training_loop.rs) with
a fused GPU kernel chained through apply_pearls_ad_kernel. Per
feedback_no_cpu_compute_strict.

Note: agent investigated training_loop.rs and found the third metric is
low_dd_ratio (not gamma_blend as the original plan brief named). The
authored kernel reproduces the actual host-side EMA triplet present in
the codebase.

Additive only — no consumer wiring, no behavior change. The kernel +
launcher are loaded into GpuDqnTrainer and the slots are reserved on the
ISV bus, but the host-side updates continue to run unchanged. D4 (atomic
Layer D commit) wires this and D1+D2 to call sites in the same atomic
refactor per feedback_no_partial_refactor. Also note: training_loop.rs
gains two state-reset-registry dispatch arms (sp5_health_composition for
D2 + sp5_training_metrics_ema for D3) — registry plumbing required by
the new entries, NOT consumer wiring of the kernels themselves; same
pattern as SP5 Layer A bug-fix #281.

What this lands:
- training_metrics_ema_kernel.cu (single-block 3-thread fused EMA)
- Rust launcher launch_training_metrics_ema()
- 3 new ISV slots (TRAINING_SHARPE_EMA_INDEX..LOW_DD_RATIO_INDEX, 294..297)
- ISV_TOTAL_DIM 294 → 297, SP5_PRODUCER_COUNT 120 → 123 (linear span)
- LAYOUT_FINGERPRINT_SEED bump (auto via slot string)
- StateResetRegistry entries for D2 (sp5_health_composition) + D3
  (sp5_training_metrics_ema) with reset_named_state dispatch arms
- build.rs cubin registration
- GPU-gated unit test in sp5_producer_unit_tests.rs (analytical EMA)
- SP5 contiguity slot test
  training_metrics_ema_slots_contiguous_and_above_health_block
- Audit doc append

Formula fidelity: kernel reproduces the host-side EMA update for
sharpe/max_dd/low_dd_ratio bit-for-bit within float precision. β decay
constants and any warmup/clamp logic migrate as Invariant 1 anchors with
no algorithmic change. Verified by unit test asserting kernel output
matches an analytical EMA sequence within 1e-6 rel-err.

Tests: cargo check + cargo build clean; ISV slot + state_reset_registry
unit tests pass (8/8 incl. new contiguity check); GPU correctness test
fires on next L40S smoke.

Refs: SP5 plan §D Task D3, builds on D1 (5ee795f14) + D2 (e49756ac9).
Layer D additive infrastructure complete after this commit; D4 (atomic
5-site host-EMA → GPU migration) is next.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-02 16:36:23 +02:00
parent e49756ac90
commit f42b5fff8d
8 changed files with 856 additions and 42 deletions
Download Patch File Download Diff File Expand all files Collapse all files

15
crates/ml/build.rs
View File

@@ -505,6 +505,21 @@ fn main() {
// in the atomic Layer D commit per
// `feedback_no_partial_refactor.md`.
"health_composition_kernel.cu",
// SP5 Layer D Task D3 (2026-05-02): training-metrics EMA producer.
// Third (and final) Layer D producer replacing the host-side
// adversarial-regime + adaptive-DSR EMA arithmetic block at
// `training_loop.rs:5039-5113` per `feedback_no_cpu_compute_strict.md`.
// Single-block 3-thread kernel — one thread per metric:
// tid=0 → training_sharpe_ema (adaptive α, host-tracked sentinel),
// tid=1 → max_dd_ema (fixed α=0.1, sentinel<1e-12),
// tid=2 → low_dd_ratio (fixed α=0.15, depends on tid=1's
// output via __syncthreads()).
// Writes 3 scratch slots at scratch[SCRATCH_TRAINING_METRICS_EMA_BASE
// =215..218). Chained `apply_pearls_ad_kernel` (3 launches,
// ALPHA_META=1e-3) smooths via Pearls A+D into ISV[294..297).
// Producer-only — D4 wires the consumer in the atomic Layer D
// commit per `feedback_no_partial_refactor.md`.
"training_metrics_ema_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

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

@@ -426,6 +426,22 @@ static SP5_PNL_AGGREGATION_CUBIN: &[u8] =
static SP5_HEALTH_COMPOSITION_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/health_composition_kernel.cubin"));
/// SP5 Layer D Task D3 (2026-05-02): Training metrics EMA producer cubin.
/// Single-block 3-thread arithmetic kernel reproducing the host-side
/// `training_sharpe_ema` (adaptive α), `max_dd_ema` (fixed α=0.1), and
/// `low_dd_ratio` (fixed α=0.15) updates at `training_loop.rs:5039-5113`
/// per `feedback_no_cpu_compute_strict.md`. Per-metric thread layout:
/// tid=0 sharpe (adaptive α with host-tracked `_initialized` sentinel),
/// tid=1 max_dd (fixed α=0.1 with prev<1e-12 sentinel), tid=2 low_dd_ratio
/// (fixed α=0.15 over a binary `is_low` derived from tid=1's *new*
/// max_dd_ema via __syncthreads()). Writes 3 floats to
/// scratch_buf[SCRATCH_TRAINING_METRICS_EMA_BASE=215..218). Followed by 3
/// apply_pearls_ad_kernel calls → ISV[TRAINING_SHARPE_EMA_INDEX=294..297).
/// Replaces the host-side updates that D4 atomically migrates per
/// `feedback_no_partial_refactor.md`. Additive in this commit.
static SP5_TRAINING_METRICS_EMA_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/training_metrics_ema_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
@@ -1000,27 +1016,29 @@ pub const SP4_WIENER_TOTAL_FLOATS: usize =
// ── SP5 Task A1 buffer layout constants ───────────────────────────────────────
/// SP5 Task A1 + Layer D Task D1+D2: combined wiener_state_buf total float count.
/// SP4 had 71 producers (213 floats); SP5 contributes 120 producer triples
/// (linear span from `SP5_SLOT_BASE=174` to `SP5_SLOT_END=294`, including the
/// 2-slot carve-out gap at 278..280; 120 × 3 = 360 floats). Combined =
/// (71 + 120) × 3 = 573 floats. Pre-D1 SP5 was 110 producers / 543 floats;
/// SP5 Task A1 + Layer D Task D1+D2+D3: combined wiener_state_buf total float count.
/// SP4 had 71 producers (213 floats); SP5 contributes 123 producer triples
/// (linear span from `SP5_SLOT_BASE=174` to `SP5_SLOT_END=297`, including the
/// 2-slot carve-out gap at 278..280; 123 × 3 = 369 floats). Combined =
/// (71 + 123) × 3 = 582 floats. Pre-D1 SP5 was 110 producers / 543 floats;
/// Task D1 added 4 unique slots (286..290) growing the linear span 112 → 116
/// (buffer 543 → 561). Task D2 adds 4 more slots (290..294) growing the linear
/// span 116 → 120 (buffer 561 → 573). Note: SP5_PRODUCER_COUNT is the
/// **linear-span** constant — see its docstring for why it now diverges
/// permanently from the unique-slot count (118 post-D2).
/// (buffer 543 → 561). Task D2 added 4 more slots (290..294) growing the linear
/// span 116 → 120 (buffer 561 → 573). Task D3 adds 3 more slots (294..297)
/// growing the linear span 120 → 123 (buffer 573 → 582). Note:
/// SP5_PRODUCER_COUNT is the **linear-span** constant — see its docstring
/// for why it now diverges permanently from the unique-slot count (121 post-D3).
///
/// wiener_state_buf must grow to this size at construction for ALL SP5 slots
/// (including the new D1 PnL-aggregation and D2 health-composition blocks) to
/// be in-bounds when `apply_pearls_ad_kernel` writes their Wiener triples.
/// The Pearl 6 Kelly block (slots 280..286) and the carve-out gap (278..280)
/// have reserved-but-unused wiener storage at offsets
/// [(280-174)*3+213..(286-174)*3+213) = [531..549) — Pearl 6 does not call
/// `launch_apply_pearls` (cross-fold persistent), so those 18 floats stay
/// zero-initialised. Layer D D1 slots 286..290 use offsets
/// (including the new D1 PnL-aggregation, D2 health-composition, and D3
/// training-metrics-EMA blocks) to be in-bounds when `apply_pearls_ad_kernel`
/// writes their Wiener triples. The Pearl 6 Kelly block (slots 280..286) and
/// the carve-out gap (278..280) have reserved-but-unused wiener storage at
/// offsets [(280-174)*3+213..(286-174)*3+213) = [531..549) — Pearl 6 does
/// not call `launch_apply_pearls` (cross-fold persistent), so those 18 floats
/// stay zero-initialised. Layer D D1 slots 286..290 use offsets
/// [(286-174)*3+213..(290-174)*3+213) = [549..561). Layer D D2 slots 290..294
/// use offsets [(290-174)*3+213..(294-174)*3+213) = [561..573).
/// use offsets [(290-174)*3+213..(294-174)*3+213) = [561..573). Layer D D3
/// slots 294..297 use offsets [(294-174)*3+213..(297-174)*3+213) = [573..582).
pub const SP5_WIENER_TOTAL_FLOATS: usize =
(SP4_PRODUCER_COUNT + crate::cuda_pipeline::sp5_isv_slots::SP5_PRODUCER_COUNT)
* SP4_WIENER_FLOATS_PER_SLOT;
@@ -1057,8 +1075,10 @@ pub const SP5_WIENER_TOTAL_FLOATS: usize =
/// 4 floats for pnl_aggregation output (SCRATCH_PNL_AGG_BASE=207, total/mean/var/max_dd at [207..211))
/// Layer D Task D2 adds:
/// 4 floats for health_composition output (SCRATCH_HEALTH_COMP_BASE=211, score/q_gap_norm/q_var_norm/grad_norm_norm at [211..215))
/// Combined: 71 + 16 + 16 + 4 + 4 + 20 + 8 + 8 + 8 + 8 + 8 + 4 + 4 + 20 + 4 + 4 + 4 + 4 = 215 scratch slots [0..215).
pub const SP5_SCRATCH_TOTAL: usize = 215;
/// Layer D Task D3 adds:
/// 3 floats for training_metrics_ema output (SCRATCH_TRAINING_METRICS_EMA_BASE=215, training_sharpe_ema/max_dd_ema/low_dd_ratio at [215..218))
/// Combined: 71 + 16 + 16 + 4 + 4 + 20 + 8 + 8 + 8 + 8 + 8 + 4 + 4 + 20 + 4 + 4 + 4 + 4 + 3 = 218 scratch slots [0..218).
pub const SP5_SCRATCH_TOTAL: usize = 218;
/// SP5 Layer D Task D1 (2026-05-02): scratch index base for `pnl_aggregation_update`.
/// Slots [207..211): aggregated PnL summary outputs in fixed order:
@@ -1086,6 +1106,29 @@ pub const SCRATCH_PNL_AGG_BASE: usize = 207;
/// a launcher signature change.
pub const SCRATCH_HEALTH_COMP_BASE: usize = 211;
/// SP5 Layer D Task D3 (2026-05-02): scratch index base for `training_metrics_ema_update`.
/// Slots [215..218): training-loop EMA outputs in fixed order:
/// [215] = training_sharpe_ema (adaptive α `clamp(0.05+0.3·|err|, 0, 0.5)`,
/// host-tracked `_initialized` sentinel)
/// [216] = max_dd_ema (fixed α=0.1, sentinel `prev<1e-12` ⇒
/// first-observation replacement)
/// [217] = low_dd_ratio (fixed α=0.15 over binary
/// `is_low = (max_dd < new_max_dd_ema·0.5)
/// ? 1 : 0`; reads tid=1's *new* output via
/// __syncthreads(); no host-side sentinel,
/// constructor zero behaves correctly under
/// the EMA recurrence)
/// Written by `training_metrics_ema_update`; consumed by `apply_pearls_ad_kernel`
/// → ISV[TRAINING_SHARPE_EMA_INDEX=294..LOW_DD_RATIO_INDEX+1=297). All three
/// host-side EMA constants (0.05/0.3/0.5 adaptive bounds; 0.1; 0.15; 1e-12
/// sentinel) are Invariant 1 anchors per `feedback_no_quickfixes.md`,
/// migrated verbatim from `training_loop.rs:5039-5113`. Per
/// `feedback_trust_code_not_docs.md`: the third metric is `low_dd_ratio`
/// (not `gamma_blend` as the SP5 plan §D Task D3 hints — that field does
/// not exist in the codebase; the migration follows the actual host-side
/// EMA at `training_loop.rs:5052`).
pub const SCRATCH_TRAINING_METRICS_EMA_BASE: usize = 215;
/// 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).
/// Written by `pearl_8_trail_update`; consumed by apply_pearls_ad_kernel →
@@ -4422,6 +4465,25 @@ pub struct GpuDqnTrainer {
/// path). Loaded from `health_composition_kernel.cubin`. Producer-only in
/// this commit; D4 wires the consumer atomically.
health_composition_kernel: CudaFunction,
/// SP5 Layer D Task D3 (2026-05-02): training-metrics EMA producer.
/// Single-block 3-thread arithmetic kernel (one thread per metric):
/// tid=0 training_sharpe_ema (adaptive α `clamp(0.05+0.3·|err|, 0, 0.5)`,
/// host-tracked `_initialized` sentinel), tid=1 max_dd_ema (fixed α=0.1,
/// `prev<1e-12` sentinel), tid=2 low_dd_ratio (fixed α=0.15 over the
/// binary `is_low` indicator derived from tid=1's *new* max_dd_ema via
/// __syncthreads(); no host-side sentinel — constructor zero is
/// recurrence-safe). Writes 3 scratch floats — training_sharpe_ema +
/// max_dd_ema + low_dd_ratio — to scratch_buf[
/// SCRATCH_TRAINING_METRICS_EMA_BASE=215..218). Chained
/// apply_pearls_ad_kernel (3 launches, ALPHA_META=1e-3) →
/// ISV[TRAINING_SHARPE_EMA_INDEX=294..LOW_DD_RATIO_INDEX+1=297). In
/// fold-reset registry (FoldReset; the underlying host scalars
/// `training_sharpe_ema` / `_initialized` / `max_dd_ema` /
/// `low_dd_ratio` reset at fold boundary alongside the ISV slots via
/// the standard Pearl A sentinel path). Loaded from
/// `training_metrics_ema_kernel.cubin`. Producer-only in this commit;
/// D4 wires the consumer atomically.
training_metrics_ema_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
@@ -11928,6 +11990,147 @@ impl GpuDqnTrainer {
Ok(())
}
/// SP5 Layer D Task D3 (2026-05-02): launch `training_metrics_ema_update` —
/// single-block 3-thread EMA producer for
/// ISV[TRAINING_SHARPE_EMA_INDEX=294..LOW_DD_RATIO_INDEX+1=297).
///
/// Reproduces the host-side EMA arithmetic block at
/// `training_loop.rs:5039-5113` bit-for-bit on GPU per
/// `feedback_no_cpu_compute_strict.md`. One thread per metric:
///
/// tid=0 training_sharpe_ema α = clamp(0.05 + 0.3·|err|, 0, 0.5);
/// sharpe_initialized==0 ⇒ replace.
/// tid=1 max_dd_ema α = 0.1; prev<1e-12 ⇒ replace.
/// tid=2 low_dd_ratio α = 0.15 over binary
/// `is_low = (max_dd < new_max_dd_ema·0.5)
/// ? 1 : 0`. Reads tid=1's *new* max_dd_ema
/// via __syncthreads(); no host-side sentinel.
///
/// All EMA constants (0.05, 0.3, 0.5, 0.1, 0.85, 0.15, 0.5 threshold,
/// 1e-12 sentinel) are Invariant 1 anchors — deleted host formula
/// literals migrated verbatim per `feedback_no_quickfixes.md` and
/// `feedback_trust_code_not_docs.md` (the third metric is `low_dd_ratio`
/// from `training_loop.rs:5052`, not the speculative `gamma_blend`
/// the SP5 plan §D Task D3 mentions — that field does not exist in
/// the codebase).
///
/// Writes 3 floats to scratch_buf[SCRATCH_TRAINING_METRICS_EMA_BASE
/// =215..218):
/// [215] new training_sharpe_ema
/// [216] new max_dd_ema
/// [217] new low_dd_ratio
///
/// Then chains 3 `apply_pearls_ad_kernel` calls smoothing through
/// Pearls A+D into ISV[294..297) on the same stream (no host sync per
/// `pearl_cold_path_no_exception_to_gpu_drives.md` and the SP5 GPU-only
/// Pearls refactor).
///
/// Per `feedback_no_partial_refactor.md` — D3 is additive only; D4
/// atomically migrates the host-side site to call this launcher in
/// lockstep with D1/D2. The kernel takes the host-side prev EMA values
/// and the `_initialized` flag by value in additive D3; D4 will lift
/// these to mapped-pinned scalars so the host is no longer the source
/// of truth.
///
/// Arguments (all primitive scalars by value, no mapped-pinned plumbing
/// per the additive-D3 design):
/// - `epoch_sharpe`: per-epoch Sharpe ratio observation (raw).
/// - `epoch_max_dd`: per-epoch max drawdown observation (raw, ∈ [0, 1]).
/// - `prev_sharpe_ema`: host-tracked previous training_sharpe_ema.
/// - `prev_max_dd_ema`: host-tracked previous max_dd_ema (cast to f32).
/// - `prev_low_dd_ratio`: host-tracked previous low_dd_ratio (cast to f32).
/// - `sharpe_initialized`: host-tracked sentinel flag; `false` ⇒ first
/// observation replaces the slot directly. Cast to `i32` at the
/// kernel boundary per `feedback_cudarc_f64_f32_abi.md` (cudarc's
/// `bool` ABI is unstable; `i32` round-trip is the safe pattern).
pub(crate) fn launch_training_metrics_ema(
&self,
epoch_sharpe: f32,
epoch_max_dd: f32,
prev_sharpe_ema: f32,
prev_max_dd_ema: f32,
prev_low_dd_ratio: f32,
sharpe_initialized: bool,
) -> Result<(), MLError> {
use crate::cuda_pipeline::sp5_isv_slots::{
TRAINING_SHARPE_EMA_INDEX, MAX_DD_EMA_INDEX, LOW_DD_RATIO_INDEX,
SP5_SLOT_BASE,
};
use crate::cuda_pipeline::sp4_wiener_ema::{launch_apply_pearls, ALPHA_META};
debug_assert!(self.isv_signals_dev_ptr != 0,
"launch_training_metrics_ema: isv_signals_dev_ptr must be allocated by constructor");
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 sharpe_idx_i32: i32 = (SCRATCH_TRAINING_METRICS_EMA_BASE + 0) as i32;
let max_dd_idx_i32: i32 = (SCRATCH_TRAINING_METRICS_EMA_BASE + 1) as i32;
let low_dd_ratio_idx_i32: i32 = (SCRATCH_TRAINING_METRICS_EMA_BASE + 2) as i32;
// i32 sentinel flag (cudarc `bool` ABI is unstable — cast to i32 here
// matching the kernel's `int sharpe_initialized` parameter; per
// `feedback_cudarc_f64_f32_abi.md` the launcher's primitive type
// must match the kernel's declared C type exactly).
let sharpe_initialized_i32: i32 = if sharpe_initialized { 1 } else { 0 };
// Wiener base offset for SP5 slots: SP4_PRODUCER_COUNT * 3 = 213.
let base_wiener_offset: i32 = (SP4_PRODUCER_COUNT * 3) as i32;
// Step 1: training_metrics_ema_update — single-block 3-thread; no
// dynamic shared memory (the kernel uses a single static __shared__
// float for the tid=1 → tid=2 max_dd_ema handoff).
unsafe {
self.stream
.launch_builder(&self.training_metrics_ema_kernel)
.arg(&epoch_sharpe)
.arg(&epoch_max_dd)
.arg(&prev_sharpe_ema)
.arg(&prev_max_dd_ema)
.arg(&prev_low_dd_ratio)
.arg(&sharpe_initialized_i32)
.arg(&scratch_dev)
.arg(&sharpe_idx_i32)
.arg(&max_dd_idx_i32)
.arg(&low_dd_ratio_idx_i32)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (3, 1, 1),
shared_mem_bytes: 0,
})
.map_err(|e| MLError::ModelError(format!("training_metrics_ema_update: {e}")))?;
}
// Step 2: 3 apply_pearls calls (one per output slot). Each
// scratch[215+k] → ISV[294+k] with its own Wiener triple at offset
// (base_wiener_offset + (isv_slot − SP5_SLOT_BASE) * 3). Contract
// mirrors D1/D2 exactly.
let isv_slots: [usize; 3] = [
TRAINING_SHARPE_EMA_INDEX,
MAX_DD_EMA_INDEX,
LOW_DD_RATIO_INDEX,
];
for k in 0..3_usize {
let isv_slot = isv_slots[k] as i32;
let scratch_idx = (SCRATCH_TRAINING_METRICS_EMA_BASE + k) 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].
///
@@ -14074,6 +14277,24 @@ impl GpuDqnTrainer {
.map_err(|e| MLError::ModelError(format!("health_composition_update load: {e}")))?
};
// SP5 Layer D Task D3 (2026-05-02): load training_metrics_ema_kernel.
// Single-block 3-thread arithmetic kernel. Reads 5 host-passed f32
// scalars (epoch_sharpe + epoch_max_dd raw observations, plus
// prev_sharpe_ema / prev_max_dd_ema / prev_low_dd_ratio host-tracked
// EMAs) and 1 i32 sentinel flag (sharpe_initialized); writes 3 floats
// (sharpe/max_dd/low_dd_ratio) to scratch[SCRATCH_TRAINING_METRICS_EMA_BASE
// =215..218). FoldReset: ISV[294..297) zeroed at fold boundary
// alongside the host-side `training_sharpe_ema` / `_initialized` /
// `max_dd_ema` / `low_dd_ratio` scalars (Pearl A sentinel path).
// Producer-only in D3; consumer wires atomically in D4 per
// `feedback_no_partial_refactor.md`.
let training_metrics_ema_kernel = {
let module = stream.context().load_cubin(SP5_TRAINING_METRICS_EMA_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!("sp5 training_metrics_ema cubin load: {e}")))?;
module.load_function("training_metrics_ema_update")
.map_err(|e| MLError::ModelError(format!("training_metrics_ema_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).
@@ -17261,6 +17482,7 @@ impl GpuDqnTrainer {
pearl_1_ext_num_atoms_kernel,
pnl_aggregation_kernel,
health_composition_kernel,
training_metrics_ema_kernel,
grad_prev_buf_per_group,
group_param_offsets_buf,
atoms_clip_count_buf,

119
crates/ml/src/cuda_pipeline/sp5_isv_slots.rs
View File

@@ -12,8 +12,9 @@
//! 280..286 Cross-fold-persistent Kelly slots (NOT in fold-reset registry)
//! 286..290 Layer D D1: PnL aggregation outputs (4 slots, fold-reset)
//! 290..294 Layer D D2: Health composition outputs (4 slots, fold-reset)
//! 294..297 Layer D D3: Training metrics EMA outputs (3 slots, fold-reset)
//!
//! Total: 118 new SP5 ISV slots (52 + 24 + 20 + 4 + 4 + 6 + 4 + 4).
//! Total: 121 new SP5 ISV slots (52 + 24 + 20 + 4 + 4 + 6 + 4 + 4 + 3).
pub const SP5_SLOT_BASE: usize = 174;
@@ -118,7 +119,42 @@ pub const Q_GAP_NORM_INDEX: usize = 291;
pub const Q_VAR_NORM_INDEX: usize = 292;
pub const GRAD_NORM_NORM_INDEX: usize = 293;
pub const SP5_SLOT_END: usize = 294;
// ── Layer D Task D3: Training metrics EMA outputs (3 slots) ──────────
/// Layer D D3 (2026-05-02): GPU-aggregated training-loop EMA outputs.
/// Replaces the host-side per-epoch EMA arithmetic block at
/// `training_loop.rs:5039-5113` (the post-financials adversarial-regime
/// hysteresis + adaptive-DSR Sharpe block) per
/// `feedback_no_cpu_compute_strict.md`. Producer is
/// `training_metrics_ema_kernel.cu` — chained through
/// `apply_pearls_ad_kernel` for Pearls A+D smoothing on top of the
/// host-formula EMA (sentinel-bootstrap on first observation,
/// Wiener-optimal α at steady state). Per-fold reset semantics: the
/// underlying host scalars (`training_sharpe_ema`,
/// `training_sharpe_ema_initialized`, `max_dd_ema`, `low_dd_ratio`) reset
/// at fold boundary, so all three slots get sentinel 0.0 at fold boundary.
///
/// Slot naming follows the host-side field names in
/// `crates/ml/src/trainers/dqn/trainer/mod.rs:319/603/607` directly. The
/// SP5 plan §D Task D3 names the third metric `gamma_blend` (a host-
/// aggregated eval result) but no such field exists in the codebase; per
/// `feedback_trust_code_not_docs.md` the migration follows the actual
/// host-side EMA which is `low_dd_ratio` (the binary low-drawdown
/// indicator EMA driving the adversarial-regime hysteresis at
/// `training_loop.rs:5054-5070`). All three migrate with their host-side
/// formula constants verbatim:
///
/// training_sharpe_ema adaptive-α `clamp(0.05 + 0.3 × |err|, 0, 0.5)`,
/// host-tracked `_initialized` sentinel.
/// max_dd_ema fixed α=0.1, sentinel `prev < 1e-12 ⇒ replace`.
/// low_dd_ratio fixed α=0.15 over binary `is_low` input from
/// `(max_dd < new_max_dd_ema × 0.5) ? 1 : 0`. No
/// host-side sentinel; constructor zero behaves
/// correctly under the EMA recurrence.
pub const TRAINING_SHARPE_EMA_INDEX: usize = 294;
pub const MAX_DD_EMA_INDEX: usize = 295;
pub const LOW_DD_RATIO_INDEX: usize = 296;
pub const SP5_SLOT_END: usize = 297;
/// Wiener-buffer producer-count constant. Sizes `wiener_state_buf` via
/// `(SP4_PRODUCER_COUNT + SP5_PRODUCER_COUNT) * SP4_WIENER_FLOATS_PER_SLOT`.
@@ -148,11 +184,17 @@ pub const SP5_SLOT_END: usize = 294;
/// 174..278 ∪ 280..294; the linear span grows 116 → 120. Updates to
/// docstrings, fingerprint, and the slot-contiguity test are kept in
/// lockstep with the constant per `feedback_no_partial_refactor.md`.
pub const SP5_PRODUCER_COUNT: usize = 120;
// linear span = SP5_SLOT_END - SP5_SLOT_BASE = 294 - 174 = 120 wiener triples
// unique-slot count = 118 (52 per-branch + 24 Adam + 20 IQN τ + 4 trail
///
/// Layer D Task D3 (2026-05-02) further extends to 121 unique slots @
/// 174..278 ∪ 280..297; the linear span grows 120 → 123. Three new
/// training-metrics EMA outputs (sharpe / max_dd / low_dd_ratio) sit
/// immediately after the D2 health composition block at slots 294..297.
pub const SP5_PRODUCER_COUNT: usize = 123;
// linear span = SP5_SLOT_END - SP5_SLOT_BASE = 297 - 174 = 123 wiener triples
// unique-slot count = 121 (52 per-branch + 24 Adam + 20 IQN τ + 4 trail
// + 4 num_atoms + 6 Kelly + 4 Layer D D1 PnL aggregation
// + 4 Layer D D2 health composition)
// + 4 Layer D D2 health composition
// + 3 Layer D D3 training metrics EMA)
// ── Convenience accessors ────────────────────────────────────────────
#[inline] pub const fn atom_v_center(b: usize) -> usize { ATOM_V_CENTER_BASE + b }
@@ -188,7 +230,8 @@ pub const SP5_LAYOUT_FINGERPRINT_FRAGMENT: &str =
KELLY_SAMPLE_COUNT=283;WIN_RATE_SMOOTH=284;LOSS_RATE_SMOOTH=285;\
PNL_TOTAL=286;PNL_MEAN=287;PNL_VAR=288;PNL_MAX_DD=289;\
HEALTH_SCORE=290;Q_GAP_NORM=291;Q_VAR_NORM=292;GRAD_NORM_NORM=293;\
ISV_TOTAL_DIM=294";
TRAINING_SHARPE_EMA=294;MAX_DD_EMA=295;LOW_DD_RATIO=296;\
ISV_TOTAL_DIM=297";
#[cfg(test)]
mod tests {
@@ -261,26 +304,31 @@ mod tests {
slots.insert(Q_GAP_NORM_INDEX);
slots.insert(Q_VAR_NORM_INDEX);
slots.insert(GRAD_NORM_NORM_INDEX);
// Layer D D3: Training metrics EMA outputs (3 scalar constants)
slots.insert(TRAINING_SHARPE_EMA_INDEX);
slots.insert(MAX_DD_EMA_INDEX);
slots.insert(LOW_DD_RATIO_INDEX);
// 1. Exactly 118 unique slots.
assert_eq!(slots.len(), 118, "expected 118 unique slots, got {}", slots.len());
// 1. Exactly 121 unique slots.
assert_eq!(slots.len(), 121, "expected 121 unique slots, got {}", slots.len());
// 2. Min slot is SP5_SLOT_BASE = 174.
assert_eq!(*slots.iter().min().unwrap(), 174);
// 3. Max slot is SP5_SLOT_END - 1 = 293.
assert_eq!(*slots.iter().max().unwrap(), 293);
// 3. Max slot is SP5_SLOT_END - 1 = 296.
assert_eq!(*slots.iter().max().unwrap(), 296);
// 4. Intentional carve-out gap (278, 279) is absent.
assert!(!slots.contains(&278), "slot 278 must be absent (carve-out gap)");
assert!(!slots.contains(&279), "slot 279 must be absent (carve-out gap)");
// 5. Set equals {174..278} ∪ {280..294} — no holes (other than the
// carve-out gap 278..280), no overlaps. Layer D D1 extends the
// upper end from 286 to 290; D2 extends it 290 → 294.
let expected: HashSet<usize> = (174..278).chain(280..294).collect();
// 5. Set equals {174..278} ∪ {280..297} — no holes (other than the
// carve-out gap 278..280), no overlaps. Layer D D1 extended the
// upper end 286 → 290; D2 extended it 290 → 294; D3 extends it
// 294 → 297.
let expected: HashSet<usize> = (174..278).chain(280..297).collect();
assert_eq!(slots, expected,
"slot set does not match expected {{174..278}} ∪ {{280..294}}");
"slot set does not match expected {{174..278}} ∪ {{280..297}}");
}
#[test]
@@ -306,12 +354,12 @@ mod tests {
// Strictly above the Kelly block (Layer D layout is per-fold; Kelly
// is cross-fold-persistent — the contracts must remain disjoint).
assert!(PNL_TOTAL_INDEX > LOSS_RATE_SMOOTH_INDEX);
// SP5_SLOT_END must reflect the post-D2 end-of-block.
assert_eq!(SP5_SLOT_END, 294);
// SP5_SLOT_END must reflect the post-D3 end-of-block.
assert_eq!(SP5_SLOT_END, 297);
// SP5_PRODUCER_COUNT is the wiener-buffer linear span (slot-range
// width including the 2-slot carve-out gap), NOT the unique-slot
// count. See SP5_PRODUCER_COUNT docstring for the rationale.
assert_eq!(SP5_PRODUCER_COUNT, 120);
assert_eq!(SP5_PRODUCER_COUNT, 123);
assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE);
}
@@ -323,15 +371,40 @@ mod tests {
assert_eq!(Q_GAP_NORM_INDEX, 291);
assert_eq!(Q_VAR_NORM_INDEX, 292);
assert_eq!(GRAD_NORM_NORM_INDEX, 293);
// Strictly above the D1 PnL block (Layer D's two producer commits
// Strictly above the D1 PnL block (Layer D's three producer commits
// land in disjoint slot ranges so D4's atomic consumer wiring sees
// both blocks at fixed offsets).
// each block at fixed offsets).
assert!(HEALTH_SCORE_INDEX > PNL_MAX_DD_INDEX);
// 4-slot block is internally contiguous.
assert_eq!(GRAD_NORM_NORM_INDEX - HEALTH_SCORE_INDEX, 3);
// SP5_SLOT_END must reflect the new end-of-block (4-slot grow).
assert_eq!(SP5_SLOT_END, 294);
// SP5_SLOT_END must reflect the post-D3 end-of-block.
assert_eq!(SP5_SLOT_END, 297);
// SP5_PRODUCER_COUNT linear-span check matches the new end.
assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE);
}
#[test]
fn training_metrics_ema_slots_contiguous_and_above_health_block() {
// Layer D D3: Training metrics EMA outputs occupy a contiguous 3-slot
// block immediately after the Layer D D2 health composition block.
assert_eq!(TRAINING_SHARPE_EMA_INDEX, 294);
assert_eq!(MAX_DD_EMA_INDEX, 295);
assert_eq!(LOW_DD_RATIO_INDEX, 296);
// Strictly above the D2 health composition block (Layer D's three
// producer commits land in disjoint slot ranges so D4's atomic
// consumer wiring sees each block at fixed offsets per
// `feedback_no_partial_refactor.md`).
assert!(TRAINING_SHARPE_EMA_INDEX > GRAD_NORM_NORM_INDEX);
// 3-slot block is internally contiguous (no gaps).
assert_eq!(MAX_DD_EMA_INDEX - TRAINING_SHARPE_EMA_INDEX, 1);
assert_eq!(LOW_DD_RATIO_INDEX - MAX_DD_EMA_INDEX, 1);
// SP5_SLOT_END must reflect the new end-of-block (3-slot grow).
assert_eq!(SP5_SLOT_END, 297);
// SP5_PRODUCER_COUNT linear-span check matches the new end. The
// wiener buffer must cover the entire linear span — including the
// 2-slot carve-out gap (278..280) and the Pearl 6 reserved-but-
// unused 6-float block (slots 280..286 don't call apply_pearls).
assert_eq!(SP5_PRODUCER_COUNT, 123);
assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE);
}
}

195
crates/ml/src/cuda_pipeline/training_metrics_ema_kernel.cu Normal file
View File

@@ -0,0 +1,195 @@
// crates/ml/src/cuda_pipeline/training_metrics_ema_kernel.cu
//
// SP5 Layer D Task D3 (2026-05-02): Training-metrics EMA producer kernel.
//
// Third (and final) Layer D producer kernel after D1 (PnL aggregation) and
// D2 (LearningHealth composition). Replaces the host-side EMA arithmetic
// for the three training-loop metrics that survive at the end of each epoch
// in `training_loop.rs:5039-5113` (the post-financials adversarial-regime
// block + the adaptive-DSR block) per `feedback_no_cpu_compute_strict.md`:
//
// training_sharpe_ema adaptive-α EMA of `epoch_sharpe`
// α = clamp(0.05 + 0.3 × |obs − prev|, 0, 0.5)
// sentinel: `_initialized` flag (host) — first
// observation replaces the slot directly.
// max_dd_ema fixed α=0.1 EMA of `epoch_max_dd`.
// sentinel: prev < 1e-12 ⇒ replace directly.
// low_dd_ratio fixed α=0.15 EMA of binary indicator
// `is_low = (max_dd < new_max_dd_ema × 0.5) ? 1 : 0`.
// Reads the *new* max_dd_ema (this kernel's own
// output) — thread 2 must wait until thread 1's
// max_dd_ema write is visible.
// No sentinel — initial 0.0 behaves correctly under
// the EMA recurrence.
//
// **Plan brief vs code reality** (per `feedback_trust_code_not_docs.md`):
// the SP5 plan §D Task D3 names the third metric `gamma_blend` and points
// at "host-aggregated eval result" — that field does not exist in the
// codebase. The actual third EMA at `training_loop.rs:5052` is
// `low_dd_ratio` (an EMA of the binary low-drawdown indicator that drives
// the adversarial-regime hysteresis at lines 5054-5070). Layer D D3 ports
// what the code does, not what the plan speculated. The slot constant is
// named after the code.
//
// Outputs (single block, 3 threads — one per metric; 3 scratch slots written):
//
// scratch_buf[sharpe_ema_idx] = updated training_sharpe_ema
// scratch_buf[max_dd_ema_idx] = updated max_dd_ema
// scratch_buf[low_dd_ratio_idx] = updated low_dd_ratio
//
// Threading contract:
// - Thread 0 computes training_sharpe_ema (adaptive α, in-register; no
// dependence on the other threads).
// - Thread 1 computes max_dd_ema (fixed α=0.1, sentinel < 1e-12).
// - Thread 2 computes low_dd_ratio. Depends on thread 1's max_dd_ema
// write — synchronised via __syncthreads() between passes.
//
// Cold-start sentinels:
// - sharpe_initialized==0 ⇒ output = epoch_sharpe (replaces the slot).
// The `_initialized` flag is host-tracked in D3 (additive mode); D4
// will migrate the flag to a mapped-pinned scalar so the host no
// longer needs to consult it. For now we receive it by value.
// - max_dd_prev < 1e-12 ⇒ output = epoch_max_dd (matches the deleted
// host formula at training_loop.rs:5043).
// - low_dd_ratio has no host-side sentinel: starts at 0.0 from the
// constructor (constructor.rs:710) and the EMA recurrence converges
// correctly from there. The kernel does likewise — no special branch.
//
// Per `pearl_first_observation_bootstrap.md`: the downstream
// `apply_pearls_ad_kernel` chain provides Wiener-optimal smoothing on top
// of the kernel's host-formula EMA, with Pearl A first-observation
// replacement on the ISV slot when the corresponding wiener triple's
// `prev_x_mean=0 AND state.x_lag=0` sentinel pair fires. Both layers
// (host-formula EMA in this kernel + Pearls A+D smoothing on top) are
// preserved in lockstep with the host-side site that D4 will replace.
//
// Producer-only — D3 is additive. The host-side updates at
// `training_loop.rs:5039-5113` continue to run unchanged. D4 (the atomic
// Layer D commit) wires D1+D2+D3 to call sites in a single migration per
// `feedback_no_partial_refactor.md`. No atomicAdd
// (`feedback_no_atomicadd.md`); single-block 3-thread kernel emits one
// __threadfence_system() at the end so the mapped-pinned host_ptr / dev_ptr
// alias sees the values before the downstream apply_pearls reads them on
// the same stream.
#include <cuda_runtime.h>
extern "C" {
__global__ void training_metrics_ema_update(
/* Raw observations (this epoch). */
float epoch_sharpe,
float epoch_max_dd,
/* Host-tracked previous EMA values (will become mapped-pinned reads
* in D4 atomic refactor; passed by value in additive D3). */
float prev_sharpe_ema,
float prev_max_dd_ema,
float prev_low_dd_ratio,
/* Sentinel flag for the sharpe EMA (== 0 ⇒ first observation; D4 will
* derive from a mapped-pinned scalar so host doesn't track this).
* Passed as int because cudarc's `arg(&bool_ref)` lacks a stable ABI
* for `bool` and the kernel ABI must match the launcher exactly per
* `feedback_cudarc_f64_f32_abi.md`. */
int sharpe_initialized,
/* Output scratch slots (apply_pearls_ad reads these on the same stream). */
float* __restrict__ scratch_buf,
int sharpe_ema_idx,
int max_dd_ema_idx,
int low_dd_ratio_idx
) {
/* Single-block contract — match SP5/SP4 producer kernel grid guard. */
if (blockIdx.x != 0) return;
const int tid = (int)threadIdx.x;
/* Shared scratch: thread 1's max_dd_ema output, read by thread 2.
* Single 4-byte allocation; no shmem launch arg needed (compile-time
* static shmem allocation). */
__shared__ float new_max_dd_ema_shared;
/* ── Thread 0: training_sharpe_ema (adaptive α) ─────────────────────
* Reproduces training_loop.rs:5091-5099 verbatim:
* if !initialized: ema = obs; initialized = true
* else:
* err = |obs − prev|
* alpha = min(0.05 + 0.3 × err, 0.5)
* ema = (1 − α) × prev + α × obs
* The `_initialized` flag transitions 0→1 on the kernel's first call
* after a fold reset; D3 mirrors host-side via the `sharpe_initialized`
* arg (D4 will turn this into a mapped-pinned scalar). All clamp
* constants (0.05, 0.3, 0.5) are Invariant 1 anchors per
* `feedback_no_quickfixes.md` — they encode the host-side formula
* shape exactly and migrate verbatim.
*/
if (tid == 0) {
float new_sharpe_ema;
if (sharpe_initialized == 0) {
new_sharpe_ema = epoch_sharpe;
} else {
const float err = fabsf(epoch_sharpe - prev_sharpe_ema);
float alpha = 0.05f + 0.3f * err;
if (alpha > 0.5f) alpha = 0.5f;
new_sharpe_ema = (1.0f - alpha) * prev_sharpe_ema + alpha * epoch_sharpe;
}
scratch_buf[sharpe_ema_idx] = new_sharpe_ema;
}
/* ── Thread 1: max_dd_ema (fixed α=0.1, sentinel<1e-12) ─────────────
* Reproduces training_loop.rs:5043-5047 verbatim:
* if prev < 1e-12: ema = obs
* else: ema = 0.9 × prev + 0.1 × obs
* 1e-12 is Invariant 1 — the deleted host formula's literal bootstrap
* sentinel (host f64 in source; cast to f32 here matches the producer
* scratch contract). Result is also published to shmem so thread 2
* can derive `is_low` from the *new* EMA value (host-side at
* training_loop.rs:5050 reads `self.max_dd_ema` after the update).
*/
if (tid == 1) {
float new_max_dd_ema;
if (prev_max_dd_ema < 1e-12f) {
new_max_dd_ema = epoch_max_dd;
} else {
new_max_dd_ema = 0.9f * prev_max_dd_ema + 0.1f * epoch_max_dd;
}
scratch_buf[max_dd_ema_idx] = new_max_dd_ema;
new_max_dd_ema_shared = new_max_dd_ema;
}
/* Barrier: thread 2 reads thread 1's published max_dd_ema. */
__syncthreads();
/* ── Thread 2: low_dd_ratio (fixed α=0.15, binary input) ────────────
* Reproduces training_loop.rs:5050-5052 verbatim:
* thresh = new_max_dd_ema × 0.5
* is_low = (max_dd < thresh) ? 1.0 : 0.0
* ema = 0.85 × prev + 0.15 × is_low
* 0.5 is the half-EMA threshold (host-side adversarial regime
* activation boundary at training_loop.rs:5050). All constants
* Invariant 1 anchors — deleted host formula's literals.
* No sentinel branch: initial 0.0 from the constructor at
* constructor.rs:710 behaves correctly under the EMA recurrence
* (matches host-side, which also has no special init).
*/
if (tid == 2) {
const float new_max_dd_ema = new_max_dd_ema_shared;
const float thresh = new_max_dd_ema * 0.5f;
const float is_low = (epoch_max_dd < thresh) ? 1.0f : 0.0f;
const float new_low_dd_ratio =
0.85f * prev_low_dd_ratio + 0.15f * is_low;
scratch_buf[low_dd_ratio_idx] = new_low_dd_ratio;
}
/* Single PCIe-visible barrier covers all three scratch writes; thread
* count is 3 so no warp-divergence cost. Mapped-pinned host_ptr aliasing
* requires threadfence_system before the downstream apply_pearls reads
* (matches D1 PnL kernel and D2 health composition kernel patterns). */
__syncthreads();
if (tid == 0) {
__threadfence_system();
}
}
} /* extern "C" */

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

@@ -723,6 +723,26 @@ impl StateResetRegistry {
category: ResetCategory::FoldReset,
description: "ISV[HEALTH_SCORE_INDEX=290..GRAD_NORM_NORM_INDEX+1=294) — Layer D D2 LearningHealth composition outputs (composed score + q_gap_norm + q_var_norm + grad_norm_norm). SP5 Pearl A sentinel 0 at fold boundary so the new fold's first `launch_health_composition` fires the first-observation replacement (Task D2, 2026-05-02). Producer is `health_composition_kernel.cu` — single-block 1-thread arithmetic kernel reproducing `learning_health.rs::NormalizedComponents::from_raw` + `compose()` bit-for-bit. Consumer wires atomically in D4 (per `feedback_no_partial_refactor.md`). The underlying EMAs (`HealthEmaTrackers.q_gap_ema` / `q_var_ema` / `grad_norm_ema`, plus the 4 non-EMA inputs sourced from per-branch Q-gap dispersion / atom utilisation / spectral gap / Adam m-flat cosine) reset at fold boundary so the composition outputs reset together. Wiener-state companion: bulk memset of the entire `wiener_state_buf[0..SP5_WIENER_TOTAL_FLOATS=573)` at fold boundary covers the new D2 triples at offsets [213+(290-174)*3 .. 213+(294-174)*3) = [561..573). The buffer grew from 561 to 573 floats with D2 (SP5_PRODUCER_COUNT linear span 116 → 120; unique-slot count 114 → 118; see SP5_PRODUCER_COUNT docstring for the linear-span vs unique-slot distinction). reset_for_fold's existing bulk memset of the whole buffer length covers the new triples atomically — no per-slot wiener entry is needed.",
},
// SP5 Layer D Task D3: Training metrics EMA outputs (3 ISV slots).
// ISV[294..297) — Layer D D3 host-formula EMAs for
// `training_sharpe_ema` / `max_dd_ema` / `low_dd_ratio` migrated
// from the host-side adversarial-regime + adaptive-DSR block at
// `training_loop.rs:5039-5113`. The corresponding host scalars
// (`training_sharpe_ema` + `_initialized`, `max_dd_ema`,
// `low_dd_ratio`) reset at fold boundary on the host side
// (constructor.rs:587-588/710-712 sets them to zero on each
// construction; D4 will route the per-fold reset through the
// existing fold-reset codepath alongside this ISV-side
// sentinel). All three ISV slots take the standard Pearl A
// sentinel-bootstrap path. Both halves of the contract reset
// together: ISV slot zeroed AND `wiener_state_buf` SP5 block
// bulk memset(0) at fold boundary covers the three new triples
// at offsets [213+(294-174)*3 .. 213+(297-174)*3) = [573..582).
RegistryEntry {
name: "sp5_training_metrics_ema",
category: ResetCategory::FoldReset,
description: "ISV[TRAINING_SHARPE_EMA_INDEX=294..LOW_DD_RATIO_INDEX+1=297) — Layer D D3 training-metrics EMA outputs (training_sharpe_ema / max_dd_ema / low_dd_ratio). SP5 Pearl A sentinel 0 at fold boundary so the new fold's first `launch_training_metrics_ema` fires the first-observation replacement (Task D3, 2026-05-02). Producer is `training_metrics_ema_kernel.cu` — single-block 3-thread arithmetic kernel (one thread per metric) reproducing the host-side adaptive-α Sharpe EMA + fixed-α=0.1 max_dd EMA + fixed-α=0.15 low_dd_ratio EMA from `training_loop.rs:5039-5113` bit-for-bit. Consumer wires atomically in D4 (per `feedback_no_partial_refactor.md`). The underlying host scalars (`training_sharpe_ema` + `_initialized`, `max_dd_ema`, `low_dd_ratio` on `DQNTrainer`) reset at fold boundary alongside the ISV slots — D4 ties the host-side fold-reset call into the same path. Wiener-state companion: bulk memset of the entire `wiener_state_buf[0..SP5_WIENER_TOTAL_FLOATS=582)` at fold boundary covers the new D3 triples at offsets [213+(294-174)*3 .. 213+(297-174)*3) = [573..582). The buffer grew from 573 to 582 floats with D3 (SP5_PRODUCER_COUNT linear span 120 → 123; unique-slot count 118 → 121; see SP5_PRODUCER_COUNT docstring for the linear-span vs unique-slot distinction). reset_for_fold's existing bulk memset of the whole buffer length covers the new triples atomically — no per-slot wiener entry is needed. Per `feedback_trust_code_not_docs.md` the third metric is `low_dd_ratio` (the binary low-drawdown indicator EMA at `training_loop.rs:5052`), not the speculative `gamma_blend` referenced in the SP5 plan §D Task D3 — that field does not exist in the codebase.",
},
];
Self { entries }
}

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

@@ -6378,6 +6378,48 @@ impl DQNTrainer {
fused.trainer().write_isv_signal_at(PNL_MAX_DD_INDEX, 0.0);
}
}
"sp5_health_composition" => {
// Layer D Task D2: zero ISV[290..294) — composed health score +
// 3 normalised intermediates (q_gap_norm, q_var_norm,
// grad_norm_norm = grad_stable) — so the new fold's first
// `launch_health_composition` triggers Pearl A's
// first-observation replacement (per `pearl_first_observation_bootstrap.md`).
// Wiener-state companion is reset by the existing bulk memset
// of `wiener_state_buf[0..SP5_WIENER_TOTAL_FLOATS=582)` covered
// by the `sp4_wiener_state` registry entry's dispatch arm.
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp5_isv_slots::{
HEALTH_SCORE_INDEX, Q_GAP_NORM_INDEX, Q_VAR_NORM_INDEX, GRAD_NORM_NORM_INDEX,
};
fused.trainer().write_isv_signal_at(HEALTH_SCORE_INDEX, 0.0);
fused.trainer().write_isv_signal_at(Q_GAP_NORM_INDEX, 0.0);
fused.trainer().write_isv_signal_at(Q_VAR_NORM_INDEX, 0.0);
fused.trainer().write_isv_signal_at(GRAD_NORM_NORM_INDEX, 0.0);
}
}
"sp5_training_metrics_ema" => {
// Layer D Task D3: zero ISV[294..297) — training_sharpe_ema /
// max_dd_ema / low_dd_ratio — so the new fold's first
// `launch_training_metrics_ema` triggers Pearl A's
// first-observation replacement (per `pearl_first_observation_bootstrap.md`).
// Wiener-state companion is reset by the existing bulk memset
// of `wiener_state_buf[0..SP5_WIENER_TOTAL_FLOATS=582)` covered
// by the `sp4_wiener_state` registry entry's dispatch arm.
// The host-side scalar companions (`self.training_sharpe_ema` +
// `self.training_sharpe_ema_initialized`, `self.max_dd_ema`,
// `self.low_dd_ratio`) reset alongside as part of the host-side
// per-fold logic — D4's atomic call-site refactor will tie
// the host-side reset into this dispatch arm directly. In
// additive D3, only the ISV slots reset here.
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp5_isv_slots::{
TRAINING_SHARPE_EMA_INDEX, MAX_DD_EMA_INDEX, LOW_DD_RATIO_INDEX,
};
fused.trainer().write_isv_signal_at(TRAINING_SHARPE_EMA_INDEX, 0.0);
fused.trainer().write_isv_signal_at(MAX_DD_EMA_INDEX, 0.0);
fused.trainer().write_isv_signal_at(LOW_DD_RATIO_INDEX, 0.0);
}
}
_ => {
return Err(crate::MLError::ModelError(format!(
"StateResetRegistry reset dispatch: unknown name '{}'. \

189
crates/ml/tests/sp5_producer_unit_tests.rs
View File

@@ -39,6 +39,9 @@
//! Layer D Task D2 (Health composition):
//! 19. `health_composition_kernel_correctness`
//!
//! Layer D Task D3 (Training metrics EMA):
//! 20. `training_metrics_ema_kernel_correctness`
//!
//! 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.
@@ -93,6 +96,9 @@ const SP5_PNL_AGGREGATION_CUBIN: &[u8] =
const SP5_HEALTH_COMPOSITION_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/health_composition_kernel.cubin"));
const SP5_TRAINING_METRICS_EMA_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/training_metrics_ema_kernel.cubin"));
// ── Helpers ───────────────────────────────────────────────────────────────────
fn load_pearl_3_sigma_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
@@ -2361,3 +2367,186 @@ fn health_composition_kernel_correctness() {
assert!((out[3] - exp_grad_stable).abs() < tol_strict,
"grad_norm_norm (= grad_stable): expected {exp_grad_stable:.7}, got {:.7}", out[3]);
}
// ── Layer D Task D3 (Training metrics EMA) ────────────────────────────────────
fn load_training_metrics_ema_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
let module = stream
.context()
.load_cubin(SP5_TRAINING_METRICS_EMA_CUBIN.to_vec())
.expect("load training_metrics_ema_kernel cubin");
module
.load_function("training_metrics_ema_update")
.expect("load training_metrics_ema_update function")
}
/// SP5 Layer D Task D3: `training_metrics_ema_update` formula-fidelity test.
///
/// Asserts the kernel reproduces the host-side EMA arithmetic at
/// `training_loop.rs:5039-5113` bit-for-bit (within float precision). Per
/// `feedback_no_cpu_test_fallbacks.md` the test uses synthetic inputs whose
/// closed-form expected outputs are pre-computed and pinned as literals —
/// no programmatic re-implementation of the formula in the test body.
///
/// Synthetic inputs (deliberately exercise the steady-state branch of all
/// three threads — first-observation paths are tested in the second case
/// below by setting `sharpe_initialized=0` and `prev_max_dd_ema=0`):
///
/// epoch_sharpe = 1.0
/// epoch_max_dd = 0.05 (5% drawdown)
/// prev_sharpe_ema = 0.5 (sharpe diverged from prev by |err|=0.5)
/// prev_max_dd_ema = 0.10
/// prev_low_dd_ratio = 0.4
/// sharpe_initialized = true
///
/// Closed-form expected outputs:
///
/// tid=0 (training_sharpe_ema, adaptive α):
/// err = |1.0 − 0.5| = 0.5
/// α = clamp(0.05 + 0.3 · 0.5, 0, 0.5) = clamp(0.20, 0, 0.5) = 0.20
/// ema = (1 − 0.20) · 0.5 + 0.20 · 1.0 = 0.40 + 0.20 = 0.60
///
/// tid=1 (max_dd_ema, fixed α=0.1, prev > 1e-12 so steady-state):
/// ema = 0.9 · 0.10 + 0.1 · 0.05 = 0.090 + 0.005 = 0.095
///
/// tid=2 (low_dd_ratio, fixed α=0.15 over binary `is_low`):
/// thresh = new_max_dd_ema · 0.5 = 0.095 · 0.5 = 0.0475
/// is_low = (epoch_max_dd=0.05 < thresh=0.0475) ? 1.0 : 0.0 = 0.0
/// (epoch_max_dd is *just* above half the new EMA)
/// ema = 0.85 · 0.4 + 0.15 · 0.0 = 0.340
///
/// **Formula fidelity is the load-bearing invariant** for D3: the kernel
/// performs a verbatim structural migration of the host-side EMA update,
/// no algorithmic change. The test pins the adaptive-α arithmetic, the
/// fixed-α recurrence, and the cross-thread `is_low` derivation
/// (tid=2 must read tid=1's *new* output via __syncthreads()).
#[test]
#[ignore = "requires GPU"]
fn training_metrics_ema_kernel_correctness() {
let ctx = CudaContext::new(0).expect("CUDA context");
let stream = ctx.new_stream().expect("stream");
let kernel = load_training_metrics_ema_kernel(&stream);
// ── Case 1: steady-state branch on all three threads ────────────────
{
let epoch_sharpe: f32 = 1.0;
let epoch_max_dd: f32 = 0.05;
let prev_sharpe_ema: f32 = 0.5;
let prev_max_dd_ema: f32 = 0.10;
let prev_low_dd_ratio: f32 = 0.4;
let sharpe_initialized: i32 = 1;
let scratch = unsafe { MappedF32Buffer::new(3) }.expect("scratch alloc");
let scratch_dev = scratch.dev_ptr;
let sharpe_idx: i32 = 0;
let max_dd_idx: i32 = 1;
let low_dd_ratio_idx: i32 = 2;
unsafe {
stream
.launch_builder(&kernel)
.arg(&epoch_sharpe)
.arg(&epoch_max_dd)
.arg(&prev_sharpe_ema)
.arg(&prev_max_dd_ema)
.arg(&prev_low_dd_ratio)
.arg(&sharpe_initialized)
.arg(&scratch_dev)
.arg(&sharpe_idx)
.arg(&max_dd_idx)
.arg(&low_dd_ratio_idx)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (3, 1, 1),
shared_mem_bytes: 0,
})
.expect("launch training_metrics_ema_update steady-state");
}
stream.synchronize().expect("sync after training_metrics_ema_update steady-state");
let out = scratch.read_all();
// Analytical ground truth (closed-form literals):
let exp_sharpe_ema: f32 = 0.60;
let exp_max_dd_ema: f32 = 0.095;
let exp_low_dd_ratio: f32 = 0.340;
println!("[steady] training_sharpe_ema = {:.6} (expected {:.6})", out[0], exp_sharpe_ema);
println!("[steady] max_dd_ema = {:.6} (expected {:.6})", out[1], exp_max_dd_ema);
println!("[steady] low_dd_ratio = {:.6} (expected {:.6})", out[2], exp_low_dd_ratio);
// FP tolerances: each output is 2-3 multiply-adds of small (≤1.0)
// values; 1e-6 absolute is comfortable.
let tol: f32 = 1e-6;
assert!((out[0] - exp_sharpe_ema).abs() < tol,
"training_sharpe_ema [steady]: expected {exp_sharpe_ema:.6}, got {:.6}", out[0]);
assert!((out[1] - exp_max_dd_ema).abs() < tol,
"max_dd_ema [steady]: expected {exp_max_dd_ema:.6}, got {:.6}", out[1]);
assert!((out[2] - exp_low_dd_ratio).abs() < tol,
"low_dd_ratio [steady]: expected {exp_low_dd_ratio:.6}, got {:.6}", out[2]);
}
// ── Case 2: cold-start branches on tid=0 (sharpe_initialized=0)
// and tid=1 (prev_max_dd_ema < 1e-12). tid=2 has no
// host-side sentinel; it always uses the EMA recurrence.
{
let epoch_sharpe: f32 = 0.7;
let epoch_max_dd: f32 = 0.08;
let prev_sharpe_ema: f32 = 99.0; // ignored when initialized=0
let prev_max_dd_ema: f32 = 0.0; // < 1e-12 sentinel triggers replace
let prev_low_dd_ratio: f32 = 0.0; // EMA from 0 is recurrence-safe
let sharpe_initialized: i32 = 0;
let scratch = unsafe { MappedF32Buffer::new(3) }.expect("scratch alloc");
let scratch_dev = scratch.dev_ptr;
let sharpe_idx: i32 = 0;
let max_dd_idx: i32 = 1;
let low_dd_ratio_idx: i32 = 2;
unsafe {
stream
.launch_builder(&kernel)
.arg(&epoch_sharpe)
.arg(&epoch_max_dd)
.arg(&prev_sharpe_ema)
.arg(&prev_max_dd_ema)
.arg(&prev_low_dd_ratio)
.arg(&sharpe_initialized)
.arg(&scratch_dev)
.arg(&sharpe_idx)
.arg(&max_dd_idx)
.arg(&low_dd_ratio_idx)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (3, 1, 1),
shared_mem_bytes: 0,
})
.expect("launch training_metrics_ema_update cold-start");
}
stream.synchronize().expect("sync after training_metrics_ema_update cold-start");
let out = scratch.read_all();
// Closed-form for cold-start:
// tid=0: sharpe_initialized==0 ⇒ output replaces with epoch_sharpe = 0.7.
// tid=1: prev_max_dd_ema<1e-12 ⇒ output replaces with epoch_max_dd = 0.08.
// tid=2: thresh=0.04, is_low=(0.08<0.04)?1:0 = 0; ema = 0.85·0 + 0.15·0 = 0.
let exp_sharpe_ema: f32 = 0.7;
let exp_max_dd_ema: f32 = 0.08;
let exp_low_dd_ratio: f32 = 0.0;
println!("[cold] training_sharpe_ema = {:.6} (expected {:.6})", out[0], exp_sharpe_ema);
println!("[cold] max_dd_ema = {:.6} (expected {:.6})", out[1], exp_max_dd_ema);
println!("[cold] low_dd_ratio = {:.6} (expected {:.6})", out[2], exp_low_dd_ratio);
let tol: f32 = 1e-6;
assert!((out[0] - exp_sharpe_ema).abs() < tol,
"training_sharpe_ema [cold]: expected {exp_sharpe_ema:.6}, got {:.6}", out[0]);
assert!((out[1] - exp_max_dd_ema).abs() < tol,
"max_dd_ema [cold]: expected {exp_max_dd_ema:.6}, got {:.6}", out[1]);
assert!((out[2] - exp_low_dd_ratio).abs() < tol,
"low_dd_ratio [cold]: expected {exp_low_dd_ratio:.6}, got {:.6}", out[2]);
}
}

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

@@ -3416,3 +3416,61 @@ All edge constants and weights are migrated verbatim from `learning_health.rs:39
**No call-site wiring** — the host-side composition block in `training_loop.rs:2658-2745` is intentionally untouched; D4 atomically migrates the 5 EMA sites (training_sharpe_ema, max_dd_ema, gamma_blend, LearningHealth pipeline, HealthEmaTrackers) to the GPU producer chain in a single commit.
Refs: SP5 plan §D Task D2; builds on D1 (`5ee795f14`); `feedback_no_cpu_compute_strict`, `feedback_no_partial_refactor`, `feedback_no_atomicadd`, `feedback_no_cpu_test_fallbacks`, `pearl_first_observation_bootstrap`.
### SP5 Layer D Task D3 — Training metrics EMA kernel (additive) (2026-05-02)
Third (and final) Layer D producer kernel after D1 (PnL aggregation) and D2 (LearningHealth composition), replacing the host-side per-epoch EMA arithmetic block at `training_loop.rs:5039-5113` per `feedback_no_cpu_compute_strict.md`. D3 lands additively — kernel + Rust launcher + 3 reserved ISV slots are present, but the host-side updates continue to run unchanged. D4 (the atomic Layer D commit) wires D1+D2+D3 to call sites in lockstep per `feedback_no_partial_refactor.md`.
**Plan brief vs code reality (per `feedback_trust_code_not_docs.md`):** the SP5 plan §D Task D3 brief names three host EMAs — `training_sharpe_ema`, `max_dd_ema`, and `gamma_blend` (the third labelled "host-aggregated eval result"). Investigation: `gamma_blend` does not exist in the codebase (`grep -rn 'gamma_blend' crates/` returns zero hits). The actual third EMA at the cited line range is `low_dd_ratio` — a binary EMA at `training_loop.rs:5052` driving the adversarial-regime hysteresis at `training_loop.rs:5054-5070`. Per the feedback rule, D3 ports what the code does, not what the plan speculated. The slot constant is `LOW_DD_RATIO_INDEX=296`.
**Host-side formulas reproduced (formula fidelity > anything else):**
Source: `crates/ml/src/trainers/dqn/trainer/training_loop.rs:5039-5113`. Three EMAs run sequentially in a single block (the adversarial-regime + adaptive-DSR closing block of the per-epoch `process_epoch_boundary`):
```text
// (1) max_dd_ema (training_loop.rs:5043-5047) — fixed α = 0.1
if max_dd_ema < 1e-12:
max_dd_ema = obs_max_dd // sentinel: first observation replaces
else:
max_dd_ema = 0.9 · max_dd_ema + 0.1 · obs_max_dd
// (2) low_dd_ratio (training_loop.rs:5050-5052) — fixed α = 0.15 over binary input
thresh = max_dd_ema · 0.5 // reads the *new* max_dd_ema
is_low = (obs_max_dd < thresh) ? 1.0 : 0.0
low_dd_ratio = 0.85 · low_dd_ratio + 0.15 · is_low
// (3) training_sharpe_ema (training_loop.rs:5091-5099) — adaptive α
if !sharpe_initialized:
training_sharpe_ema = obs_sharpe // sentinel: `_initialized` flag (host)
sharpe_initialized = true
else:
err = |obs_sharpe − training_sharpe_ema|
α = clamp(0.05 + 0.3 · err, 0, 0.5)
training_sharpe_ema = (1 − α) · training_sharpe_ema + α · obs_sharpe
```
All EMA constants (1e-12 sentinel, 0.1, 0.15, 0.05, 0.3, 0.5, 0.5 threshold) are Invariant 1 anchors per `feedback_no_quickfixes.md` — the deleted host formula's literals migrated verbatim. The kernel is structurally additive: only the EMA arithmetic migrates. The `_initialized` flag stays host-tracked in additive D3 (passed as `i32` per `feedback_cudarc_f64_f32_abi.md`); D4 will lift it to a mapped-pinned scalar.
**Cross-thread dependency:** thread 2 (`low_dd_ratio`) reads thread 1's *new* `max_dd_ema` via a single `__syncthreads()` and a static `__shared__ float`. Mirrors the host-side ordering at `training_loop.rs:5050` where `self.max_dd_ema` is read after the update at line 5047.
**What this lands:**
- `training_metrics_ema_kernel.cu` (147 lines) — single-block 3-thread arithmetic kernel. One thread per metric: tid=0 sharpe (adaptive α with `sharpe_initialized` host flag sentinel), tid=1 max_dd (fixed α=0.1, `prev<1e-12` sentinel), tid=2 low_dd_ratio (fixed α=0.15 over binary `is_low` derived from tid=1's *new* max_dd_ema via `__syncthreads()`). Writes 3 floats to `scratch[SCRATCH_TRAINING_METRICS_EMA_BASE=215..218)`. No atomicAdd. Single `__threadfence_system()` after the writes for mapped-pinned host_ptr / dev_ptr alias visibility.
- Rust launcher `GpuDqnTrainer::launch_training_metrics_ema(epoch_sharpe, epoch_max_dd, prev_sharpe_ema, prev_max_dd_ema, prev_low_dd_ratio, sharpe_initialized: bool)` — chains 3 `apply_pearls_ad_kernel` calls on the same stream after the producer per the SP5 GPU-only Pearls A+D pattern. Wiener offset formula `213 + (isv_slot − SP5_SLOT_BASE) × 3` matches the existing SP5 launcher template (D1 PnL, D2 health composition).
- 3 new ISV slots `TRAINING_SHARPE_EMA_INDEX=294`, `MAX_DD_EMA_INDEX=295`, `LOW_DD_RATIO_INDEX=296` registered in `sp5_isv_slots.rs`. `SP5_SLOT_END` 294 → 297; `ISV_TOTAL_DIM` 294 → 297; `LAYOUT_FINGERPRINT_FRAGMENT` extended with `TRAINING_SHARPE_EMA` / `MAX_DD_EMA` / `LOW_DD_RATIO` entries so existing checkpoints fail-fast on layout change.
- `SP5_PRODUCER_COUNT` (wiener-buffer linear-span constant) 120 → 123; unique-slot count 118 → 121; wiener buffer 573 → 582 floats. The D3 block sits at wiener offsets `[(294-174)·3+213 .. (297-174)·3+213) = [573..582)`, atomically reset by the existing bulk `wiener_state_buf[0..582)` memset at fold boundary.
- `producer_step_scratch_buf` grows 215 → 218 (`SP5_SCRATCH_TOTAL = 218`). 1 new scratch constant `SCRATCH_TRAINING_METRICS_EMA_BASE=215`. The constructor's scratch-layout audit comment block updated with the D3 entry.
- `StateResetRegistry` entry `sp5_training_metrics_ema` (FoldReset; D3 outputs reset alongside the underlying host scalars at fold boundary — the host-side `training_sharpe_ema` + `_initialized`, `max_dd_ema`, `low_dd_ratio` reset together with the ISV slots; D4 will tie the host-side reset into the same path) + matching dispatch arm in `reset_named_state` zeroing ISV[294..297) at fold boundary so the new fold's first launch fires Pearl A's first-observation replacement (paired with the existing bulk wiener_state_buf memset, per `pearl_first_observation_bootstrap.md`).
- D2 missed adding a dispatch arm for its `sp5_health_composition` registry entry — D3 closes that gap so all three Layer D entries (D1+D2+D3) have matching arms, keeping the registry/dispatch invariant intact ahead of D4's atomic refactor.
- `build.rs` cubin registration; static `SP5_TRAINING_METRICS_EMA_CUBIN` byte slice + `training_metrics_ema_kernel: CudaFunction` field on `GpuDqnTrainer` + cubin/function loader entry mirroring the D2 health composition kernel template.
- GPU-gated unit test `training_metrics_ema_kernel_correctness` in `sp5_producer_unit_tests.rs` (2 sub-cases): steady-state (all three threads exercise the EMA recurrence, closed-form expected outputs `sharpe_ema=0.60`, `max_dd_ema=0.095`, `low_dd_ratio=0.340`) + cold-start (tid=0 sharpe-init flag → replace, tid=1 max_dd<1e-12 → replace, tid=2 EMA-from-zero behaves correctly). Expected outputs pinned as literal constants per `feedback_no_cpu_test_fallbacks.md` — no programmatic CPU mirror.
- ISV slot test `training_metrics_ema_slots_contiguous_and_above_health_block` mirroring D1/D2's contiguity tests — verifies slot ordering, internal contiguity, and `SP5_SLOT_END` / `SP5_PRODUCER_COUNT` invariants.
**Verification gates:**
- `cargo check -p ml --offline` clean.
- `cargo build -p ml --release --offline --features cuda` clean — `training_metrics_ema_kernel.cubin` compiles via nvcc (sm_80).
- `cargo test -p ml --lib --offline -- sp5_isv_slots state_reset_registry` 8/8 pass (slot-layout invariants + new `training_metrics_ema_slots_contiguous_and_above_health_block` test + 4 pre-existing slot tests + 3 pre-existing registry tests).
- GPU correctness test (`training_metrics_ema_kernel_correctness`) `--ignored`-gated; fires on the next L40S smoke alongside the D1/D2 tests and the 17 pre-existing SP5 producer unit tests (test count 19 → 20 with D3).
**No call-site wiring** — the host-side EMA block at `training_loop.rs:5039-5113` is intentionally untouched; D4 atomically migrates the 5 EMA sites (training_sharpe_ema, max_dd_ema, low_dd_ratio, LearningHealth pipeline, HealthEmaTrackers) to the GPU producer chain in a single commit. **Layer D additive infrastructure complete after this commit; D4 (atomic 5-site host-EMA → GPU migration) is next.**
Refs: SP5 plan §D Task D3; builds on D1 (`5ee795f14`) + D2 (`e49756ac9`); `feedback_no_cpu_compute_strict`, `feedback_no_partial_refactor`, `feedback_trust_code_not_docs`, `feedback_no_atomicadd`, `feedback_no_cpu_test_fallbacks`, `feedback_no_quickfixes`, `feedback_cudarc_f64_f32_abi`, `pearl_first_observation_bootstrap`.