feat(dqn-v2): E.5 Mode A attention-focus ISV — fully GPU-driven
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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)
|
||||
|
||||
93
crates/ml/src/cuda_pipeline/attention_focus_ema_kernel.cu
Normal file
93
crates/ml/src/cuda_pipeline/attention_focus_ema_kernel.cu
Normal file
@@ -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 <cuda_runtime.h>
|
||||
|
||||
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;
|
||||
}
|
||||
}
|
||||
@@ -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::<f32>();
|
||||
let dir_end = padded_byte_offset(¶m_sizes, 28) as usize / std::mem::size_of::<f32>();
|
||||
let mag_start = padded_byte_offset(¶m_sizes, 28) as usize / std::mem::size_of::<f32>();
|
||||
let mag_end = padded_byte_offset(¶m_sizes, 30) as usize / std::mem::size_of::<f32>();
|
||||
(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.
|
||||
|
||||
@@ -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`.
|
||||
|
||||
@@ -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
|
||||
|
||||
115
crates/ml/src/trainers/dqn/monitors/attention_monitor.rs
Normal file
115
crates/ml/src/trainers/dqn/monitors/attention_monitor.rs
Normal file
@@ -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<f32> {
|
||||
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}");
|
||||
}
|
||||
}
|
||||
@@ -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;
|
||||
|
||||
@@ -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 }
|
||||
}
|
||||
|
||||
@@ -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.
|
||||
|
||||
@@ -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.
|
||||
|
||||
@@ -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 |
|
||||
|
||||
Reference in New Issue
Block a user