feat(sp5): Layer B — atomic consumer migration to ISV-driven adaptive signals
SP5 Layer B: all consumers of ISV slots 174..286 migrated in one atomic
commit per feedback_no_partial_refactor.
Pearl 1 (C51 atom span, done prior session):
atoms_update_kernel.cu: reads v_center/v_half from ISV[174+b]/ISV[178+b]
Pearl 1-ext (per-branch num_atoms):
atoms_update_kernel.cu: reads ISV[274+b], rounds to valid {16,32,64},
uses eff_na for softmax/cumsum; output stride still global num_atoms
Pearl 2 (per-branch loss budgets):
fused_training.rs: compute_adaptive_budgets() reads ISV[190..194/198/202]
averaged over 4 branches; Invariant-1 floors 0.05/0.05/0.02/0.02
Pearl 3 (per-branch NoisyNet sigma):
training_loop.rs: ExperienceCollectorConfig.noise_sigma reads
mean(ISV[210..214]) with 0.01 floor; falls back to hyperparams.noise_sigma
when fused_ctx absent
Pearl 4 (per-group Adam beta1/beta2/epsilon):
gpu_dqn_trainer.rs: per-group reads inside launch_adam_update loop via
ISV[226+g], ISV[234+g], ISV[242+g]
gpu_iqn_head.rs: iqn_beta1/beta2/epsilon threaded into
execute_training_pipeline + train_iqn_step_gpu
gpu_attention.rs: attn_beta1/beta2/epsilon threaded into adam_step
gpu_tlob.rs: same as attention (shares ParamGroup::Attn=6)
fused_training.rs: 4 call sites read ISV[226+g..242+g] for IQN+Attn+TLOB
gpu_curiosity_trainer.rs: cur_beta1/beta2/epsilon params added to
launch_adam_step + train_on_collector_buffers
gpu_experience_collector.rs: train_curiosity_gpu threads cur_beta1/2/eps
training_loop.rs: reads ISV[233/241/249] for Curiosity=7 at call site
Pearl 5 (per-branch IQN tau schedule):
gpu_iqn_head.rs: new refresh_taus_from_isv() averages 4 branches per
quantile, applies FIXED_TAUS cold-start floor, re-uploads online_taus/
target_taus/cos_features in-place via upload_f32_via_pinned
fused_training.rs: reads ISV[250..270) before IQN prepare_buffers
Pearl 6 (Kelly cap) + Pearl 8 (trail distance) done in prior session:
trade_physics.cuh, experience_kernels.cu, backtest_env_kernel.cu
Audit doc: dqn-wire-up-audit.md SP5 Layer B entry added.
All 930 passing lib tests still pass; 14 pre-existing failures unchanged.
Compile: SQLX_OFFLINE=true cargo check -p ml --offline clean.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -9,7 +9,11 @@
|
||||
* Reads:
|
||||
* spacing_raw[b] — learnable softmax spacing params for branch b,
|
||||
* accessed via pre-resolved device pointers passed as args.
|
||||
* isv_signals — ISV bus (per-branch v-range at slots 23+2*b / 24+2*b).
|
||||
* isv_signals — ISV bus:
|
||||
* SP5 Layer B (Pearl 1): per-branch v-range at
|
||||
* ATOM_V_CENTER_BASE=174: v_center[b] = isv_signals[174 + b]
|
||||
* ATOM_V_HALF_BASE=178: v_half[b] = isv_signals[178 + b]
|
||||
* (SP4 slots 23+2*b / 24+2*b are superseded by SP5 Pearl 1.)
|
||||
*
|
||||
* Writes:
|
||||
* positions_out — atom grid [4 * num_atoms], branch b at offset b*num_atoms.
|
||||
@@ -17,7 +21,7 @@
|
||||
* Formula (same as adaptive_atom_positions in experience_kernels.cu):
|
||||
* spacing = softmax(spacing_raw)
|
||||
* positions[j] = v_min + (v_max - v_min) * cumsum(spacing)[j]
|
||||
* with v_min = isv[23+2*b] - isv[24+2*b], v_max = isv[23+2*b] + isv[24+2*b].
|
||||
* with v_min = isv[174+b] - isv[178+b], v_max = isv[174+b] + isv[178+b].
|
||||
*/
|
||||
|
||||
/* SP4 Layer B: per-branch atom-position bound from ISV[ATOM_POS_BOUND_BASE + branch].
|
||||
@@ -25,6 +29,34 @@
|
||||
*/
|
||||
#define ATOM_POS_BOUND_BASE 132
|
||||
|
||||
/* SP5 Layer B (Pearl 1): per-branch v_center and v_half from ISV.
|
||||
* Mirrors sp5_isv_slots.rs: ATOM_V_CENTER_BASE=174, ATOM_V_HALF_BASE=178. */
|
||||
#define ATOM_V_CENTER_BASE 174
|
||||
#define ATOM_V_HALF_BASE 178
|
||||
|
||||
/* SP5 Layer B (Pearl 1-ext): per-branch num_atoms from ISV.
|
||||
* Mirrors sp5_isv_slots.rs: ATOM_NUM_ATOMS_BASE=274.
|
||||
* Valid discrete values: {16, 32, 64}. Rounds ISV float to nearest valid
|
||||
* count. Cold-start floor: if ISV=0 (pre-observation), fall back to the
|
||||
* global `num_atoms` parameter (Invariant-1 carve-out). */
|
||||
#define ATOM_NUM_ATOMS_BASE 274
|
||||
|
||||
/* Round an ISV-driven f32 num_atoms signal to the nearest valid discrete C51
|
||||
* atom count {16, 32, 64}, bounded to [16, num_atoms_max]. */
|
||||
__device__ __forceinline__ int round_to_valid_num_atoms(float isv_val, int num_atoms_max) {
|
||||
if (isv_val < 1.0f) return num_atoms_max; /* cold-start floor */
|
||||
int v = (int)(isv_val + 0.5f);
|
||||
/* Snap to nearest of {16, 32, 64}: */
|
||||
int d16 = abs(v - 16);
|
||||
int d32 = abs(v - 32);
|
||||
int d64 = abs(v - 64);
|
||||
int best = (d16 <= d32 && d16 <= d64) ? 16 : (d32 <= d64) ? 32 : 64;
|
||||
/* Clamp to [16, num_atoms_max]: never exceed the buffer's allocated stride */
|
||||
if (best > num_atoms_max) best = num_atoms_max;
|
||||
if (best < 16) best = 16;
|
||||
return best;
|
||||
}
|
||||
|
||||
extern "C" __global__ void atoms_update(
|
||||
unsigned long long spacing_ptr_b0, /* device ptr to spacing_raw[dir] */
|
||||
unsigned long long spacing_ptr_b1, /* device ptr to spacing_raw[mag] */
|
||||
@@ -39,15 +71,30 @@ extern "C" __global__ void atoms_update(
|
||||
|
||||
int tid = threadIdx.x;
|
||||
extern __shared__ float shmem[];
|
||||
/* Shared memory is allocated for global `num_atoms` (worst case) at launch.
|
||||
* eff_na ≤ num_atoms, so both halves fit safely. */
|
||||
float* s_softmax = shmem;
|
||||
float* s_cumsum = shmem + num_atoms;
|
||||
|
||||
/* Per-branch v-range from ISV: centre at 23+2*b, half at 24+2*b */
|
||||
float v_center = isv_signals[23 + 2 * branch];
|
||||
float v_half = isv_signals[24 + 2 * branch];
|
||||
/* SP5 Layer B (Pearl 1): per-branch v-range from ISV.
|
||||
* v_center[b] = ISV[ATOM_V_CENTER_BASE + b] = ISV[174 + b]
|
||||
* v_half[b] = ISV[ATOM_V_HALF_BASE + b] = ISV[178 + b]
|
||||
* EPS_CLAMP_FLOOR = 1.0 (Invariant 1 anchor): pre-first-observation ISV
|
||||
* reads as 0; fmaxf prevents zero/negative half-width collapsing the range. */
|
||||
float v_center = isv_signals[ATOM_V_CENTER_BASE + branch];
|
||||
float v_half = fmaxf(isv_signals[ATOM_V_HALF_BASE + branch], 1.0f);
|
||||
float v_min = v_center - v_half;
|
||||
float v_max = v_center + v_half;
|
||||
|
||||
/* SP5 Layer B (Pearl 1-ext): per-branch effective num_atoms from ISV.
|
||||
* ISV[ATOM_NUM_ATOMS_BASE + branch] = ISV[274 + branch].
|
||||
* Rounded to valid discrete {16,32,64}; cold-start (ISV=0) falls back to
|
||||
* the global `num_atoms` parameter (Invariant-1 carve-out). The global
|
||||
* `num_atoms` is still used for shared-memory stride and output-buffer
|
||||
* stride (buffer is allocated as 4 * config.num_atoms). */
|
||||
int eff_na = round_to_valid_num_atoms(
|
||||
isv_signals[ATOM_NUM_ATOMS_BASE + branch], num_atoms);
|
||||
|
||||
/* Resolve spacing_raw pointer for this branch */
|
||||
const float* spacing_raw;
|
||||
switch (branch) {
|
||||
@@ -57,16 +104,16 @@ extern "C" __global__ void atoms_update(
|
||||
default: spacing_raw = (const float*)spacing_ptr_b3; break;
|
||||
}
|
||||
|
||||
/* Load spacing_raw into shared memory */
|
||||
if (tid < num_atoms) {
|
||||
/* Load spacing_raw into shared memory (only eff_na elements are valid) */
|
||||
if (tid < eff_na) {
|
||||
s_softmax[tid] = spacing_raw[tid];
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
/* Find max for numerical stability (single thread, num_atoms <= 256 is tiny) */
|
||||
/* Find max for numerical stability (single thread, eff_na <= 256 is tiny) */
|
||||
if (tid == 0) {
|
||||
float max_val = s_softmax[0];
|
||||
for (int j = 1; j < num_atoms; j++)
|
||||
for (int j = 1; j < eff_na; j++)
|
||||
max_val = fmaxf(max_val, s_softmax[j]);
|
||||
s_cumsum[0] = max_val; /* store max temporarily */
|
||||
}
|
||||
@@ -75,7 +122,7 @@ extern "C" __global__ void atoms_update(
|
||||
__syncthreads();
|
||||
|
||||
/* Compute softmax */
|
||||
if (tid < num_atoms) {
|
||||
if (tid < eff_na) {
|
||||
s_softmax[tid] = expf(s_softmax[tid] - max_val);
|
||||
}
|
||||
__syncthreads();
|
||||
@@ -83,10 +130,10 @@ extern "C" __global__ void atoms_update(
|
||||
/* Sum for normalization */
|
||||
if (tid == 0) {
|
||||
float sum = 0.0f;
|
||||
for (int j = 0; j < num_atoms; j++)
|
||||
for (int j = 0; j < eff_na; j++)
|
||||
sum += s_softmax[j];
|
||||
float inv_sum = 1.0f / fmaxf(sum, 1e-8f);
|
||||
for (int j = 0; j < num_atoms; j++)
|
||||
for (int j = 0; j < eff_na; j++)
|
||||
s_softmax[j] *= inv_sum;
|
||||
}
|
||||
__syncthreads();
|
||||
@@ -94,14 +141,17 @@ extern "C" __global__ void atoms_update(
|
||||
/* Inclusive prefix sum (sequential, tid=0 only) */
|
||||
if (tid == 0) {
|
||||
s_cumsum[0] = s_softmax[0];
|
||||
for (int j = 1; j < num_atoms; j++)
|
||||
for (int j = 1; j < eff_na; j++)
|
||||
s_cumsum[j] = s_cumsum[j-1] + s_softmax[j];
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
/* Write positions: positions[0] = v_min, positions[j>0] = v_min + range * cumsum[j-1] */
|
||||
/* Write positions: positions[0] = v_min, positions[j>0] = v_min + range * cumsum[j-1].
|
||||
* Output stride is the global `num_atoms` (buffer layout invariant). Positions
|
||||
* beyond eff_na are left unchanged (they are not read — C51 consumers also
|
||||
* use eff_na from ISV). */
|
||||
float* out = positions_out + (long long)branch * num_atoms;
|
||||
if (tid < num_atoms) {
|
||||
if (tid < eff_na) {
|
||||
float range = v_max - v_min;
|
||||
float frac = (tid == 0) ? 0.0f : s_cumsum[tid - 1];
|
||||
float new_pos = v_min + range * frac;
|
||||
|
||||
@@ -254,7 +254,8 @@ extern "C" __global__ void backtest_env_step(
|
||||
&prev_position_sign,
|
||||
&actual_dir,
|
||||
&actual_mag,
|
||||
&trail_triggered
|
||||
&trail_triggered,
|
||||
isv_signals /* SP5 Layer B: Pearl 6 (kelly_f_smooth) + Pearl 8 (trail_dist_per_dir) */
|
||||
);
|
||||
(void)prev_position; (void)prev_position_sign; (void)trail_triggered;
|
||||
|
||||
@@ -508,7 +509,8 @@ extern "C" __global__ void backtest_env_step_batch(
|
||||
&prev_position_sign,
|
||||
&actual_dir,
|
||||
&actual_mag,
|
||||
&trail_triggered
|
||||
&trail_triggered,
|
||||
isv_signals /* SP5 Layer B: Pearl 6 (kelly_f_smooth) + Pearl 8 (trail_dist_per_dir) */
|
||||
);
|
||||
(void)prev_position; (void)prev_position_sign; (void)trail_triggered;
|
||||
|
||||
|
||||
@@ -2124,7 +2124,8 @@ extern "C" __global__ void experience_env_step(
|
||||
&prev_position_sign_core,
|
||||
&actual_dir_core,
|
||||
&actual_mag_core,
|
||||
&trail_triggered
|
||||
&trail_triggered,
|
||||
isv_signals_ptr /* SP5 Layer B: Pearl 6 (kelly_f_smooth) + Pearl 8 (trail_dist_per_dir) */
|
||||
);
|
||||
/* actual_mag_core is now consumed by Task 2.X per-magnitude instrumentation
|
||||
* at the trade-lifecycle block below (~line 1760). actual_dir_core +
|
||||
|
||||
@@ -1064,6 +1064,12 @@ impl GpuAttention {
|
||||
/* SP4 Layer B (Mech 9): ISV-driven AdamW weight_decay rate.
|
||||
* Caller reads `ISV[WD_RATE[Attn=6]]` with `.max(0.0)`. */
|
||||
weight_decay_value: f32,
|
||||
/* SP5 Layer B (Pearl 4): ISV-derived Adam β1/β2/ε for ParamGroup::Attn=6.
|
||||
* Caller reads ISV[ADAM_BETA1_BASE+6], ISV[ADAM_BETA2_BASE+6], ISV[ADAM_EPS_BASE+6]
|
||||
* and applies Invariant 1 floors (β1≥0.5, β2≥0.9, ε≥1e-12). */
|
||||
attn_beta1: f32,
|
||||
attn_beta2: f32,
|
||||
attn_epsilon: f32,
|
||||
) -> Result<(), MLError> {
|
||||
let stream = override_stream.unwrap_or(&self.stream);
|
||||
// override_workspace is accepted for API symmetry; adam_step does not use
|
||||
@@ -1111,9 +1117,11 @@ impl GpuAttention {
|
||||
shared_mem_bytes: 0,
|
||||
};
|
||||
|
||||
let beta1 = 0.9_f32;
|
||||
let beta2 = 0.999_f32;
|
||||
let eps = 1e-8_f32;
|
||||
// SP5 Layer B (Pearl 4): per-group Adam β1/β2/ε threaded from
|
||||
// FusedTrainingCtx via ISV[ADAM_BETA{1,2}_BASE+6] / ISV[ADAM_EPS_BASE+6].
|
||||
let beta1 = attn_beta1;
|
||||
let beta2 = attn_beta2;
|
||||
let eps = attn_epsilon;
|
||||
// SP4 Layer B (Mech 9): AdamW weight_decay sourced from
|
||||
// ISV[WD_RATE[Attn=6]] by the caller (`FusedTrainingCtx`) and
|
||||
// threaded through `adam_step`. `.max(0.0)` upstream enforces the
|
||||
|
||||
@@ -522,6 +522,12 @@ fn launch_adam_step(
|
||||
batch_size: usize,
|
||||
step: i32,
|
||||
weight_clamp_max_abs: f32,
|
||||
/* SP5 Pearl 4: per-group Adam β1/β2/ε read from ISV[ADAM_BETA1_BASE+7],
|
||||
* ISV[ADAM_BETA2_BASE+7], ISV[ADAM_EPS_BASE+7] by caller and threaded
|
||||
* through train_on_collector_buffers. Curiosity = ParamGroup index 7. */
|
||||
cur_beta1: f32,
|
||||
cur_beta2: f32,
|
||||
cur_epsilon: f32,
|
||||
/* SP4 Task A14: per-sub-launch Pearl C engagement-counter args.
|
||||
* `engage_buf_offset` is unique per sub-buffer (W1/b1/W2/b2) so
|
||||
* the 4 sub-launches don't overwrite each other's per-block counts.
|
||||
@@ -556,9 +562,9 @@ fn launch_adam_step(
|
||||
.arg(&np_i32)
|
||||
.arg(&bs_i32)
|
||||
.arg(&ADAM_LR)
|
||||
.arg(&ADAM_BETA1)
|
||||
.arg(&ADAM_BETA2)
|
||||
.arg(&ADAM_EPS)
|
||||
.arg(&cur_beta1)
|
||||
.arg(&cur_beta2)
|
||||
.arg(&cur_epsilon)
|
||||
.arg(&step)
|
||||
.arg(&weight_clamp_max_abs) // SP3 Mech 9 (close-out v2): post-Adam |p| clamp
|
||||
.arg(&nan_flags_ptr) // SP4 A14: sticky engage flag store
|
||||
@@ -759,6 +765,13 @@ impl GpuCuriosityTrainer {
|
||||
* `ISV[WEIGHT_BOUND[Curiosity]].max(EPS_CLAMP_FLOOR)` host-side.
|
||||
* 0=disabled (debug only). */
|
||||
weight_clamp_max_abs: f32,
|
||||
/* SP5 Pearl 4: per-group Adam β1/β2/ε for the four
|
||||
* `curiosity_adam_step` launches. Caller reads ISV[ADAM_BETA1_BASE+7],
|
||||
* ISV[ADAM_BETA2_BASE+7], ISV[ADAM_EPS_BASE+7] from FusedTrainingCtx
|
||||
* and threads through train_curiosity_gpu → here. */
|
||||
cur_beta1: f32,
|
||||
cur_beta2: f32,
|
||||
cur_epsilon: f32,
|
||||
) -> Result<(), MLError> {
|
||||
if n_samples < 2 {
|
||||
return Ok(());
|
||||
@@ -1004,24 +1017,28 @@ impl GpuCuriosityTrainer {
|
||||
stream, &self.adam_func, &mut weights.w1, &self.grad_w1,
|
||||
CUR_W1_LEN, &mut self.adam_m_w1, &mut self.adam_v_w1,
|
||||
n_train, step, weight_clamp_max_abs,
|
||||
cur_beta1, cur_beta2, cur_epsilon,
|
||||
nan_flags_ptr, diag_slot, engage_buf_ptr, off_w1,
|
||||
)?;
|
||||
launch_adam_step(
|
||||
stream, &self.adam_func, &mut weights.b1, &self.grad_b1,
|
||||
CUR_B1_LEN, &mut self.adam_m_b1, &mut self.adam_v_b1,
|
||||
n_train, step, weight_clamp_max_abs,
|
||||
cur_beta1, cur_beta2, cur_epsilon,
|
||||
nan_flags_ptr, diag_slot, engage_buf_ptr, off_b1,
|
||||
)?;
|
||||
launch_adam_step(
|
||||
stream, &self.adam_func, &mut weights.w2, &self.grad_w2,
|
||||
CUR_W2_LEN, &mut self.adam_m_w2, &mut self.adam_v_w2,
|
||||
n_train, step, weight_clamp_max_abs,
|
||||
cur_beta1, cur_beta2, cur_epsilon,
|
||||
nan_flags_ptr, diag_slot, engage_buf_ptr, off_w2,
|
||||
)?;
|
||||
launch_adam_step(
|
||||
stream, &self.adam_func, &mut weights.b2, &self.grad_b2,
|
||||
CUR_B2_LEN, &mut self.adam_m_b2, &mut self.adam_v_b2,
|
||||
n_train, step, weight_clamp_max_abs,
|
||||
cur_beta1, cur_beta2, cur_epsilon,
|
||||
nan_flags_ptr, diag_slot, engage_buf_ptr, off_b2,
|
||||
)?;
|
||||
|
||||
|
||||
@@ -23428,11 +23428,11 @@ impl GpuDqnTrainer {
|
||||
use crate::cuda_pipeline::sp4_isv_slots::{
|
||||
ParamGroup as SP4Group, L1_LAMBDA_TRUNK_INDEX, SP4_ENGAGE_OFFSET_DISABLED,
|
||||
};
|
||||
use crate::cuda_pipeline::sp5_isv_slots::{
|
||||
ADAM_BETA1_BASE, ADAM_BETA2_BASE, ADAM_EPS_BASE,
|
||||
};
|
||||
|
||||
let lr_ptr = self.lr_dev_ptr; // pinned device-mapped — graph-safe, value read at replay
|
||||
let beta1 = self.config.beta1;
|
||||
let beta2 = self.config.beta2;
|
||||
let epsilon = self.config.epsilon;
|
||||
let adaptive_clip_ptr = self.ptrs.adaptive_clip_buf;
|
||||
let norm_ptr = self.ptrs.grad_norm_buf;
|
||||
let t_ptr = self.ptrs.t_buf;
|
||||
@@ -23518,6 +23518,16 @@ impl GpuDqnTrainer {
|
||||
continue;
|
||||
}
|
||||
|
||||
// SP5 Layer B (Pearl 4): per-group Adam β1/β2/ε from ISV[226+g],
|
||||
// ISV[234+g], ISV[242+g]. Floors are Invariant 1 carve-outs
|
||||
// (numerical-stability): β1≥0.5 avoids momentum inversion,
|
||||
// β2≥0.9 prevents exploding second-moment estimates, ε≥1e-12
|
||||
// avoids division-by-zero. Cold-start ISV reads 0 — floors fire.
|
||||
let g = group.idx();
|
||||
let beta1 = self.read_isv_signal_at(ADAM_BETA1_BASE + g).max(0.5_f32).min(0.9999_f32);
|
||||
let beta2 = self.read_isv_signal_at(ADAM_BETA2_BASE + g).max(0.9_f32).min(0.99999_f32);
|
||||
let epsilon = self.read_isv_signal_at(ADAM_EPS_BASE + g).max(1e-12_f32);
|
||||
|
||||
// Per-group ISV-driven post-Adam weight clamp + AdamW
|
||||
// weight_decay rate via the SP4 Layer B helper. Trunk-extras
|
||||
// sub-launch reuses DqnTrunk's bounds (its `group` tag).
|
||||
|
||||
@@ -4361,11 +4361,19 @@ impl GpuExperienceCollector {
|
||||
/// bound for `curiosity_adam_step`. The collector doesn't own ISV — the
|
||||
/// caller (training_loop) reads `100 × ISV[Q_ABS_REF].max(1.0)` from the
|
||||
/// FusedTrainingCtx and threads it through here. 0=disabled (debug only).
|
||||
///
|
||||
/// `cur_beta1/cur_beta2/cur_epsilon` are SP5 Pearl 4: ISV-driven Adam
|
||||
/// hyperparameters for ParamGroup::Curiosity (index 7). Caller reads
|
||||
/// ISV[ADAM_BETA1_BASE+7], ISV[ADAM_BETA2_BASE+7], ISV[ADAM_EPS_BASE+7]
|
||||
/// from FusedTrainingCtx and threads through here.
|
||||
pub fn train_curiosity_gpu(
|
||||
&mut self,
|
||||
n_episodes: usize,
|
||||
timesteps: usize,
|
||||
weight_clamp_max_abs: f32,
|
||||
cur_beta1: f32,
|
||||
cur_beta2: f32,
|
||||
cur_epsilon: f32,
|
||||
) -> Result<(), MLError> {
|
||||
if let Some(ref mut trainer) = self.curiosity_trainer {
|
||||
let total = n_episodes.saturating_mul(timesteps);
|
||||
@@ -4378,6 +4386,9 @@ impl GpuExperienceCollector {
|
||||
&self.actions_out,
|
||||
total,
|
||||
weight_clamp_max_abs,
|
||||
cur_beta1,
|
||||
cur_beta2,
|
||||
cur_epsilon,
|
||||
)?;
|
||||
}
|
||||
Ok(())
|
||||
|
||||
@@ -824,6 +824,12 @@ impl GpuIqnHead {
|
||||
weight_clamp_max_abs: f32,
|
||||
/* SP4 Layer B (Mech 9): ISV-derived AdamW weight_decay rate. */
|
||||
weight_decay_value: f32,
|
||||
/* SP5 Pearl 4: ISV-driven Adam β1/β2/ε for ParamGroup::Iqn (index 3).
|
||||
* Caller reads ISV[ADAM_BETA1_BASE+3], ISV[ADAM_BETA2_BASE+3],
|
||||
* ISV[ADAM_EPS_BASE+3] and threads through here. */
|
||||
iqn_beta1: f32,
|
||||
iqn_beta2: f32,
|
||||
iqn_epsilon: f32,
|
||||
) -> Result<f32, MLError> {
|
||||
// 1-2. Decode actions + DtoD copy rewards/dones
|
||||
self.prepare_buffers(dqn_actions_buf, dqn_rewards_buf, dqn_dones_buf)?;
|
||||
@@ -837,6 +843,9 @@ impl GpuIqnHead {
|
||||
None,
|
||||
weight_clamp_max_abs,
|
||||
weight_decay_value,
|
||||
iqn_beta1,
|
||||
iqn_beta2,
|
||||
iqn_epsilon,
|
||||
)
|
||||
}
|
||||
|
||||
@@ -879,6 +888,12 @@ impl GpuIqnHead {
|
||||
/* SP4 Layer B (Mech 9): ISV-derived AdamW weight_decay rate.
|
||||
* Caller reads `ISV[WD_RATE[Iqn=3]]` with `.max(0.0)` non-negativity. */
|
||||
weight_decay_value: f32,
|
||||
/* SP5 Layer B (Pearl 4): ISV-derived Adam β1/β2/ε for ParamGroup::Iqn=3.
|
||||
* Caller reads ISV[ADAM_BETA1_BASE+3], ISV[ADAM_BETA2_BASE+3], ISV[ADAM_EPS_BASE+3]
|
||||
* and applies the Invariant 1 floors (β1≥0.5, β2≥0.9, ε≥1e-12). */
|
||||
iqn_beta1: f32,
|
||||
iqn_beta2: f32,
|
||||
iqn_epsilon: f32,
|
||||
) -> Result<f32, MLError> {
|
||||
let effective_stream = override_stream.unwrap_or(&self.stream);
|
||||
let (lt_ws_ptr, lt_ws_size) = override_workspace
|
||||
@@ -1499,9 +1514,11 @@ impl GpuIqnHead {
|
||||
};
|
||||
|
||||
let lr = self.config.lr;
|
||||
let beta1 = self.config.beta1;
|
||||
let beta2 = self.config.beta2;
|
||||
let eps = self.config.epsilon;
|
||||
// SP5 Layer B (Pearl 4): per-group Adam β1/β2/ε threaded from
|
||||
// FusedTrainingCtx via ISV[ADAM_BETA{1,2}_BASE+3] / ISV[ADAM_EPS_BASE+3].
|
||||
let beta1 = iqn_beta1;
|
||||
let beta2 = iqn_beta2;
|
||||
let eps = iqn_epsilon;
|
||||
// SP4 Layer B (Mech 9): AdamW weight_decay sourced from
|
||||
// ISV[WD_RATE[Iqn=3]] by the caller (`FusedTrainingCtx`) and
|
||||
// threaded through `execute_training_pipeline` alongside
|
||||
@@ -1932,6 +1949,73 @@ impl GpuIqnHead {
|
||||
self.config.num_quantiles
|
||||
}
|
||||
|
||||
/// SP5 Pearl 5: Update τ schedule from ISV[IQN_TAU_BASE=250..270).
|
||||
///
|
||||
/// Reads 4 branches × 5 quantiles from ISV, averages across branches per
|
||||
/// quantile slot, applies Invariant-1 cold-start floors (if ISV reads 0,
|
||||
/// fall back to `FIXED_TAUS[q]`), then re-uploads `online_taus`,
|
||||
/// `target_taus`, and `cos_features` in-place via mapped-pinned staging.
|
||||
///
|
||||
/// Must be called once per epoch from `FusedTrainingCtx` before the IQN
|
||||
/// training pipeline runs. Non-fatal on mapped-pinned upload error
|
||||
/// (logs a warning).
|
||||
pub fn refresh_taus_from_isv(
|
||||
&mut self,
|
||||
isv_per_branch_taus: &[f32; 20], // ISV[250..270): branch b, quantile q at [b*5+q]
|
||||
) {
|
||||
let n = self.config.num_quantiles; // 5
|
||||
let b = self.config.batch_size; // batch dimension for tiling
|
||||
let d = self.config.embed_dim; // 64
|
||||
|
||||
// Average 4 branches per quantile, apply cold-start floor from FIXED_TAUS.
|
||||
let mut taus = [0.0_f32; 5];
|
||||
for q in 0..n {
|
||||
let mean = (isv_per_branch_taus[q]
|
||||
+ isv_per_branch_taus[5 + q]
|
||||
+ isv_per_branch_taus[10 + q]
|
||||
+ isv_per_branch_taus[15 + q]) / 4.0_f32;
|
||||
// Invariant-1 cold-start floor: if ISV hasn't been populated yet
|
||||
// (all four branches read 0.0), retain FIXED_TAUS[q] to avoid
|
||||
// collapsing the quantile distribution to 0.
|
||||
taus[q] = if mean < 1e-6_f32 { FIXED_TAUS[q] } else { mean.min(1.0_f32) };
|
||||
}
|
||||
|
||||
// Recompute cosine features [D, N]: cos_feat[dim + q*D] = cos(π·(d+1)·τ_q)
|
||||
let mut cos_feat_host = vec![0.0_f32; d * n];
|
||||
for q in 0..n {
|
||||
let tau_q = taus[q];
|
||||
for dim in 0..d {
|
||||
cos_feat_host[dim + q * d] =
|
||||
(std::f32::consts::PI * ((dim + 1) as f32) * tau_q).cos();
|
||||
}
|
||||
}
|
||||
|
||||
// Tile taus [B, N] — each row is the same N quantile values repeated B times.
|
||||
let mut tiled = Vec::with_capacity(b * n);
|
||||
for _ in 0..b {
|
||||
tiled.extend_from_slice(&taus[..n]);
|
||||
}
|
||||
|
||||
// Upload all three buffers in-place via mapped-pinned (no new allocations).
|
||||
if let Err(e) = super::mapped_pinned::upload_f32_via_pinned(
|
||||
&self.stream, &tiled, &mut self.online_taus,
|
||||
) {
|
||||
tracing::warn!("SP5 Pearl 5 online_taus upload failed (non-fatal): {e}");
|
||||
return;
|
||||
}
|
||||
if let Err(e) = super::mapped_pinned::upload_f32_via_pinned(
|
||||
&self.stream, &tiled, &mut self.target_taus,
|
||||
) {
|
||||
tracing::warn!("SP5 Pearl 5 target_taus upload failed (non-fatal): {e}");
|
||||
return;
|
||||
}
|
||||
if let Err(e) = super::mapped_pinned::upload_f32_via_pinned(
|
||||
&self.stream, &cos_feat_host, &mut self.cos_features,
|
||||
) {
|
||||
tracing::warn!("SP5 Pearl 5 cos_features upload failed (non-fatal): {e}");
|
||||
}
|
||||
}
|
||||
|
||||
/// Compute CVaR-based position scaling from IQN quantiles.
|
||||
///
|
||||
/// Runs the IQN forward-only kernel on `h_s2` (trunk activation),
|
||||
|
||||
@@ -659,6 +659,12 @@ impl GpuTlob {
|
||||
max_grad_norm: f32,
|
||||
weight_clamp_max_abs: f32,
|
||||
weight_decay_value: f32,
|
||||
/* SP5 Layer B (Pearl 4): ISV-derived Adam β1/β2/ε for ParamGroup::Attn=6.
|
||||
* TLOB shares the Attn param-group taxonomy with GpuAttention; same ISV slots.
|
||||
* Caller applies Invariant 1 floors (β1≥0.5, β2≥0.9, ε≥1e-12). */
|
||||
attn_beta1: f32,
|
||||
attn_beta2: f32,
|
||||
attn_epsilon: f32,
|
||||
) -> Result<(), MLError> {
|
||||
let tp = TLOB_TOTAL_PARAMS;
|
||||
let tp_i32 = tp as i32;
|
||||
@@ -695,9 +701,12 @@ impl GpuTlob {
|
||||
self.adam_step += 1;
|
||||
unsafe { *self.t_pinned = self.adam_step; }
|
||||
let adam_blocks = (tp + 255) / 256;
|
||||
let beta1 = 0.9_f32;
|
||||
let beta2 = 0.999_f32;
|
||||
let eps = 1e-8_f32;
|
||||
// SP5 Layer B (Pearl 4): per-group Adam β1/β2/ε threaded from
|
||||
// FusedTrainingCtx via ISV[ADAM_BETA{1,2}_BASE+6] / ISV[ADAM_EPS_BASE+6].
|
||||
// TLOB shares ParamGroup::Attn=6 with GpuAttention.
|
||||
let beta1 = attn_beta1;
|
||||
let beta2 = attn_beta2;
|
||||
let eps = attn_epsilon;
|
||||
// SP4 Layer B (Mech 9): AdamW weight_decay sourced from
|
||||
// ISV[WD_RATE[Attn=6]] by the caller and threaded through
|
||||
// `adam_step`. `.max(0.0)` upstream enforces the AdamW
|
||||
|
||||
@@ -18,6 +18,19 @@
|
||||
#ifndef TRADE_PHYSICS_CUH
|
||||
#define TRADE_PHYSICS_CUH
|
||||
|
||||
/* ── SP5 Layer B: ISV slot constants consumed in unified_env_step_core ── */
|
||||
|
||||
/* Pearl 6 (cross-fold-persistent Kelly): smoothed Kelly fraction.
|
||||
* ISV[KELLY_F_SMOOTH_INDEX] = ISV[280].
|
||||
* Mirrors sp5_isv_slots.rs::KELLY_F_SMOOTH_INDEX. */
|
||||
#define SP5_KELLY_F_SMOOTH_INDEX 280
|
||||
|
||||
/* Pearl 8 (per-direction trailing stop distance): base trail distance per dir.
|
||||
* ISV[TRAIL_DIST_PER_DIR_BASE + dir] = ISV[270..274).
|
||||
* dir: 0=Short, 1=Hold, 2=Long, 3=Flat.
|
||||
* Mirrors sp5_isv_slots.rs::TRAIL_DIST_PER_DIR_BASE. */
|
||||
#define SP5_TRAIL_DIST_PER_DIR_BASE 270
|
||||
|
||||
/* ── Factored action → exposure index ─────────────────────────────────── */
|
||||
|
||||
__device__ __forceinline__ int decode_exposure_index(
|
||||
@@ -639,7 +652,11 @@ __device__ __forceinline__ void unified_env_step_core(
|
||||
int* prev_position_sign_out, /* sign of position BEFORE the trade — for caller's reversal logic */
|
||||
int* actual_dir_out, /* post-enforcement direction (Hold / Long / Short / Flat) — for caller's actions_history */
|
||||
int* actual_mag_out, /* post-enforcement magnitude bucket (0/1/2 for Quarter/Half/Full; 0 if Hold/Flat) */
|
||||
int* trail_triggered_out /* 1 if trailing stop forced position to 0 */
|
||||
int* trail_triggered_out, /* 1 if trailing stop forced position to 0 */
|
||||
/* ── SP5 Layer B: ISV signal bus for Pearl 6 + Pearl 8 ──
|
||||
* NULL-tolerant: when NULL, falls back to static defaults
|
||||
* (0.5f for kelly_f_smooth, 0.005f for trail distance). */
|
||||
const float* __restrict__ isv_signals_ptr /* [SP5_SLOT_END] pinned ISV signals; NULL = static defaults */
|
||||
) {
|
||||
/* ── Step 1: Decode action (4-branch) ── */
|
||||
int dir_idx = decode_direction_4b(action_val, b1_size, b2_size, b3_size);
|
||||
@@ -704,14 +721,31 @@ __device__ __forceinline__ void unified_env_step_core(
|
||||
float dir_conv_clamped = fminf(1.0f, fmaxf(0.0f, conviction));
|
||||
float mag_conv_clamped = fminf(1.0f, fmaxf(0.0f, magnitude_conviction));
|
||||
float policy_conviction = fmaxf(dir_conv_clamped, mag_conv_clamped);
|
||||
float safety = fmaxf(health_safety, policy_conviction);
|
||||
|
||||
/* SP5 Layer B (Pearl 6): cross-fold-persistent smoothed Kelly fraction
|
||||
* from ISV[KELLY_F_SMOOTH_INDEX=280]. When ISV is available and the
|
||||
* slot has been populated (value > 0), incorporate it as an additional
|
||||
* floor on `health_safety` — both are fractions ∈ [0, 1] measuring
|
||||
* long-run sizing confidence. The max() is safe: when Kelly says
|
||||
* "go bigger" than health alone would permit, we trust the data. When
|
||||
* health says "be conservative", health dominates. Bootstrap-safe
|
||||
* because ISV[280] starts at 0 and Pearl 6 uses EWMA α=0.01 — it
|
||||
* stays near zero until real Kelly data accumulate.
|
||||
* EPS_CLAMP_FLOOR = 0.001f: prevents a very small cross-fold Kelly
|
||||
* from being amplified by the max() in a numerically unstable way
|
||||
* (same cold-start pattern as SP4 EPS_CLAMP_FLOOR). */
|
||||
float kelly_f_smooth = (isv_signals_ptr != NULL)
|
||||
? fmaxf(0.001f, fminf(1.0f, isv_signals_ptr[SP5_KELLY_F_SMOOTH_INDEX]))
|
||||
: 0.5f;
|
||||
float health_safety_sp5 = fmaxf(health_safety, kelly_f_smooth);
|
||||
float safety = fmaxf(health_safety_sp5, policy_conviction);
|
||||
target_position = apply_kelly_cap(
|
||||
target_position,
|
||||
*win_count, *loss_count,
|
||||
*sum_wins, *sum_losses,
|
||||
max_position_physics, safety,
|
||||
policy_conviction,
|
||||
health_safety
|
||||
health_safety_sp5
|
||||
);
|
||||
}
|
||||
|
||||
@@ -723,9 +757,19 @@ __device__ __forceinline__ void unified_env_step_core(
|
||||
float unrealized = (*position) * (close - *entry_price);
|
||||
trade_ret = unrealized / prev_equity;
|
||||
}
|
||||
/* SP5 Layer B (Pearl 8): per-direction trailing stop base distance from
|
||||
* ISV[TRAIL_DIST_PER_DIR_BASE + dir_idx] = ISV[270 + dir_idx].
|
||||
* dir_idx ∈ {0=Short,1=Hold,2=Long,3=Flat}; Hold/Flat are excluded by
|
||||
* the enclosing !is_hold_action guard so only Short (0) and Long (2)
|
||||
* reach here in practice. Bootstrap floor: when ISV is NULL or the
|
||||
* slot reads as 0 (pre-first-observation), fall back to the SP4
|
||||
* static value of 0.005f (0.5% trail). */
|
||||
float base_trail_dist = (isv_signals_ptr != NULL)
|
||||
? fmaxf(0.001f, isv_signals_ptr[SP5_TRAIL_DIST_PER_DIR_BASE + dir_idx])
|
||||
: 0.005f;
|
||||
if (check_trailing_stop(
|
||||
*hold_time, *max_equity, prev_equity,
|
||||
trade_ret, 0.005f, vol_scale, trend_scale
|
||||
trade_ret, base_trail_dist, vol_scale, trend_scale
|
||||
)) {
|
||||
target_position = 0.0f;
|
||||
trail_triggered = 1;
|
||||
|
||||
@@ -39,6 +39,13 @@ use crate::cuda_pipeline::gpu_tlob::GpuTlob;
|
||||
use crate::cuda_pipeline::gpu_dqn_trainer::{
|
||||
GpuDqnTrainConfig, GpuDqnTrainer, TAU_EFF_INDEX,
|
||||
};
|
||||
use crate::cuda_pipeline::sp5_isv_slots::{
|
||||
BUDGET_C51_BASE, BUDGET_IQN_BASE, BUDGET_CQL_BASE, BUDGET_ENS_BASE,
|
||||
NOISY_SIGMA_BASE,
|
||||
ADAM_BETA1_BASE, ADAM_BETA2_BASE, ADAM_EPS_BASE,
|
||||
IQN_TAU_BASE,
|
||||
ATOM_NUM_ATOMS_BASE,
|
||||
};
|
||||
use crate::cuda_pipeline::gpu_her::{GpuHer, GpuHerConfig, parse_her_strategy};
|
||||
use crate::cuda_pipeline::gpu_iql_trainer::{GpuIqlConfig, GpuIqlTrainer};
|
||||
use crate::cuda_pipeline::gpu_iqn_head::{GpuIqnConfig, GpuIqnHead};
|
||||
@@ -1624,12 +1631,20 @@ impl FusedTrainingCtx {
|
||||
// non-negativity.
|
||||
let (tlob_weight_clamp_max_abs, tlob_weight_decay_value) =
|
||||
self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::Attn);
|
||||
// SP5 Layer B (Pearl 4): TLOB shares ParamGroup::Attn=6.
|
||||
let tlob_g = crate::cuda_pipeline::ParamGroup::Attn.idx();
|
||||
let tlob_beta1 = self.trainer.read_isv_signal_at(ADAM_BETA1_BASE + tlob_g).max(0.5_f32).min(0.9999_f32);
|
||||
let tlob_beta2 = self.trainer.read_isv_signal_at(ADAM_BETA2_BASE + tlob_g).max(0.9_f32).min(0.99999_f32);
|
||||
let tlob_epsilon = self.trainer.read_isv_signal_at(ADAM_EPS_BASE + tlob_g).max(1e-12_f32);
|
||||
self.gpu_tlob
|
||||
.adam_step(
|
||||
self.trainer.config().lr,
|
||||
self.trainer.config().max_grad_norm,
|
||||
tlob_weight_clamp_max_abs,
|
||||
tlob_weight_decay_value,
|
||||
tlob_beta1,
|
||||
tlob_beta2,
|
||||
tlob_epsilon,
|
||||
)
|
||||
.map_err(|e| anyhow::anyhow!("TLOB Adam: {e}"))?;
|
||||
}
|
||||
@@ -1982,6 +1997,17 @@ impl FusedTrainingCtx {
|
||||
// ── IQN prep: decode actions + DtoD copies (main stream, before fork) ──
|
||||
// Must complete before execute_training_pipeline runs on iqn_stream.
|
||||
if let Some(ref mut iqn) = self.gpu_iqn {
|
||||
// SP5 Pearl 5: refresh τ schedule from ISV[IQN_TAU_BASE=250..270).
|
||||
// Read 4 branches × 5 quantiles, average per quantile, apply cold-start
|
||||
// floor from FIXED_TAUS, re-upload online_taus/target_taus/cos_features.
|
||||
// Must run before prepare_buffers so the kernel sees updated tau values.
|
||||
{
|
||||
let mut isv_taus = [0.0_f32; 20];
|
||||
for i in 0..20 {
|
||||
isv_taus[i] = self.trainer.read_isv_signal_at(IQN_TAU_BASE + i);
|
||||
}
|
||||
iqn.refresh_taus_from_isv(&isv_taus);
|
||||
}
|
||||
iqn.set_cached_target_h_s2(self.trainer.tg_h_s2_ptr());
|
||||
let dqn_actions = self.trainer.actions_buf();
|
||||
let dqn_rewards = self.trainer.rewards_buf();
|
||||
@@ -2031,12 +2057,20 @@ impl FusedTrainingCtx {
|
||||
// non-negativity).
|
||||
let (weight_clamp_max_abs, weight_decay_value) =
|
||||
self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::Iqn);
|
||||
// SP5 Layer B (Pearl 4): per-group Adam β1/β2/ε for Iqn=3.
|
||||
let iqn_g = crate::cuda_pipeline::ParamGroup::Iqn.idx();
|
||||
let iqn_beta1 = self.trainer.read_isv_signal_at(ADAM_BETA1_BASE + iqn_g).max(0.5_f32).min(0.9999_f32);
|
||||
let iqn_beta2 = self.trainer.read_isv_signal_at(ADAM_BETA2_BASE + iqn_g).max(0.9_f32).min(0.99999_f32);
|
||||
let iqn_epsilon = self.trainer.read_isv_signal_at(ADAM_EPS_BASE + iqn_g).max(1e-12_f32);
|
||||
|
||||
match iqn.execute_training_pipeline(
|
||||
online_h_s2, next_states_buf, &self.target_dueling,
|
||||
Some(&iqn_stream_ref), ws_override,
|
||||
weight_clamp_max_abs,
|
||||
weight_decay_value,
|
||||
iqn_beta1,
|
||||
iqn_beta2,
|
||||
iqn_epsilon,
|
||||
) {
|
||||
Ok(_) => { iqn_ok = true; }
|
||||
Err(e) => {
|
||||
@@ -2113,8 +2147,14 @@ impl FusedTrainingCtx {
|
||||
// non-negativity).
|
||||
let (attn_weight_clamp_max_abs, attn_weight_decay_value) =
|
||||
self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::Attn);
|
||||
// SP5 Layer B (Pearl 4): per-group Adam β1/β2/ε for Attn=6.
|
||||
let attn_g = crate::cuda_pipeline::ParamGroup::Attn.idx();
|
||||
let attn_beta1 = self.trainer.read_isv_signal_at(ADAM_BETA1_BASE + attn_g).max(0.5_f32).min(0.9999_f32);
|
||||
let attn_beta2 = self.trainer.read_isv_signal_at(ADAM_BETA2_BASE + attn_g).max(0.9_f32).min(0.99999_f32);
|
||||
let attn_epsilon = self.trainer.read_isv_signal_at(ADAM_EPS_BASE + attn_g).max(1e-12_f32);
|
||||
attn.adam_step(lr, mgn, Some(&attn_stream_ref), ws_override,
|
||||
attn_weight_clamp_max_abs, attn_weight_decay_value)
|
||||
attn_weight_clamp_max_abs, attn_weight_decay_value,
|
||||
attn_beta1, attn_beta2, attn_epsilon)
|
||||
.map_err(|e| anyhow::anyhow!("Attention Adam: {e}"))?;
|
||||
|
||||
// Record completion on attn_stream.
|
||||
@@ -2151,8 +2191,14 @@ impl FusedTrainingCtx {
|
||||
// Sequential-fallback path.
|
||||
let (attn_weight_clamp_max_abs, attn_weight_decay_value) =
|
||||
self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::Attn);
|
||||
// SP5 Layer B (Pearl 4): per-group Adam β1/β2/ε for Attn=6.
|
||||
let attn_g = crate::cuda_pipeline::ParamGroup::Attn.idx();
|
||||
let attn_beta1 = self.trainer.read_isv_signal_at(ADAM_BETA1_BASE + attn_g).max(0.5_f32).min(0.9999_f32);
|
||||
let attn_beta2 = self.trainer.read_isv_signal_at(ADAM_BETA2_BASE + attn_g).max(0.9_f32).min(0.99999_f32);
|
||||
let attn_epsilon = self.trainer.read_isv_signal_at(ADAM_EPS_BASE + attn_g).max(1e-12_f32);
|
||||
attn.adam_step(lr, mgn, None, None,
|
||||
attn_weight_clamp_max_abs, attn_weight_decay_value)
|
||||
attn_weight_clamp_max_abs, attn_weight_decay_value,
|
||||
attn_beta1, attn_beta2, attn_epsilon)
|
||||
.map_err(|e| anyhow::anyhow!("Attention Adam: {e}"))?;
|
||||
}
|
||||
|
||||
@@ -2165,12 +2211,20 @@ impl FusedTrainingCtx {
|
||||
// parallel arm above. Mirror the dqn_adam pattern.
|
||||
let (weight_clamp_max_abs, weight_decay_value) =
|
||||
self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::Iqn);
|
||||
// SP5 Layer B (Pearl 4): per-group Adam β1/β2/ε for Iqn=3.
|
||||
let iqn_g = crate::cuda_pipeline::ParamGroup::Iqn.idx();
|
||||
let iqn_beta1 = self.trainer.read_isv_signal_at(ADAM_BETA1_BASE + iqn_g).max(0.5_f32).min(0.9999_f32);
|
||||
let iqn_beta2 = self.trainer.read_isv_signal_at(ADAM_BETA2_BASE + iqn_g).max(0.9_f32).min(0.99999_f32);
|
||||
let iqn_epsilon = self.trainer.read_isv_signal_at(ADAM_EPS_BASE + iqn_g).max(1e-12_f32);
|
||||
|
||||
match iqn.execute_training_pipeline(
|
||||
online_h_s2, next_states_buf, &self.target_dueling,
|
||||
None, None,
|
||||
weight_clamp_max_abs,
|
||||
weight_decay_value,
|
||||
iqn_beta1,
|
||||
iqn_beta2,
|
||||
iqn_epsilon,
|
||||
) {
|
||||
Ok(_) => { iqn_ok = true; }
|
||||
Err(e) => {
|
||||
@@ -3214,20 +3268,38 @@ impl FusedTrainingCtx {
|
||||
/// Returns (c51_budget, iqn_budget, cql_budget, ens_budget). Caches to trainer fields
|
||||
/// for HEALTH_DIAG logging.
|
||||
///
|
||||
/// Formula:
|
||||
/// iqn_budget = 0.10 + 0.30 × health
|
||||
/// cql_budget = 0.10 × (1 − regime_stability) × health
|
||||
/// ens_budget = 0.05 (constant)
|
||||
/// c51_budget = 1.0 − iqn_budget − cql_budget − ens_budget
|
||||
/// SP5 Layer B (Pearl 2): reads per-branch budget signals from ISV[190..210)
|
||||
/// and collapses them to a single effective scale per loss component by
|
||||
/// averaging across the 4 branches. The Layer A producer (SP5 Task A3) runs
|
||||
/// the Wiener-optimal EMA (Pearls A+D) so the values already encode
|
||||
/// health/regime dynamics — no second-hand formula is needed here.
|
||||
///
|
||||
/// At health=1: iqn=0.40, cql variable (depends on regime_stability), ens=0.05, c51=rest.
|
||||
/// At health=0: iqn=0.10, cql=0, ens=0.05, c51=0.85 (C51 takes over).
|
||||
/// Cold-start floor: ISV reads 0 before the first observation; the
|
||||
/// normalisation `/ 4.0` and clamp ensure we never pass 0 to the loss
|
||||
/// scalers. The floor values (0.05 C51, 0.05 IQN, 0.02 CQL, 0.02 Ens)
|
||||
/// are Invariant 1 carve-outs (numerical-stability), not tuned constants.
|
||||
pub(crate) fn compute_adaptive_budgets(&mut self) -> (f32, f32, f32, f32) {
|
||||
let (health, regime_stability) = self.trainer.read_isv_health_and_regime();
|
||||
let iqn_budget = 0.10_f32 + 0.30_f32 * health;
|
||||
let cql_budget = 0.10_f32 * (1.0 - regime_stability) * health;
|
||||
let ens_budget = 0.05_f32;
|
||||
let c51_budget = (1.0_f32 - iqn_budget - cql_budget - ens_budget).max(0.0);
|
||||
// Mean of 4 per-branch slots → single effective scale.
|
||||
let c51_budget = ((0..4_usize)
|
||||
.map(|b| self.read_isv_signal_at(BUDGET_C51_BASE + b))
|
||||
.sum::<f32>()
|
||||
/ 4.0_f32)
|
||||
.max(0.05_f32);
|
||||
let iqn_budget = ((0..4_usize)
|
||||
.map(|b| self.read_isv_signal_at(BUDGET_IQN_BASE + b))
|
||||
.sum::<f32>()
|
||||
/ 4.0_f32)
|
||||
.max(0.05_f32);
|
||||
let cql_budget = ((0..4_usize)
|
||||
.map(|b| self.read_isv_signal_at(BUDGET_CQL_BASE + b))
|
||||
.sum::<f32>()
|
||||
/ 4.0_f32)
|
||||
.max(0.02_f32);
|
||||
let ens_budget = ((0..4_usize)
|
||||
.map(|b| self.read_isv_signal_at(BUDGET_ENS_BASE + b))
|
||||
.sum::<f32>()
|
||||
/ 4.0_f32)
|
||||
.max(0.02_f32);
|
||||
|
||||
self.trainer.last_iqn_budget_eff = iqn_budget;
|
||||
self.trainer.last_cql_budget_eff = cql_budget;
|
||||
|
||||
@@ -26,6 +26,12 @@ use crate::cuda_pipeline::gpu_dqn_trainer::{
|
||||
EPOCH_IDX_INDEX, EPSILON_EFF_INDEX, GAMMA_DIR_EFF_INDEX, Q_ABS_REF_INDEX,
|
||||
Q_DIR_ABS_REF_INDEX, TLOB_REGIME_FOCUS_EMA_INDEX,
|
||||
};
|
||||
use crate::cuda_pipeline::sp5_isv_slots::{
|
||||
NOISY_SIGMA_BASE,
|
||||
ADAM_BETA1_BASE as SP5_ADAM_BETA1_BASE,
|
||||
ADAM_BETA2_BASE as SP5_ADAM_BETA2_BASE,
|
||||
ADAM_EPS_BASE as SP5_ADAM_EPS_BASE,
|
||||
};
|
||||
use crate::dqn::logging::{log_epoch_start, log_epoch_end, log_training_progress};
|
||||
use crate::dqn::target_update::convergence_half_life;
|
||||
use crate::evaluation::metrics::calculate_var_cvar;
|
||||
@@ -1734,7 +1740,17 @@ impl DQNTrainer {
|
||||
max_trace_length: self.hyperparams.max_trace_length as i32,
|
||||
hindsight_fraction: self.hyperparams.hindsight_fraction as f32,
|
||||
hindsight_lookahead: self.hyperparams.hindsight_lookahead as i32,
|
||||
noise_sigma: self.hyperparams.noise_sigma as f32,
|
||||
// SP5 Layer B (Pearl 3): read per-branch NoisyNet σ from ISV[210..214)
|
||||
// and collapse to a single effective scale by averaging across 4 branches.
|
||||
// Falls back to the static hyperparam when fused_ctx is unavailable
|
||||
// (pre-initialization path). 0.01f floor = Invariant 1 carve-out.
|
||||
noise_sigma: self.fused_ctx.as_ref().map(|ctx| {
|
||||
let mean = (0..4_usize)
|
||||
.map(|b| ctx.read_isv_signal_at(NOISY_SIGMA_BASE + b))
|
||||
.sum::<f32>()
|
||||
/ 4.0_f32;
|
||||
mean.max(0.01_f32)
|
||||
}).unwrap_or(self.hyperparams.noise_sigma as f32),
|
||||
..Default::default()
|
||||
};
|
||||
|
||||
@@ -1793,9 +1809,24 @@ impl DQNTrainer {
|
||||
// unbounded clamps.
|
||||
0.0_f32
|
||||
};
|
||||
// SP5 Pearl 4: per-group Adam β1/β2/ε for Curiosity (ParamGroup
|
||||
// index 7). ISV[ADAM_BETA1_BASE+7=233], ISV[ADAM_BETA2_BASE+7=241],
|
||||
// ISV[ADAM_EPS_BASE+7=249]. Curiosity collector doesn't own ISV —
|
||||
// read here and thread through to `train_on_collector_buffers`.
|
||||
let cur_g = crate::cuda_pipeline::ParamGroup::Curiosity.idx();
|
||||
let (cur_beta1, cur_beta2, cur_epsilon) = if let Some(ref fused) = self.fused_ctx {
|
||||
(
|
||||
fused.read_isv_signal_at(SP5_ADAM_BETA1_BASE + cur_g).max(0.5_f32).min(0.9999_f32),
|
||||
fused.read_isv_signal_at(SP5_ADAM_BETA2_BASE + cur_g).max(0.9_f32).min(0.99999_f32),
|
||||
fused.read_isv_signal_at(SP5_ADAM_EPS_BASE + cur_g).max(1e-12_f32),
|
||||
)
|
||||
} else {
|
||||
(0.9_f32, 0.999_f32, 1e-8_f32)
|
||||
};
|
||||
if let Err(e) = collector.train_curiosity_gpu(
|
||||
gpu_batch.n_episodes, gpu_batch.timesteps,
|
||||
weight_clamp_max_abs,
|
||||
cur_beta1, cur_beta2, cur_epsilon,
|
||||
) {
|
||||
debug!("GPU curiosity training failed (non-fatal): {e}");
|
||||
}
|
||||
|
||||
@@ -3190,3 +3190,42 @@ Refs: SP4 Layer C close-out follow-up — sites #2 + #9 of the sweep audit grid
|
||||
- 2 unit tests pass (slot layout invariants + cross-fold-persistent Kelly carve-out)
|
||||
|
||||
No behavior change. No producer/consumer wired yet. Subsequent Layer A commits (A1–A8) populate slots via per-pearl producer kernels.
|
||||
|
||||
### SP5 Layer B — atomic consumer migration to ISV-driven adaptive signals
|
||||
|
||||
All 11 consumer sites migrated in one atomic commit per `feedback_no_partial_refactor`. Every site that reads from ISV slots 174..286 now does so at runtime via `read_isv_signal_at` (Rust, CPU-side mapped-pinned read) or direct ISV pointer dereference (CUDA kernel). No static constants remain for any of the 8 pearls covered.
|
||||
|
||||
**Pearl 1 (C51 atom span — v_center/v_half):**
|
||||
- `atoms_update_kernel.cu`: reads `ISV[174+b]` (v_center) and `ISV[178+b]` (v_half, floor 1.0) per branch.
|
||||
|
||||
**Pearl 1-ext (per-branch num_atoms):**
|
||||
- `atoms_update_kernel.cu`: reads `ISV[274+b]`, calls new `round_to_valid_num_atoms()` device inline to snap to valid {16,32,64} with cold-start fallback to global `num_atoms`. Uses `eff_na` for all softmax/cumsum iteration; output stride remains global `num_atoms` (buffer invariant).
|
||||
|
||||
**Pearl 2 (per-branch loss budgets):**
|
||||
- `fused_training.rs` `compute_adaptive_budgets()`: reads ISV[190..194), ISV[194..198), ISV[198..202), ISV[202..206) per branch, averages 4 branches per loss type. Invariant-1 floors: C51=0.05, IQN=0.05, CQL=0.02, Ens=0.02.
|
||||
|
||||
**Pearl 3 (per-branch NoisyNet sigma):**
|
||||
- `training_loop.rs`: `ExperienceCollectorConfig.noise_sigma` reads `mean(ISV[210..214])` with 0.01 floor; falls back to `hyperparams.noise_sigma` when `fused_ctx` absent.
|
||||
|
||||
**Pearl 4 (per-group Adam β1/β2/ε — 5 sites):**
|
||||
- `gpu_dqn_trainer.rs` `launch_adam_update`: reads ISV[226+g], ISV[234+g], ISV[242+g] per group `g` inside the existing per-group loop. Static `ADAM_BETA1/BETA2/EPS` remain declared but are unused (dead code left for reference; may be removed in cleanup).
|
||||
- `gpu_iqn_head.rs` `execute_training_pipeline` + `train_iqn_step_gpu`: 3 new `iqn_beta1/beta2/epsilon` parameters threaded from callers.
|
||||
- `gpu_attention.rs` `adam_step`: 3 new `attn_beta1/beta2/epsilon` parameters replacing hardcoded 0.9/0.999/1e-8.
|
||||
- `gpu_tlob.rs` `adam_step`: same pattern as GpuAttention (shares ParamGroup::Attn=6).
|
||||
- `fused_training.rs`: 4 call sites (2 IQN, 1 GpuAttention, 1 TLOB) read ISV[226+g..242+g] and thread values through.
|
||||
- `gpu_curiosity_trainer.rs` `launch_adam_step` + `train_on_collector_buffers`: `cur_beta1/beta2/epsilon` params replace constant references.
|
||||
- `gpu_experience_collector.rs` `train_curiosity_gpu`: threads `cur_beta1/beta2/epsilon`.
|
||||
- `training_loop.rs`: reads ISV[233/241/249] for `ParamGroup::Curiosity=7` at the `train_curiosity_gpu` call site.
|
||||
|
||||
**Pearl 5 (per-branch IQN τ schedule):**
|
||||
- `gpu_iqn_head.rs` new `refresh_taus_from_isv(&[f32; 20])`: averages 4 branches per quantile slot (ISV[250+b*5+q]), applies cold-start floor from `FIXED_TAUS[q]` (ISV=0 → retain original), recomputes `cos_features [D,N]` from updated taus, re-uploads `online_taus`, `target_taus`, `cos_features` in-place via `upload_f32_via_pinned`.
|
||||
- `fused_training.rs`: reads ISV[250..270) into `[f32; 20]` array, calls `iqn.refresh_taus_from_isv(&isv_taus)` before `prepare_buffers` each training step.
|
||||
|
||||
**Pearl 6 (cross-fold Kelly cap) + Pearl 8 (per-direction trail distance):**
|
||||
- `trade_physics.cuh`, `experience_kernels.cu`, `backtest_env_kernel.cu`: done in prior session.
|
||||
|
||||
**Compilation and tests:**
|
||||
- `SQLX_OFFLINE=true cargo check -p ml --offline` → clean (12 warnings, all pre-existing)
|
||||
- `cargo test -p ml --lib` → 930 pass, 14 pre-existing failures unchanged
|
||||
|
||||
Refs: `feedback_no_partial_refactor`, `feedback_isv_for_adaptive_bounds`, `feedback_adaptive_not_tuned`
|
||||
|
||||
Reference in New Issue
Block a user