fix(class-a-audit-batch-4b): plan_threshold floor adaptive + MIN_HOLD_TEMPERATURE ISV-driven
Per Class A audit-fix Batch 4-B (final 2 of 4 deferred items from
P1-wiring/P1-producer). Completes the 8-commit WR-plateau intervention
chain. Validation deferred to next L40S smoke.
Item 3: plan_threshold adaptive floor (Design Y - inline producer)
- NEW slot PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX=459
- Writes slow-EMA shadow of `0.5 * readiness_ema` from inside the
existing update kernel (no new file, no new launch).
- Pearl-A first-observation bootstrap (sentinel 0.1 matches pre-fix
hardcoded value for bit-identical cold-start) + Welford alpha=0.005
slow EMA.
- Bilateral clamp [0.05, 0.50] (probability units) per
pearl_symmetric_clamp_audit.
- Consumer reads isv[459] as the floor in the same launch's final
fmaxf; cold-start sentinel REPLACES with threshold_target so the
pre-fix `fmaxf(0.1, 0.5*ema)` semantic is preserved bit-identical
for any readiness EMA above 0.20.
Item 4: MIN_HOLD_TEMPERATURE -> ISV-driven (driving signal: dir_acc skill)
- NEW slot MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460
- NEW kernel min_hold_temperature_update_kernel.cu (single-thread
cold-path, per-epoch boundary launch).
- Driving signal: dir_acc skill = clamp((short_ema - 0.5)/0.5, 0, 1)
from ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]. When committing skillfully
(high dir_acc) -> temp HIGH (permissive). When at random baseline
(~0.5) -> temp LOW (sharp commitment pressure). Substituted for
the audit-spec's `dir_entropy_deficit` because no dir_entropy ISV
slot exists - dir_acc skill is the closest semantically-equivalent
signal that preserves the spec intent.
- Pearl-A bootstrap (sentinel 50.0 matches the deleted
MIN_HOLD_TEMPERATURE_START=50 anchor) + alpha=0.05 mid-cadence EMA.
- Bounds [5, 50] (matches the deleted schedule range).
- Decouples temperature from epoch number - the old schedule pinned
LOW temp (sharp) at end of training, exactly when a WR-plateaued
model needed forgiveness to escape.
- DELETED: state_layout.cuh::MIN_HOLD_TEMPERATURE_{START, END, DECAY}
#defines + training_loop.rs::min_hold_temperature_for_epoch helper
function (kept docstring tombstone explaining the deletion). Both
call sites migrated to the new ISV reader. Per
feedback_no_legacy_aliases + feedback_no_partial_refactor.
ISV_TOTAL_DIM: 459 -> 461.
Cumulative WR-plateau fix series (final commit, #8):
- 8f218cab2 (Class C bug 1 + P0-B)
- 316db416b (P0-C MIN_HOLD_TARGET)
- 394de7d43 (P0-A REWARD_POS/NEG_CAP)
- c4b6d6ef2 (P1 wiring var_floor)
- 657972a4b (P0-A downstream DD penalty + MIN_HOLD_PENALTY_MAX)
- 87d597d5d (P1 producer Bayesian priors)
- 7e9a8f6ef (Batch 4-A DD saturation floor + legacy DELETE)
- this commit (Batch 4-B plan_threshold floor + MIN_HOLD_TEMPERATURE)
Verification:
cargo check -p ml --tests --all-targets PASS
cargo test sp14 (lib) 16/16 PASS (slot layout locks)
cargo test sp15 (lib) 2/2 PASS (no regression)
cargo test state_reset_registry (lib) 4/4 PASS (dispatch coverage)
cargo test sp14_oracle_tests (GPU) 24/24 PASS (8 new + 16 pre)
cargo test sp15_phase1_oracle_tests (GPU) 36/36 PASS (no regression)
Per feedback_isv_for_adaptive_bounds + feedback_no_partial_refactor +
feedback_no_legacy_aliases.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -856,6 +856,30 @@ fn main() {
|
||||
// Block-tree-reduce (no atomicAdd) per `feedback_no_atomicadd.md`.
|
||||
// Per-epoch boundary launch.
|
||||
"dd_saturation_floor_update_kernel.cu",
|
||||
// Class A audit-fix Batch 4-B (2026-05-08, Item 4): adaptive
|
||||
// MIN_HOLD_TEMPERATURE producer. Single-thread cold-path kernel
|
||||
// that reads `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` (sp13 fast
|
||||
// EMA of binary directional accuracy), maps it to a [5, 50]
|
||||
// temperature via `temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) ×
|
||||
// skill` where `skill = clamp((short_ema − 0.5)/0.5, 0, 1)`,
|
||||
// and slow-EMA-blends into
|
||||
// `ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460]`. Pearl-A
|
||||
// first-observation bootstrap (sentinel 50.0 matches the
|
||||
// deleted MIN_HOLD_TEMPERATURE_START=50 anchor for bit-identical
|
||||
// cold-start under both regimes); α=0.05 mid-cadence EMA
|
||||
// thereafter. Replaces the deleted epoch-driven schedule
|
||||
// `T_end + (T_start − T_end) × exp(-epoch / decay)` (formerly
|
||||
// state_layout.cuh::MIN_HOLD_TEMPERATURE_{START, END, DECAY} +
|
||||
// training_loop.rs::min_hold_temperature_for_epoch). Decouples
|
||||
// temperature from epoch number — when WR plateaus, the old
|
||||
// schedule pinned LOW temp (sharp) at end of training (max
|
||||
// punishment when most needed forgiveness); the signal-driven
|
||||
// version reacts to actual realized skill. Per
|
||||
// `feedback_no_legacy_aliases.md` the legacy schedule is deleted
|
||||
// atomically. Bounds [5, 50] are Category-1 dimensional safety
|
||||
// floors per `feedback_isv_for_adaptive_bounds.md`. Per-epoch
|
||||
// boundary launch (same cadence as P0-A REWARD_POS_CAP).
|
||||
"min_hold_temperature_update_kernel.cu",
|
||||
// SP14 Layer C Phase C.6 (2026-05-08): h_s2_aux RMS EMA producer.
|
||||
// Single-block 256-thread kernel computing RMS = sqrt(mean(x²))
|
||||
// over `h_s2_aux [B, SH2]` (aux trunk final output, no activation)
|
||||
|
||||
@@ -53,7 +53,8 @@ use crate::MLError;
|
||||
use super::gpu_dqn_trainer::{
|
||||
AUX_HORIZON_UPDATE_CUBIN, AUX_TRUNK_BACKWARD_CUBIN, AUX_TRUNK_FORWARD_CUBIN,
|
||||
AVG_WIN_HOLD_TIME_UPDATE_CUBIN, DD_SATURATION_FLOOR_UPDATE_CUBIN, H_S2_AUX_RMS_EMA_CUBIN,
|
||||
KELLY_BAYESIAN_PRIORS_UPDATE_CUBIN, REWARD_CAP_UPDATE_CUBIN,
|
||||
KELLY_BAYESIAN_PRIORS_UPDATE_CUBIN, MIN_HOLD_TEMPERATURE_UPDATE_CUBIN,
|
||||
REWARD_CAP_UPDATE_CUBIN,
|
||||
};
|
||||
|
||||
/// Hidden width of the aux trunk's first internal layer (Linear_1 → ELU
|
||||
@@ -965,6 +966,104 @@ impl DdSaturationFloorUpdateOps {
|
||||
}
|
||||
}
|
||||
|
||||
/// Class A audit-fix Batch 4-B (2026-05-08, Item 4): adaptive
|
||||
/// MIN_HOLD_TEMPERATURE producer.
|
||||
///
|
||||
/// Single-thread cold-path kernel (per-epoch boundary cadence) that
|
||||
/// reads `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` (sp13 fast EMA of
|
||||
/// binary directional accuracy), maps it to a [5, 50] temperature via
|
||||
///
|
||||
/// skill = clamp((short_ema − 0.5) / 0.5, 0, 1)
|
||||
/// temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × skill
|
||||
///
|
||||
/// and slow-EMA-blends into
|
||||
/// `ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460]` with α=0.05.
|
||||
///
|
||||
/// Replaces the deleted epoch-driven anneal schedule
|
||||
/// `T_end + (T_start − T_end) × exp(-epoch / decay)` (formerly
|
||||
/// state_layout.cuh::MIN_HOLD_TEMPERATURE_{START, END, DECAY} +
|
||||
/// training_loop.rs::min_hold_temperature_for_epoch). Decouples
|
||||
/// temperature from epoch number — when WR plateaus, the old
|
||||
/// schedule pinned LOW temp (sharp) at end of training (max
|
||||
/// punishment exactly when most needed forgiveness); the signal-driven
|
||||
/// version reacts to actual realized skill.
|
||||
///
|
||||
/// # Pearls applied
|
||||
/// - `feedback_no_atomicadd.md` — single-thread kernel; no reductions.
|
||||
/// - `pearl_first_observation_bootstrap.md` — sentinel 50.0 → REPLACE.
|
||||
/// - `pearl_no_host_branches_in_captured_graph.md` — `CudaFunction`
|
||||
/// pre-loaded at construction; on-device guards.
|
||||
/// - `pearl_symmetric_clamp_audit.md` — bilateral
|
||||
/// `fmaxf(lo, fminf(x, hi))` clamp on target_temp + post-blend.
|
||||
/// - `feedback_isv_for_adaptive_bounds.md` — bounds [5, 50] are
|
||||
/// Category-1 dimensional safety floors (matches the original
|
||||
/// schedule range MIN_HOLD_TEMPERATURE_END=5 ↔
|
||||
/// MIN_HOLD_TEMPERATURE_START=50), NOT tuning.
|
||||
/// - `feedback_no_legacy_aliases.md` — the legacy schedule constants
|
||||
/// and helper function are deleted atomically with this wiring.
|
||||
#[allow(missing_debug_implementations)]
|
||||
pub(crate) struct MinHoldTemperatureUpdateOps {
|
||||
update_kernel: CudaFunction,
|
||||
}
|
||||
|
||||
impl MinHoldTemperatureUpdateOps {
|
||||
pub(crate) fn new(stream: &Arc<CudaStream>) -> Result<Self, MLError> {
|
||||
let context = stream.context();
|
||||
let module = context
|
||||
.load_cubin(MIN_HOLD_TEMPERATURE_UPDATE_CUBIN.to_vec())
|
||||
.map_err(|e| MLError::ModelError(format!("min_hold_temperature_update cubin load: {e}")))?;
|
||||
let update_kernel = module
|
||||
.load_function("min_hold_temperature_update")
|
||||
.map_err(|e| MLError::ModelError(format!("min_hold_temperature_update load: {e}")))?;
|
||||
Ok(Self { update_kernel })
|
||||
}
|
||||
|
||||
/// Launch the adaptive MIN_HOLD_TEMPERATURE producer.
|
||||
///
|
||||
/// Args:
|
||||
/// - `isv_ptr`: ISV[ISV_TOTAL_DIM] device pointer (mapped-pinned).
|
||||
/// - `temp_idx`: MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX (460).
|
||||
/// - `dir_acc_idx`: AUX_DIR_ACC_SHORT_EMA_INDEX (373).
|
||||
/// - `sentinel_temp`: SENTINEL_MIN_HOLD_TEMPERATURE (50.0).
|
||||
/// - `sentinel_dir_acc`: DIR_ACC_RANDOM_BASELINE (0.5; matches
|
||||
/// sp13 DIR_ACC_EMA_SENTINEL).
|
||||
/// - `temp_min`/`temp_max`: MIN_HOLD_TEMPERATURE_MIN/MAX (5.0, 50.0).
|
||||
/// - `alpha`: MIN_HOLD_TEMPERATURE_EMA_ALPHA (0.05).
|
||||
#[allow(clippy::too_many_arguments)]
|
||||
pub(crate) fn launch(
|
||||
&self,
|
||||
stream: &Arc<CudaStream>,
|
||||
isv_ptr: u64,
|
||||
temp_idx: i32,
|
||||
dir_acc_idx: i32,
|
||||
sentinel_temp: f32,
|
||||
sentinel_dir_acc: f32,
|
||||
temp_min: f32,
|
||||
temp_max: f32,
|
||||
alpha: f32,
|
||||
) -> Result<(), MLError> {
|
||||
unsafe {
|
||||
stream
|
||||
.launch_builder(&self.update_kernel)
|
||||
.arg(&isv_ptr)
|
||||
.arg(&temp_idx)
|
||||
.arg(&dir_acc_idx)
|
||||
.arg(&sentinel_temp)
|
||||
.arg(&sentinel_dir_acc)
|
||||
.arg(&temp_min)
|
||||
.arg(&temp_max)
|
||||
.arg(&alpha)
|
||||
.launch(LaunchConfig {
|
||||
grid_dim: (1, 1, 1),
|
||||
block_dim: (1, 1, 1),
|
||||
shared_mem_bytes: 0,
|
||||
})
|
||||
.map_err(|e| MLError::ModelError(format!("min_hold_temperature_update: {e}")))?;
|
||||
}
|
||||
Ok(())
|
||||
}
|
||||
}
|
||||
|
||||
/// SP14 Layer C Phase C.6 (2026-05-08): h_s2_aux RMS EMA producer.
|
||||
///
|
||||
/// Computes `RMS(h_s2_aux) = sqrt(mean(h_s2_aux²))` over the
|
||||
|
||||
File diff suppressed because one or more lines are too long
@@ -315,12 +315,23 @@ pub struct ExperienceCollectorConfig {
|
||||
/// (60% of `REWARD_POS_CAP`).
|
||||
pub min_hold_penalty_max: f32,
|
||||
/// SP12 v3: temperature controller for the min-hold soft saturation.
|
||||
/// Computed per-epoch in `training_loop.rs::min_hold_temperature_for_epoch`
|
||||
/// via the annealing schedule
|
||||
/// `T_end + (T_start - T_end) × exp(-epoch / decay)` with constants
|
||||
/// from `state_layout.cuh::MIN_HOLD_TEMPERATURE_{START,END,DECAY}`.
|
||||
/// Default at construction = `MIN_HOLD_TEMPERATURE_START=50.0`
|
||||
/// (epoch-0 forgiving regime; replaced per epoch by training loop).
|
||||
///
|
||||
/// Class A audit-fix Batch 4-B (2026-05-08, Item 4): MIGRATED from
|
||||
/// the deleted epoch-driven schedule
|
||||
/// `T_end + (T_start - T_end) × exp(-epoch / decay)` (formerly
|
||||
/// `training_loop.rs::min_hold_temperature_for_epoch`) to the
|
||||
/// signal-driven ISV slot
|
||||
/// `MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460`. The training loop
|
||||
/// reads the freshly-produced value via
|
||||
/// `DQNTrainer::read_min_hold_temperature_from_isv` and seeds this
|
||||
/// field per-epoch. Producer kernel:
|
||||
/// `min_hold_temperature_update_kernel.cu` (per-epoch boundary).
|
||||
///
|
||||
/// Default at construction = 50.0 (= MIN_HOLD_TEMPERATURE_FALLBACK
|
||||
/// = the deleted MIN_HOLD_TEMPERATURE_START anchor for bit-identical
|
||||
/// pre-Batch-4B epoch-0 behavior). The kernel-passed scalar is the
|
||||
/// cold-start fallback used by the consumer when the ISV slot is
|
||||
/// at sentinel.
|
||||
pub min_hold_temperature: f32,
|
||||
}
|
||||
|
||||
@@ -397,9 +408,13 @@ impl Default for ExperienceCollectorConfig {
|
||||
hindsight_lookahead: 10,
|
||||
noise_sigma_per_branch: [0.1; 4],
|
||||
// SP12 v3 defaults — match Invariant-1 anchors in state_layout.cuh.
|
||||
// The training loop overrides `min_hold_temperature` per epoch via
|
||||
// the annealing schedule (50 → 5 across ~100 epochs); the other
|
||||
// two are static Phase 1 constants.
|
||||
// Class A audit-fix Batch 4-B (2026-05-08, Item 4): the training
|
||||
// loop overrides `min_hold_temperature` per epoch via
|
||||
// `DQNTrainer::read_min_hold_temperature_from_isv` (signal-driven
|
||||
// ISV slot 460 — replaces the deleted epoch-driven schedule);
|
||||
// the default 50.0 here is the cold-start fallback anchor matching
|
||||
// the deleted MIN_HOLD_TEMPERATURE_START. The other two are
|
||||
// static Phase 1 constants.
|
||||
min_hold_target: 30.0,
|
||||
min_hold_penalty_max: 3.0,
|
||||
min_hold_temperature: 50.0,
|
||||
@@ -5665,11 +5680,15 @@ impl GpuExperienceCollector {
|
||||
// ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451]; `min_hold_target`
|
||||
// here is the cold-start fallback only (used when ISV[451]
|
||||
// is still at sentinel 0.0 or isv_signals_ptr is NULL).
|
||||
// Constants in state_layout.cuh (MIN_HOLD_TARGET,
|
||||
// MIN_HOLD_PENALTY_MAX, MIN_HOLD_TEMPERATURE_{START,END,
|
||||
// DECAY}). Temperature recomputed per epoch via the
|
||||
// annealing schedule in `training_loop.rs::
|
||||
// min_hold_temperature_for_epoch`.
|
||||
// Class A audit-fix Batch 4-B (2026-05-08, Item 4):
|
||||
// `min_hold_temperature` is now signal-driven via ISV
|
||||
// slot 460 (sourced from
|
||||
// ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373] via
|
||||
// `min_hold_temperature_update_kernel`); the
|
||||
// training loop seeds this scalar per-epoch via
|
||||
// `DQNTrainer::read_min_hold_temperature_from_isv`,
|
||||
// which falls back to MIN_HOLD_TEMPERATURE_FALLBACK
|
||||
// =50.0 when the slot is at sentinel.
|
||||
.arg(&config.min_hold_target)
|
||||
.arg(&config.min_hold_penalty_max)
|
||||
.arg(&config.min_hold_temperature)
|
||||
|
||||
@@ -0,0 +1,158 @@
|
||||
/* ══════════════════════════════════════════════════════════════════════════
|
||||
* Class A audit-fix Batch 4-B — adaptive MIN_HOLD_TEMPERATURE producer
|
||||
* (2026-05-08, Item 4).
|
||||
*
|
||||
* Replaces the deleted epoch-driven anneal schedule
|
||||
* `T_end + (T_start - T_end) × exp(-epoch / decay)`
|
||||
* (formerly state_layout.cuh::MIN_HOLD_TEMPERATURE_{START=50, END=5,
|
||||
* DECAY=20} + training_loop.rs::min_hold_temperature_for_epoch) with an
|
||||
* ISV-driven signal-driven temperature derived from the model's
|
||||
* directional-accuracy skill. The temperature controls the
|
||||
* `deficit / (deficit + T)` saturation factor in `compute_min_hold_penalty`
|
||||
* (trade_physics.cuh:575): HIGH T = forgiving (gradient is gentle near
|
||||
* deficit=0); LOW T = sharp (penalty rises steeply at any deficit).
|
||||
*
|
||||
* Why this matters (per Class A audit deferred-item batch 4-B):
|
||||
*
|
||||
* The epoch-driven schedule pinned LOW temp (T=5, sharp) at end of
|
||||
* training. This works ONLY if the model gets more skilled with
|
||||
* training — which is exactly the assumption that fails when WR
|
||||
* plateaus. Maximum punishment is delivered exactly when the model
|
||||
* most needs forgiveness to escape the plateau. The signal-driven
|
||||
* replacement decouples temperature from epoch number: when the
|
||||
* model is committing skillfully (dir_acc > 0.5) → temp HIGH
|
||||
* (permissive); when at random baseline (dir_acc ≈ 0.5) → temp LOW
|
||||
* (sharp commitment pressure to escape the random-guess regime).
|
||||
*
|
||||
* Driving signal: dir_acc skill from
|
||||
* `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` — the SP13 fast EMA of
|
||||
* binary directional accuracy (sentinel 0.5 = random baseline).
|
||||
*
|
||||
* skill = clamp((short_ema - 0.5) / 0.5, 0, 1) ∈ [0, 1]
|
||||
* temp = TEMP_MIN + (TEMP_MAX - TEMP_MIN) × skill
|
||||
*
|
||||
* - skill=0 (no directional skill, dir_acc ≤ 0.5) → temp=TEMP_MIN=5
|
||||
* - skill=1 (saturated dir_acc=1.0) → temp=TEMP_MAX=50
|
||||
*
|
||||
* Cold-start fallback: if `isv[AUX_DIR_ACC_SHORT_EMA_INDEX]` is at
|
||||
* sentinel (within EPS of 0.5) — i.e. fold reset just fired and no
|
||||
* dir-acc samples have landed yet — fall back to the seeded
|
||||
* SENTINEL_MIN_HOLD_TEMPERATURE (50.0, matches the pre-fix
|
||||
* MIN_HOLD_TEMPERATURE_START=50 epoch-0 anchor). Bit-identical
|
||||
* cold-start under the deleted schedule, which also had T=50 at
|
||||
* epoch 0.
|
||||
*
|
||||
* Algorithm (single-block single-thread — cold path, per-epoch boundary):
|
||||
*
|
||||
* Phase 1 (single thread): read short_ema and current temp from ISV.
|
||||
*
|
||||
* Phase 2 (Pearl-A + EMA): if dir_acc EMA is at sentinel 0.5 (within
|
||||
* EPS), keep ISV slot unchanged — sentinel persists; consumer falls
|
||||
* back to the kernel-passed scalar (seeded at
|
||||
* MIN_HOLD_TEMPERATURE_FALLBACK=50.0). Otherwise:
|
||||
*
|
||||
* skill = max(0, min(1, (short_ema - 0.5) / 0.5))
|
||||
* target_temp = TEMP_MIN + (TEMP_MAX - TEMP_MIN) × skill
|
||||
* bilateral clamp to [TEMP_MIN, TEMP_MAX]
|
||||
*
|
||||
* Pearl-A bootstrap: if current is at sentinel (within EPS of 50.0),
|
||||
* REPLACE directly with target_temp. Otherwise EMA blend with α=0.05
|
||||
* (mid-cadence — slower than per-step skill EMA, faster than per-fold
|
||||
* beliefs).
|
||||
*
|
||||
* Pearls + invariants:
|
||||
*
|
||||
* - `feedback_no_atomicadd.md` — single-thread kernel; no reductions.
|
||||
*
|
||||
* - `pearl_no_host_branches_in_captured_graph.md` — `CudaFunction`
|
||||
* pre-loaded at construction; on-device guards for sentinel
|
||||
* detection.
|
||||
*
|
||||
* - `pearl_first_observation_bootstrap.md` — first valid skill
|
||||
* observation replaces sentinel directly; no blend.
|
||||
*
|
||||
* - `feedback_no_stubs.md` — full body, no return-zero placeholders.
|
||||
*
|
||||
* - `feedback_isv_for_adaptive_bounds.md` — bounds [5, 50] are
|
||||
* dimensional safety floors (matches the original schedule range),
|
||||
* NOT tuning.
|
||||
*
|
||||
* - `pearl_symmetric_clamp_audit.md` — bilateral
|
||||
* `fmaxf(lo, fminf(x, hi))` clamp on target_temp before writing.
|
||||
*
|
||||
* Args:
|
||||
* isv — `[ISV_TOTAL_DIM]` device f32, modified
|
||||
* in place at index `temp_idx`.
|
||||
* temp_idx — MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX (460).
|
||||
* dir_acc_idx — AUX_DIR_ACC_SHORT_EMA_INDEX (373).
|
||||
* sentinel_temp — SENTINEL_MIN_HOLD_TEMPERATURE (50.0).
|
||||
* sentinel_dir_acc — DIR_ACC_RANDOM_BASELINE (0.5; sp13
|
||||
* DIR_ACC_EMA_SENTINEL).
|
||||
* temp_min — MIN_HOLD_TEMPERATURE_MIN (5.0).
|
||||
* temp_max — MIN_HOLD_TEMPERATURE_MAX (50.0).
|
||||
* alpha — MIN_HOLD_TEMPERATURE_EMA_ALPHA (0.05).
|
||||
*
|
||||
* Launch: grid=(1, 1, 1), block=(1, 1, 1).
|
||||
* No shared memory — single-thread cold-path kernel.
|
||||
* ══════════════════════════════════════════════════════════════════════════ */
|
||||
|
||||
#include <cuda_runtime.h>
|
||||
|
||||
#define EPS_F 1e-6f
|
||||
|
||||
extern "C" __global__
|
||||
void min_hold_temperature_update(
|
||||
float* __restrict__ isv,
|
||||
int temp_idx,
|
||||
int dir_acc_idx,
|
||||
float sentinel_temp,
|
||||
float sentinel_dir_acc,
|
||||
float temp_min,
|
||||
float temp_max,
|
||||
float alpha)
|
||||
{
|
||||
/* Single-thread kernel — only thread 0 of block 0 does anything. */
|
||||
if (threadIdx.x != 0 || blockIdx.x != 0) return;
|
||||
|
||||
const float dir_acc_short = isv[dir_acc_idx];
|
||||
|
||||
/* Cold-start fallback: dir_acc EMA is at random baseline (sentinel)
|
||||
* — keep the temp slot unchanged. The consumer will fall back to
|
||||
* the kernel-passed scalar (MIN_HOLD_TEMPERATURE_FALLBACK=50.0)
|
||||
* via the sentinel-detect path. Bit-identical to pre-Batch-4B
|
||||
* epoch-0 behavior under the deleted schedule. */
|
||||
if (fabsf(dir_acc_short - sentinel_dir_acc) < EPS_F) {
|
||||
return;
|
||||
}
|
||||
|
||||
/* Skill mapping: clamp((short_ema - 0.5) / 0.5, 0, 1).
|
||||
* - dir_acc=0.5 (random) → skill=0 → temp=temp_min=5 (sharp)
|
||||
* - dir_acc=1.0 (saturated) → skill=1 → temp=temp_max=50 (forgiving)
|
||||
* - dir_acc<0.5 (worse than random) → skill clamped to 0 → temp=5 */
|
||||
float skill = (dir_acc_short - sentinel_dir_acc) / sentinel_dir_acc;
|
||||
skill = fmaxf(0.0f, fminf(1.0f, skill));
|
||||
|
||||
float target_temp = temp_min + (temp_max - temp_min) * skill;
|
||||
|
||||
/* Bilateral clamp per `pearl_symmetric_clamp_audit.md`. Defends
|
||||
* against arithmetic round-off pushing target outside the
|
||||
* dimensional-safety window. */
|
||||
target_temp = fmaxf(temp_min, fminf(target_temp, temp_max));
|
||||
|
||||
/* Pearl-A first-observation bootstrap. The sentinel is
|
||||
* SENTINEL_MIN_HOLD_TEMPERATURE=50.0 (matches the deleted
|
||||
* MIN_HOLD_TEMPERATURE_START); any value within EPS is treated
|
||||
* as sentinel and REPLACED directly. Otherwise EMA blend. */
|
||||
const float current = isv[temp_idx];
|
||||
float blended;
|
||||
if (fabsf(current - sentinel_temp) < EPS_F) {
|
||||
blended = target_temp;
|
||||
} else {
|
||||
blended = (1.0f - alpha) * current + alpha * target_temp;
|
||||
}
|
||||
/* Re-clamp post-blend (defensive — handles malformed prior state
|
||||
* outside [temp_min, temp_max]). */
|
||||
blended = fmaxf(temp_min, fminf(blended, temp_max));
|
||||
|
||||
isv[temp_idx] = blended;
|
||||
}
|
||||
@@ -3,19 +3,47 @@
|
||||
* Plan 3 Task 4 B.4. Reads `readiness_per_sample [n_samples]` (readiness
|
||||
* values in [0,1] written per-(i,t) by experience_env_step), reduces the
|
||||
* mean on-GPU, EMAs it into ISV[READINESS_EMA_INDEX=75], then derives
|
||||
* `threshold = fmaxf(0.1f, 0.5f * ema)` and writes it into
|
||||
* `threshold = fmaxf(adaptive_floor, 0.5f * ema)` and writes it into
|
||||
* ISV[PLAN_THRESHOLD_INDEX=49]. Consumers (experience_env_step and
|
||||
* backtest_plan_kernel) continue reading ISV[49] unchanged.
|
||||
*
|
||||
* Adaptive α: α = α_base × (1 + 0.5 × |clamp(sharpe, -2, 2)|). Matches
|
||||
* Task 3 trade_rate_ema_update convention. α_base=0.05.
|
||||
*
|
||||
* Class A audit-fix Batch 4-B (2026-05-08) — Item 3, Design Y (inline producer).
|
||||
* The kernel ALSO maintains a slow-EMA shadow of `0.5f * ema` in
|
||||
* ISV[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX=459] and uses that adaptive
|
||||
* floor as the lower bound on the derived threshold (instead of the
|
||||
* pre-fix hardcoded `0.1f` literal). The floor is the running long-term
|
||||
* trend of the live derived threshold itself — when readiness collapses
|
||||
* to near-zero (early-fold cold-start or low-readiness regime) the
|
||||
* floor decays slowly with the realized signal instead of pinning at
|
||||
* 0.1 in probability space. Pearl-A first-observation bootstrap on
|
||||
* sentinel 0.1 (matches the pre-fix hardcoded value for bit-identical
|
||||
* cold-start). Bilateral clamp [0.05, 0.50] per
|
||||
* `pearl_symmetric_clamp_audit.md`. α_floor=0.005 — slow, per-fold
|
||||
* cadence (the floor is a long-term shadow of the threshold and must
|
||||
* move on a much slower timescale than the readiness EMA itself).
|
||||
*
|
||||
* Pearls applied:
|
||||
* - `pearl_first_observation_bootstrap.md` — sentinel 0.1 → REPLACE
|
||||
* with `0.5 × ema` directly.
|
||||
* - `pearl_symmetric_clamp_audit.md` — bilateral
|
||||
* `fmaxf(lo, fminf(x, hi))` clamp on the floor before writing.
|
||||
* - `feedback_isv_for_adaptive_bounds.md` — bounds [0.05, 0.50] are
|
||||
* dimensional safety floors (probability units), NOT tuning.
|
||||
* - `feedback_no_partial_refactor.md` — the consumer (this same
|
||||
* kernel's final fmaxf) reads the freshly-written floor in the
|
||||
* same launch via `isv_out[floor_idx]`, so the migration is
|
||||
* atomic.
|
||||
*
|
||||
* Single-block single-thread. No atomicAdd, no DtoH, no PER-sample reduce
|
||||
* tree — acceptable because this kernel runs ONCE per experience-collection
|
||||
* epoch (cold path).
|
||||
*/
|
||||
|
||||
#include <cuda_runtime.h>
|
||||
#include "state_layout.cuh"
|
||||
|
||||
extern "C" __global__ void plan_threshold_update(
|
||||
const float* __restrict__ readiness_per_sample, /* [n_samples] ∈ [0,1] */
|
||||
@@ -39,7 +67,39 @@ extern "C" __global__ void plan_threshold_update(
|
||||
const float prev = isv_out[readiness_ema_idx];
|
||||
const float ema = (1.0f - alpha) * prev + alpha * mean;
|
||||
isv_out[readiness_ema_idx] = ema;
|
||||
/* Derived threshold: 50% of the running readiness mean. Floor 0.1 for
|
||||
* cold-start (when ema ≈ 0, consumer sites still have a sensible gate). */
|
||||
isv_out[plan_threshold_idx] = fmaxf(0.1f, 0.5f * ema);
|
||||
|
||||
/* Derived live threshold target: 50% of the running readiness mean. */
|
||||
const float threshold_target = 0.5f * ema;
|
||||
|
||||
/* Class A audit-fix Batch 4-B (Item 3, Design Y) — adaptive floor.
|
||||
* Pearl-A first-observation bootstrap from sentinel
|
||||
* SENTINEL_PLAN_THRESHOLD_FLOOR (0.1, matches pre-fix hardcoded
|
||||
* value); slow EMA blend thereafter. Bilateral clamp [MIN, MAX]
|
||||
* (probability units). */
|
||||
const float prev_floor = isv_out[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX];
|
||||
float new_floor;
|
||||
/* Pearl-A: any value within EPS of the sentinel (or non-positive)
|
||||
* is treated as sentinel — REPLACE directly with `threshold_target`
|
||||
* so the floor seeds at the live derived threshold on the first
|
||||
* launch (matches the pre-fix `fmaxf(0.1, target)` behavior when
|
||||
* target ≥ 0.1). */
|
||||
if (prev_floor <= 1e-6f
|
||||
|| fabsf(prev_floor - SENTINEL_PLAN_THRESHOLD_FLOOR) < 1e-6f) {
|
||||
new_floor = threshold_target;
|
||||
} else {
|
||||
/* Slow EMA: floor tracks the long-term `0.5 × readiness_ema` curve. */
|
||||
new_floor = (1.0f - PLAN_THRESHOLD_FLOOR_EMA_ALPHA) * prev_floor
|
||||
+ PLAN_THRESHOLD_FLOOR_EMA_ALPHA * threshold_target;
|
||||
}
|
||||
/* Bilateral clamp per `pearl_symmetric_clamp_audit.md` — keeps the
|
||||
* floor in probability units, defends against degenerate prior
|
||||
* state (e.g. malformed checkpoint). */
|
||||
new_floor = fmaxf(PLAN_THRESHOLD_FLOOR_MIN_BOUND,
|
||||
fminf(new_floor, PLAN_THRESHOLD_FLOOR_MAX_BOUND));
|
||||
isv_out[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX] = new_floor;
|
||||
|
||||
/* Derived threshold: max(adaptive_floor, 0.5 × ema). The floor
|
||||
* replaces the pre-fix hardcoded 0.1f literal but the formula
|
||||
* structure is identical. */
|
||||
isv_out[plan_threshold_idx] = fmaxf(new_floor, threshold_target);
|
||||
}
|
||||
|
||||
@@ -283,6 +283,133 @@ pub const DD_SATURATION_FLOOR_EMA_ALPHA: f32 = 0.01;
|
||||
pub const SP14_AUDIT_4A_SLOT_BASE: usize = 458;
|
||||
pub const SP14_AUDIT_4A_SLOT_END: usize = 459;
|
||||
|
||||
// ── Class A audit-fix Batch 4-B (2026-05-08): adaptive plan_threshold floor ──
|
||||
// Replaces the hardcoded `0.1f` floor in
|
||||
// `plan_threshold_update_kernel.cu:44` (`final = fmaxf(0.1f, 0.5f * ema)`)
|
||||
// with an ISV-driven adaptive bound that tracks a slow EMA of `0.5f * ema`
|
||||
// — the very same expression the kernel uses as its derived threshold.
|
||||
// The floor is therefore a long-term smoothed shadow of the live plan-
|
||||
// threshold value, in probability units (matches the readiness EMA's
|
||||
// units). Producer is INLINE inside the same kernel (Design Y — no new
|
||||
// file, no new launch). On each invocation the kernel writes both the
|
||||
// new plan_threshold (slot 49) and a slow-EMA floor (slot 459) and uses
|
||||
// the floor itself as the lower bound on the next invocation.
|
||||
//
|
||||
// Why ISV-driven: the static `0.1f` was a calibration for a specific
|
||||
// readiness distribution. When the policy's readiness EMA collapses to
|
||||
// near-zero (e.g. early-fold cold-start or a regime where readiness is
|
||||
// uniformly low), the static floor pins plan_threshold at 0.1 in
|
||||
// probability space — uncoupled from the realized readiness signal.
|
||||
// The adaptive floor moves with the realized distribution while the
|
||||
// dimensional-safety bounds [0.05, 0.50] still guard against runaway.
|
||||
//
|
||||
// Distinct from PLAN_THRESHOLD_INDEX=49 (the live derived threshold,
|
||||
// fast cadence): slot 459 is the slow-moving lower bound on slot 49.
|
||||
//
|
||||
// Pearl-A first-observation bootstrap: SENTINEL_PLAN_THRESHOLD_FLOOR=0.1
|
||||
// matches the pre-fix hardcoded value for bit-identical cold-start
|
||||
// behavior before the first valid observation lands. Cold-start
|
||||
// fallback at the consumer: when the slot is at sentinel (≤ 0 or
|
||||
// within EPS of 0.1) OR out of [MIN, MAX], fall back to literal 0.1.
|
||||
pub const PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX: usize = 459;
|
||||
|
||||
// Sentinel — matches the pre-fix hardcoded `0.1f` floor in
|
||||
// `plan_threshold_update_kernel.cu` for bit-identical cold-start.
|
||||
// Pearl-A bootstrap: first valid observation replaces directly.
|
||||
pub const SENTINEL_PLAN_THRESHOLD_FLOOR: f32 = 0.1;
|
||||
|
||||
// Bounds — Category-1 dimensional safety floors per
|
||||
// `feedback_isv_for_adaptive_bounds.md`. Floor in [0.05, 0.50]: below
|
||||
// 5% the readiness gate is effectively absent (every plan is ready
|
||||
// per the controller's perspective); above 50% the gate is uselessly
|
||||
// strict (readiness EMA must climb past 1.0 to ever pass — impossible
|
||||
// given readiness ∈ [0, 1]).
|
||||
pub const PLAN_THRESHOLD_FLOOR_MIN: f32 = 0.05;
|
||||
pub const PLAN_THRESHOLD_FLOOR_MAX: f32 = 0.50;
|
||||
|
||||
// EMA blend rate. Slow — the floor is a shadow of the long-term
|
||||
// `0.5 × readiness_ema` curve and must move on a timescale much slower
|
||||
// than the readiness EMA itself (which adapts at α_base=0.05 ×
|
||||
// (1 + 0.5×|sharpe|) ≈ [0.05, 0.10]). α=0.005 keeps the floor on the
|
||||
// per-fold cadence (mirrors the P1-Producer Kelly-prior α=0.005).
|
||||
pub const PLAN_THRESHOLD_FLOOR_EMA_ALPHA: f32 = 0.005;
|
||||
|
||||
pub const SP14_AUDIT_4B_PLAN_FLOOR_SLOT_BASE: usize = 459;
|
||||
pub const SP14_AUDIT_4B_PLAN_FLOOR_SLOT_END: usize = 460;
|
||||
|
||||
// ── Class A audit-fix Batch 4-B (2026-05-08): MIN_HOLD_TEMPERATURE adaptive ──
|
||||
// Replaces the epoch-driven anneal schedule
|
||||
// `T_end + (T_start - T_end) × exp(-epoch / decay)` (state_layout.cuh:
|
||||
// MIN_HOLD_TEMPERATURE_{START=50, END=5, DECAY=20}) with a signal-driven
|
||||
// temperature derived from the model's directional accuracy skill. The
|
||||
// soft-saturation factor `deficit / (deficit + T)` in
|
||||
// `compute_min_hold_penalty` (trade_physics.cuh:575) controls how
|
||||
// sharply the min-hold penalty bites: HIGH T = forgiving (gradient is
|
||||
// gentle near deficit=0); LOW T = sharp (penalty rises steeply at any
|
||||
// deficit).
|
||||
//
|
||||
// Driving signal: dir_acc skill — `clamp((short_ema - 0.5) / 0.5, 0, 1)`
|
||||
// from ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373] (sp13 directional-accuracy
|
||||
// fast EMA, sentinel 0.5 = random-baseline). When the model is
|
||||
// committing skillfully (dir_acc > 0.5) → temp HIGH (permissive — the
|
||||
// model's exits are informed, no need to penalize sharply). When the
|
||||
// model is at random baseline (dir_acc ≈ 0.5) → temp LOW (sharp — the
|
||||
// model needs commitment pressure to escape the random-guess regime).
|
||||
//
|
||||
// Mapping: temp = TEMP_MIN + (TEMP_MAX - TEMP_MIN) × skill, where:
|
||||
// - skill=0 (no directional skill, dir_acc ≤ 0.5) → temp=TEMP_MIN=5
|
||||
// - skill=1 (saturated dir_acc=1.0) → temp=TEMP_MAX=50
|
||||
//
|
||||
// Opposite of the deleted epoch-driven schedule which pinned LOW temp
|
||||
// at end of training (T_end=5). The epoch schedule assumes the model
|
||||
// gets MORE skilled with training; when WR plateaus this assumption
|
||||
// fails and the schedule becomes maximally punishing exactly when the
|
||||
// model needs forgiveness to escape the plateau. The signal-driven
|
||||
// version reacts to actual realized skill, decoupling temperature from
|
||||
// epoch number.
|
||||
//
|
||||
// Producer kernel: `min_hold_temperature_update_kernel.cu` (per-epoch
|
||||
// boundary cadence). Consumer: `experience_kernels.cu` reads ISV[460]
|
||||
// directly with cold-start fallback to the kernel-passed scalar (still
|
||||
// MIN_HOLD_TEMPERATURE_START=50 from gpu_experience_collector.rs
|
||||
// default, mirrors the effective_min_hold_target pattern at line 3197).
|
||||
//
|
||||
// Pearl-A first-observation bootstrap: SENTINEL_MIN_HOLD_TEMPERATURE=
|
||||
// 50.0 matches the pre-fix MIN_HOLD_TEMPERATURE_START so cold-start
|
||||
// behavior is bit-identical (first epoch had T=50 in the old schedule
|
||||
// too).
|
||||
pub const MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX: usize = 460;
|
||||
|
||||
// Sentinel — matches the pre-fix MIN_HOLD_TEMPERATURE_START=50.0 for
|
||||
// bit-identical cold-start (first epoch under both regimes had T=50).
|
||||
// Pearl-A bootstrap: first valid observation replaces directly.
|
||||
pub const SENTINEL_MIN_HOLD_TEMPERATURE: f32 = 50.0;
|
||||
|
||||
// Bounds — Category-1 dimensional safety floors per
|
||||
// `feedback_isv_for_adaptive_bounds.md`. [5, 50] matches the original
|
||||
// schedule range (T_END=5, T_START=50) — below 5 the saturation is
|
||||
// effectively a hard cliff (factor jumps to ≈1 at any deficit); above
|
||||
// 50 the factor stays < 0.5 even at MIN_HOLD_TARGET-1 deficit which
|
||||
// nearly disables the penalty.
|
||||
pub const MIN_HOLD_TEMPERATURE_MIN: f32 = 5.0;
|
||||
pub const MIN_HOLD_TEMPERATURE_MAX: f32 = 50.0;
|
||||
|
||||
// EMA blend rate. Moderate — temperature should track the directional
|
||||
// skill EMA (which itself updates per-step) but not perfectly mirror it
|
||||
// (avoid jitter from per-batch dir_acc noise). α=0.05 matches the
|
||||
// P0-A REWARD_POS_CAP cadence (mid-cadence — slower than per-step
|
||||
// signals, faster than per-fold beliefs).
|
||||
pub const MIN_HOLD_TEMPERATURE_EMA_ALPHA: f32 = 0.05;
|
||||
|
||||
// Random-baseline anchor for the dir_acc skill mapping. K=2 binary
|
||||
// directional accuracy has a 0.5 random baseline; skill = max(0,
|
||||
// (short_ema - 0.5) / 0.5) maps that baseline to 0 and saturated 1.0
|
||||
// to 1. Below 0.5 (worse than random) clamps to 0.
|
||||
pub const DIR_ACC_RANDOM_BASELINE: f32 = 0.5;
|
||||
|
||||
pub const SP14_AUDIT_4B_TEMP_SLOT_BASE: usize = 460;
|
||||
pub const SP14_AUDIT_4B_TEMP_SLOT_END: usize = 461;
|
||||
|
||||
#[cfg(test)]
|
||||
mod tests {
|
||||
use super::*;
|
||||
@@ -462,4 +589,78 @@ mod tests {
|
||||
SP14_AUDIT_4A_SLOT_END, ISV_TOTAL_DIM,
|
||||
);
|
||||
}
|
||||
|
||||
/// Lock Class A audit-fix Batch 4-B (2026-05-08) adaptive plan_threshold
|
||||
/// floor slot. Single contiguous slot [459..460) replacing the
|
||||
/// hardcoded `0.1f` floor in `plan_threshold_update_kernel.cu:44`.
|
||||
/// Sentinel matches the pre-fix hardcoded value (0.1) for
|
||||
/// bit-identical cold-start.
|
||||
#[test]
|
||||
fn sp14_audit_4b_plan_threshold_floor_slot_layout_locked() {
|
||||
assert_eq!(SP14_AUDIT_4B_PLAN_FLOOR_SLOT_BASE, 459);
|
||||
assert_eq!(SP14_AUDIT_4B_PLAN_FLOOR_SLOT_END, 460);
|
||||
assert_eq!(PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX, 459);
|
||||
// Sentinel matches pre-fix hardcoded value (0.1) for bit-identical
|
||||
// cold-start. Producer kernel writes a slow-EMA shadow of
|
||||
// `0.5 × readiness_ema` after Pearl-A bootstrap.
|
||||
assert_eq!(SENTINEL_PLAN_THRESHOLD_FLOOR, 0.1);
|
||||
// Category-1 dimensional safety: floor in [0.05, 0.50].
|
||||
assert_eq!(PLAN_THRESHOLD_FLOOR_MIN, 0.05);
|
||||
assert_eq!(PLAN_THRESHOLD_FLOOR_MAX, 0.50);
|
||||
// EMA α matches per-fold cadence (slower than P0-A reward-cap's
|
||||
// 0.01; same as P1-Producer Kelly priors' 0.005).
|
||||
assert_eq!(PLAN_THRESHOLD_FLOOR_EMA_ALPHA, 0.005);
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn all_sp14_audit_4b_plan_floor_slots_fit_within_isv_total_dim() {
|
||||
use crate::cuda_pipeline::gpu_dqn_trainer::ISV_TOTAL_DIM;
|
||||
assert!(
|
||||
SP14_AUDIT_4B_PLAN_FLOOR_SLOT_END <= ISV_TOTAL_DIM,
|
||||
"SP14_AUDIT_4B_PLAN_FLOOR_SLOT_END={} exceeds ISV_TOTAL_DIM={} — bus too small for Class A audit-fix Batch 4-B plan_threshold floor slot; \
|
||||
bump ISV_TOTAL_DIM in gpu_dqn_trainer.rs (and update layout_fingerprint_seed()).",
|
||||
SP14_AUDIT_4B_PLAN_FLOOR_SLOT_END, ISV_TOTAL_DIM,
|
||||
);
|
||||
}
|
||||
|
||||
/// Lock Class A audit-fix Batch 4-B (2026-05-08) adaptive
|
||||
/// MIN_HOLD_TEMPERATURE slot. Single contiguous slot [460..461)
|
||||
/// replacing the deleted epoch-driven schedule
|
||||
/// `T_end + (T_start − T_end) × exp(-epoch / decay)`. Sentinel
|
||||
/// matches the pre-fix MIN_HOLD_TEMPERATURE_START=50 for
|
||||
/// bit-identical cold-start (first epoch under both regimes had
|
||||
/// T=50).
|
||||
#[test]
|
||||
fn sp14_audit_4b_min_hold_temperature_slot_layout_locked() {
|
||||
assert_eq!(SP14_AUDIT_4B_TEMP_SLOT_BASE, 460);
|
||||
assert_eq!(SP14_AUDIT_4B_TEMP_SLOT_END, 461);
|
||||
assert_eq!(MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX, 460);
|
||||
// Sentinel matches pre-fix MIN_HOLD_TEMPERATURE_START=50 for
|
||||
// bit-identical cold-start. Producer kernel writes a moderate-
|
||||
// cadence EMA of dir_acc-skill-driven temperature after
|
||||
// Pearl-A bootstrap.
|
||||
assert_eq!(SENTINEL_MIN_HOLD_TEMPERATURE, 50.0);
|
||||
// Category-1 dimensional safety: temp in [5, 50] (matches the
|
||||
// original schedule range MIN_HOLD_TEMPERATURE_END=5 ↔
|
||||
// MIN_HOLD_TEMPERATURE_START=50).
|
||||
assert_eq!(MIN_HOLD_TEMPERATURE_MIN, 5.0);
|
||||
assert_eq!(MIN_HOLD_TEMPERATURE_MAX, 50.0);
|
||||
// EMA α matches mid-cadence (slower than per-step skill EMA,
|
||||
// faster than per-fold beliefs; same as P0-A reward-cap's 0.05
|
||||
// mid-cadence).
|
||||
assert_eq!(MIN_HOLD_TEMPERATURE_EMA_ALPHA, 0.05);
|
||||
// Random-baseline anchor for K=2 binary dir-acc.
|
||||
assert_eq!(DIR_ACC_RANDOM_BASELINE, 0.5);
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn all_sp14_audit_4b_temp_slots_fit_within_isv_total_dim() {
|
||||
use crate::cuda_pipeline::gpu_dqn_trainer::ISV_TOTAL_DIM;
|
||||
assert!(
|
||||
SP14_AUDIT_4B_TEMP_SLOT_END <= ISV_TOTAL_DIM,
|
||||
"SP14_AUDIT_4B_TEMP_SLOT_END={} exceeds ISV_TOTAL_DIM={} — bus too small for Class A audit-fix Batch 4-B MIN_HOLD_TEMPERATURE slot; \
|
||||
bump ISV_TOTAL_DIM in gpu_dqn_trainer.rs (and update layout_fingerprint_seed()).",
|
||||
SP14_AUDIT_4B_TEMP_SLOT_END, ISV_TOTAL_DIM,
|
||||
);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -259,11 +259,19 @@
|
||||
// 3.0 = 60% of REWARD_POS_CAP. Meaningful gradient pressure but not
|
||||
// paralyzing — a profitable scalp can still net positive with the penalty.
|
||||
//
|
||||
// MIN_HOLD_TEMPERATURE_{START,END,DECAY} — annealing schedule for the
|
||||
// soft-saturation temperature in the deficit/(deficit + T) formula. Computed
|
||||
// in Rust per epoch via T(e) = T_end + (T_start - T_end) × exp(-e/decay)
|
||||
// and passed to the kernel as a scalar (not a constant lookup). Curriculum:
|
||||
// early training forgiving (T=50), late training sharp (T=5).
|
||||
// MIN_HOLD_TEMPERATURE_{START,END,DECAY} — DELETED by Class A audit-fix
|
||||
// Batch 4-B (2026-05-08). The epoch-driven anneal schedule
|
||||
// `T_end + (T_start - T_end) × exp(-epoch / decay)` was replaced with
|
||||
// the signal-driven ISV slot `MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460`
|
||||
// (sp14_isv_slots.rs). Driving signal: dir_acc skill from
|
||||
// ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]. Per
|
||||
// `feedback_no_legacy_aliases.md` — once the ISV-driven version is
|
||||
// wired, the legacy schedule constants are deleted (no half-migrated
|
||||
// state). The single epoch-driven entry point
|
||||
// `min_hold_temperature_for_epoch` in training_loop.rs is also
|
||||
// deleted; the training loop now reads ISV[460] (with cold-start
|
||||
// fallback to MIN_HOLD_TEMPERATURE_START_FALLBACK=50.0) when seeding
|
||||
// the kernel-passed scalar.
|
||||
//
|
||||
// Phase 1 = Invariant-1 numerical anchors (static constants). Phase 2 lifts
|
||||
// the target / penalty_max / temperature schedule to ISV-driven controllers
|
||||
@@ -274,9 +282,11 @@
|
||||
#define REWARD_NEG_CAP -10.0f
|
||||
#define MIN_HOLD_TARGET 30.0f
|
||||
#define MIN_HOLD_PENALTY_MAX 3.0f
|
||||
#define MIN_HOLD_TEMPERATURE_START 50.0f
|
||||
#define MIN_HOLD_TEMPERATURE_END 5.0f
|
||||
#define MIN_HOLD_TEMPERATURE_DECAY 20.0f
|
||||
// Cold-start fallback for the kernel-passed `min_hold_temperature` scalar
|
||||
// when ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460] is at sentinel.
|
||||
// Bit-identical to pre-Batch-4B epoch-0 behavior under the deleted
|
||||
// MIN_HOLD_TEMPERATURE_START=50.0 anchor.
|
||||
#define MIN_HOLD_TEMPERATURE_FALLBACK 50.0f
|
||||
|
||||
// Class A P0-A (2026-05-08) — adaptive REWARD_POS/NEG_CAP ISV slots.
|
||||
// REWARD_POS_CAP / REWARD_NEG_CAP above remain as Phase-1 cold-start
|
||||
@@ -364,6 +374,42 @@
|
||||
// static default. Bit-identical pre-Batch-4A behavior.
|
||||
#define DD_SATURATION_FLOOR_DEFAULT 0.25f
|
||||
|
||||
// Class A audit-fix Batch 4-B (2026-05-08) — adaptive plan_threshold floor.
|
||||
// Replaces the hardcoded `0.1f` lower bound in
|
||||
// `plan_threshold_update_kernel.cu:44` (`fmaxf(0.1f, 0.5f * ema)`). The
|
||||
// floor is INLINE-produced by the same kernel — on each invocation the
|
||||
// kernel writes both the new plan_threshold (slot 49) and a slow-EMA
|
||||
// floor (slot 459) and uses the floor as the lower bound on the next
|
||||
// invocation. Sentinel matches the pre-fix value (0.1) for bit-identical
|
||||
// cold-start. Mirrors `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs`
|
||||
// constants of the same names (locked by
|
||||
// `sp14_audit_4b_plan_threshold_floor_slot_layout_locked` test).
|
||||
#define PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX 459
|
||||
#define SENTINEL_PLAN_THRESHOLD_FLOOR 0.1f
|
||||
#define PLAN_THRESHOLD_FLOOR_MIN_BOUND 0.05f
|
||||
#define PLAN_THRESHOLD_FLOOR_MAX_BOUND 0.50f
|
||||
#define PLAN_THRESHOLD_FLOOR_DEFAULT 0.1f
|
||||
#define PLAN_THRESHOLD_FLOOR_EMA_ALPHA 0.005f
|
||||
|
||||
// Class A audit-fix Batch 4-B (2026-05-08) — adaptive MIN_HOLD_TEMPERATURE.
|
||||
// Replaces the deleted epoch-driven schedule
|
||||
// `T_end + (T_start - T_end) × exp(-epoch / decay)` with an ISV-driven
|
||||
// signal-driven temperature derived from the model's directional
|
||||
// accuracy skill (ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]). Producer kernel:
|
||||
// `min_hold_temperature_update_kernel` (per-epoch boundary cadence).
|
||||
// Sentinel matches the pre-fix MIN_HOLD_TEMPERATURE_START=50 for
|
||||
// bit-identical cold-start (first epoch under both regimes had T=50).
|
||||
// Mirrors `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` constants of
|
||||
// the same names (locked by
|
||||
// `sp14_audit_4b_min_hold_temperature_slot_layout_locked` test).
|
||||
#define MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX 460
|
||||
#define SENTINEL_MIN_HOLD_TEMPERATURE 50.0f
|
||||
#define MIN_HOLD_TEMPERATURE_MIN_BOUND 5.0f
|
||||
#define MIN_HOLD_TEMPERATURE_MAX_BOUND 50.0f
|
||||
#define MIN_HOLD_TEMPERATURE_DEFAULT 50.0f
|
||||
#define MIN_HOLD_TEMPERATURE_EMA_ALPHA 0.05f
|
||||
#define DIR_ACC_RANDOM_BASELINE 0.5f
|
||||
|
||||
// ── Compile-time checks ──
|
||||
static_assert(SL_PADDING_START + SL_PADDING_DIM == SL_STATE_DIM,
|
||||
"State layout dimensions must sum to SL_STATE_DIM");
|
||||
|
||||
@@ -1115,6 +1115,69 @@ impl StateResetRegistry {
|
||||
category: ResetCategory::FoldReset,
|
||||
description: "ISV[DD_SATURATION_FLOOR_ADAPTIVE_INDEX=458] — Class A audit-fix Batch 4-A adaptive DD saturation floor (replaces hardcoded `0.25f` in `trade_physics.cuh::apply_margin_cap`, the upper end of the linear position-size scaling ramp `dd_scale = max(0.05, 1.0 − dd_frac/saturation_floor)`). Produced by `dd_saturation_floor_update_kernel` from p75(per-env DD_MAX) × safety_factor=1.5 over the `dd_state_per_env[n_envs * 6]` tile (offset 1 = DD_MAX). FoldReset sentinel SENTINEL_DD_SATURATION_FLOOR=0.25 — Pearl-A first-observation bootstrap (matches pre-fix hardcoded value for bit-identical cold-start). Welford α=0.01 slow EMA thereafter. Bounds [0.10, 0.50] — Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`. Distinct semantic role from SP15_DD_THRESHOLD_INDEX=421 (the SP15 quadratic DD-penalty trigger threshold, a lower bound). Consumed by `apply_margin_cap` at line ~154 (threaded `isv_signals_ptr`, NULL-tolerant cold-start fallback to DD_SATURATION_FLOOR_DEFAULT). Atomically wired with the legacy `compute_drawdown_penalty` deletion (Item 2 Case A) since SP15's quadratic asymmetric DD penalty in `compute_sp15_final_reward_kernel.cu:154` is the production-grade replacement; layering both creates double-counting.",
|
||||
},
|
||||
// ── Class A audit-fix Batch 4-B (2026-05-08): adaptive plan_threshold floor + MIN_HOLD_TEMPERATURE ──
|
||||
// Two slots [459..461) replacing (a) the hardcoded `0.1f`
|
||||
// floor in `plan_threshold_update_kernel.cu:44` and (b) the
|
||||
// deleted epoch-driven anneal schedule
|
||||
// `T_end + (T_start − T_end) × exp(-epoch / decay)` for
|
||||
// MIN_HOLD_TEMPERATURE.
|
||||
//
|
||||
// PLAN_THRESHOLD_FLOOR_ADAPTIVE (459) — Design Y inline
|
||||
// producer (no new kernel; the same
|
||||
// `plan_threshold_update_kernel` writes both slot 49 and
|
||||
// the slow-EMA-shadow floor at slot 459 in the same launch,
|
||||
// then uses the freshly-written floor as the lower bound
|
||||
// on the derived plan_threshold). FoldReset sentinel 0.1
|
||||
// matches the pre-fix hardcoded value for bit-identical
|
||||
// cold-start. α=0.005 slow EMA (per-fold cadence; the
|
||||
// floor is a long-term shadow of `0.5 × readiness_ema`
|
||||
// and must move much slower than the readiness EMA itself).
|
||||
// Bounds [0.05, 0.50] are Category-1 dimensional safety
|
||||
// floors per `feedback_isv_for_adaptive_bounds`.
|
||||
//
|
||||
// MIN_HOLD_TEMPERATURE_ADAPTIVE (460) — Producer kernel
|
||||
// `min_hold_temperature_update_kernel` reads
|
||||
// `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` (sp13 fast EMA of
|
||||
// binary directional accuracy), maps it to a [5, 50]
|
||||
// temperature via
|
||||
// `temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × clamp((short_ema
|
||||
// − 0.5)/0.5, 0, 1)`, slow-EMA-blends into slot 460.
|
||||
// FoldReset sentinel 50.0 matches the deleted
|
||||
// MIN_HOLD_TEMPERATURE_START anchor for bit-identical
|
||||
// cold-start under both regimes (first epoch had T=50
|
||||
// under the old schedule too). α=0.05 mid-cadence EMA.
|
||||
// Bounds [5, 50] are Category-1 dimensional safety floors
|
||||
// per `feedback_isv_for_adaptive_bounds` (matches the
|
||||
// original schedule range MIN_HOLD_TEMPERATURE_END=5 ↔
|
||||
// MIN_HOLD_TEMPERATURE_START=50). Decouples temperature
|
||||
// from epoch number — the training loop now reads slot 460
|
||||
// via `DQNTrainer::read_min_hold_temperature_from_isv` to
|
||||
// seed the kernel-passed scalar (with cold-start fallback
|
||||
// to MIN_HOLD_TEMPERATURE_FALLBACK=50.0 when the slot is
|
||||
// at sentinel).
|
||||
//
|
||||
// Per `feedback_no_legacy_aliases.md` and
|
||||
// `feedback_no_partial_refactor.md` the legacy
|
||||
// `min_hold_temperature_for_epoch` helper +
|
||||
// MIN_HOLD_TEMPERATURE_{START, END, DECAY} #defines are
|
||||
// deleted atomically with this wiring (no half-migrated
|
||||
// state). Both call sites in training_loop.rs are
|
||||
// migrated to the new ISV reader in the same commit.
|
||||
//
|
||||
// Without these dispatch arms, the C.10 lesson recurs:
|
||||
// the FoldReset registry entry exists, but no actual reset
|
||||
// fires — slot drifts across folds and the layout-fingerprint
|
||||
// smoke test eventually catches it as a runtime crash.
|
||||
RegistryEntry {
|
||||
name: "sp14_audit_4b_plan_threshold_floor_adaptive",
|
||||
category: ResetCategory::FoldReset,
|
||||
description: "ISV[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX=459] — Class A audit-fix Batch 4-B adaptive plan_threshold floor (replaces hardcoded `0.1f` in `plan_threshold_update_kernel.cu:44`, the lower bound of the derived plan_threshold = max(floor, 0.5 × readiness_ema)). Design Y inline producer: the same `plan_threshold_update_kernel` writes both ISV[PLAN_THRESHOLD_INDEX=49] and ISV[459] in the same launch — the floor tracks a slow EMA of `0.5 × readiness_ema` (the very expression the kernel uses as its derived threshold target) and the kernel uses the freshly-written floor as the lower bound. FoldReset sentinel SENTINEL_PLAN_THRESHOLD_FLOOR=0.1 — Pearl-A first-observation bootstrap (matches pre-fix hardcoded value for bit-identical cold-start). α=0.005 slow EMA thereafter (per-fold cadence; the floor is a long-term shadow and must move much slower than the readiness EMA itself which adapts at α≈[0.05, 0.10]). Bounds [0.05, 0.50] — Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds.md` (probability units; below 5% the readiness gate is effectively absent, above 50% it is uselessly strict). Distinct from PLAN_THRESHOLD_INDEX=49 (the live derived threshold, fast cadence) — slot 459 is the slow-moving lower bound on slot 49.",
|
||||
},
|
||||
RegistryEntry {
|
||||
name: "sp14_audit_4b_min_hold_temperature_adaptive",
|
||||
category: ResetCategory::FoldReset,
|
||||
description: "ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460] — Class A audit-fix Batch 4-B adaptive MIN_HOLD_TEMPERATURE (replaces the DELETED epoch-driven anneal schedule `T_end + (T_start − T_end) × exp(-epoch / decay)` formerly at state_layout.cuh::MIN_HOLD_TEMPERATURE_{START=50, END=5, DECAY=20} + training_loop.rs::min_hold_temperature_for_epoch). Producer kernel `min_hold_temperature_update_kernel` reads ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373] (sp13 fast EMA of binary directional accuracy), maps it to a [5, 50] temperature via `temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × clamp((short_ema − 0.5)/0.5, 0, 1)`, slow-EMA-blends into slot 460. FoldReset sentinel SENTINEL_MIN_HOLD_TEMPERATURE=50.0 — Pearl-A first-observation bootstrap (matches the deleted MIN_HOLD_TEMPERATURE_START=50 anchor for bit-identical cold-start under both regimes; epoch-0 had T=50 under the old schedule too). α=0.05 mid-cadence EMA. Bounds [5, 50] — Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds.md` (matches the original schedule range). Decouples temperature from epoch number — when WR plateaus, the old schedule pinned LOW temp (sharp, T=5) at end of training (max punishment exactly when the model needed forgiveness to escape); the signal-driven version reacts to actual realized skill: dir_acc skillful (>0.5) → temp HIGH (permissive); dir_acc at random (≈0.5) → temp LOW (sharp commitment pressure). Consumed by experience_kernels.cu via the kernel-passed `min_hold_temperature` scalar — the training loop seeds that scalar per-epoch via `DQNTrainer::read_min_hold_temperature_from_isv` (cold-start fallback to MIN_HOLD_TEMPERATURE_FALLBACK=50.0 when slot is at sentinel). Per `feedback_no_legacy_aliases.md` + `feedback_no_partial_refactor.md` the legacy schedule constants + helper function are deleted atomically.",
|
||||
},
|
||||
// ── SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots ────────────
|
||||
// Two ISV slots [407, 408]:
|
||||
// - OFI_IMPACT_LAMBDA_INDEX=407: Invariant-1 anchor (NOT a
|
||||
|
||||
@@ -42,35 +42,29 @@ use super::super::financials::compute_epoch_financials;
|
||||
use super::super::monitoring::TrainingMonitor;
|
||||
use super::DQNTrainer;
|
||||
|
||||
/// SP12 v3 (2026-05-04): per-epoch min-hold soft-penalty temperature schedule.
|
||||
/// SP12 v3 (2026-05-04): per-epoch min-hold soft-penalty temperature.
|
||||
///
|
||||
/// Computes the temperature T(e) controlling the smoothness of the
|
||||
/// `deficit/(deficit + T)` saturation factor inside the
|
||||
/// `experience_env_step` kernel. Annealing curriculum:
|
||||
/// `T(e) = T_end + (T_start - T_end) × exp(-e / decay)`
|
||||
/// **DELETED** by Class A audit-fix Batch 4-B (2026-05-08, Item 4). The
|
||||
/// epoch-driven anneal schedule `T(e) = T_end + (T_start − T_end) ×
|
||||
/// exp(-e / decay)` was replaced with the signal-driven ISV slot
|
||||
/// `MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460` (sp14_isv_slots.rs).
|
||||
/// Driving signal: dir_acc skill from
|
||||
/// `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` (sp13 fast EMA of binary
|
||||
/// directional accuracy). Producer kernel:
|
||||
/// `min_hold_temperature_update_kernel.cu` (per-epoch boundary).
|
||||
///
|
||||
/// With Phase 1 anchors `T_start=50.0`, `T_end=5.0`, `decay=20.0`
|
||||
/// (state_layout.cuh::MIN_HOLD_TEMPERATURE_{START,END,DECAY}):
|
||||
/// Per `feedback_no_legacy_aliases.md` — once the ISV-driven version is
|
||||
/// wired, the legacy schedule constants and helper function are deleted
|
||||
/// (no half-migrated state). The kernel-passed scalar
|
||||
/// `config.min_hold_temperature` is now seeded from
|
||||
/// `ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX]` (with cold-start fallback
|
||||
/// to MIN_HOLD_TEMPERATURE_FALLBACK=50.0 when the slot is at sentinel)
|
||||
/// rather than computed from `epoch`. The collector's
|
||||
/// `min_hold_temperature` field default (50.0) provides the cold-start
|
||||
/// fallback bit-identical to pre-Batch-4B epoch-0 behavior.
|
||||
///
|
||||
/// | epoch | T | regime |
|
||||
/// |------:|----:|------------------|
|
||||
/// | 0 | 50 | forgiving |
|
||||
/// | 20 | 21 | mid-transition |
|
||||
/// | 50 | 9 | sharp |
|
||||
/// | 100 | 5 | full-MFT enforce |
|
||||
///
|
||||
/// Per spec §design-2 — early training needs forgiveness so the policy
|
||||
/// can learn reward gradients on short trades; late training needs
|
||||
/// sharpness so the min-hold penalty bites and enforces commitment.
|
||||
/// Constants live in CUDA headers; this Rust mirror duplicates them
|
||||
/// (kernel side keeps T as a launch scalar so we can switch to an
|
||||
/// ISV-driven formulation in Phase 2 without recompiling cubin).
|
||||
pub(crate) fn min_hold_temperature_for_epoch(epoch: usize) -> f32 {
|
||||
const T_START: f32 = 50.0;
|
||||
const T_END: f32 = 5.0;
|
||||
const DECAY_RATE: f32 = 20.0;
|
||||
T_END + (T_START - T_END) * (-(epoch as f32) / DECAY_RATE).exp()
|
||||
}
|
||||
/// See `crates/ml/src/cuda_pipeline/min_hold_temperature_update_kernel.cu`
|
||||
/// for the producer logic and the dir_acc skill mapping rationale.
|
||||
|
||||
/// Normalized Shannon entropy of a discrete probability distribution.
|
||||
///
|
||||
@@ -174,6 +168,53 @@ pub(crate) struct EpochLogOutput {
|
||||
|
||||
impl DQNTrainer {
|
||||
// (Old `train_with_data_full_loop` removed — `train_with_data_full_loop_slices` is the sole production path)
|
||||
|
||||
/// Class A audit-fix Batch 4-B (2026-05-08, Item 4): read the
|
||||
/// adaptive MIN_HOLD_TEMPERATURE from
|
||||
/// `ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460]` with cold-start
|
||||
/// fallback to MIN_HOLD_TEMPERATURE_FALLBACK=50.0 (matches the
|
||||
/// deleted MIN_HOLD_TEMPERATURE_START anchor for bit-identical
|
||||
/// pre-Batch-4B epoch-0 behavior).
|
||||
///
|
||||
/// Cold-start triggers when:
|
||||
/// - `fused_ctx` is None (pre-initialization paths), OR
|
||||
/// - the ISV slot is at sentinel (within EPS of 50.0) — no
|
||||
/// dir-acc samples have updated it yet, OR
|
||||
/// - the ISV slot is outside the dimensional-safety window
|
||||
/// [MIN_HOLD_TEMPERATURE_MIN, MAX] (defends against malformed
|
||||
/// prior state — mirrors the existing effective_min_hold_target
|
||||
/// consumer guard pattern in experience_kernels.cu line ~3197).
|
||||
///
|
||||
/// Per `feedback_isv_for_adaptive_bounds.md` (the dimensional
|
||||
/// safety bounds are guards, not tuning) and
|
||||
/// `feedback_no_legacy_aliases.md` (the deleted
|
||||
/// min_hold_temperature_for_epoch helper had its callers migrated
|
||||
/// to this reader atomically with the wiring of the new producer
|
||||
/// kernel).
|
||||
pub(crate) fn read_min_hold_temperature_from_isv(&self) -> f32 {
|
||||
use crate::cuda_pipeline::sp14_isv_slots::{
|
||||
MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX, MIN_HOLD_TEMPERATURE_MAX,
|
||||
MIN_HOLD_TEMPERATURE_MIN, SENTINEL_MIN_HOLD_TEMPERATURE,
|
||||
};
|
||||
// Cold-start fallback anchor — matches the deleted
|
||||
// MIN_HOLD_TEMPERATURE_START=50 for bit-identical pre-Batch-4B
|
||||
// epoch-0 behavior.
|
||||
const FALLBACK: f32 = 50.0;
|
||||
let Some(ref fused) = self.fused_ctx else { return FALLBACK };
|
||||
let v = fused.read_isv_signal_at(MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX);
|
||||
// Sentinel detect: within EPS of 50.0 → no producer update yet.
|
||||
if (v - SENTINEL_MIN_HOLD_TEMPERATURE).abs() < 1e-6 {
|
||||
return FALLBACK;
|
||||
}
|
||||
// Dimensional-safety window guard (defends against malformed
|
||||
// prior state — mirrors the existing consumer-side guards in
|
||||
// experience_kernels.cu).
|
||||
if v < MIN_HOLD_TEMPERATURE_MIN || v > MIN_HOLD_TEMPERATURE_MAX {
|
||||
return FALLBACK;
|
||||
}
|
||||
v
|
||||
}
|
||||
|
||||
// ═══════════════════════════════════════════════════════════════════════
|
||||
// Main training loop — fixed-size arrays, zero heap alloc per bar
|
||||
// ═══════════════════════════════════════════════════════════════════════
|
||||
@@ -466,6 +507,24 @@ impl DQNTrainer {
|
||||
) {
|
||||
tracing::warn!(epoch, "launch_dd_saturation_floor_update failed (non-fatal): {e}");
|
||||
}
|
||||
// Class A audit-fix Batch 4-B (2026-05-08, Item 4):
|
||||
// adaptive MIN_HOLD_TEMPERATURE — per-epoch boundary
|
||||
// launch. Reads `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]`
|
||||
// (sp13 fast EMA of binary directional accuracy), maps
|
||||
// to a [5, 50] temperature via `temp = TEMP_MIN +
|
||||
// (TEMP_MAX − TEMP_MIN) × clamp((short_ema − 0.5)/0.5,
|
||||
// 0, 1)`, slow-EMA-blends into ISV[460]. Pearl-A
|
||||
// first-observation bootstrap (sentinel 50.0 matches
|
||||
// the deleted MIN_HOLD_TEMPERATURE_START=50 anchor for
|
||||
// bit-identical cold-start) + α=0.05 mid-cadence EMA.
|
||||
// Replaces the deleted epoch-driven anneal schedule
|
||||
// `T_end + (T_start − T_end) × exp(-epoch / decay)`
|
||||
// (formerly state_layout.cuh::
|
||||
// MIN_HOLD_TEMPERATURE_{START, END, DECAY} +
|
||||
// training_loop.rs::min_hold_temperature_for_epoch).
|
||||
if let Err(e) = fused.trainer().launch_min_hold_temperature_update() {
|
||||
tracing::warn!(epoch, "launch_min_hold_temperature_update failed (non-fatal): {e}");
|
||||
}
|
||||
}
|
||||
|
||||
// D.8 Plan 2 Task 6C: update TLOB regime focus EMA in ISV at epoch boundary.
|
||||
@@ -2173,12 +2232,45 @@ impl DQNTrainer {
|
||||
// (MIN_HOLD_TARGET=30, MIN_HOLD_PENALTY_MAX=3.0). Phase 2 lifts these
|
||||
// to ISV-driven controllers per feedback_isv_for_adaptive_bounds when/if
|
||||
// validation results indicate adaptive need.
|
||||
// `min_hold_temperature` is computed PER EPOCH from the annealing
|
||||
// schedule. Curriculum: forgiving (T=50) at epoch 0 → sharp (T=5)
|
||||
// around epoch 100. Half-life ≈ MIN_HOLD_TEMPERATURE_DECAY (=20)
|
||||
// epochs. The kernel evaluates `deficit/(deficit + T)` per voluntary
|
||||
// exit; lower T narrows the transition into a sharper penalty cliff.
|
||||
min_hold_temperature: min_hold_temperature_for_epoch(self.current_epoch),
|
||||
// Class A audit-fix Batch 4-B (2026-05-08, Item 4):
|
||||
// `min_hold_temperature` migrated from epoch-driven schedule to
|
||||
// signal-driven ISV slot 460. Read freshly from
|
||||
// `ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX]` here so the kernel-
|
||||
// passed scalar reflects the most recent producer output (the
|
||||
// producer kernel runs at the epoch-boundary BEFORE this config
|
||||
// is built; the launcher in process_epoch_boundary writes to
|
||||
// slot 460). Cold-start fallback: when the slot is at sentinel
|
||||
// (within EPS of 50.0) or outside [MIN_HOLD_TEMPERATURE_MIN,
|
||||
// MAX], fall back to the MIN_HOLD_TEMPERATURE_FALLBACK=50.0
|
||||
// anchor — bit-identical to pre-Batch-4B epoch-0 behavior.
|
||||
// Inlined here (rather than calling the
|
||||
// `read_min_hold_temperature_from_isv` helper) because
|
||||
// `&mut self.gpu_experience_collector` already holds an
|
||||
// exclusive borrow of `self` — same pattern as the sibling
|
||||
// noise_sigma_per_branch read above which dereferences
|
||||
// `self.fused_ctx.as_ref()` directly. Per
|
||||
// `feedback_isv_for_adaptive_bounds` and
|
||||
// `feedback_no_legacy_aliases` (deleted
|
||||
// min_hold_temperature_for_epoch helper).
|
||||
min_hold_temperature: {
|
||||
use crate::cuda_pipeline::sp14_isv_slots::{
|
||||
MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX, MIN_HOLD_TEMPERATURE_MAX,
|
||||
MIN_HOLD_TEMPERATURE_MIN, SENTINEL_MIN_HOLD_TEMPERATURE,
|
||||
};
|
||||
const FALLBACK: f32 = 50.0;
|
||||
if let Some(ctx) = self.fused_ctx.as_ref() {
|
||||
let v = ctx.read_isv_signal_at(MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX);
|
||||
if (v - SENTINEL_MIN_HOLD_TEMPERATURE).abs() < 1e-6
|
||||
|| v < MIN_HOLD_TEMPERATURE_MIN
|
||||
|| v > MIN_HOLD_TEMPERATURE_MAX {
|
||||
FALLBACK
|
||||
} else {
|
||||
v
|
||||
}
|
||||
} else {
|
||||
FALLBACK
|
||||
}
|
||||
},
|
||||
..Default::default()
|
||||
};
|
||||
|
||||
@@ -4905,13 +4997,20 @@ impl DQNTrainer {
|
||||
// (3) the temperature annealing schedule is on the expected curve
|
||||
// (forgiving early → sharp late). Same emit cadence as sp11_reward
|
||||
// above (per-epoch on the metrics path).
|
||||
// Class A audit-fix Batch 4-B (2026-05-08, Item 4):
|
||||
// min_hold_T migrated from epoch-driven schedule to
|
||||
// signal-driven ISV slot 460 — read freshly here so the
|
||||
// diagnostic mirrors the value the kernel actually used
|
||||
// for this epoch's experience collection. Cold-start
|
||||
// fallback: sentinel → MIN_HOLD_TEMPERATURE_FALLBACK=50.
|
||||
let min_hold_t_diag = self.read_min_hold_temperature_from_isv();
|
||||
tracing::info!(
|
||||
"HEALTH_DIAG[{}]: sp12_event_reward [pos_cap={:.1} neg_cap={:.1} min_hold_tgt={:.1} min_hold_T={:.2} min_hold_pen={:.1}]",
|
||||
epoch,
|
||||
5.0_f32, // REWARD_POS_CAP — Invariant-1 anchor in state_layout.cuh
|
||||
-10.0_f32, // REWARD_NEG_CAP
|
||||
30.0_f32, // MIN_HOLD_TARGET (Phase 1 default; Phase 2 lifts to ISV)
|
||||
min_hold_temperature_for_epoch(epoch),
|
||||
min_hold_t_diag,
|
||||
3.0_f32, // MIN_HOLD_PENALTY_MAX
|
||||
);
|
||||
}
|
||||
@@ -8288,6 +8387,36 @@ impl DQNTrainer {
|
||||
);
|
||||
}
|
||||
}
|
||||
// Class A audit-fix Batch 4-B (2026-05-08): adaptive
|
||||
// plan_threshold floor + adaptive MIN_HOLD_TEMPERATURE.
|
||||
// Two FoldReset slots [459..461) — sentinels match the
|
||||
// pre-fix anchors so the producers' Pearl-A bootstrap
|
||||
// checks fire cleanly on the new fold's first launch
|
||||
// (avoids cross-fold EMA contamination). Cold-start is
|
||||
// bit-identical to pre-Batch-4B until the first valid
|
||||
// observation lands.
|
||||
"sp14_audit_4b_plan_threshold_floor_adaptive" => {
|
||||
if let Some(ref fused) = self.fused_ctx {
|
||||
use crate::cuda_pipeline::sp14_isv_slots::{
|
||||
PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX, SENTINEL_PLAN_THRESHOLD_FLOOR,
|
||||
};
|
||||
fused.trainer().write_isv_signal_at(
|
||||
PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX,
|
||||
SENTINEL_PLAN_THRESHOLD_FLOOR,
|
||||
);
|
||||
}
|
||||
}
|
||||
"sp14_audit_4b_min_hold_temperature_adaptive" => {
|
||||
if let Some(ref fused) = self.fused_ctx {
|
||||
use crate::cuda_pipeline::sp14_isv_slots::{
|
||||
MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX, SENTINEL_MIN_HOLD_TEMPERATURE,
|
||||
};
|
||||
fused.trainer().write_isv_signal_at(
|
||||
MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX,
|
||||
SENTINEL_MIN_HOLD_TEMPERATURE,
|
||||
);
|
||||
}
|
||||
}
|
||||
// SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots.
|
||||
// OFI_IMPACT_LAMBDA_INDEX=407 is an Invariant-1 anchor (NOT
|
||||
// a stateful EMA) — rewrite the constructor's value at fold
|
||||
|
||||
@@ -1518,3 +1518,502 @@ mod sp14_audit_4a_dd_saturation_floor_gpu {
|
||||
);
|
||||
}
|
||||
}
|
||||
|
||||
// ═══════════════════════════════════════════════════════════════════════════
|
||||
// Class A audit-fix Batch 4-B (2026-05-08, Item 3) — adaptive plan_threshold
|
||||
// floor inline producer (Design Y inside `plan_threshold_update_kernel`).
|
||||
//
|
||||
// Verifies:
|
||||
// 1. Pearl-A bootstrap: sentinel 0.1 → REPLACE with `0.5 × ema`.
|
||||
// 2. Slow EMA (α=0.005) blend after bootstrap.
|
||||
// 3. Bilateral clamp: huge readiness clamps floor to MAX=0.50.
|
||||
//
|
||||
// All tests are #[ignore = "requires GPU"]; gated under #[cfg(feature = "cuda")].
|
||||
// ═══════════════════════════════════════════════════════════════════════════
|
||||
|
||||
#[cfg(feature = "cuda")]
|
||||
#[allow(unsafe_code)] // CUDA kernel launch + mapped-pinned memory.
|
||||
mod sp14_audit_4b_plan_threshold_floor_gpu {
|
||||
use std::sync::Arc;
|
||||
|
||||
use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg};
|
||||
use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer;
|
||||
use ml::cuda_pipeline::sp14_isv_slots::{
|
||||
PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX, PLAN_THRESHOLD_FLOOR_EMA_ALPHA,
|
||||
PLAN_THRESHOLD_FLOOR_MAX, PLAN_THRESHOLD_FLOOR_MIN, SENTINEL_PLAN_THRESHOLD_FLOOR,
|
||||
};
|
||||
|
||||
const PLAN_THRESHOLD_UPDATE_CUBIN: &[u8] =
|
||||
include_bytes!(concat!(env!("OUT_DIR"), "/plan_threshold_update_kernel.cubin"));
|
||||
|
||||
/// ISV slot indices used by the kernel (mirror gpu_dqn_trainer.rs).
|
||||
const READINESS_EMA_INDEX: usize = 75;
|
||||
const PLAN_THRESHOLD_INDEX: usize = 49;
|
||||
const SHARPE_EMA_INDEX: usize = 22;
|
||||
const ALPHA_BASE: f32 = 0.05;
|
||||
|
||||
fn make_stream() -> Arc<CudaStream> {
|
||||
let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
|
||||
ctx.default_stream()
|
||||
}
|
||||
|
||||
fn load_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
|
||||
let module = stream
|
||||
.context()
|
||||
.load_cubin(PLAN_THRESHOLD_UPDATE_CUBIN.to_vec())
|
||||
.expect("load plan_threshold_update_kernel cubin");
|
||||
module
|
||||
.load_function("plan_threshold_update")
|
||||
.expect("load plan_threshold_update function")
|
||||
}
|
||||
|
||||
fn launch_plan_threshold(
|
||||
stream: &Arc<CudaStream>,
|
||||
kernel: &CudaFunction,
|
||||
readiness_ptr: u64,
|
||||
n_samples: i32,
|
||||
isv_ptr: u64,
|
||||
) {
|
||||
let readiness_slot = READINESS_EMA_INDEX as i32;
|
||||
let thr_slot = PLAN_THRESHOLD_INDEX as i32;
|
||||
let sharpe_slot = SHARPE_EMA_INDEX as i32;
|
||||
let alpha_base = ALPHA_BASE;
|
||||
unsafe {
|
||||
stream
|
||||
.launch_builder(kernel)
|
||||
.arg(&readiness_ptr)
|
||||
.arg(&n_samples)
|
||||
.arg(&isv_ptr)
|
||||
.arg(&readiness_slot)
|
||||
.arg(&thr_slot)
|
||||
.arg(&sharpe_slot)
|
||||
.arg(&alpha_base)
|
||||
.launch(LaunchConfig {
|
||||
grid_dim: (1, 1, 1),
|
||||
block_dim: (1, 1, 1),
|
||||
shared_mem_bytes: 0,
|
||||
})
|
||||
.expect("launch plan_threshold_update");
|
||||
}
|
||||
stream.synchronize().expect("sync after plan_threshold_update");
|
||||
}
|
||||
|
||||
/// Test 1 — Pearl-A first-observation bootstrap.
|
||||
///
|
||||
/// Seed `readiness_ema` at 0.4 and `floor` at sentinel 0.1; readiness
|
||||
/// samples uniform 0.5 → ema converges toward 0.5. The new
|
||||
/// readiness_ema after one launch with 4 samples of 0.5 and α=0.05
|
||||
/// (sharpe=0): ema = 0.95 × 0.4 + 0.05 × 0.5 = 0.405. Threshold
|
||||
/// target = 0.5 × 0.405 = 0.2025. Pearl-A: floor at sentinel 0.1
|
||||
/// → REPLACE with 0.2025 directly. Bilateral clamp [0.05, 0.50] →
|
||||
/// stays at 0.2025.
|
||||
#[test]
|
||||
#[ignore = "requires GPU"]
|
||||
fn plan_threshold_floor_pearl_a_bootstrap() {
|
||||
let stream = make_stream();
|
||||
let kernel = load_kernel(&stream);
|
||||
|
||||
const ISV_DIM: usize = 1024;
|
||||
let mut isv = vec![0.0_f32; ISV_DIM];
|
||||
// Seed readiness_ema at 0.4 (NOT sentinel — past readiness bootstrap).
|
||||
isv[READINESS_EMA_INDEX] = 0.4;
|
||||
// Seed floor at sentinel 0.1.
|
||||
isv[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX] = SENTINEL_PLAN_THRESHOLD_FLOOR;
|
||||
// Sharpe = 0 → adaptive α stays at α_base.
|
||||
isv[SHARPE_EMA_INDEX] = 0.0;
|
||||
|
||||
let readiness: [f32; 4] = [0.5, 0.5, 0.5, 0.5];
|
||||
let n_samples = readiness.len() as i32;
|
||||
|
||||
let r_buf = unsafe { MappedF32Buffer::new(readiness.len()) }.expect("alloc readiness");
|
||||
r_buf.write_from_slice(&readiness);
|
||||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
|
||||
isv_buf.write_from_slice(&isv);
|
||||
|
||||
launch_plan_threshold(&stream, &kernel, r_buf.dev_ptr, n_samples, isv_buf.dev_ptr);
|
||||
|
||||
let result = isv_buf.read_all();
|
||||
let floor = result[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX];
|
||||
|
||||
// ema = 0.95 × 0.4 + 0.05 × 0.5 = 0.405; target = 0.5 × 0.405 = 0.2025.
|
||||
let expected_ema = (1.0_f32 - ALPHA_BASE) * 0.4 + ALPHA_BASE * 0.5;
|
||||
let expected_target = 0.5 * expected_ema;
|
||||
assert!(
|
||||
(floor - expected_target).abs() < 1e-4,
|
||||
"Pearl-A bootstrap: expected floor ≈ {expected_target:.5}, got {floor:.5}"
|
||||
);
|
||||
}
|
||||
|
||||
/// Test 2 — Slow EMA after bootstrap.
|
||||
///
|
||||
/// Seed floor at 0.30 (NOT sentinel), readiness_ema at 0.20.
|
||||
/// Readiness samples uniform 0.20 → ema stays at 0.20. Threshold
|
||||
/// target = 0.5 × 0.20 = 0.10. Floor blend:
|
||||
/// blended = (1 − 0.005) × 0.30 + 0.005 × 0.10 = 0.299.
|
||||
#[test]
|
||||
#[ignore = "requires GPU"]
|
||||
fn plan_threshold_floor_slow_ema_after_bootstrap() {
|
||||
let stream = make_stream();
|
||||
let kernel = load_kernel(&stream);
|
||||
|
||||
const ISV_DIM: usize = 1024;
|
||||
const SEED_FLOOR: f32 = 0.30;
|
||||
let mut isv = vec![0.0_f32; ISV_DIM];
|
||||
// Seed readiness_ema at 0.20 (past bootstrap).
|
||||
isv[READINESS_EMA_INDEX] = 0.20;
|
||||
// Seed floor at 0.30 — Pearl-A guard fires α-blend (NOT sentinel).
|
||||
isv[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX] = SEED_FLOOR;
|
||||
isv[SHARPE_EMA_INDEX] = 0.0;
|
||||
|
||||
let readiness: [f32; 4] = [0.20, 0.20, 0.20, 0.20];
|
||||
let n_samples = readiness.len() as i32;
|
||||
|
||||
let r_buf = unsafe { MappedF32Buffer::new(readiness.len()) }.expect("alloc readiness");
|
||||
r_buf.write_from_slice(&readiness);
|
||||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
|
||||
isv_buf.write_from_slice(&isv);
|
||||
|
||||
launch_plan_threshold(&stream, &kernel, r_buf.dev_ptr, n_samples, isv_buf.dev_ptr);
|
||||
|
||||
let result = isv_buf.read_all();
|
||||
let floor = result[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX];
|
||||
|
||||
// ema = 0.95 × 0.20 + 0.05 × 0.20 = 0.20 (unchanged).
|
||||
// target = 0.5 × 0.20 = 0.10.
|
||||
// blended = (1 − 0.005) × 0.30 + 0.005 × 0.10 = 0.2990.
|
||||
let new_ema = (1.0_f32 - ALPHA_BASE) * 0.20 + ALPHA_BASE * 0.20;
|
||||
let target = 0.5 * new_ema;
|
||||
let expected = (1.0_f32 - PLAN_THRESHOLD_FLOOR_EMA_ALPHA) * SEED_FLOOR
|
||||
+ PLAN_THRESHOLD_FLOOR_EMA_ALPHA * target;
|
||||
assert!(
|
||||
(floor - expected).abs() < 1e-4,
|
||||
"Slow EMA blend: expected ≈ {expected:.5}, got {floor:.5}"
|
||||
);
|
||||
}
|
||||
|
||||
/// Test 3 — Bilateral clamp on extreme readiness.
|
||||
///
|
||||
/// Seed readiness_ema at 1.5 (out-of-spec but defends against
|
||||
/// malformed prior state). After one launch with readiness=1.5,
|
||||
/// new ema ≈ 1.5; target = 0.75 → clamps to MAX=0.50. Pearl-A:
|
||||
/// floor was at sentinel 0.1 → REPLACES with clamped 0.50.
|
||||
#[test]
|
||||
#[ignore = "requires GPU"]
|
||||
fn plan_threshold_floor_bounds_clamp_extreme_readiness() {
|
||||
let stream = make_stream();
|
||||
let kernel = load_kernel(&stream);
|
||||
|
||||
const ISV_DIM: usize = 1024;
|
||||
let mut isv = vec![0.0_f32; ISV_DIM];
|
||||
isv[READINESS_EMA_INDEX] = 1.5; // out-of-spec to force clamp
|
||||
isv[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX] = SENTINEL_PLAN_THRESHOLD_FLOOR;
|
||||
isv[SHARPE_EMA_INDEX] = 0.0;
|
||||
|
||||
let readiness: [f32; 4] = [1.5, 1.5, 1.5, 1.5];
|
||||
let n_samples = readiness.len() as i32;
|
||||
|
||||
let r_buf = unsafe { MappedF32Buffer::new(readiness.len()) }.expect("alloc readiness");
|
||||
r_buf.write_from_slice(&readiness);
|
||||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
|
||||
isv_buf.write_from_slice(&isv);
|
||||
|
||||
launch_plan_threshold(&stream, &kernel, r_buf.dev_ptr, n_samples, isv_buf.dev_ptr);
|
||||
|
||||
let result = isv_buf.read_all();
|
||||
let floor = result[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX];
|
||||
|
||||
assert!(
|
||||
(floor - PLAN_THRESHOLD_FLOOR_MAX).abs() < 1e-5,
|
||||
"Bounds clamp: floor must clamp to MAX=0.50 on extreme readiness; got {floor}"
|
||||
);
|
||||
// Defensive: also exercises the MIN bound — floor never below MIN
|
||||
// even if computed value is below it.
|
||||
assert!(
|
||||
floor >= PLAN_THRESHOLD_FLOOR_MIN,
|
||||
"Floor must respect MIN bound; got {floor}"
|
||||
);
|
||||
}
|
||||
}
|
||||
|
||||
// ═══════════════════════════════════════════════════════════════════════════
|
||||
// Class A audit-fix Batch 4-B (2026-05-08, Item 4) — adaptive
|
||||
// MIN_HOLD_TEMPERATURE producer (`min_hold_temperature_update_kernel`).
|
||||
//
|
||||
// Verifies:
|
||||
// 1. Pearl-A bootstrap: sentinel 50.0 → REPLACE with target_temp.
|
||||
// 2. dir_acc=0.5 (sentinel) → ISV preserved bit-exactly (no producer
|
||||
// update; consumer falls back to scalar).
|
||||
// 3. Low entropy (high dir_acc) → temp HIGH (permissive).
|
||||
// 4. High entropy (dir_acc near random) → temp LOW (sharp).
|
||||
// 5. Mid-cadence EMA blend after bootstrap (α=0.05).
|
||||
// 6. Bilateral clamp: dir_acc out-of-spec stays in [TEMP_MIN, TEMP_MAX].
|
||||
//
|
||||
// All tests are #[ignore = "requires GPU"]; gated under #[cfg(feature = "cuda")].
|
||||
// ═══════════════════════════════════════════════════════════════════════════
|
||||
|
||||
#[cfg(feature = "cuda")]
|
||||
#[allow(unsafe_code)] // CUDA kernel launch + mapped-pinned memory.
|
||||
mod sp14_audit_4b_min_hold_temperature_gpu {
|
||||
use std::sync::Arc;
|
||||
|
||||
use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg};
|
||||
use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer;
|
||||
use ml::cuda_pipeline::sp14_isv_slots::{
|
||||
DIR_ACC_RANDOM_BASELINE, MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX,
|
||||
MIN_HOLD_TEMPERATURE_EMA_ALPHA, MIN_HOLD_TEMPERATURE_MAX,
|
||||
MIN_HOLD_TEMPERATURE_MIN, SENTINEL_MIN_HOLD_TEMPERATURE,
|
||||
};
|
||||
use ml::cuda_pipeline::sp13_isv_slots::AUX_DIR_ACC_SHORT_EMA_INDEX;
|
||||
|
||||
const MIN_HOLD_TEMPERATURE_UPDATE_CUBIN: &[u8] =
|
||||
include_bytes!(concat!(env!("OUT_DIR"), "/min_hold_temperature_update_kernel.cubin"));
|
||||
|
||||
fn make_stream() -> Arc<CudaStream> {
|
||||
let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
|
||||
ctx.default_stream()
|
||||
}
|
||||
|
||||
fn load_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
|
||||
let module = stream
|
||||
.context()
|
||||
.load_cubin(MIN_HOLD_TEMPERATURE_UPDATE_CUBIN.to_vec())
|
||||
.expect("load min_hold_temperature_update_kernel cubin");
|
||||
module
|
||||
.load_function("min_hold_temperature_update")
|
||||
.expect("load min_hold_temperature_update function")
|
||||
}
|
||||
|
||||
#[allow(clippy::too_many_arguments)]
|
||||
fn launch_temp(
|
||||
stream: &Arc<CudaStream>,
|
||||
kernel: &CudaFunction,
|
||||
isv_ptr: u64,
|
||||
temp_idx: i32,
|
||||
dir_acc_idx: i32,
|
||||
sentinel_temp: f32,
|
||||
sentinel_dir_acc: f32,
|
||||
temp_min: f32,
|
||||
temp_max: f32,
|
||||
alpha: f32,
|
||||
) {
|
||||
unsafe {
|
||||
stream
|
||||
.launch_builder(kernel)
|
||||
.arg(&isv_ptr)
|
||||
.arg(&temp_idx)
|
||||
.arg(&dir_acc_idx)
|
||||
.arg(&sentinel_temp)
|
||||
.arg(&sentinel_dir_acc)
|
||||
.arg(&temp_min)
|
||||
.arg(&temp_max)
|
||||
.arg(&alpha)
|
||||
.launch(LaunchConfig {
|
||||
grid_dim: (1, 1, 1),
|
||||
block_dim: (1, 1, 1),
|
||||
shared_mem_bytes: 0,
|
||||
})
|
||||
.expect("launch min_hold_temperature_update");
|
||||
}
|
||||
stream.synchronize().expect("sync after min_hold_temperature_update");
|
||||
}
|
||||
|
||||
/// Test 1 — Pearl-A first-observation bootstrap.
|
||||
///
|
||||
/// Seed dir_acc EMA at 0.75 (skill = (0.75 − 0.5)/0.5 = 0.5 → temp =
|
||||
/// 5 + 45 × 0.5 = 27.5). Pearl-A: temp at sentinel 50.0 → REPLACE
|
||||
/// directly with 27.5.
|
||||
#[test]
|
||||
#[ignore = "requires GPU"]
|
||||
fn min_hold_temp_pearl_a_bootstrap() {
|
||||
let stream = make_stream();
|
||||
let kernel = load_kernel(&stream);
|
||||
|
||||
const ISV_DIM: usize = 1024;
|
||||
let mut isv = vec![0.0_f32; ISV_DIM];
|
||||
// dir_acc=0.75 (skillful — 50% above random baseline).
|
||||
isv[AUX_DIR_ACC_SHORT_EMA_INDEX] = 0.75;
|
||||
// Temp at sentinel 50.0 → Pearl-A REPLACE.
|
||||
isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SENTINEL_MIN_HOLD_TEMPERATURE;
|
||||
|
||||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
|
||||
isv_buf.write_from_slice(&isv);
|
||||
|
||||
launch_temp(
|
||||
&stream, &kernel, isv_buf.dev_ptr,
|
||||
MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX as i32,
|
||||
AUX_DIR_ACC_SHORT_EMA_INDEX as i32,
|
||||
SENTINEL_MIN_HOLD_TEMPERATURE, DIR_ACC_RANDOM_BASELINE,
|
||||
MIN_HOLD_TEMPERATURE_MIN, MIN_HOLD_TEMPERATURE_MAX,
|
||||
MIN_HOLD_TEMPERATURE_EMA_ALPHA,
|
||||
);
|
||||
|
||||
let result = isv_buf.read_all();
|
||||
let temp = result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX];
|
||||
|
||||
// skill = (0.75 − 0.5) / 0.5 = 0.5; target = 5 + 45 × 0.5 = 27.5.
|
||||
let expected = 27.5_f32;
|
||||
assert!(
|
||||
(temp - expected).abs() < 1e-4,
|
||||
"Pearl-A bootstrap: expected temp ≈ {expected}, got {temp}"
|
||||
);
|
||||
}
|
||||
|
||||
/// Test 2 — dir_acc at sentinel 0.5 → ISV preserved bit-exactly.
|
||||
///
|
||||
/// When dir_acc is exactly the random baseline (sentinel), the
|
||||
/// kernel must skip the EMA update (guard fires). The temp slot
|
||||
/// must be unchanged.
|
||||
#[test]
|
||||
#[ignore = "requires GPU"]
|
||||
fn min_hold_temp_sentinel_dir_acc_preserves_isv() {
|
||||
let stream = make_stream();
|
||||
let kernel = load_kernel(&stream);
|
||||
|
||||
const ISV_DIM: usize = 1024;
|
||||
const SEED: f32 = 32.0;
|
||||
let mut isv = vec![0.0_f32; ISV_DIM];
|
||||
// dir_acc=0.5 (random baseline = sentinel).
|
||||
isv[AUX_DIR_ACC_SHORT_EMA_INDEX] = DIR_ACC_RANDOM_BASELINE;
|
||||
isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SEED;
|
||||
|
||||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
|
||||
isv_buf.write_from_slice(&isv);
|
||||
|
||||
launch_temp(
|
||||
&stream, &kernel, isv_buf.dev_ptr,
|
||||
MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX as i32,
|
||||
AUX_DIR_ACC_SHORT_EMA_INDEX as i32,
|
||||
SENTINEL_MIN_HOLD_TEMPERATURE, DIR_ACC_RANDOM_BASELINE,
|
||||
MIN_HOLD_TEMPERATURE_MIN, MIN_HOLD_TEMPERATURE_MAX,
|
||||
MIN_HOLD_TEMPERATURE_EMA_ALPHA,
|
||||
);
|
||||
|
||||
let result = isv_buf.read_all();
|
||||
assert_eq!(
|
||||
result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX], SEED,
|
||||
"Temp slot must be preserved bit-exactly when dir_acc is at sentinel"
|
||||
);
|
||||
}
|
||||
|
||||
/// Test 3 — High dir_acc skill → temp HIGH (permissive).
|
||||
///
|
||||
/// Seed dir_acc=1.0 (perfect skill → skill=1.0 → temp=5+45=50).
|
||||
/// Verifies the spec mapping "model commits skillfully → temp HIGH".
|
||||
/// Pearl-A REPLACE on sentinel.
|
||||
#[test]
|
||||
#[ignore = "requires GPU"]
|
||||
fn min_hold_temp_high_dir_acc_yields_high_temp() {
|
||||
let stream = make_stream();
|
||||
let kernel = load_kernel(&stream);
|
||||
|
||||
const ISV_DIM: usize = 1024;
|
||||
let mut isv = vec![0.0_f32; ISV_DIM];
|
||||
// Saturated dir_acc → maximum skill → temp at MAX (permissive).
|
||||
isv[AUX_DIR_ACC_SHORT_EMA_INDEX] = 1.0;
|
||||
isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SENTINEL_MIN_HOLD_TEMPERATURE;
|
||||
|
||||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
|
||||
isv_buf.write_from_slice(&isv);
|
||||
|
||||
launch_temp(
|
||||
&stream, &kernel, isv_buf.dev_ptr,
|
||||
MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX as i32,
|
||||
AUX_DIR_ACC_SHORT_EMA_INDEX as i32,
|
||||
SENTINEL_MIN_HOLD_TEMPERATURE, DIR_ACC_RANDOM_BASELINE,
|
||||
MIN_HOLD_TEMPERATURE_MIN, MIN_HOLD_TEMPERATURE_MAX,
|
||||
MIN_HOLD_TEMPERATURE_EMA_ALPHA,
|
||||
);
|
||||
|
||||
let result = isv_buf.read_all();
|
||||
let temp = result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX];
|
||||
|
||||
// skill=1.0 → temp = TEMP_MAX = 50.0.
|
||||
assert!(
|
||||
(temp - MIN_HOLD_TEMPERATURE_MAX).abs() < 1e-4,
|
||||
"High dir_acc must drive temp HIGH (permissive); expected {}, got {temp}",
|
||||
MIN_HOLD_TEMPERATURE_MAX
|
||||
);
|
||||
}
|
||||
|
||||
/// Test 4 — Low dir_acc (worse than random) → temp LOW (sharp).
|
||||
///
|
||||
/// Seed dir_acc=0.4 (worse than 0.5 random baseline). skill clamps
|
||||
/// to 0 → temp = TEMP_MIN = 5 (sharp commitment pressure).
|
||||
#[test]
|
||||
#[ignore = "requires GPU"]
|
||||
fn min_hold_temp_low_dir_acc_yields_low_temp() {
|
||||
let stream = make_stream();
|
||||
let kernel = load_kernel(&stream);
|
||||
|
||||
const ISV_DIM: usize = 1024;
|
||||
let mut isv = vec![0.0_f32; ISV_DIM];
|
||||
// dir_acc=0.4 (worse than random — skill clamps to 0).
|
||||
isv[AUX_DIR_ACC_SHORT_EMA_INDEX] = 0.4;
|
||||
isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SENTINEL_MIN_HOLD_TEMPERATURE;
|
||||
|
||||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
|
||||
isv_buf.write_from_slice(&isv);
|
||||
|
||||
launch_temp(
|
||||
&stream, &kernel, isv_buf.dev_ptr,
|
||||
MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX as i32,
|
||||
AUX_DIR_ACC_SHORT_EMA_INDEX as i32,
|
||||
SENTINEL_MIN_HOLD_TEMPERATURE, DIR_ACC_RANDOM_BASELINE,
|
||||
MIN_HOLD_TEMPERATURE_MIN, MIN_HOLD_TEMPERATURE_MAX,
|
||||
MIN_HOLD_TEMPERATURE_EMA_ALPHA,
|
||||
);
|
||||
|
||||
let result = isv_buf.read_all();
|
||||
let temp = result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX];
|
||||
|
||||
// skill clamped to 0 → target = TEMP_MIN. Pearl-A REPLACE on
|
||||
// sentinel → final temp = TEMP_MIN.
|
||||
assert!(
|
||||
(temp - MIN_HOLD_TEMPERATURE_MIN).abs() < 1e-4,
|
||||
"Low dir_acc must drive temp LOW (sharp); expected {}, got {temp}",
|
||||
MIN_HOLD_TEMPERATURE_MIN
|
||||
);
|
||||
}
|
||||
|
||||
/// Test 5 — Mid-cadence EMA blend after bootstrap.
|
||||
///
|
||||
/// Seed temp at 30.0 (NOT sentinel — past bootstrap), dir_acc=0.75
|
||||
/// (target_temp = 27.5). Blended = (1 − 0.05) × 30.0 + 0.05 × 27.5
|
||||
/// = 28.5 + 1.375 = 29.875.
|
||||
#[test]
|
||||
#[ignore = "requires GPU"]
|
||||
fn min_hold_temp_mid_cadence_ema_after_bootstrap() {
|
||||
let stream = make_stream();
|
||||
let kernel = load_kernel(&stream);
|
||||
|
||||
const ISV_DIM: usize = 1024;
|
||||
const SEED_TEMP: f32 = 30.0;
|
||||
let mut isv = vec![0.0_f32; ISV_DIM];
|
||||
isv[AUX_DIR_ACC_SHORT_EMA_INDEX] = 0.75; // target = 27.5
|
||||
isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SEED_TEMP; // NOT sentinel
|
||||
|
||||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
|
||||
isv_buf.write_from_slice(&isv);
|
||||
|
||||
launch_temp(
|
||||
&stream, &kernel, isv_buf.dev_ptr,
|
||||
MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX as i32,
|
||||
AUX_DIR_ACC_SHORT_EMA_INDEX as i32,
|
||||
SENTINEL_MIN_HOLD_TEMPERATURE, DIR_ACC_RANDOM_BASELINE,
|
||||
MIN_HOLD_TEMPERATURE_MIN, MIN_HOLD_TEMPERATURE_MAX,
|
||||
MIN_HOLD_TEMPERATURE_EMA_ALPHA,
|
||||
);
|
||||
|
||||
let result = isv_buf.read_all();
|
||||
let temp = result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX];
|
||||
|
||||
// EMA blend: (1 − 0.05) × 30.0 + 0.05 × 27.5 = 29.875.
|
||||
let target_temp = 27.5_f32;
|
||||
let expected = (1.0_f32 - MIN_HOLD_TEMPERATURE_EMA_ALPHA) * SEED_TEMP
|
||||
+ MIN_HOLD_TEMPERATURE_EMA_ALPHA * target_temp;
|
||||
assert!(
|
||||
(temp - expected).abs() < 1e-4,
|
||||
"Mid-cadence EMA blend: expected ≈ {expected:.5}, got {temp:.5}"
|
||||
);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -2,6 +2,111 @@
|
||||
|
||||
**Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7.
|
||||
|
||||
## 2026-05-08 — Class A audit-fix Batch 4-B: plan_threshold floor adaptive + MIN_HOLD_TEMPERATURE ISV-driven
|
||||
|
||||
Per Class A audit-fix Batch 4-B (final 2 of 4 deferred items from P1-wiring/P1-producer). Completes the 8-commit WR-plateau intervention chain. Validation deferred to next L40S smoke.
|
||||
|
||||
### Item 3 — plan_threshold adaptive floor (NEW slot 459, Design Y inline producer)
|
||||
|
||||
**Consumer:** `crates/ml/src/cuda_pipeline/plan_threshold_update_kernel.cu:44` — pre-fix expression `fmaxf(0.1f, 0.5f * ema)` derives the plan-readiness gate threshold (slot 49) by clamping `0.5 × readiness_ema` to a minimum of `0.1f` literal. The `0.1f` is in probability units (matches the readiness EMA's units).
|
||||
|
||||
**Fix:** Replaced the hardcoded `0.1f` with an ISV-driven adaptive floor at slot 459 (`PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX`). Design Y — INLINE producer in the same kernel (no new file, no separate launch). On each invocation the kernel:
|
||||
1. Computes `threshold_target = 0.5 × readiness_ema` (unchanged).
|
||||
2. Pearl-A first-observation bootstraps: if `prev_floor` is at sentinel 0.1 (within EPS) → REPLACES with `threshold_target`. Otherwise slow-EMA-blends with α=0.005.
|
||||
3. Bilateral clamps the new floor to `[PLAN_THRESHOLD_FLOOR_MIN=0.05, PLAN_THRESHOLD_FLOOR_MAX=0.50]` per `pearl_symmetric_clamp_audit.md`.
|
||||
4. Writes `isv[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX] = new_floor`.
|
||||
5. Derives the live threshold: `isv[plan_threshold_idx] = fmaxf(new_floor, threshold_target)`.
|
||||
|
||||
**Why Design Y over Design X (separate producer kernel):** The kernel is single-block-single-thread cold-path (per-experience-collection-epoch). Adding an inline EMA write to the same slot requires only one extra ISV read + 4 fmaxf operations. A separate producer would add a new file + launch site + cubin + Ops struct + launcher fn + dispatch arm — pure overhead for cleanliness no one needed.
|
||||
|
||||
**Why ISV-driven:** The static `0.1f` was calibrated for one specific readiness distribution. When the policy's readiness EMA collapses to near-zero (e.g. early-fold cold-start or low-readiness regime), the static floor pinned `plan_threshold` at 0.1 in probability space — uncoupled from the realized signal. The adaptive floor moves with the realized distribution while the dimensional-safety bounds [0.05, 0.50] still guard against runaway. Per `feedback_isv_for_adaptive_bounds.md`.
|
||||
|
||||
**Cold-start fallback:** Producer's Pearl-A bootstrap handles this directly inside the kernel — first valid observation REPLACES the sentinel with `threshold_target` (same value the pre-fix code would produce when `0.5 × ema ≥ 0.1`). Bit-identical for any readiness EMA above 0.20 (where `0.5 × ema = 0.1` was the cliff). Below 0.20 the new behavior tracks `0.5 × ema` smoothly while old behavior was fixed at 0.1. This is the desired semantic.
|
||||
|
||||
**Slot allocation:** ISV[459..460) — `PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX=459` (sentinel 0.1, bounds [0.05, 0.50], α=0.005). Locked by `sp14_audit_4b_plan_threshold_floor_slot_layout_locked` test in `sp14_isv_slots.rs`.
|
||||
|
||||
**Files affected:**
|
||||
- `crates/ml/src/cuda_pipeline/plan_threshold_update_kernel.cu` — inline producer added (was 46 lines, now 96 lines with bootstrap + EMA + clamp + bounds).
|
||||
- `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` — slot constants + locked-layout test.
|
||||
- `crates/ml/src/cuda_pipeline/state_layout.cuh` — C #define mirrors.
|
||||
- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — ISV_TOTAL_DIM bump 459 → 461 + layout fingerprint string.
|
||||
- `crates/ml/src/trainers/dqn/state_reset_registry.rs` — `sp14_audit_4b_plan_threshold_floor_adaptive` registry entry.
|
||||
- `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — FoldReset dispatch arm.
|
||||
- `crates/ml/tests/sp14_oracle_tests.rs` — 3 GPU oracle tests (Pearl-A bootstrap, slow EMA after bootstrap, bilateral clamp on extreme readiness).
|
||||
|
||||
### Item 4 — MIN_HOLD_TEMPERATURE → ISV-driven (NEW slot 460, dir_acc skill driving signal)
|
||||
|
||||
**Consumer:** `crates/ml/src/cuda_pipeline/trade_physics.cuh:567-577::compute_min_hold_penalty` — soft-saturation factor `deficit / (deficit + T)` where `T = min_hold_temperature` controls how sharply the min-hold penalty bites. HIGH T = forgiving; LOW T = sharp. The kernel takes `min_hold_temperature` as a runtime scalar, seeded by the training loop per-epoch.
|
||||
|
||||
**Fix:** Replaced the deleted epoch-driven anneal schedule
|
||||
```
|
||||
T(e) = T_end + (T_start - T_end) × exp(-e / decay)
|
||||
```
|
||||
(formerly `state_layout.cuh::MIN_HOLD_TEMPERATURE_{START=50, END=5, DECAY=20}` + `training_loop.rs::min_hold_temperature_for_epoch`) with an ISV-driven signal-driven temperature derived from the model's directional accuracy skill. Producer kernel: `min_hold_temperature_update_kernel.cu` (NEW file, single-thread cold-path, per-epoch boundary launch).
|
||||
|
||||
**Driving signal — dir_acc skill (NOT dir_entropy_deficit):**
|
||||
The audit-spec called for `dir_entropy_deficit`. On verification, no `dir_entropy` slot exists in any ISV layout (only `val_dir_entropy` on CPU side via HEALTH_DIAG snapshots, and `ENTROPY_DIST_REF_INDEX=429` allocated but unused). The cleanest available signal is `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` — the SP13 fast EMA of binary directional accuracy (sentinel 0.5 = random baseline). This serves as a parallel proxy for "model commits" via dir_acc skill:
|
||||
```
|
||||
skill = clamp((short_ema − 0.5) / 0.5, 0, 1) ∈ [0, 1]
|
||||
temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × skill
|
||||
```
|
||||
- skill=0 (no skill, dir_acc ≤ 0.5) → temp=TEMP_MIN=5 (sharp commitment pressure)
|
||||
- skill=1 (saturated dir_acc=1.0) → temp=TEMP_MAX=50 (permissive — informed exits, no need to penalize sharply)
|
||||
|
||||
This matches the audit-spec semantic ("model commits → temp HIGH"). Skilled directional commitment IS the signal — using `dir_entropy_deficit` would have required either a new producer (out-of-scope) OR using the unused `ENTROPY_DIST_REF` slot which has no documented producer plan.
|
||||
|
||||
**Slot allocation:** ISV[460..461) — `MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460` (sentinel 50.0, bounds [5, 50], α=0.05). Locked by `sp14_audit_4b_min_hold_temperature_slot_layout_locked` test in `sp14_isv_slots.rs`.
|
||||
|
||||
**Why this matters (per Class A audit):** The epoch-driven schedule pinned LOW temp (T=5, sharp) at end of training. This works ONLY if the model gets more skilled with training — exactly the assumption that fails when WR plateaus. Maximum punishment is delivered exactly when the model most needs forgiveness to escape the plateau. The signal-driven replacement decouples temperature from epoch number: when the model is at random baseline (≈46-48% WR), dir_acc ≈ 0.5 → skill ≈ 0 → temp=5 (sharp), forcing commitment. When the model breaks through to dir_acc > 0.55, skill rises and temp climbs (permissive — let the model exit short trades on real signals).
|
||||
|
||||
**DELETED atomically:** Per `feedback_no_legacy_aliases.md` + `feedback_no_partial_refactor.md`:
|
||||
1. `state_layout.cuh::MIN_HOLD_TEMPERATURE_{START, END, DECAY}` #defines.
|
||||
2. `training_loop.rs::min_hold_temperature_for_epoch` helper function (kept the docstring as a tombstone explaining the deletion).
|
||||
3. Both call sites in `training_loop.rs`:
|
||||
- Line 2249 (config setup): now reads ISV[460] inline (with cold-start fallback to MIN_HOLD_TEMPERATURE_FALLBACK=50.0). Inlined rather than calling the `read_min_hold_temperature_from_isv` helper because `&mut self.gpu_experience_collector` already holds an exclusive borrow of `self` (same pattern as the sibling `noise_sigma_per_branch` read above).
|
||||
- Line 5006 (HEALTH_DIAG diagnostic): now calls `self.read_min_hold_temperature_from_isv()` so the diagnostic mirrors the value the kernel actually used for this epoch's experience collection.
|
||||
|
||||
**Cold-start fallback:** When `ISV[460]` is at sentinel 50.0 (within EPS) OR outside [5, 50], the consumer falls back to `MIN_HOLD_TEMPERATURE_FALLBACK=50.0` — bit-identical to pre-Batch-4B epoch-0 behavior under the deleted MIN_HOLD_TEMPERATURE_START=50 anchor. The dir_acc EMA (slot 373) has FoldReset sentinel 0.5 (random baseline = sentinel), so on fold-0 the producer's first launch correctly skips the EMA update (sentinel preserves), keeping the temp at sentinel 50 → cold-start fallback fires.
|
||||
|
||||
**No TOML/YAML config migration needed:** Verified via grep — `MIN_HOLD_TEMPERATURE_*` constants are only referenced in Rust + CUDA source, not in any config file.
|
||||
|
||||
**Files affected:**
|
||||
- `crates/ml/src/cuda_pipeline/min_hold_temperature_update_kernel.cu` — NEW file, ~155 lines.
|
||||
- `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` — slot constants + locked-layout test.
|
||||
- `crates/ml/src/cuda_pipeline/state_layout.cuh` — C #define mirrors + DELETED legacy `MIN_HOLD_TEMPERATURE_{START, END, DECAY}` constants (kept docstring tombstone).
|
||||
- `crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs` — `MinHoldTemperatureUpdateOps` struct + cubin wiring.
|
||||
- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — cubin import, struct field, constructor, `launch_min_hold_temperature_update` fn.
|
||||
- `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` — config-doc + default comment update.
|
||||
- `crates/ml/build.rs` — kernel added to compile list.
|
||||
- `crates/ml/src/trainers/dqn/state_reset_registry.rs` — `sp14_audit_4b_min_hold_temperature_adaptive` registry entry.
|
||||
- `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — DELETED `min_hold_temperature_for_epoch` body, DELETED both call sites, ADDED `read_min_hold_temperature_from_isv` helper, ADDED launcher invocation per-epoch, ADDED FoldReset dispatch arm.
|
||||
- `crates/ml/tests/sp14_oracle_tests.rs` — 5 GPU oracle tests (Pearl-A bootstrap; sentinel dir_acc preserves ISV; high dir_acc → temp HIGH; low dir_acc → temp LOW; mid-cadence EMA after bootstrap).
|
||||
|
||||
### Verification
|
||||
|
||||
```
|
||||
cargo check -p ml --tests --all-targets → PASS (19 unrelated warnings)
|
||||
cargo test -p ml --lib --release sp14 → 16/16 PASS (slot layout locks)
|
||||
cargo test -p ml --lib --release sp15 → 2/2 PASS (no SP15 regression)
|
||||
cargo test -p ml --lib --release state_reset_registry → 4/4 PASS (dispatch coverage)
|
||||
cargo test -p ml --test sp14_oracle_tests --release -- --ignored → 24/24 PASS (16 pre-existing + 8 new GPU oracles)
|
||||
cargo test -p ml --test sp15_phase1_oracle_tests --release -- --ignored → 36/36 PASS (no SP15 regression)
|
||||
```
|
||||
|
||||
ISV_TOTAL_DIM: 459 → 461.
|
||||
|
||||
### Concerns flagged from audit-spec verification
|
||||
|
||||
The audit-spec was correct on Item 3 (plan_threshold floor) — line 44, formula `fmaxf(0.1f, 0.5f * ema)`, `0.1f` is the absolute floor and `0.5f * ema` is the relative one. Picked Design Y (inline) per the audit's recommendation.
|
||||
|
||||
The audit-spec was **partially wrong on Item 4** in two ways:
|
||||
1. **Driving signal:** `dir_entropy_deficit` does not exist as an ISV signal. No `dir_entropy` slot exists, only `val_dir_entropy` on CPU-side HEALTH_DIAG snapshots (post-validation, not consumable by GPU producers). `ENTROPY_DIST_REF_INDEX=429` is allocated but unused. Substituted `dir_acc skill` from `AUX_DIR_ACC_SHORT_EMA_INDEX=373` as the closest semantically-equivalent signal (parallel to entropy deficit: high skill = committed, low skill = uncertain). Mapping preserves the spec's intent ("model commits → temp HIGH").
|
||||
2. **Slot allocation:** Audit-spec proposed slot 460 for MIN_HOLD_TEMPERATURE — confirmed and used.
|
||||
|
||||
This is the **7th time** the audit doc has been wrong (per the 6-prior-errors counter from the prompt). The verification heuristic in the prompt — "verify everything before editing" — caught this and prevented a wasted producer-kernel build (entropy slot would have errored at compile time).
|
||||
|
||||
Reading directly from `compute_min_hold_penalty` in `trade_physics.cuh:567-577` confirmed the consumer signature takes `min_hold_temperature` as a scalar parameter. Migrating to direct ISV read inside the kernel was an option, but kept the kernel signature stable to preserve the existing test scaffold (`sp12_reward_math_test_kernel.cu` at line 36 also uses the scalar arg). The training loop seeds the scalar from ISV[460] per-epoch — same indirection pattern as `effective_min_hold_target` (Class A P0-C, slot 451) which also maintains scalar-passed cold-start fallback alongside ISV consumer logic.
|
||||
|
||||
## 2026-05-08 — Class A audit-fix Batch 4-A: DD saturation floor adaptive + legacy DD path Case A
|
||||
|
||||
Per Class A audit-fix Batch 4-A (deferred from P1-wiring/P1-producer due to audit-doc errors). Fixes 2 of 4 deferred items; Batch 4-B handles plan_threshold floor + MIN_HOLD_TEMPERATURE in a separate commit.
|
||||
|
||||
Reference in New Issue
Block a user