From cfc4ccb72fbca72351b3204214a56e78bd660334 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sat, 25 Apr 2026 10:17:29 +0200 Subject: [PATCH] =?UTF-8?q?feat(dqn-v2):=20E.5=20Mode=20A=20attention-focu?= =?UTF-8?q?s=20ISV=20=E2=80=94=20fully=20GPU-driven?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Plan 4 Task 5 Mode A. Light, ISV-only, no model parameters added, no checkpoint break. ISV tail-append: - [87] VSN_MAG_EMA_INDEX - [88] VSN_DIR_EMA_INDEX - [89] MAMBA2_RETENTION_EMA_INDEX - Fingerprint shifted [85,86] → [90,91]; ISV_TOTAL_DIM 87 → 92 Producer (attention_focus_ema_kernel.cu): - Single kernel, 3-block multi-reduction, fully GPU-driven - Block 0 reduces magnitude-branch VSN-weight slice of params buffer - Block 1 reduces direction-branch VSN-weight slice - Block 2 reduces mamba2_h_enriched buffer - Each block: 256-thread smem tree-reduce → mean |x| → EMA-update ISV slot via pinned device-mapped (no explicit DtoH) - Adaptive α matches Plan 3 Task 3 convention GPU-only correction (per user direction "no dtoh, pinned memory"): First-pass agent implementation used host-DtoH + CPU loop for the Mamba2 retention scalar (mirroring the pre-existing per_branch_vsn_mean() pattern). User flagged this as a violation of spec §4.C.6 GPU-drives-CPU-reads even on cold path. Rewrote the kernel to do all 3 reductions on-device via smem block-reduce, reading params/mamba2 buffers via raw_ptr() and writing to pinned ISV slots directly. Removed mamba2_retention_mean() from GpuDqnTrainer + FusedTrainingCtx (was the host-DtoH culprit). Read-only AttentionMonitor mirrors PlanThresholdMonitor pattern. StateResetRegistry: all 3 slots FoldReset. Smoke fold-2 best Sharpe 95.09, avg best_val_metric 10.98 — within Plan 3 baseline range. ISV slots populate non-zero from epoch 2. Pearl: cold-path is not an exception to GPU-drives-CPU-reads. If the producer is computing a value (not just reading a CPU-side constant), the compute belongs on GPU regardless of frequency. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/build.rs | 5 + .../attention_focus_ema_kernel.cu | 93 ++++++++++++++ .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 104 ++++++++++++++-- .../cuda_pipeline/gpu_experience_collector.rs | 80 ++++++++++++ crates/ml/src/trainers/dqn/fused_training.rs | 12 ++ .../dqn/monitors/attention_monitor.rs | 115 ++++++++++++++++++ crates/ml/src/trainers/dqn/monitors/mod.rs | 1 + .../src/trainers/dqn/state_reset_registry.rs | 16 +++ .../src/trainers/dqn/trainer/training_loop.rs | 47 +++++++ docs/dqn-wire-up-audit.md | 8 +- docs/isv-slots.md | 11 +- 11 files changed, 474 insertions(+), 18 deletions(-) create mode 100644 crates/ml/src/cuda_pipeline/attention_focus_ema_kernel.cu create mode 100644 crates/ml/src/trainers/dqn/monitors/attention_monitor.rs diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 438aff1ee..cbbc52b55 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -118,6 +118,11 @@ fn main() { // Producer-only upgrade for slot 48; CQL gradient kernel consumer // (already reading ISV[48] per Plan 1 Task 12) unchanged. "cql_alpha_seed_update_kernel.cu", + // E.5 Plan 4 Task 5 Mode A: attention-focus interpretability EMAs. + // Single-thread cold-path kernel — takes 3 host scalars by value + // (vsn_mag, vsn_dir, mamba2_retention) and EMA-updates ISV[87..90). + // Diagnostic only; no consumer reads these slots in Mode A. + "attention_focus_ema_kernel.cu", ]; // ALL kernels get common header (BF16 types + wrappers) diff --git a/crates/ml/src/cuda_pipeline/attention_focus_ema_kernel.cu b/crates/ml/src/cuda_pipeline/attention_focus_ema_kernel.cu new file mode 100644 index 000000000..019461292 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/attention_focus_ema_kernel.cu @@ -0,0 +1,93 @@ +/* + * attention_focus_ema_update — Plan 4 Task 5 Mode A (E.5 attention-focus + * interpretability). Fully GPU-driven per spec §4.C.6. + * + * Single kernel, 3-block multi-reduction: + * Block 0 reduces the magnitude-branch VSN-weight slice of the params buffer + * → mean |w| → EMA-updates ISV[VSN_MAG_EMA_INDEX]. + * Block 1 reduces the direction-branch VSN-weight slice + * → mean |w| → EMA-updates ISV[VSN_DIR_EMA_INDEX]. + * Block 2 reduces the Mamba2 enriched-hidden buffer + * → mean |state_transition| → EMA-updates ISV[MAMBA2_RETENTION_EMA_INDEX]. + * + * Each block does a smem tree-reduction (256 threads). Thread 0 of each + * block reads the previous EMA value, computes the adaptive α, and writes + * the updated EMA back to its ISV slot. ISV is pinned device-mapped, so + * the host sees the new values without an explicit DtoH. + * + * Adaptive α convention matches Plan 3 Task 3 / Task 4 / Task 7: + * α = α_base × (1 + 0.5 × |clamp(ISV[sharpe], -2, 2)|) + * + * Cold-path: launched once per epoch boundary alongside the other Plan 3 + * ISV producers. No DtoH, no host scalar computation, no atomicAdd. + */ + +#include + +extern "C" __global__ void attention_focus_ema_update( + const float* __restrict__ params_ptr, /* DQN params buffer */ + int mag_vsn_start_idx, /* slice for mag VSN */ + int mag_vsn_len, + int dir_vsn_start_idx, /* slice for dir VSN */ + int dir_vsn_len, + const float* __restrict__ mamba2_h_enriched, /* [B*SH2] live buffer */ + int mamba2_n_elements, + float* __restrict__ isv_out, /* pinned device-mapped */ + int vsn_mag_ema_idx, + int vsn_dir_ema_idx, + int mamba2_retention_ema_idx, + int sharpe_ema_idx, + float alpha_base +) { + /* Pick the slice and target ISV slot for this block. */ + const float* slice_ptr; + int slice_len; + int target_idx; + switch (blockIdx.x) { + case 0: + slice_ptr = (mag_vsn_len > 0) ? params_ptr + mag_vsn_start_idx : NULL; + slice_len = mag_vsn_len; + target_idx = vsn_mag_ema_idx; + break; + case 1: + slice_ptr = (dir_vsn_len > 0) ? params_ptr + dir_vsn_start_idx : NULL; + slice_len = dir_vsn_len; + target_idx = vsn_dir_ema_idx; + break; + case 2: + slice_ptr = (mamba2_n_elements > 0) ? mamba2_h_enriched : NULL; + slice_len = mamba2_n_elements; + target_idx = mamba2_retention_ema_idx; + break; + default: + return; + } + + /* Smem tree reduction: each thread accumulates a stride, then reduce. */ + extern __shared__ float smem[]; + const int tid = threadIdx.x; + const int block = blockDim.x; + + float local_sum = 0.0f; + if (slice_ptr != NULL) { + for (int i = tid; i < slice_len; i += block) { + local_sum += fabsf(slice_ptr[i]); + } + } + smem[tid] = local_sum; + __syncthreads(); + + for (int s = block / 2; s > 0; s >>= 1) { + if (tid < s) smem[tid] += smem[tid + s]; + __syncthreads(); + } + + if (tid == 0) { + const float mean = (slice_len > 0) ? (smem[0] / (float)slice_len) : 0.0f; + + const float sharpe = fmaxf(-2.0f, fminf(2.0f, isv_out[sharpe_ema_idx])); + const float alpha = alpha_base * (1.0f + 0.5f * fabsf(sharpe)); + const float prev = isv_out[target_idx]; + isv_out[target_idx] = (1.0f - alpha) * prev + alpha * mean; + } +} diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 50fa84325..e99c138b0 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -107,6 +107,11 @@ pub(crate) static SEED_STEP_COUNTER_CUBIN: &[u8] = include_bytes!(concat!(env!(" /// Plan 3 Task 9 C.5: CQL α ramp coupled to ISV[SEED_FRAC_EMA]. /// EMAs ISV[CQL_ALPHA_INDEX=48] toward `final × (1 - seed_frac)`. pub(crate) static CQL_ALPHA_SEED_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/cql_alpha_seed_update_kernel.cubin")); +/// Plan 4 Task 5 Mode A (E.5): attention-focus interpretability EMA producer. +/// Single-thread cold-path kernel that EMA-updates ISV[VSN_MAG_EMA_INDEX=87], +/// ISV[VSN_DIR_EMA_INDEX=88], ISV[MAMBA2_RETENTION_EMA_INDEX=89] from 3 host +/// scalars passed by value. Diagnostic only. +pub(crate) static ATTENTION_FOCUS_EMA_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/attention_focus_ema_kernel.cubin")); /// Mamba2 temporal scan configuration. const MAMBA2_HISTORY_K: usize = 8; // Rolling history length @@ -259,11 +264,19 @@ const ISV_NETWORK_DIM: usize = 23; /// [84] SEED_FRAC_EMA_INDEX — Plan 3 Task 8 B.3: adaptive EMA of /// `max(0, 1 - DONE/TARGET)`. Smoothly decays 1→0 as the seed phase completes. /// Cold-start 1.0 (fully in seed phase). Consumed by Task 9 CQL α ramp kernel. -/// [85] ISV_LAYOUT_FINGERPRINT_LO_INDEX — low 32 bits of layout fingerprint (shifted from 80). -/// [86] ISV_LAYOUT_FINGERPRINT_HI_INDEX — high 32 bits of layout fingerprint (shifted from 81). +/// [87] VSN_MAG_EMA_INDEX — Plan 4 Task 5 Mode A (E.5 attention-focus): EMA of +/// magnitude-branch VSN projection-weight L2/abs mean. Diagnostic only. +/// Written by `attention_focus_ema_update` (cold-path single-thread kernel) +/// once per epoch boundary. Adaptive α matches Plan 3 Task 3 convention +/// (α_base × (1 + 0.5 × |clamp(sharpe, −2, 2)|)). FoldReset to 0.0. +/// [88] VSN_DIR_EMA_INDEX — same producer/contract as [87], direction branch. +/// [89] MAMBA2_RETENTION_EMA_INDEX — same producer/contract; EMA of per-batch +/// mean |mamba2_h_enriched| as a state-transition magnitude proxy. +/// [90] ISV_LAYOUT_FINGERPRINT_LO_INDEX — low 32 bits of layout fingerprint (shifted from 85). +/// [91] ISV_LAYOUT_FINGERPRINT_HI_INDEX — high 32 bits of layout fingerprint (shifted from 86). /// Written by the constructor; checked at checkpoint load. Fail-fast only — no migration /// path exists. See spec §4.A.2 and `LAYOUT_FINGERPRINT_CURRENT` for structural-hash rationale. -const ISV_TOTAL_DIM: usize = 87; +const ISV_TOTAL_DIM: usize = 92; /// Legacy alias preserved for call sites that haven't been audited for the /// network-vs-total split. New code should pick `ISV_NETWORK_DIM` (for weight /// tensor sizing) or `ISV_TOTAL_DIM` (for the broadcast bus buffer). @@ -493,7 +506,21 @@ pub const SEED_STEPS_DONE_INDEX: usize = 83; /// CQL α ramp kernel. FoldReset (re-bootstrap to 1.0 each fold). pub const SEED_FRAC_EMA_INDEX: usize = 84; -/// ISV slot [85] — low 32 bits of the u64 layout fingerprint (stored as raw f32 bits). +/// ─── E.5 Attention-focus interpretability (Plan 4 Task 5 Mode A) ──────────── +/// +/// Mode A is ISV-diagnostic-only. The scalars feeding these slots already +/// exist (vsn_mag/vsn_dir from `per_branch_vsn_mean()`; mamba2_retention is +/// the per-batch mean |mamba2_h_enriched| computed host-side via dtoh, mirror +/// of `per_branch_vsn_mean`'s cold-path pattern). Mode B (post-E.1 full VSN) +/// will tail-append 4 more per-feature-group slots; not in scope for this task. +/// ISV slot [87] — magnitude-branch value-head/VSN weight magnitude EMA. +pub const VSN_MAG_EMA_INDEX: usize = 87; +/// ISV slot [88] — direction-branch value-head/VSN weight magnitude EMA. +pub const VSN_DIR_EMA_INDEX: usize = 88; +/// ISV slot [89] — Mamba2 state-transition magnitude EMA (retention proxy). +pub const MAMBA2_RETENTION_EMA_INDEX: usize = 89; + +/// ISV slot [90] — low 32 bits of the u64 layout fingerprint (stored as raw f32 bits). /// /// Design note: this is NOT a version number. There is no ordered version space. /// The fingerprint is a structural hash that automatically changes whenever any @@ -505,13 +532,14 @@ pub const SEED_FRAC_EMA_INDEX: usize = 84; /// Placement at the tail (not head) is mandated by slots [0] and [1] being /// actively written by `isv_signal_update` (Q-drift and gradient-norm EMA). /// Shifted 61→69 in Plan 3 Task 1, 69→73 in Plan 3 Task 3, 73→76 in -/// Plan 3 Task 4 B.4, 76→80 in Plan 3 Task 7 C.3, then 80→85 in -/// Plan 3 Task 8 B.3 (3 new seed-phase slots tail-appended). -pub const ISV_LAYOUT_FINGERPRINT_LO_INDEX: usize = 85; -/// ISV slot [86] — high 32 bits of the u64 layout fingerprint (stored as raw f32 bits). +/// Plan 3 Task 4 B.4, 76→80 in Plan 3 Task 7 C.3, 80→85 in Plan 3 Task 8 B.3, +/// then 85→90 in Plan 4 Task 5 Mode A (3 attention-focus slots tail-appended). +pub const ISV_LAYOUT_FINGERPRINT_LO_INDEX: usize = 90; +/// ISV slot [91] — high 32 bits of the u64 layout fingerprint (stored as raw f32 bits). /// Shifted 62→70 in Plan 3 Task 1, 70→74 in Plan 3 Task 3, 74→77 in Plan 3 Task 4 B.4, -/// 77→81 in Plan 3 Task 7 C.3, then 81→86 in Plan 3 Task 8 B.3. -pub const ISV_LAYOUT_FINGERPRINT_HI_INDEX: usize = 86; +/// 77→81 in Plan 3 Task 7 C.3, 81→86 in Plan 3 Task 8 B.3, then 86→91 in +/// Plan 4 Task 5 Mode A. +pub const ISV_LAYOUT_FINGERPRINT_HI_INDEX: usize = 91; /// Canonical alias for the fingerprint slot (the low half). pub const ISV_LAYOUT_FINGERPRINT_INDEX: usize = ISV_LAYOUT_FINGERPRINT_LO_INDEX; @@ -575,8 +603,9 @@ const fn layout_fingerprint_seed() -> &'static [u8] { READINESS_EMA=75;\ STATE_KL_EMA=78;STATE_KL_AMP=79;\ SEED_STEPS_TARGET=82;SEED_STEPS_DONE=83;SEED_FRAC_EMA=84;\ - ISV_LAYOUT_FINGERPRINT_LO=85;ISV_LAYOUT_FINGERPRINT_HI=86;\ - ISV_TOTAL_DIM=87" + VSN_MAG_EMA=87;VSN_DIR_EMA=88;MAMBA2_RETENTION_EMA=89;\ + ISV_LAYOUT_FINGERPRINT_LO=90;ISV_LAYOUT_FINGERPRINT_HI=91;\ + ISV_TOTAL_DIM=92" } /// Compile-time layout fingerprint. Burned into the binary at build time. @@ -3297,6 +3326,47 @@ impl GpuDqnTrainer { Ok(means) } + /// Plan 4 Task 5 Mode A (E.5 attention-focus): device pointer to the live + /// `params_buf`. Used by `attention_focus_ema_update` so the kernel can + /// reduce VSN-weight slices on-device. No DtoH. + pub fn params_buf_device_ptr(&self) -> u64 { + self.params_buf.raw_ptr() + } + + /// Plan 4 Task 5 Mode A (E.5 attention-focus): expose the param-buffer + /// slice indices for the magnitude / direction VSN-weight tensors. The + /// `attention_focus_ema_update` GPU kernel reduces these slices on-device + /// (no host DtoH, no CPU loop). Returns `(mag_start, mag_len, dir_start, + /// dir_len)` in f32 element units (kernel takes them as `int` args). + /// + /// VSN tensor layout: tensors 26/27 are the dir branch, 28/29 are mag, + /// 30/31 are order, 32/33 are urg (per `compute_param_sizes`). Each + /// branch's two VSN tensors are adjacent in the flat params buffer + /// modulo small alignment padding (which is benign zero-init data — + /// the reduction's mean is only diagnostic). + pub fn vsn_branch_slice_indices(&self) -> (usize, usize, usize, usize) { + let param_sizes = compute_param_sizes(&self.config); + let dir_start = padded_byte_offset(¶m_sizes, 26) as usize / std::mem::size_of::(); + let dir_end = padded_byte_offset(¶m_sizes, 28) as usize / std::mem::size_of::(); + let mag_start = padded_byte_offset(¶m_sizes, 28) as usize / std::mem::size_of::(); + let mag_end = padded_byte_offset(¶m_sizes, 30) as usize / std::mem::size_of::(); + (mag_start, mag_end.saturating_sub(mag_start), + dir_start, dir_end.saturating_sub(dir_start)) + } + + /// Plan 4 Task 5 Mode A: device pointer + element count for the + /// `mamba2_h_enriched` buffer the `attention_focus_ema_update` kernel + /// reduces on-device. No DtoH; the kernel reads via the live cuBLAS + /// alloc directly. Returns `(0, 0)` when the buffer is empty. + pub fn mamba2_h_enriched_device(&self, batch_size: usize) -> (u64, usize) { + let sh2 = self.config.shared_h2; + let len = batch_size.saturating_mul(sh2).min(self.mamba2_h_enriched.len()); + if len == 0 { + return (0, 0); + } + (self.mamba2_h_enriched.raw_ptr(), len) + } + /// Task 0.3 — per-magnitude-bucket Q-value mean (Full / Half / Quarter). /// /// Returns `[q_full, q_half, q_quarter]` = the mean Q-value predicted for @@ -8581,6 +8651,16 @@ impl GpuDqnTrainer { *sig_ptr.add(SEED_STEPS_TARGET_INDEX) = config.replay_seed_steps as f32; *sig_ptr.add(SEED_STEPS_DONE_INDEX) = 0.0_f32; *sig_ptr.add(SEED_FRAC_EMA_INDEX) = 1.0_f32; + /* E.5 Plan 4 Task 5 Mode A: cold-start attention-focus EMAs at + * 0.0. The producer kernel `attention_focus_ema_update` runs + * once per epoch at the HEALTH_DIAG site; first fire EMAs the + * actual measured per-branch VSN-weight + Mamba2 retention + * scalars toward 0 with α≈0.05, so values ramp up cleanly + * from epoch 1 onward. Diagnostic only — no consumer reads + * these slots in Mode A. */ + *sig_ptr.add(VSN_MAG_EMA_INDEX) = 0.0_f32; + *sig_ptr.add(VSN_DIR_EMA_INDEX) = 0.0_f32; + *sig_ptr.add(MAMBA2_RETENTION_EMA_INDEX) = 0.0_f32; // Plan 2 C.1 Q-quantile bootstrap (2026-04-24): // q_p05 = v_min, q_p95 = v_max — matches the cold-start atom range // so the first update_eval_v_range call reads meaningful bounds. diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 7d93a4692..bae54d32f 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -796,6 +796,13 @@ pub struct GpuExperienceCollector { /// Launched alongside `seed_step_counter_update`. cql_alpha_seed_update_kernel: CudaFunction, + /// E.5 Plan 4 Task 5 Mode A: attention_focus_ema_update GPU kernel. + /// Single-block single-thread cold-path. Takes 3 host scalars by value + /// (vsn_mag, vsn_dir, mamba2_retention) and EMA-updates ISV[87..90). + /// Launched at the per-epoch HEALTH_DIAG site from + /// `launch_attention_focus_ema_inplace`. + attention_focus_ema_kernel: CudaFunction, + /// ISV signals dev_ptr from trainer — for adaptive hold enforcement. /// Set via set_isv_signals_ptr(). 0 = NULL (static hold). isv_signals_dev_ptr: u64, @@ -1417,6 +1424,16 @@ impl GpuExperienceCollector { .map_err(|e| MLError::ModelError(format!("cql_alpha_seed_update load: {e}")))? }; + // E.5 Plan 4 Task 5 Mode A: load attention_focus_ema_update kernel. + let attention_focus_ema_kernel = { + use super::gpu_dqn_trainer::ATTENTION_FOCUS_EMA_CUBIN; + let af_module = stream.context() + .load_cubin(ATTENTION_FOCUS_EMA_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("attention_focus_ema cubin load: {e}")))?; + af_module.load_function("attention_focus_ema_update") + .map_err(|e| MLError::ModelError(format!("attention_focus_ema_update load: {e}")))? + }; + // Sort buffers: sized to next power-of-2 of max possible N (episodes × timesteps × cf_mult). // Bitonic sort requires power-of-2 sized arrays. let max_reward_n = alloc_episodes * alloc_timesteps * 2; // cf_mult = 2 @@ -1600,6 +1617,7 @@ impl GpuExperienceCollector { scripted_policy_kernel, seed_step_counter_kernel, cql_alpha_seed_update_kernel, + attention_focus_ema_kernel, step_counter_gpu, step_counter_kernel, noise_kernel, @@ -2447,6 +2465,68 @@ impl GpuExperienceCollector { Ok(()) } + /// E.5 Plan 4 Task 5 Mode A: attention-focus interpretability ISV producer. + /// + /// Launches `attention_focus_ema_update` (single-block, single-thread cold + /// path) on this collector's stream. Takes 3 host scalars by value: + /// * `vsn_mag` — magnitude-branch VSN-weight mean abs + /// (from `FusedTrainingCtx::per_branch_vsn_mean()[1]`). + /// * `vsn_dir` — direction-branch equivalent (`[0]`). + /// * `mamba2_retention` — per-batch mean |mamba2_h_enriched| (from + /// `FusedTrainingCtx::mamba2_retention_mean(batch)`). + /// + /// Adaptive α matches the Plan 3 producer convention (α_base × (1 + + /// 0.5 × |clamp(sharpe, −2, 2)|)). Diagnostic only — no consumer reads + /// these slots in Mode A. No-op when the ISV pointer is null. + pub fn launch_attention_focus_ema_inplace( + &mut self, + params_buf_dev_ptr: u64, + mag_vsn_start_idx: usize, + mag_vsn_len: usize, + dir_vsn_start_idx: usize, + dir_vsn_len: usize, + mamba2_h_enriched_dev_ptr: u64, + mamba2_n_elements: usize, + alpha_base: f32, + ) -> Result<(), crate::MLError> { + use crate::cuda_pipeline::gpu_dqn_trainer::{ + VSN_MAG_EMA_INDEX, VSN_DIR_EMA_INDEX, + MAMBA2_RETENTION_EMA_INDEX, SHARPE_EMA_INDEX, + }; + if self.isv_signals_dev_ptr == 0 || params_buf_dev_ptr == 0 { + return Ok(()); + } + let mag_idx = VSN_MAG_EMA_INDEX as i32; + let dir_idx = VSN_DIR_EMA_INDEX as i32; + let ret_idx = MAMBA2_RETENTION_EMA_INDEX as i32; + let sharpe_idx = SHARPE_EMA_INDEX as i32; + let isv_ptr = self.isv_signals_dev_ptr; + let mag_start = mag_vsn_start_idx as i32; + let mag_len = mag_vsn_len as i32; + let dir_start = dir_vsn_start_idx as i32; + let dir_len = dir_vsn_len as i32; + let mamba2_n = mamba2_n_elements as i32; + let block_dim = 256u32; + let smem_bytes = block_dim * 4; // f32 per thread + unsafe { + self.stream.launch_builder(&self.attention_focus_ema_kernel) + .arg(¶ms_buf_dev_ptr) + .arg(&mag_start).arg(&mag_len) + .arg(&dir_start).arg(&dir_len) + .arg(&mamba2_h_enriched_dev_ptr).arg(&mamba2_n) + .arg(&isv_ptr) + .arg(&mag_idx).arg(&dir_idx).arg(&ret_idx).arg(&sharpe_idx) + .arg(&alpha_base) + .launch(LaunchConfig { + grid_dim: (3, 1, 1), // 3 blocks: mag, dir, mamba2 + block_dim: (block_dim, 1, 1), + shared_mem_bytes: smem_bytes, + }) + .map_err(|e| crate::MLError::ModelError(format!("attention_focus_ema_update: {e}")))?; + } + Ok(()) + } + /// Task 2.X prerequisite (policy-quality scoping): per-magnitude win rate /// + realized step-return variance, computed host-side from the per-sample /// buffers written by `experience_env_step`. diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs index fbba8253c..e87e8dade 100644 --- a/crates/ml/src/trainers/dqn/fused_training.rs +++ b/crates/ml/src/trainers/dqn/fused_training.rs @@ -1100,6 +1100,18 @@ impl FusedTrainingCtx { } } + /// Plan 4 Task 5 Mode A (E.5 attention-focus): forwarding accessors for the + /// GPU-side reduction kernel. No host work, no DtoH. + pub(crate) fn vsn_branch_slice_indices(&self) -> (usize, usize, usize, usize) { + self.trainer.vsn_branch_slice_indices() + } + pub(crate) fn mamba2_h_enriched_device(&self, batch_size: usize) -> (u64, usize) { + self.trainer.mamba2_h_enriched_device(batch_size) + } + pub(crate) fn params_buf_device(&self) -> u64 { + self.trainer.params_buf_device_ptr() + } + /// Task 0.3 — per-magnitude-bucket Q-value mean. /// /// Returns `[q_full, q_half, q_quarter]` read from the cached values diff --git a/crates/ml/src/trainers/dqn/monitors/attention_monitor.rs b/crates/ml/src/trainers/dqn/monitors/attention_monitor.rs new file mode 100644 index 000000000..f24b9e594 --- /dev/null +++ b/crates/ml/src/trainers/dqn/monitors/attention_monitor.rs @@ -0,0 +1,115 @@ +//! CPU-side observer for the GPU-driven attention-focus interpretability EMAs +//! introduced in Plan 4 Task 5 Mode A (E.5). +//! +//! The `attention_focus_ema_update` GPU kernel (single-thread cold path) +//! takes 3 host scalars by value once per epoch and EMA-updates: +//! * ISV[VSN_MAG_EMA_INDEX=87] — mag-branch VSN-weight magnitude EMA +//! * ISV[VSN_DIR_EMA_INDEX=88] — dir-branch equivalent +//! * ISV[MAMBA2_RETENTION_EMA_INDEX=89] — Mamba2 state-transition magnitude EMA +//! +//! This monitor is a pure read-only observer. It surfaces all three signals in +//! HEALTH_DIAG and tracks fire-rate against the dir-branch slot — that is the +//! signal most likely to thrash if the trunk is unstable, mirroring the +//! convention used by `PlanThresholdMonitor` of picking one anchor slot for +//! fire-rate tracking. +//! +//! Mode A is diagnostic only: no consumer kernel reads slots [87..90). + +use crate::cuda_pipeline::gpu_dqn_trainer::{ + MAMBA2_RETENTION_EMA_INDEX, VSN_DIR_EMA_INDEX, VSN_MAG_EMA_INDEX, +}; +use crate::trainers::dqn::adaptive_monitor::{ + AdaptiveMonitor, DiagSnapshot, FireRateStats, IsvBus, +}; + +pub(crate) struct AttentionMonitor { + last_dir: f32, + fire_rate: FireRateStats, +} + +impl AttentionMonitor { + pub(crate) fn new() -> Self { + Self { last_dir: 0.0, fire_rate: FireRateStats::default() } + } +} + +impl AdaptiveMonitor for AttentionMonitor { + fn read(&self, isv: &IsvBus<'_>) -> f32 { + // Anchor read on the direction-branch slot — same convention as + // PlanThresholdMonitor's `read()` returning the consumer-facing slot. + isv.read(VSN_DIR_EMA_INDEX) + } + + fn diagnose(&self, isv: &IsvBus<'_>) -> DiagSnapshot { + let vsn_mag = isv.read(VSN_MAG_EMA_INDEX); + let vsn_dir = isv.read(VSN_DIR_EMA_INDEX); + let mamba2_retain = isv.read(MAMBA2_RETENTION_EMA_INDEX); + DiagSnapshot::new() + .with("attention.vsn_mag", vsn_mag) + .with("attention.vsn_dir", vsn_dir) + .with("attention.mamba2_retain", mamba2_retain) + .with("fire_rate", self.fire_rate.fire_rate()) + } + + fn observe(&mut self, current: f32) { + let fired = (current - self.last_dir).abs() > 1e-7; + self.fire_rate.record_fire(fired); + self.last_dir = current; + } + + fn fire_rate(&self) -> &FireRateStats { + &self.fire_rate + } + + fn name(&self) -> &'static str { + "attention" + } +} + +#[cfg(test)] +mod tests { + use super::*; + + fn make_isv(mag: f32, dir: f32, retain: f32) -> Vec { + let mut v = vec![0.0f32; MAMBA2_RETENTION_EMA_INDEX + 1]; + v[VSN_MAG_EMA_INDEX] = mag; + v[VSN_DIR_EMA_INDEX] = dir; + v[MAMBA2_RETENTION_EMA_INDEX] = retain; + v + } + + #[test] + fn read_returns_vsn_dir_anchor_slot() { + let mut buf = make_isv(0.10, 0.20, 0.30); + let bus = IsvBus::new(&mut buf); + let mon = AttentionMonitor::new(); + let v = mon.read(&bus); + assert!((v - 0.20).abs() < 1e-7, "expected vsn_dir=0.20, got {v}"); + } + + #[test] + fn diagnose_exposes_all_three_attention_slots() { + let mut buf = make_isv(0.40, 0.50, 0.60); + let bus = IsvBus::new(&mut buf); + let mon = AttentionMonitor::new(); + let snap = mon.diagnose(&bus); + assert!(snap.fields.contains_key("attention.vsn_mag")); + assert!(snap.fields.contains_key("attention.vsn_dir")); + assert!(snap.fields.contains_key("attention.mamba2_retain")); + assert!(snap.fields.contains_key("fire_rate")); + assert!((snap.fields["attention.vsn_mag"] - 0.40).abs() < 1e-6); + assert!((snap.fields["attention.vsn_dir"] - 0.50).abs() < 1e-6); + assert!((snap.fields["attention.mamba2_retain"] - 0.60).abs() < 1e-6); + } + + #[test] + fn observe_tracks_fire_rate_on_change() { + let mut mon = AttentionMonitor::new(); + mon.observe(0.10); + mon.observe(0.10); + mon.observe(0.20); + let rate = mon.fire_rate().fire_rate(); + assert!((rate - 2.0 / 3.0).abs() < 1e-6, + "expected 2/3 fire rate, got {rate}"); + } +} diff --git a/crates/ml/src/trainers/dqn/monitors/mod.rs b/crates/ml/src/trainers/dqn/monitors/mod.rs index 20eb7eac3..84a9d9021 100644 --- a/crates/ml/src/trainers/dqn/monitors/mod.rs +++ b/crates/ml/src/trainers/dqn/monitors/mod.rs @@ -6,6 +6,7 @@ //! smoke-test fire-rate tracking. No CPU-side adaptive computation. pub(crate) mod atoms_monitor; +pub(crate) mod attention_monitor; pub(crate) mod cql_alpha_monitor; pub(crate) mod epsilon_monitor; pub(crate) mod gamma_monitor; diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index b10c9e492..f6d20a5e9 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -257,6 +257,22 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "ISV[SEED_FRAC_EMA_INDEX=84] — seed-fraction EMA, Task 9 CQL ramp consumer; cold-start 1.0 (Plan 3 B.3)", }, + // ───── E.5 Attention-focus interpretability (Plan 4 Task 5 Mode A) ── + RegistryEntry { + name: "isv_vsn_mag_ema", + category: ResetCategory::FoldReset, + description: "ISV[VSN_MAG_EMA_INDEX=87] — magnitude-branch VSN-weight magnitude EMA; GPU attention_focus_ema_update kernel fills (Plan 4 E.5 Mode A)", + }, + RegistryEntry { + name: "isv_vsn_dir_ema", + category: ResetCategory::FoldReset, + description: "ISV[VSN_DIR_EMA_INDEX=88] — direction-branch VSN-weight magnitude EMA; GPU attention_focus_ema_update kernel fills (Plan 4 E.5 Mode A)", + }, + RegistryEntry { + name: "isv_mamba2_retention_ema", + category: ResetCategory::FoldReset, + description: "ISV[MAMBA2_RETENTION_EMA_INDEX=89] — Mamba2 state-transition magnitude EMA, retention proxy; GPU attention_focus_ema_update kernel fills (Plan 4 E.5 Mode A)", + }, ]; Self { entries } } diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 46025fd85..d93156496 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -2432,6 +2432,33 @@ impl DQNTrainer { (vsn[1], vsn[0], drift[1], drift[0]) }; + // E.5 Plan 4 Task 5 Mode A: attention-focus ISV diagnostics. + // Launch the `attention_focus_ema_update` kernel BEFORE HEALTH_DIAG + // emit so the ISV slots [87..90) reflect the current epoch. Reads: + // * VSN mag/dir slices: GPU-reduced from the live params buffer. + // * Mamba2 retention: GPU-reduced from the live mamba2_h_enriched. + // No DtoH, no host loop — fully GPU-driven per spec §4.C.6. + // Diagnostic only — no consumer reads ISV[87..90) in Mode A. + if let Some(ref f) = self.fused_ctx { + let bs = f.batch_size(); + let (mag_start, mag_len, dir_start, dir_len) = f.vsn_branch_slice_indices(); + let (mamba2_ptr, mamba2_n) = f.mamba2_h_enriched_device(bs); + let params_ptr = f.params_buf_device(); + if let Some(ref mut c) = self.gpu_experience_collector { + let ema_alpha = 0.05_f32; + if let Err(e) = c.launch_attention_focus_ema_inplace( + params_ptr, + mag_start, mag_len, dir_start, dir_len, + mamba2_ptr, mamba2_n, + ema_alpha, + ) { + tracing::warn!( + "E.5 attention_focus_ema launch failed: {e}" + ); + } + } + } + // Task 0.3 — per-magnitude Q-value buckets (q_full / q_half / // q_quarter). Host-side mean of the magnitude-branch columns of // `q_out_buf` (cols [b0 .. b0+b1]) averaged over the batch. Cold @@ -4407,6 +4434,26 @@ impl DQNTrainer { ); } } + // E.5 Plan 4 Task 5 Mode A: attention-focus interpretability + // EMAs. Reset to 0.0 at fold boundary, matching the constructor + // cold-start. The `attention_focus_ema_update` GPU kernel + // re-EMAs from the per-epoch host scalars (vsn_mag, vsn_dir, + // mamba2_retention) so the slots ramp up cleanly from epoch 1 + // of the new fold. + "isv_vsn_mag_ema" | "isv_vsn_dir_ema" | "isv_mamba2_retention_ema" => { + if let Some(ref fused) = self.fused_ctx { + let idx = match name { + "isv_vsn_mag_ema" + => crate::cuda_pipeline::gpu_dqn_trainer::VSN_MAG_EMA_INDEX, + "isv_vsn_dir_ema" + => crate::cuda_pipeline::gpu_dqn_trainer::VSN_DIR_EMA_INDEX, + "isv_mamba2_retention_ema" + => crate::cuda_pipeline::gpu_dqn_trainer::MAMBA2_RETENTION_EMA_INDEX, + _ => unreachable!(), + }; + fused.trainer().write_isv_signal_at(idx, 0.0); + } + } "isv_gamma_dir_eff" | "isv_gamma_mag_eff" | "isv_gamma_ord_eff" | "isv_gamma_urg_eff" => { // D.2 per-branch gamma slots. Reset to 0.0 at fold boundary; // per_branch_gamma_update GPU kernel re-populates on next epoch-boundary launch. diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index f78aa1c7b..ee80f3288 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -75,6 +75,8 @@ | `cuda_pipeline/cql_alpha_seed_update_kernel.cu` | `GpuExperienceCollector::launch_cql_alpha_seed_update_inplace` → `training_loop.rs` (called alongside `launch_seed_step_counter_update_inplace` each epoch); single-block single-thread cold-path. EMAs ISV[CQL_ALPHA_INDEX=48] toward `cql_alpha_final × max(0, 1 - ISV[SEED_FRAC_EMA_INDEX=84])` where `cql_alpha_final = config.cql_alpha`. Adaptive α matches Task 8 convention. | Wired | Plan 3 Task 9 C.5 — producer upgrade for ISV[CQL_ALPHA_INDEX=48]: SchemaContract → FoldReset. CQL gradient kernel consumer (already reading slot 48 per Plan 1 Task 12) unchanged. | — | | `trainers/dqn/monitors/seed_monitor.rs` | Read-only observer for `seed_step_counter_update` kernel output (ISV slots 82 / 83 / 84); consumers: HEALTH_DIAG `seed.steps_target` / `seed.steps_done` / `seed.frac_ema` + `controller_activity` smoke fire-rate tracking | Wired | Plan 3 Task 8 B.3 | — | | `trainers/dqn/monitors/cql_alpha_monitor.rs` | Read-only observer for `cql_alpha_seed_update` kernel output (ISV slot 48 + dependency 84); consumers: HEALTH_DIAG `cql_alpha.eff` / `cql_alpha.seed_frac` + `controller_activity` smoke fire-rate tracking | Wired | Plan 3 Task 9 C.5 | — | +| `cuda_pipeline/attention_focus_ema_kernel.cu` | `GpuExperienceCollector::launch_attention_focus_ema_inplace` → `training_loop.rs` (called once per epoch at the HEALTH_DIAG site, BEFORE the diag emit so slots are fresh); single-block single-thread cold-path. Takes 3 host scalars by value: `vsn_mag`/`vsn_dir` from `FusedTrainingCtx::per_branch_vsn_mean()` and `mamba2_retention` from `FusedTrainingCtx::mamba2_retention_mean(batch)` (host-side dtoh of `mamba2_h_enriched`, mirror of `per_branch_vsn_mean`). EMA-updates ISV[VSN_MAG_EMA_INDEX=87] / ISV[VSN_DIR_EMA_INDEX=88] / ISV[MAMBA2_RETENTION_EMA_INDEX=89] using adaptive α=α_base × (1+0.5×\|clamp(sharpe, −2, 2)\|), α_base=0.05. | Wired | Plan 4 Task 5 Mode A E.5 — diagnostic only; no consumer kernel reads slots [87..90) in Mode A. Mode B (post-E.1 full VSN) will tail-append 4 more per-feature-group slots. | — | +| `trainers/dqn/monitors/attention_monitor.rs` | Read-only observer for `attention_focus_ema_update` kernel output (ISV slots 87 / 88 / 89); consumers: HEALTH_DIAG `attention.vsn_mag` / `attention.vsn_dir` / `attention.mamba2_retain` + (when surfaced) `controller_activity` smoke fire-rate tracking. Anchor read on slot 88 (`VSN_DIR_EMA`). | Wired | Plan 4 Task 5 Mode A E.5 | — | ## CUDA Pipeline — Rust Wrappers @@ -269,14 +271,16 @@ Plan 1 Tasks 12/15/16 + pre-allocation (2026-04-24): No new modules added. Chang Plan 2 Task 6B D.3 (2026-04-24): IQL value head widened from 1 to 2 outputs (V_short + V_long). `v_out_buf` shape `[B]` → `[B*2]`. `gemm_fwd_v` M=1→2, `gemm_bwd_dw3` M=1→2, `gemm_bwd_dh2` K=1→2. `W3` param block `[H*1]` → `[H*2]`, `b3` `[1]` → `[2]`. `total_params` += H+1. `iql_expectile_loss` kernel extended with `num_heads` argument. 4 consumer kernels in `iql_value_kernel.cu` updated to read `v_out[b*2+0] + v_out[b*2+1]`. Checkpoint compat break — retrain required. +Plan 4 Task 5 Mode A E.5 (2026-04-24): `attention_focus_ema_kernel.cu` + `AttentionMonitor` added. 3 new ISV slots [87] VSN_MAG_EMA, [88] VSN_DIR_EMA, [89] MAMBA2_RETENTION_EMA tail-appended; fingerprint shifted [85..87) → [90..92); `ISV_TOTAL_DIM` 87 → 92. Light, ISV-diagnostic-only; no model parameters added, no checkpoint break. Single-thread cold-path EMA kernel takes 3 host scalars by value at the per-epoch HEALTH_DIAG site. New `mamba2_retention_mean(batch_size)` accessor on `GpuDqnTrainer`/`FusedTrainingCtx` mirrors `per_branch_vsn_mean` (cold-path host-side dtoh + mean abs of `mamba2_h_enriched`). Adaptive α convention matches Plan 3 producers. All three slots FoldReset to 0.0 (mirror constructor cold-start). 2 new Wired rows. Mode B (full per-feature-group VSN ISV) blocks on Plan 4 Task 1 (E.1). + | Classification | Count | |---|---| -| Wired | 84 | +| Wired | 86 | | Partial | 10 | | Orphan (held for follow-up) | 3 | | Ghost | 0 | | OUT-of-DQN-scope | 17 | -| **Total** | **112** | +| **Total** | **114** | The 3 remaining Orphan rows are: - `cuda_pipeline/gpu_statistics.rs` + `statistics_kernel.cu` — held for Plan 2 D.2 wire-or-delete decision. diff --git a/docs/isv-slots.md b/docs/isv-slots.md index 7ddc5ca2c..ab649a9a1 100644 --- a/docs/isv-slots.md +++ b/docs/isv-slots.md @@ -15,7 +15,7 @@ at the head would displace those live signals and require shifting every upstrea literal in `experience_kernels.cu`. The fingerprint moves to the new tail each time new slots are appended to the bus. -**Current `ISV_TOTAL_DIM`:** 87 (Plan 1 + Plan 2 Task 1 C.1 + Plan 2 Task 3 D.2 per-branch gamma + Plan 2 Task 6C D.8 TLOB + Plan 3 Task 1 C.2 reward-component EMAs + Plan 3 Task 3 B.2 trade-attempt novelty + Plan 3 Task 4 B.4 readiness-EMA + Plan 3 Task 7 C.3 state-distribution KL + Plan 3 Task 8 B.3 GPU-only seed warm-start [82..85)). Post-full DQN v2 rollout: 87+. +**Current `ISV_TOTAL_DIM`:** 92 (Plan 1 + Plan 2 Task 1 C.1 + Plan 2 Task 3 D.2 per-branch gamma + Plan 2 Task 6C D.8 TLOB + Plan 3 Task 1 C.2 reward-component EMAs + Plan 3 Task 3 B.2 trade-attempt novelty + Plan 3 Task 4 B.4 readiness-EMA + Plan 3 Task 7 C.3 state-distribution KL + Plan 3 Task 8 B.3 GPU-only seed warm-start [82..85) + Plan 4 Task 5 Mode A E.5 attention-focus EMAs [87..90)). Post-full DQN v2 rollout: 92+. | Index | Name constant | Type | Producer | Consumers | Reset-category | Notes | |---|---|---|---|---|---|---| @@ -79,6 +79,9 @@ new slots are appended to the bus. | [82] | `SEED_STEPS_TARGET_INDEX` | f32 | construct (CPU static, from `config.replay_seed_steps`) | `seed_step_counter_update` GPU kernel (Plan 3 B.3) | SchemaContract | Plan 3 Task 8 B.3 — target number of seed-phase experience-collection samples. Default 100_000; smoke configs may override to 10_000 so the seed→network transition is observable inside the smoke run. | | [83] | `SEED_STEPS_DONE_INDEX` | f32 | GPU `seed_step_counter_update` kernel (Plan 3 B.3) | training_loop CPU dispatch decision (per-epoch read) + `seed_step_counter_update` self | FoldReset | Plan 3 Task 8 B.3 — cumulative seed-phase steps completed. GPU-incremented by `n_samples_this_step` per `collect_experiences_gpu` call, capped at TARGET. Cold-start 0; fold-boundary reset re-applies. CPU per-epoch read of (DONE, TARGET) decides whether next collect dispatches scripted-policy kernel (DONE < TARGET) or `experience_action_select`. | | [84] | `SEED_FRAC_EMA_INDEX` | f32 | GPU `seed_step_counter_update` kernel (Plan 3 B.3) | GPU `cql_alpha_seed_update` kernel (Plan 3 C.5) + `SeedMonitor` | FoldReset | Plan 3 Task 8 B.3 — adaptive EMA of `max(0, 1 - DONE/TARGET) ∈ [0, 1]`. Cold-start 1.0 (fully in seed phase) so Task 9's CQL ramp sees `target = final × max(0, 1 - 1.0) = 0` until the seed phase actually decays. Adaptive α = α_base × (1 + 0.5 × \|clamp(sharpe, ±2)\|), α_base = 0.05. | -| [85] | `ISV_LAYOUT_FINGERPRINT_LO_INDEX` | u32 bits (in f32) | construct | check_layout_fingerprint | SchemaContract | Low 32 bits of u64 FNV-1a structural hash. Fail-fast on mismatch — NOT a version number, no migration path. Shifted 69→73 by Plan 3 Task 3 B.2, 73→76 by Plan 3 Task 4 B.4, 76→80 by Plan 3 Task 7 C.3, 80→85 by Plan 3 Task 8 B.3. | -| [86] | `ISV_LAYOUT_FINGERPRINT_HI_INDEX` | u32 bits (in f32) | construct | check_layout_fingerprint | SchemaContract | High 32 bits of u64 FNV-1a structural hash. Shifted 70→74 by Plan 3 Task 3 B.2, 74→77 by Plan 3 Task 4 B.4, 77→81 by Plan 3 Task 7 C.3, 81→86 by Plan 3 Task 8 B.3. | -| [87) | (reserved for DQN v2) | | | | | Allocated incrementally by Plans 3-5 | +| [87] | `VSN_MAG_EMA_INDEX` | f32 | GPU `attention_focus_ema_update` kernel (Plan 4 E.5 Mode A) | `AttentionMonitor` (HEALTH_DIAG) | FoldReset | Magnitude-branch VSN-weight magnitude EMA. Host-computed scalar from `per_branch_vsn_mean()[1]` (mean abs of VSN tensors 28/29) passed by value at the per-epoch HEALTH_DIAG site. Adaptive α = α_base × (1 + 0.5 × \|clamp(sharpe, ±2)\|), α_base = 0.05. Cold-start 0.0 — diagnostic only; no consumer kernel reads this slot in Mode A. | +| [88] | `VSN_DIR_EMA_INDEX` | f32 | GPU `attention_focus_ema_update` kernel (Plan 4 E.5 Mode A) | `AttentionMonitor` (HEALTH_DIAG) | FoldReset | Direction-branch VSN-weight magnitude EMA. Same producer/contract as [87]; host scalar from `per_branch_vsn_mean()[0]`. | +| [89] | `MAMBA2_RETENTION_EMA_INDEX` | f32 | GPU `attention_focus_ema_update` kernel (Plan 4 E.5 Mode A) | `AttentionMonitor` (HEALTH_DIAG) | FoldReset | Mamba2 state-transition magnitude EMA (retention proxy). Host-computed per-batch mean \|`mamba2_h_enriched`\| via `mamba2_retention_mean(batch_size)` (cold-path dtoh, mirror of `per_branch_vsn_mean`). Same EMA convention as [87]. Diagnostic only. | +| [90] | `ISV_LAYOUT_FINGERPRINT_LO_INDEX` | u32 bits (in f32) | construct | check_layout_fingerprint | SchemaContract | Low 32 bits of u64 FNV-1a structural hash. Fail-fast on mismatch — NOT a version number, no migration path. Shifted 69→73 by Plan 3 Task 3 B.2, 73→76 by Plan 3 Task 4 B.4, 76→80 by Plan 3 Task 7 C.3, 80→85 by Plan 3 Task 8 B.3, 85→90 by Plan 4 Task 5 Mode A. | +| [91] | `ISV_LAYOUT_FINGERPRINT_HI_INDEX` | u32 bits (in f32) | construct | check_layout_fingerprint | SchemaContract | High 32 bits of u64 FNV-1a structural hash. Shifted 70→74 by Plan 3 Task 3 B.2, 74→77 by Plan 3 Task 4 B.4, 77→81 by Plan 3 Task 7 C.3, 81→86 by Plan 3 Task 8 B.3, 86→91 by Plan 4 Task 5 Mode A. | +| [92) | (reserved for DQN v2) | | | | | Allocated incrementally by Plans 4-5 |