feat(dqn-v2): C.6 static ISV writes (Tasks 12, 15, 16) + slot pre-allocation
Lands the static-configuration branch of Plan 1 C.6: cql_alpha, conviction_floor (no-op if IQL_BRANCH_SCALE_FLOOR already serves), plan_threshold written to dedicated ISV slots at constructor. Consumer kernels (CQL, backtest_plan, experience) read from ISV instead of config fields / hardcoded literals. Also pre-allocates ISV slots for the 6 upcoming GPU-kernel tasks (atoms, gamma, kelly_cap, tau, epsilon outputs + CPU-born epoch inputs). Those slots start at 0; GPU kernels in follow-up commits fill them. New ISV slots: - EPOCH_IDX_INDEX=39, TOTAL_EPOCHS_INDEX=40 (CPU-born inputs) - EPSILON_EFF_INDEX=41, TAU_EFF_INDEX=42, GAMMA_EFF_INDEX=43, KELLY_CAP_EFF_INDEX=44 (GPU-written in follow-up tasks) - CQL_ALPHA_INDEX=45, PLAN_THRESHOLD_INDEX=46 (static config; this commit) - Task 15 confirmed no-op: IQL_BRANCH_SCALE_FLOOR_INDEX=36 already serves conviction-floor role (constructor + ISV read in iql kernel). Layout fingerprint auto-updated via seed-byte edits; fingerprint re-tail at [47..49). ISV_TOTAL_DIM 39 -> 49. GpuDqnTrainConfig gains total_epochs field; fused_training.rs passes hyperparams.epochs at construction; default 0 for smoke tests. write_isv_signal_at bound extended from ISV_DIM(23) to ISV_TOTAL_DIM(49) so tail slots are writable by CPU. cql_alpha consumer: compute_cql_logit_gradients reads base from ISV[CQL_ALPHA_INDEX] instead of config.cql_alpha; adaptive formula (base x health x (1-regime_stability)) unchanged. plan_threshold consumers: experience_kernels.cu (experience_state_gather, experience_action_select, experience_env_step) and backtest_plan_kernel.cu (backtest_plan_state_isv) read plan_thr from ISV[ISV_PLAN_THRESHOLD_IDX=46] via isv_signals pointer already present in both kernels; null-guard defaults to 0.5f for smoke-scale runs without ISV warmup. ISV_PLAN_THRESHOLD_IDX=46 defined in state_layout.cuh (included by both plan kernel files); value must match PLAN_THRESHOLD_INDEX in gpu_dqn_trainer.rs. StateResetRegistry entries added for all new slots: SchemaContract for TOTAL_EPOCHS/CQL_ALPHA/PLAN_THRESHOLD; FoldReset for EPOCH_IDX and GPU-written per-fold slots. No behavioural change: all threshold/base values remain at their prior defaults; consumers now read adaptive values from ISV instead of baked-in config or hardcoded literals. Plan 1 Tasks 12, 15, 16 + pre-allocation for 9, 10, 11, 13, 14. Spec §4.C.6 (2026-04-24 GPU-drives-CPU-reads revision). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -115,8 +115,9 @@ extern "C" __global__ void backtest_plan_state_isv(
|
||||
* prior direction. In practice trade flips are rare and the plan
|
||||
* gets re-populated on the NEXT Flat→Position transition after
|
||||
* the flip closes. ── */
|
||||
float plan_thr = (isv_signals != nullptr) ? isv_signals[ISV_PLAN_THRESHOLD_IDX] : 0.5f;
|
||||
float active_tgt = plan_state[w * 7 + 0];
|
||||
if (in_position && active_tgt < 0.5f) {
|
||||
if (in_position && active_tgt < plan_thr) {
|
||||
const float* pp = plan_params + w * 6;
|
||||
plan_state[w * 7 + 0] = pp[PLAN_PARAM_TARGET_BARS]; /* target_bars */
|
||||
plan_state[w * 7 + 1] = pp[PLAN_PARAM_PROFIT_TARGET]; /* profit_target */
|
||||
@@ -146,7 +147,7 @@ extern "C" __global__ void backtest_plan_state_isv(
|
||||
float pisv[6];
|
||||
|
||||
/* [0] plan progress [0, 2]: hold_time/target_bars */
|
||||
pisv[PLAN_ISV_PROGRESS] = (plan_tgt_bars > 0.5f)
|
||||
pisv[PLAN_ISV_PROGRESS] = (plan_tgt_bars > plan_thr)
|
||||
? fminf(hold_time / plan_tgt_bars, 2.0f) : 0.0f;
|
||||
|
||||
/* [1] PnL vs profit target [bounded 2.0] */
|
||||
@@ -161,11 +162,11 @@ extern "C" __global__ void backtest_plan_state_isv(
|
||||
pisv[PLAN_ISV_ENTRY_CONVICTION] = plan_conv_ent;
|
||||
|
||||
/* [4] Conviction drift: current MLP conviction - entry conviction */
|
||||
pisv[PLAN_ISV_CONVICTION_DRIFT] = (plan_tgt_bars > 0.5f)
|
||||
pisv[PLAN_ISV_CONVICTION_DRIFT] = (plan_tgt_bars > plan_thr)
|
||||
? (plan_params[w * 6 + PLAN_PARAM_CONVICTION] - plan_conv_ent) : 0.0f;
|
||||
|
||||
/* [5] Regime shift magnitude since entry */
|
||||
pisv[PLAN_ISV_REGIME_SHIFT] = (isv_signals != nullptr && plan_tgt_bars > 0.5f)
|
||||
pisv[PLAN_ISV_REGIME_SHIFT] = (isv_signals != nullptr && plan_tgt_bars > plan_thr)
|
||||
? fabsf(isv_signals[11] - entry_stab) : 0.0f;
|
||||
|
||||
#pragma unroll
|
||||
|
||||
@@ -593,7 +593,11 @@ extern "C" __global__ void experience_state_gather(
|
||||
float plan_stop = ps[PS_PLAN_STOP_LOSS];
|
||||
float plan_conv = ps[PS_PLAN_CONVICTION];
|
||||
|
||||
plan_isv[PLAN_ISV_PROGRESS] = (plan_tgt_bars > 0.5f)
|
||||
/* Plan activation threshold from ISV[ISV_PLAN_THRESHOLD_IDX] (Plan 1 Task 16).
|
||||
* Default 0.5 when ISV not yet available (smoke-scale runs). */
|
||||
float plan_thr_sg = (isv_signals_ptr != NULL) ? isv_signals_ptr[ISV_PLAN_THRESHOLD_IDX] : 0.5f;
|
||||
|
||||
plan_isv[PLAN_ISV_PROGRESS] = (plan_tgt_bars > plan_thr_sg)
|
||||
? fminf(f_hold_time / plan_tgt_bars, 2.0f) /* plan progress [0, 2] */
|
||||
: 0.0f;
|
||||
plan_isv[PLAN_ISV_PNL_VS_TARGET] = (plan_profit > 1e-6f)
|
||||
@@ -606,15 +610,15 @@ extern "C" __global__ void experience_state_gather(
|
||||
|
||||
/* Conviction drift — current conviction minus entry conviction.
|
||||
* Negative drift = model's thesis is weakening → exit signal.
|
||||
* Only meaningful when a plan is active (ps[PS_PLAN_TARGET_BARS] > 0.5). */
|
||||
plan_isv[PLAN_ISV_CONVICTION_DRIFT] = (plan_params_ptr != NULL && ps[PS_PLAN_TARGET_BARS] > 0.5f)
|
||||
* Only meaningful when a plan is active (ps[PS_PLAN_TARGET_BARS] > plan_thr_sg). */
|
||||
plan_isv[PLAN_ISV_CONVICTION_DRIFT] = (plan_params_ptr != NULL && ps[PS_PLAN_TARGET_BARS] > plan_thr_sg)
|
||||
? (plan_params_ptr[i * 6 + PLAN_PARAM_CONVICTION] - ps[PS_PLAN_CONVICTION]) /* conviction drift */
|
||||
: 0.0f;
|
||||
|
||||
/* Regime shift since entry — |regime_stability_now - regime_stability_at_entry|.
|
||||
* Large shift = regime changed mid-trade, plan may be invalid.
|
||||
* Entry regime stored in ps[PS_PLAN_ENTRY_REGIME] during plan activation. */
|
||||
plan_isv[PLAN_ISV_REGIME_SHIFT] = (isv_signals_ptr != NULL && ps[PS_PLAN_TARGET_BARS] > 0.5f)
|
||||
plan_isv[PLAN_ISV_REGIME_SHIFT] = (isv_signals_ptr != NULL && ps[PS_PLAN_TARGET_BARS] > plan_thr_sg)
|
||||
? fabsf(isv_signals_ptr[11] - ps[PS_PLAN_ENTRY_REGIME]) /* regime shift */
|
||||
: 0.0f;
|
||||
|
||||
@@ -1145,7 +1149,8 @@ extern "C" __global__ void experience_action_select(
|
||||
/* Plan direction lock: during active plan, force current direction */
|
||||
if (portfolio_states != NULL) {
|
||||
int ps_base_plan = i * PORTFOLIO_STRIDE;
|
||||
int has_plan_active = (portfolio_states[ps_base_plan + PS_PLAN_TARGET_BARS] > 0.5f);
|
||||
float plan_thr_as = (isv_signals_ptr != NULL) ? isv_signals_ptr[ISV_PLAN_THRESHOLD_IDX] : 0.5f;
|
||||
int has_plan_active = (portfolio_states[ps_base_plan + PS_PLAN_TARGET_BARS] > plan_thr_as);
|
||||
if (has_plan_active) {
|
||||
float cur_pos_plan = portfolio_states[ps_base_plan + PS_POSITION];
|
||||
if (cur_pos_plan > 0.001f) dir_idx = DIR_LONG;
|
||||
@@ -1771,16 +1776,20 @@ extern "C" __global__ void experience_env_step(
|
||||
|
||||
/* ── Trade Plan: readiness-gated activation + enforcement ──
|
||||
* Plan head is randomly initialized → garbage plans early in training.
|
||||
* Gate by model readiness: readiness < 0.5 → plan disabled (model trades freely).
|
||||
* readiness ≥ 0.5 → plan activates and enforces. Smooth transition.
|
||||
* Gate by model readiness: readiness < plan_thr → plan disabled (model trades freely).
|
||||
* readiness >= plan_thr → plan activates and enforces. Smooth transition.
|
||||
*
|
||||
* Phase 3 unified-env: also gated by exploration_scale > 0.5. Validation
|
||||
* Phase 3 unified-env: also gated by exploration_scale > plan_thr. Validation
|
||||
* runs with exploration_scale=0 → plan disabled → position is the raw
|
||||
* Kelly/margin-capped target, matching the backtest kernel. */
|
||||
* Kelly/margin-capped target, matching the backtest kernel.
|
||||
*
|
||||
* plan_thr reads from ISV[ISV_PLAN_THRESHOLD_IDX] (Plan 1 Task 16 / spec §4.C.6).
|
||||
* Default 0.5 when ISV not yet available (smoke-scale runs). */
|
||||
float plan_thr_env = (isv_signals_ptr != NULL) ? isv_signals_ptr[ISV_PLAN_THRESHOLD_IDX] : 0.5f;
|
||||
float readiness = (readiness_ptr != NULL) ? readiness_ptr[0] : 1.0f;
|
||||
int plan_enabled = (readiness >= 0.5f
|
||||
int plan_enabled = (readiness >= plan_thr_env
|
||||
&& plan_params_ptr != NULL
|
||||
&& exploration_scale >= 0.5f);
|
||||
&& exploration_scale >= plan_thr_env);
|
||||
|
||||
/* Plan activation on Flat → Positioned (only when model is ready) */
|
||||
int was_flat_plan = (fabsf(ps0_f) < 0.001f);
|
||||
@@ -1801,7 +1810,7 @@ extern "C" __global__ void experience_env_step(
|
||||
position *= fmaxf(ps[PS_PLAN_CONVICTION] * readiness, 0.1f);
|
||||
}
|
||||
|
||||
int has_plan = (ps[PS_PLAN_TARGET_BARS] > 0.5f);
|
||||
int has_plan = (ps[PS_PLAN_TARGET_BARS] > plan_thr_env);
|
||||
|
||||
/* Recursive plan revision (N17): living plan document.
|
||||
* When model is highly confident (readiness ≥ 0.7) and conviction changes
|
||||
@@ -2391,12 +2400,12 @@ extern "C" __global__ void experience_env_step(
|
||||
}
|
||||
|
||||
/* Plan conviction as reward scaling — ONLY when plan head is mature.
|
||||
* Before readiness ≥ 0.5, conviction defaults to 1.0 (identity) so the
|
||||
* Before readiness >= plan_thr_env, conviction defaults to 1.0 (identity) so the
|
||||
* reward signal flows undampened. Without this gate, random init conviction
|
||||
* (~0.1) kills the gradient signal and prevents the model from learning. */
|
||||
if (plan_params_ptr != NULL && fabsf(reward) > 1e-8f) {
|
||||
float readiness_val = (readiness_ptr != NULL) ? readiness_ptr[0] : 0.0f;
|
||||
float conviction = (readiness_val >= 0.5f)
|
||||
float conviction = (readiness_val >= plan_thr_env)
|
||||
? plan_params_ptr[i * 6 + PLAN_PARAM_CONVICTION] /* [0, 1] from plan head */
|
||||
: 1.0f; /* identity until plan matures */
|
||||
reward *= conviction;
|
||||
|
||||
@@ -164,17 +164,25 @@ const ISV_NETWORK_DIM: usize = 23;
|
||||
/// to ensure no (sample, branch) pair collapses to zero gradient
|
||||
/// contribution.
|
||||
///
|
||||
/// Slots [37..39) carry the layout fingerprint (ISV_LAYOUT_FINGERPRINT_LO_INDEX
|
||||
/// Slots [39..47) are the DQN v2 Plan 1 expansion (spec §4.C.6, 2026-04-24):
|
||||
/// [39] EPOCH_IDX_INDEX — CPU writes current epoch index at epoch boundary.
|
||||
/// [40] TOTAL_EPOCHS_INDEX — CPU writes total epoch count at constructor (static).
|
||||
/// [41] EPSILON_EFF_INDEX — GPU-written by epsilon adaptive kernel (follow-up task).
|
||||
/// [42] TAU_EFF_INDEX — GPU-written by tau adaptive kernel (follow-up task).
|
||||
/// [43] GAMMA_EFF_INDEX — GPU-written by gamma adaptive kernel (follow-up task).
|
||||
/// [44] KELLY_CAP_EFF_INDEX — GPU-written by kelly_cap adaptive kernel (follow-up task).
|
||||
/// [45] CQL_ALPHA_INDEX — constructor writes `config.cql_alpha`; CQL formula reads
|
||||
/// base from here instead of config field (Plan 1 Task 12).
|
||||
/// [46] PLAN_THRESHOLD_INDEX — constructor writes 0.5f; plan kernels read instead of
|
||||
/// hardcoded literal (Plan 1 Task 16).
|
||||
///
|
||||
/// Slots [47..49) carry the layout fingerprint (ISV_LAYOUT_FINGERPRINT_LO_INDEX
|
||||
/// and ISV_LAYOUT_FINGERPRINT_HI_INDEX). These are placed at the tail because
|
||||
/// slots [0] and [1] are actively written by `isv_signal_update` (Q-drift and
|
||||
/// gradient-norm EMA respectively). Written by the constructor; checked at
|
||||
/// checkpoint load. Fail-fast only — no migration path exists. See spec §4.A.2
|
||||
/// and `LAYOUT_FINGERPRINT_CURRENT` for the structural-hash design rationale.
|
||||
///
|
||||
/// Extending past 39 requires updating the pinned allocation in the
|
||||
/// constructor and any kernel that takes an ISV pointer expecting a
|
||||
/// specific length.
|
||||
const ISV_TOTAL_DIM: usize = 39;
|
||||
const ISV_TOTAL_DIM: usize = 49;
|
||||
/// Legacy alias preserved for call sites that haven't been audited for the
|
||||
/// network-vs-total split. New code should pick `ISV_NETWORK_DIM` (for weight
|
||||
/// tensor sizing) or `ISV_TOTAL_DIM` (for the broadcast bus buffer).
|
||||
@@ -263,7 +271,36 @@ pub const GRAD_SCALE_LIMIT_INDEX: usize = 35;
|
||||
/// `feedback_isv_for_adaptive_bounds.md` carve-out.
|
||||
pub const IQL_BRANCH_SCALE_FLOOR_INDEX: usize = 36;
|
||||
|
||||
/// ISV slot [37] — low 32 bits of the u64 layout fingerprint (stored as raw f32 bits).
|
||||
/// ISV slot [39] — current epoch index. CPU writes at each epoch boundary.
|
||||
/// Zero at construction; updated by the epoch loop in a follow-up task.
|
||||
pub const EPOCH_IDX_INDEX: usize = 39;
|
||||
/// ISV slot [40] — total epoch count for this training run. CPU writes at
|
||||
/// constructor from `config`; constant for the run's lifetime.
|
||||
pub const TOTAL_EPOCHS_INDEX: usize = 40;
|
||||
/// ISV slot [41] — effective epsilon written by the GPU epsilon adaptive
|
||||
/// kernel (follow-up task). Starts at 0; kernel fills it each step.
|
||||
pub const EPSILON_EFF_INDEX: usize = 41;
|
||||
/// ISV slot [42] — effective tau written by the GPU tau adaptive kernel
|
||||
/// (follow-up task). Starts at 0; kernel fills it each step.
|
||||
pub const TAU_EFF_INDEX: usize = 42;
|
||||
/// ISV slot [43] — effective gamma written by the GPU gamma adaptive kernel
|
||||
/// (follow-up task). Starts at 0; kernel fills it each step.
|
||||
pub const GAMMA_EFF_INDEX: usize = 43;
|
||||
/// ISV slot [44] — effective Kelly cap written by the GPU kelly_cap adaptive
|
||||
/// kernel (follow-up task). Starts at 0; kernel fills it each step.
|
||||
pub const KELLY_CAP_EFF_INDEX: usize = 44;
|
||||
/// ISV slot [45] — CQL pessimism base coefficient. Constructor writes
|
||||
/// `config.cql_alpha`; the CQL launch formula reads from here instead of
|
||||
/// the config field (Plan 1 Task 12 / spec §4.C.6). Plan 3 B.3 may later
|
||||
/// make this reactive by writing an adaptive value each epoch.
|
||||
pub const CQL_ALPHA_INDEX: usize = 45;
|
||||
/// ISV slot [46] — plan-MLP activation threshold. Constructor writes 0.5f;
|
||||
/// plan kernels (`experience_kernels.cu`, `backtest_plan_kernel.cu`) read
|
||||
/// from here instead of the hardcoded literal (Plan 1 Task 16 / spec §4.C.6).
|
||||
/// Plan 3 B.4 may later make this reactive.
|
||||
pub const PLAN_THRESHOLD_INDEX: usize = 46;
|
||||
|
||||
/// ISV slot [47] — low 32 bits of the u64 layout fingerprint (stored as raw f32 bits).
|
||||
///
|
||||
/// Design note: this is NOT a version number. There is no ordered version space.
|
||||
/// The fingerprint is a structural hash that automatically changes whenever any
|
||||
@@ -274,9 +311,9 @@ pub const IQL_BRANCH_SCALE_FLOOR_INDEX: usize = 36;
|
||||
///
|
||||
/// Placement at the tail (not head) is mandated by slots [0] and [1] being
|
||||
/// actively written by `isv_signal_update` (Q-drift and gradient-norm EMA).
|
||||
pub const ISV_LAYOUT_FINGERPRINT_LO_INDEX: usize = 37;
|
||||
/// ISV slot [38] — high 32 bits of the u64 layout fingerprint (stored as raw f32 bits).
|
||||
pub const ISV_LAYOUT_FINGERPRINT_HI_INDEX: usize = 38;
|
||||
pub const ISV_LAYOUT_FINGERPRINT_LO_INDEX: usize = 47;
|
||||
/// ISV slot [48] — high 32 bits of the u64 layout fingerprint (stored as raw f32 bits).
|
||||
pub const ISV_LAYOUT_FINGERPRINT_HI_INDEX: usize = 48;
|
||||
/// Canonical alias for the fingerprint slot (the low half).
|
||||
pub const ISV_LAYOUT_FINGERPRINT_INDEX: usize = ISV_LAYOUT_FINGERPRINT_LO_INDEX;
|
||||
|
||||
@@ -327,8 +364,11 @@ const fn layout_fingerprint_seed() -> &'static [u8] {
|
||||
V_CENTER_ORD=27;V_HALF_ORD=28;V_CENTER_URG=29;V_HALF_URG=30;\
|
||||
GRAD_NORM_TARGET_DIR=31;GRAD_NORM_TARGET_MAG=32;GRAD_NORM_TARGET_ORD=33;GRAD_NORM_TARGET_URG=34;\
|
||||
GRAD_SCALE_LIMIT=35;IQL_BRANCH_SCALE_FLOOR=36;\
|
||||
ISV_LAYOUT_FINGERPRINT_LO=37;ISV_LAYOUT_FINGERPRINT_HI=38;\
|
||||
ISV_TOTAL_DIM=39"
|
||||
EPOCH_IDX=39;TOTAL_EPOCHS=40;\
|
||||
EPSILON_EFF=41;TAU_EFF=42;GAMMA_EFF=43;KELLY_CAP_EFF=44;\
|
||||
CQL_ALPHA=45;PLAN_THRESHOLD=46;\
|
||||
ISV_LAYOUT_FINGERPRINT_LO=47;ISV_LAYOUT_FINGERPRINT_HI=48;\
|
||||
ISV_TOTAL_DIM=49"
|
||||
}
|
||||
|
||||
/// Compile-time layout fingerprint. Burned into the binary at build time.
|
||||
@@ -564,6 +604,11 @@ pub struct GpuDqnTrainConfig {
|
||||
/// market features from portfolio+MTF+OFI features in the state vector.
|
||||
/// OFI features (8 dims, when enabled) bypass the bottleneck via portfolio_dim.
|
||||
pub market_dim: usize,
|
||||
/// Total number of training epochs for this run. Written to ISV[TOTAL_EPOCHS_INDEX]
|
||||
/// at constructor time so GPU kernels can compute epoch-fraction progress without
|
||||
/// a CPU→GPU round-trip. Default 0 means "unknown / single fold"; set this from
|
||||
/// `DQNHyperparameters::epochs` at construction.
|
||||
pub total_epochs: usize,
|
||||
}
|
||||
|
||||
impl Default for GpuDqnTrainConfig {
|
||||
@@ -610,6 +655,7 @@ impl Default for GpuDqnTrainConfig {
|
||||
causal_intervention_interval: 50,
|
||||
bottleneck_dim: 16,
|
||||
market_dim: 42, // Default: 42 base features. Overridden to 50 when OFI (MBP-10) enabled.
|
||||
total_epochs: 0,
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -2024,11 +2070,12 @@ pub struct GpuDqnTrainer {
|
||||
q_dir_bin_means_reduce_kernel: CudaFunction,
|
||||
|
||||
// ── ISV core buffers (pinned device-mapped, GPU read/write) ──
|
||||
// isv_signals_pinned is sized for ISV_TOTAL_DIM (39) — the extra slots
|
||||
// slots [23..31] carry per-branch Q-support centres/half-widths. The
|
||||
// network (w_isv_fc1) and history rotation still consume only the first
|
||||
// isv_signals_pinned is sized for ISV_TOTAL_DIM (49) — the extra slots
|
||||
// slots [23..31] carry per-branch Q-support centres/half-widths; [39..49)
|
||||
// are the Plan 1 C.6 expansion (static configs + pre-allocated GPU slots).
|
||||
// The network (w_isv_fc1) and history rotation still consume only the first
|
||||
// ISV_NETWORK_DIM (23) slots; the tail is broadcast-bus scratchpad only.
|
||||
isv_signals_pinned: *mut f32, // [ISV_TOTAL_DIM = 39]
|
||||
isv_signals_pinned: *mut f32, // [ISV_TOTAL_DIM = 49]
|
||||
isv_signals_dev_ptr: u64,
|
||||
isv_history_pinned: *mut f32, // [ISV_K * 12 = 48] — history rotates slots [0..11] only
|
||||
isv_history_dev_ptr: u64,
|
||||
@@ -5858,7 +5905,15 @@ impl GpuDqnTrainer {
|
||||
Some(k) => k.clone(),
|
||||
None => return Ok(false),
|
||||
};
|
||||
if self.config.cql_alpha <= 0.0 {
|
||||
// Read cql_alpha base from ISV[CQL_ALPHA_INDEX] (Plan 1 Task 12 / spec §4.C.6).
|
||||
// When ISV is null (smoke scale without ISV warmup), fall back to the config
|
||||
// field so smoke tests that skip ISV bootstrap still work.
|
||||
let cql_alpha_base = if self.isv_signals_pinned.is_null() {
|
||||
self.config.cql_alpha
|
||||
} else {
|
||||
unsafe { *self.isv_signals_pinned.add(CQL_ALPHA_INDEX) }
|
||||
};
|
||||
if cql_alpha_base <= 0.0 {
|
||||
return Ok(false);
|
||||
}
|
||||
|
||||
@@ -5908,14 +5963,14 @@ impl GpuDqnTrainer {
|
||||
should never hit this path."
|
||||
);
|
||||
});
|
||||
self.config.cql_alpha
|
||||
cql_alpha_base
|
||||
} else {
|
||||
let (health, regime_stability) = unsafe {
|
||||
let h = (*self.isv_signals_pinned.add(LEARNING_HEALTH_INDEX)).clamp(0.0, 1.0);
|
||||
let s = (*self.isv_signals_pinned.add(11)).clamp(0.0, 1.0);
|
||||
(h, s)
|
||||
};
|
||||
self.config.cql_alpha * (1.0 - regime_stability) * health
|
||||
cql_alpha_base * (1.0 - regime_stability) * health
|
||||
};
|
||||
|
||||
// Cache for logging via FusedTrainingCtx / DQNTrainer.
|
||||
@@ -6160,7 +6215,16 @@ impl GpuDqnTrainer {
|
||||
|
||||
/// Whether CQL is enabled and the kernel was compiled.
|
||||
pub fn has_cql(&self) -> bool {
|
||||
self.cql_logit_grad_kernel.is_some() && self.config.cql_alpha > 0.0
|
||||
if !self.cql_logit_grad_kernel.is_some() {
|
||||
return false;
|
||||
}
|
||||
// Read base from ISV[CQL_ALPHA_INDEX] when available; fall back to config.
|
||||
let base = if self.isv_signals_pinned.is_null() {
|
||||
self.config.cql_alpha
|
||||
} else {
|
||||
unsafe { *self.isv_signals_pinned.add(CQL_ALPHA_INDEX) }
|
||||
};
|
||||
base > 0.0
|
||||
}
|
||||
|
||||
/// F8/G5: Scale the C51-contributed portion of `grad_buf` by the adaptive c51_budget.
|
||||
@@ -8167,10 +8231,10 @@ impl GpuDqnTrainer {
|
||||
};
|
||||
|
||||
// ── ISV core buffers (pinned device-mapped) ──────────────────────
|
||||
// Allocation uses `ISV_TOTAL_DIM` so all scratchpad slots [23..39) have
|
||||
// Allocation uses `ISV_TOTAL_DIM` so all scratchpad slots [23..49) have
|
||||
// backing storage. The network-facing path still consumes only the first
|
||||
// `ISV_NETWORK_DIM` slots — the scratchpad tail is invisible to `w_isv_fc1`.
|
||||
// Slots [37..39) hold the layout fingerprint (tail placement because [0..2)
|
||||
// Slots [47..49) hold the layout fingerprint (tail placement because [0..2)
|
||||
// are actively written by `isv_signal_update`).
|
||||
let (isv_signals_pinned, isv_signals_dev_ptr) = {
|
||||
let num_bytes = ISV_TOTAL_DIM * std::mem::size_of::<f32>();
|
||||
@@ -8212,7 +8276,14 @@ impl GpuDqnTrainer {
|
||||
// IQL readiness→1 and the losing branch has near-zero
|
||||
// advantage. Safety bound, not tuning.
|
||||
*sig_ptr.add(IQL_BRANCH_SCALE_FLOOR_INDEX) = 0.1_f32;
|
||||
// Layout fingerprint (ISV[37..39)). Compile-time structural hash of
|
||||
// Plan 1 C.6 static-config slots (2026-04-24):
|
||||
// EPOCH_IDX starts at 0; epoch loop updates in a follow-up task.
|
||||
// EPSILON_EFF / TAU_EFF / GAMMA_EFF / KELLY_CAP_EFF start at 0;
|
||||
// GPU kernels in follow-up tasks fill them.
|
||||
*sig_ptr.add(TOTAL_EPOCHS_INDEX) = config.total_epochs as f32;
|
||||
*sig_ptr.add(CQL_ALPHA_INDEX) = config.cql_alpha;
|
||||
*sig_ptr.add(PLAN_THRESHOLD_INDEX) = 0.5_f32;
|
||||
// Layout fingerprint (ISV[47..49)). Compile-time structural hash of
|
||||
// the slot layout; checkpoint load fails-fast on mismatch.
|
||||
// Stored as a u64 split across two f32 lanes using raw bit-cast so
|
||||
// no precision is lost. See LAYOUT_FINGERPRINT_CURRENT docs for why
|
||||
@@ -9880,10 +9951,12 @@ impl GpuDqnTrainer {
|
||||
// ── A3: LearningHealth signal writers / readers ───────────────────────
|
||||
|
||||
/// Write a scalar into the ISV pinned buffer at the given index.
|
||||
/// No-op if the pinned pointer is null or index >= ISV_DIM.
|
||||
/// No-op if the pinned pointer is null or index >= ISV_TOTAL_DIM.
|
||||
/// Accepts any slot in the full bus [0..ISV_TOTAL_DIM), including the
|
||||
/// scratchpad tail [ISV_NETWORK_DIM..ISV_TOTAL_DIM).
|
||||
pub fn write_isv_signal_at(&self, index: usize, value: f32) {
|
||||
if self.isv_signals_pinned.is_null() { return; }
|
||||
if index >= ISV_DIM { return; }
|
||||
if index >= ISV_TOTAL_DIM { return; }
|
||||
unsafe { *self.isv_signals_pinned.add(index) = value; }
|
||||
}
|
||||
|
||||
|
||||
@@ -124,6 +124,16 @@
|
||||
#define MAG_FULL 2 // 1.00× max_position
|
||||
#define NUM_MAGNITUDES 3
|
||||
|
||||
// ────────────────────────────────────────────────────────────────────────────
|
||||
// ISV bus slot indices referenced by kernels that receive the ISV pointer.
|
||||
// Kept in state_layout.cuh because both experience_kernels.cu and
|
||||
// backtest_plan_kernel.cu include this header and need these indices.
|
||||
// The authoritative named constants live in gpu_dqn_trainer.rs; the values
|
||||
// here must match exactly. Any renumber requires updating both files in the
|
||||
// same commit (per feedback_no_partial_refactor.md).
|
||||
// ────────────────────────────────────────────────────────────────────────────
|
||||
#define ISV_PLAN_THRESHOLD_IDX 46 // == PLAN_THRESHOLD_INDEX — plan activation threshold
|
||||
|
||||
// ── Compile-time checks ──
|
||||
static_assert(SL_PADDING_START + SL_PADDING_DIM == SL_STATE_DIM,
|
||||
"State layout dimensions must sum to SL_STATE_DIM");
|
||||
|
||||
@@ -400,6 +400,7 @@ impl FusedTrainingCtx {
|
||||
causal_intervention_interval: hyperparams.causal_intervention_interval,
|
||||
bottleneck_dim: hyperparams.bottleneck_dim,
|
||||
market_dim: 42, // Always 42 base market features — OFI features bypass bottleneck via portfolio_dim
|
||||
total_epochs: hyperparams.epochs,
|
||||
};
|
||||
|
||||
// Create weight set pointer views AFTER GpuDqnTrainer is constructed below.
|
||||
|
||||
@@ -144,13 +144,54 @@ impl StateResetRegistry {
|
||||
RegistryEntry {
|
||||
name: "ISV_LAYOUT_FINGERPRINT",
|
||||
category: ResetCategory::SchemaContract,
|
||||
description: "ISV[37..39) — compile-time structural hash of slot layout; fail-fast only, no migration path",
|
||||
description: "ISV[47..49) — compile-time structural hash of slot layout; fail-fast only, no migration path",
|
||||
},
|
||||
RegistryEntry {
|
||||
name: "isv_iql_branch_scale_floor",
|
||||
category: ResetCategory::SchemaContract,
|
||||
description: "ISV[36] — IQL branch_scales floor (safety bound, bootstrap-only)",
|
||||
},
|
||||
RegistryEntry {
|
||||
name: "isv_total_epochs",
|
||||
category: ResetCategory::SchemaContract,
|
||||
description: "ISV[TOTAL_EPOCHS_INDEX=40] — total epoch count for this run; constant after construction",
|
||||
},
|
||||
RegistryEntry {
|
||||
name: "isv_cql_alpha",
|
||||
category: ResetCategory::SchemaContract,
|
||||
description: "ISV[CQL_ALPHA_INDEX=45] — CQL pessimism base coefficient; Plan 3 B.3 may make reactive",
|
||||
},
|
||||
RegistryEntry {
|
||||
name: "isv_plan_threshold",
|
||||
category: ResetCategory::SchemaContract,
|
||||
description: "ISV[PLAN_THRESHOLD_INDEX=46] — plan-MLP activation threshold; Plan 3 B.4 may make reactive",
|
||||
},
|
||||
// ───── Fold-reset state (pre-allocated GPU-written slots) ───
|
||||
RegistryEntry {
|
||||
name: "isv_epoch_idx",
|
||||
category: ResetCategory::FoldReset,
|
||||
description: "ISV[EPOCH_IDX_INDEX=39] — current epoch index; CPU writes at epoch boundary (follow-up task)",
|
||||
},
|
||||
RegistryEntry {
|
||||
name: "isv_epsilon_eff",
|
||||
category: ResetCategory::FoldReset,
|
||||
description: "ISV[EPSILON_EFF_INDEX=41] — effective epsilon; GPU epsilon-adaptive kernel fills (follow-up task)",
|
||||
},
|
||||
RegistryEntry {
|
||||
name: "isv_tau_eff",
|
||||
category: ResetCategory::FoldReset,
|
||||
description: "ISV[TAU_EFF_INDEX=42] — effective tau; GPU tau-adaptive kernel fills (follow-up task)",
|
||||
},
|
||||
RegistryEntry {
|
||||
name: "isv_gamma_eff",
|
||||
category: ResetCategory::FoldReset,
|
||||
description: "ISV[GAMMA_EFF_INDEX=43] — effective gamma; GPU gamma-adaptive kernel fills (follow-up task)",
|
||||
},
|
||||
RegistryEntry {
|
||||
name: "isv_kelly_cap_eff",
|
||||
category: ResetCategory::FoldReset,
|
||||
description: "ISV[KELLY_CAP_EFF_INDEX=44] — effective Kelly cap; GPU kelly-cap-adaptive kernel fills (follow-up task)",
|
||||
},
|
||||
];
|
||||
Self { entries }
|
||||
}
|
||||
|
||||
@@ -219,6 +219,8 @@
|
||||
|
||||
Updated after Task 6 cleanup (2026-04-24): 5 confirmed-orphan files deleted, 3 Orphan rows reclassified Partial, 1 Orphan reclassified with crate-level follow-up action. Plan 1 Task 8 (revised, 2026-04-24): `adaptive_controller.rs` renamed → `adaptive_monitor.rs`; `AdaptiveController` replaced with read-only `AdaptiveMonitor` per spec §4.C.6 (GPU drives, CPU reads).
|
||||
|
||||
Plan 1 Tasks 12/15/16 + pre-allocation (2026-04-24): No new modules added. Changes are ISV slot allocation + consumer migration only. Task 15 confirmed no-op (`IQL_BRANCH_SCALE_FLOOR_INDEX` already serves conviction-floor role). Tasks 12 and 16 migrate `cql_alpha` and plan-threshold consumers from config fields / hardcoded literals to ISV slots. 8 new ISV slots allocated ([39..47)); fingerprint tail moves from [37..39) to [47..49); `ISV_TOTAL_DIM` 39 → 49. `GpuDqnTrainConfig` gains `total_epochs` field (written to `TOTAL_EPOCHS_INDEX` at construction). `write_isv_signal_at` bound extended from `ISV_DIM` to `ISV_TOTAL_DIM` to allow writes beyond slot 22.
|
||||
|
||||
| Classification | Count |
|
||||
|---|---|
|
||||
| Wired | 75 |
|
||||
|
||||
@@ -8,13 +8,14 @@ fail-fast; no migration functions exist. Backward compat is structurally
|
||||
unwritable (there is no ordered version space to pair migrations against).
|
||||
See spec §4.A.2.
|
||||
|
||||
**Tail placement rationale:** The fingerprint occupies the tail (`[37..39)`) rather
|
||||
**Tail placement rationale:** The fingerprint occupies the tail (`[47..49)`) rather
|
||||
than the head because `isv_signals[0]` and `isv_signals[1]` are actively written
|
||||
by `isv_signal_update` (Q-drift EMA and gradient-norm EMA respectively). Inserting
|
||||
at the head would displace those live signals and require shifting every upstream
|
||||
literal in `experience_kernels.cu`.
|
||||
literal in `experience_kernels.cu`. The fingerprint moves to the new tail each time
|
||||
new slots are appended to the bus.
|
||||
|
||||
**Current `ISV_TOTAL_DIM`:** 39. Post-full DQN v2 rollout: 72.
|
||||
**Current `ISV_TOTAL_DIM`:** 49 (Plan 1 Tasks 12/15/16 + pre-allocation for Tasks 9-11/13/14). Post-full DQN v2 rollout: 72.
|
||||
|
||||
| Index | Name constant | Type | Producer | Consumers | Reset-category | Notes |
|
||||
|---|---|---|---|---|---|---|
|
||||
@@ -38,6 +39,15 @@ literal in `experience_kernels.cu`.
|
||||
| [31..35) | `GRAD_NORM_TARGET_*_INDEX` | f32 | grad_balance_isv_update | branch_grad_rescale | SoftReset(decay_bars=500) | Per-branch grad-norm target |
|
||||
| [35] | `GRAD_SCALE_LIMIT_INDEX` | f32 | grad_balance_isv_update | branch_grad_rescale | SoftReset(decay_bars=500) | Scale clamp limit |
|
||||
| [36] | `IQL_BRANCH_SCALE_FLOOR_INDEX` | f32 | construct (static) | iql_per_branch_advantage | SchemaContract | Per-sample branch_scales floor; safety bound |
|
||||
| [37] | `ISV_LAYOUT_FINGERPRINT_LO_INDEX` | u32 bits (in f32) | construct | check_layout_fingerprint | SchemaContract | Low 32 bits of u64 FNV-1a structural hash. Fail-fast on mismatch — NOT a version number, no migration path. |
|
||||
| [38] | `ISV_LAYOUT_FINGERPRINT_HI_INDEX` | u32 bits (in f32) | construct | check_layout_fingerprint | SchemaContract | High 32 bits of u64 FNV-1a structural hash. |
|
||||
| [39..72) | (reserved for DQN v2) | | | | | Allocated incrementally by Plans 2-5 |
|
||||
| [37..39) | (gap — fingerprint moved to tail) | — | — | — | — | Previously [37..39); fingerprint promoted to [47..49) by Plan 1 C.6 expansion. Slots unused; zero-filled. |
|
||||
| [39] | `EPOCH_IDX_INDEX` | f32 (int cast) | CPU epoch-loop (follow-up task) | GPU adaptive kernels | FoldReset | Current epoch index; 0 at construction, CPU writes at each epoch boundary |
|
||||
| [40] | `TOTAL_EPOCHS_INDEX` | f32 (int cast) | construct (CPU static) | GPU adaptive kernels | SchemaContract | Total epoch count for this run; written at construction from `config.total_epochs` |
|
||||
| [41] | `EPSILON_EFF_INDEX` | f32 | GPU epsilon-adaptive kernel (follow-up) | epsilon-greedy action select | FoldReset | Effective epsilon; 0 at construction; GPU fills each step |
|
||||
| [42] | `TAU_EFF_INDEX` | f32 | GPU tau-adaptive kernel (follow-up) | target-net Polyak update | FoldReset | Effective tau; 0 at construction; GPU fills each step |
|
||||
| [43] | `GAMMA_EFF_INDEX` | f32 | GPU gamma-adaptive kernel (follow-up) | Bellman target kernel | FoldReset | Effective gamma; 0 at construction; GPU fills each step |
|
||||
| [44] | `KELLY_CAP_EFF_INDEX` | f32 | GPU kelly-cap-adaptive kernel (follow-up) | experience_env_step | FoldReset | Effective Kelly cap; 0 at construction; GPU fills each step |
|
||||
| [45] | `CQL_ALPHA_INDEX` | f32 | construct (CPU static) | CQL adaptive formula in Rust | SchemaContract | CQL pessimism base coefficient; written from `config.cql_alpha`; read in `compute_cql_logit_gradients` |
|
||||
| [46] | `PLAN_THRESHOLD_INDEX` | f32 | construct (CPU static) | `experience_kernels.cu`, `backtest_plan_kernel.cu` | SchemaContract | Plan-MLP activation threshold; 0.5f at construction; read via `ISV_PLAN_THRESHOLD_IDX` in both kernels |
|
||||
| [47] | `ISV_LAYOUT_FINGERPRINT_LO_INDEX` | u32 bits (in f32) | construct | check_layout_fingerprint | SchemaContract | Low 32 bits of u64 FNV-1a structural hash. Fail-fast on mismatch — NOT a version number, no migration path. |
|
||||
| [48] | `ISV_LAYOUT_FINGERPRINT_HI_INDEX` | u32 bits (in f32) | construct | check_layout_fingerprint | SchemaContract | High 32 bits of u64 FNV-1a structural hash. |
|
||||
| [49..72) | (reserved for DQN v2) | | | | | Allocated incrementally by Plans 2-5 |
|
||||
|
||||
Reference in New Issue
Block a user