gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

fix(class-a-p1-producer): adaptive Bayesian Kelly priors per-fold-end

Browse Source 
Replace 4 hardcoded Bayesian priors in Kelly cap calculations
(prior_wins=2.0, prior_losses=2.0, prior_sum_wins=0.01,
prior_sum_losses=0.01) with ISV-driven slow-EMA values fed by a new
producer kernel that aggregates realized PS_KELLY_* fields across envs
from the same portfolio_state buffer kelly_cap_update_kernel reads
from at the same per-epoch boundary.

Slots [454..458):
- KELLY_PRIOR_WINS, KELLY_PRIOR_LOSSES,
  KELLY_PRIOR_SUM_WINS, KELLY_PRIOR_SUM_LOSSES

ISV_TOTAL_DIM bumps 454 -> 458.

Producer:
- kelly_bayesian_priors_update_kernel.cu (per-fold-end / per-epoch
  boundary, block-tree-reduce in shmem, no atomicAdd). Pearl-A
  first-observation bootstrap (sentinels 2.0/2.0/0.01/0.01 match
  pre-P1-Producer hardcoded values for bit-identical cold-start) +
  alpha=0.005 slow EMA. Bounds counts in [0.5, 100], sums in [0.001,
  1.0] per feedback_isv_for_adaptive_bounds. Launched RIGHT BEFORE
  launch_kelly_cap_update so that kernel sees the freshly-blended
  priors.

Consumers migrated atomically (same commit):
- kelly_cap_update_kernel.cu:39-42 -> ISV[454..458) with cold-start
  fallback via kelly_prior_or_default helper (range guard).
- trade_physics.cuh::kelly_position_cap:304-307 -> NULL-tolerant
  isv_signals_ptr threaded through apply_kelly_cap (single caller in
  unified_env_step_core line 898 already had the bus pointer; mirrors
  the existing kelly_f_smooth / kelly_warmup_floor_sp9 patterns).

State reset:
- 4 FoldReset registry entries (sp14_p1_kelly_prior_*) +
  4 dispatch arms in training_loop.rs::reset_named_state.

Oracle tests (sp14_oracle_tests.rs, GPU-gated #[ignore]):
- Pearl-A bootstrap (sentinel -> REPLACE with aggregated targets)
- No realized trades -> ISV preserved bit-exactly
- Bounds clamp on extreme aggregates
- Slow EMA blend after bootstrap

CPU tests passing:
- sp14_p1_kelly_prior_slot_layout_locked
- all_sp14_p1_slots_fit_within_isv_total_dim
- every_fold_and_soft_reset_entry_has_dispatch_arm (C.10 lesson)
- layout_fingerprint_bumps_after_sp14_wire

DEFERRED — Item 2 (MIN_HOLD_TEMPERATURE EMA): the audit-spec said
"MIN_HOLD_TEMPERATURE = 0.5f hardcoded somewhere" but the actual code
has MIN_HOLD_TEMPERATURE_{START=50.0f, END=5.0f, DECAY=20.0f} as a
PER-EPOCH ANNEALING SCHEDULE driven by min_hold_temperature_for_epoch
in training_loop.rs:68-73. The kernel takes T as a runtime scalar
specifically to enable Phase 2 ISV-driven lift "without recompiling
cubin" per the SP12 v3 design comment. The audit's claim of "0.5f
hardcoded" does not match reality; the right Phase 2 signal is a
separate spec decision and should not be guessed at per
feedback_no_quickfixes. Reporting back per the prompt's hard rule #8.

Per feedback_isv_for_adaptive_bounds + feedback_no_partial_refactor +
pearl_controller_anchors_isv_driven.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-08 09:38:51 +02:00
parent 657972a4b5
commit 87d597d5d7
12 changed files with 1376 additions and 23 deletions
Download Patch File Download Diff File Expand all files Collapse all files

21
crates/ml/build.rs
View File

@@ -806,6 +806,27 @@ fn main() {
// WR-stuck-at-46-48% plateau across 11 superprojects. Per-epoch
// boundary launch.
"reward_cap_update_kernel.cu",
// Class A P1-Producer (2026-05-08): adaptive Bayesian Kelly priors
// producer. Single-block kernel that sweeps the per-fold-end
// `portfolio_state[n_envs, PS_STRIDE]` buffer (same source the
// `kelly_cap_update_kernel` reads from at the same boundary),
// aggregates the realized PS_KELLY_{WIN_COUNT, LOSS_COUNT,
// SUM_WINS, SUM_LOSSES} fields across envs, and slow-EMA-blends
// into ISV[454..458). Pearl-A first-observation bootstrap
// (sentinels 2.0/2.0/0.01/0.01 match pre-P1-Producer hardcoded
// values for bit-identical cold-start); α=0.005 slow EMA
// thereafter (per-fold cadence — the Bayesian prior is a *prior
// belief* and should change slowly across folds; a fast
// "prior" would just be a noisy Kelly estimate). Bounds
// counts∈[0.5, 100], sums∈[0.001, 1.0] are Category-1 dimensional
// safety floors per `feedback_isv_for_adaptive_bounds`. Replaces
// hardcoded `prior_wins=2.0f / prior_losses=2.0f /
// prior_sum_wins=0.01f / prior_sum_losses=0.01f` in
// kelly_cap_update_kernel.cu:39-42 + trade_physics.cuh::
// kelly_position_cap:304-307 — Class A audit P1-Producer batch.
// Block-tree-reduce (no atomicAdd) per
// `feedback_no_atomicadd.md`. Per-epoch boundary launch.
"kelly_bayesian_priors_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)

125
crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs
View File

@@ -52,7 +52,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, H_S2_AUX_RMS_EMA_CUBIN, REWARD_CAP_UPDATE_CUBIN,
AVG_WIN_HOLD_TIME_UPDATE_CUBIN, H_S2_AUX_RMS_EMA_CUBIN, KELLY_BAYESIAN_PRIORS_UPDATE_CUBIN,
REWARD_CAP_UPDATE_CUBIN,
};
/// Hidden width of the aux trunk's first internal layer (Linear_1 → ELU
@@ -734,6 +735,128 @@ impl RewardCapUpdateOps {
}
}
/// Class A P1-Producer (2026-05-08): adaptive Bayesian Kelly priors producer.
///
/// Replaces the previously-hardcoded `prior_wins=2.0f, prior_losses=2.0f,
/// prior_sum_wins=0.01f, prior_sum_losses=0.01f` constants used in
/// `kelly_cap_update_kernel.cu` and `trade_physics.cuh::kelly_position_cap`
/// with ISV-driven slow-EMA values fed by the
/// `kelly_bayesian_priors_update_kernel` producer at per-fold-end (per-epoch
/// boundary in the current scheduling). Aggregates the realized
/// PS_KELLY_{WIN_COUNT, LOSS_COUNT, SUM_WINS, SUM_LOSSES} fields across all
/// envs from the same `portfolio_state[n_envs, PS_STRIDE]` buffer that
/// `kelly_cap_update_kernel` reads from at the same boundary, then
/// slow-EMA-blends into ISV[KELLY_PRIOR_*_INDEX] = ISV[454..458).
///
/// Single-block 256-thread kernel; shmem block-tree-reduce
/// (no atomicAdd per `feedback_no_atomicadd.md`). `CudaFunction` pre-loaded
/// at construction per `pearl_no_host_branches_in_captured_graph.md`.
///
/// Pearls applied:
/// - `feedback_no_atomicadd.md` — block-tree-reduce in shmem, single
/// global write per slot from thread 0.
/// - `pearl_first_observation_bootstrap.md` — sentinel match (within 1e-6
/// of pre-P1-Producer hardcoded value) → REPLACE; otherwise EMA blend.
/// - `pearl_no_host_branches_in_captured_graph.md` — `CudaFunction`
/// pre-loaded at construction; on-device guard for "no realized trades".
/// - `feedback_no_stubs.md` — full body, no placeholder.
/// - `feedback_isv_for_adaptive_bounds.md` — count bounds [0.5, 100],
/// sum bounds [0.001, 1.0] are dimensional safety floors, NOT tuning.
/// - `pearl_symmetric_clamp_audit.md` — bilateral
/// `fmaxf(lo, fminf(x, hi))` clamp on each slot before writing.
/// - `pearl_controller_anchors_isv_driven.md` — every controller anchor
/// (priors are anchors for the Kelly cap formula) is ISV-driven.
#[allow(missing_debug_implementations)]
pub(crate) struct KellyBayesianPriorsUpdateOps {
update_kernel: CudaFunction,
}
impl KellyBayesianPriorsUpdateOps {
/// Block dim used by the producer kernel — must match `BLK_DIM` in
/// `kelly_bayesian_priors_update_kernel.cu`.
const BLK_DIM: u32 = 256;
pub(crate) fn new(stream: &Arc<CudaStream>) -> Result<Self, MLError> {
let context = stream.context();
let module = context
.load_cubin(KELLY_BAYESIAN_PRIORS_UPDATE_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!("kelly_bayesian_priors_update cubin load: {e}")))?;
let update_kernel = module
.load_function("kelly_bayesian_priors_update")
.map_err(|e| MLError::ModelError(format!("kelly_bayesian_priors_update load: {e}")))?;
Ok(Self { update_kernel })
}
/// Launch the adaptive Bayesian-priors producer.
///
/// Args:
/// - `portfolio_state_ptr`: f32 device ptr `[n_envs * ps_stride]` —
/// the same buffer `kelly_cap_update_kernel` reads from. We use
/// the four PS_KELLY_* fields per env.
/// - `n_envs`: number of envs (alloc_episodes from the collector).
/// - `ps_stride`: PS_STRIDE from state_layout.cuh (43 post-Plan-3-D.4c).
/// - `isv_ptr`: ISV[ISV_TOTAL_DIM] device pointer (mapped-pinned).
/// - `wins_idx`/`losses_idx`/`sum_wins_idx`/`sum_losses_idx`:
/// KELLY_PRIOR_*_INDEX (454..458).
/// - `sentinel_*`: SENTINEL_KELLY_PRIOR_* (2.0/2.0/0.01/0.01).
/// - `count_min`/`count_max`: KELLY_PRIOR_COUNT_MIN/MAX (0.5, 100).
/// - `sum_min`/`sum_max`: KELLY_PRIOR_SUM_MIN/MAX (0.001, 1.0).
/// - `alpha`: KELLY_PRIOR_EMA_ALPHA (0.005, slow per-fold cadence).
#[allow(clippy::too_many_arguments)]
pub(crate) fn launch(
&self,
stream: &Arc<CudaStream>,
portfolio_state_ptr: u64,
n_envs: i32,
ps_stride: i32,
isv_ptr: u64,
wins_idx: i32,
losses_idx: i32,
sum_wins_idx: i32,
sum_losses_idx: i32,
sentinel_wins: f32,
sentinel_losses: f32,
sentinel_sum_wins: f32,
sentinel_sum_losses: f32,
count_min: f32,
count_max: f32,
sum_min: f32,
sum_max: f32,
alpha: f32,
) -> Result<(), MLError> {
// 4 arrays × BLK_DIM × sizeof(f32) — see the kernel comment block.
let smem_bytes = 4 * Self::BLK_DIM * std::mem::size_of::<f32>() as u32;
unsafe {
stream
.launch_builder(&self.update_kernel)
.arg(&portfolio_state_ptr)
.arg(&n_envs)
.arg(&ps_stride)
.arg(&isv_ptr)
.arg(&wins_idx)
.arg(&losses_idx)
.arg(&sum_wins_idx)
.arg(&sum_losses_idx)
.arg(&sentinel_wins)
.arg(&sentinel_losses)
.arg(&sentinel_sum_wins)
.arg(&sentinel_sum_losses)
.arg(&count_min)
.arg(&count_max)
.arg(&sum_min)
.arg(&sum_max)
.arg(&alpha)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (Self::BLK_DIM, 1, 1),
shared_mem_bytes: smem_bytes,
})
.map_err(|e| MLError::ModelError(format!("kelly_bayesian_priors_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

109
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
View File

File diff suppressed because one or more lines are too long

250
crates/ml/src/cuda_pipeline/kelly_bayesian_priors_update_kernel.cu Normal file
View File

@@ -0,0 +1,250 @@
/* ══════════════════════════════════════════════════════════════════════════
* Class A P1-Producer — adaptive Bayesian Kelly priors producer (2026-05-08).
*
* Replaces the hardcoded Bayesian priors used in Kelly cap calculations:
* - `kelly_cap_update_kernel.cu:39-42` (epoch-boundary effective Kelly cap)
* - `trade_physics.cuh::kelly_position_cap:304-307` (per-step env cap)
*
* Both consumers ran with `prior_wins=2.0f, prior_losses=2.0f,
* prior_sum_wins=0.01f, prior_sum_losses=0.01f` — a static "I don't know"
* Bayesian prior frozen at sprint-1 defaults. The cold-start of every new
* fold ignored what prior folds had learned about the realized Kelly
* distribution; the prior was *literally* the prior (in the Bayesian sense),
* but it was a fixed scalar, not an adaptive belief.
*
* This producer aggregates realized Kelly stats from the existing
* `portfolio_state[n_envs, PS_STRIDE]` buffer — same source the
* `kelly_cap_update_kernel` already reads from at the same per-epoch
* boundary. The four PS slots (`PS_KELLY_WIN_COUNT`,
* `PS_KELLY_LOSS_COUNT`, `PS_KELLY_SUM_WINS`, `PS_KELLY_SUM_LOSSES`) are
* realized counts/sums per env; we aggregate across envs and slow-EMA-blend
* into ISV[454..458). On the next fold's cold-start path, the consumer
* reads the EMA prior (a learned belief about the distribution shape) instead
* of the fixed sprint-1 default.
*
* Ordering: launched per-fold-end (currently per-epoch boundary in the
* scheduling chain in `training_loop.rs`) BEFORE `kelly_cap_update` so the
* effective-Kelly-cap kernel sees the freshly-blended priors. The launch
* site lives at the same epoch-boundary block as `kelly_cap_update` and
* `reward_cap_update`.
*
* Algorithm:
*
* Phase 1 (block-tree-reduce): each thread of the single block sweeps
* its tile of `n_envs` portfolio_state slots, accumulating
* local_win_count, local_loss_count, local_sum_wins, local_sum_losses
* from the four PS_KELLY_* fields. Block-tree-reduce in shmem (no
* atomicAdd per `feedback_no_atomicadd.md`) yields the four sums in
* thread 0.
*
* Phase 2 (single thread, EMA blend): if no realized trades observed
* (total_win_count + total_loss_count == 0), keep ISV slots unchanged
* (cold-start sentinel persists; consumers fall back to the static
* defaults via the NULL-tolerant guard in `kelly_position_cap` /
* `kelly_cap_update_kernel`). Otherwise:
*
* target_wins = total_win_count
* target_losses = total_loss_count
* target_sum_wins = total_sum_wins
* target_sum_losses = total_sum_losses
*
* Pearl-A first-observation bootstrap: each slot independently checks
* `|current - sentinel| < EPS_F`; if so, REPLACE with target. Otherwise
* slow EMA blend with α=0.005 (per-fold cadence — the Bayesian prior
* is a *prior belief* and should change slowly across folds; a fast
* "prior" would just be a noisy Kelly estimate). Per
* `pearl_first_observation_bootstrap.md` and
* `pearl_controller_anchors_isv_driven.md`.
*
* Bilateral clamp on each slot per `pearl_symmetric_clamp_audit.md`:
* - counts in [0.5, 100]: below 0.5 the prior is effectively absent
* (kelly_f swings on first real trade); above 100 the prior
* dominates 10+ real trades (the maturity threshold for warmup_floor
* blend in kelly_position_cap).
* - sums in [0.001, 1.0]: below 1e-3 hits the same numerical floor
* as `fmaxf(avg_loss, 0.0001f)` in the Kelly formula; above 1.0
* implies a per-trade win/loss > 100% of equity which would be
* risk-management-failure territory anyway.
*
* Pearls + invariants:
*
* - `feedback_no_atomicadd.md` — single-block kernel; block-tree-reduce
* in shmem; no atomics across blocks.
*
* - `pearl_no_host_branches_in_captured_graph.md` — `CudaFunction`
* pre-loaded at construction; on-device guards for "no realized
* trades" and Pearl-A sentinel detection.
*
* - `pearl_first_observation_bootstrap.md` — first valid observation
* replaces sentinel directly; no blend.
*
* - `feedback_no_stubs.md` — full body, no return-zero placeholders.
*
* - `feedback_isv_for_adaptive_bounds.md` — count bounds [0.5, 100],
* sum bounds [0.001, 1.0] are dimensional safety floors, NOT tuning.
*
* - `pearl_symmetric_clamp_audit.md` — bilateral
* `fmaxf(lo, fminf(x, hi))` clamp on each slot before writing.
*
* - `pearl_controller_anchors_isv_driven.md` — every controller anchor
* (the priors are anchors for the Kelly cap formula) is ISV-driven.
*
* Args:
* portfolio_state — `[n_envs * ps_stride]` device f32; the
* same buffer `kelly_cap_update_kernel`
* already reads. We use the four
* PS_KELLY_* fields per env.
* n_envs — number of envs (alloc_episodes from the
* collector).
* ps_stride — PS_STRIDE from state_layout.cuh (43
* post-Plan-3-D.4c).
* isv — `[ISV_TOTAL_DIM]` device f32, modified
* in place at indices `*_idx`.
* wins_idx — KELLY_PRIOR_WINS_INDEX (454).
* losses_idx — KELLY_PRIOR_LOSSES_INDEX (455).
* sum_wins_idx — KELLY_PRIOR_SUM_WINS_INDEX (456).
* sum_losses_idx — KELLY_PRIOR_SUM_LOSSES_INDEX (457).
* sentinel_wins — SENTINEL_KELLY_PRIOR_WINS (2.0).
* sentinel_losses — SENTINEL_KELLY_PRIOR_LOSSES (2.0).
* sentinel_sum_wins — SENTINEL_KELLY_PRIOR_SUM_WINS (0.01).
* sentinel_sum_losses — SENTINEL_KELLY_PRIOR_SUM_LOSSES (0.01).
* count_min, count_max — KELLY_PRIOR_COUNT_MIN/MAX (0.5, 100).
* sum_min, sum_max — KELLY_PRIOR_SUM_MIN/MAX (0.001, 1.0).
* alpha — KELLY_PRIOR_EMA_ALPHA (0.005, slow).
*
* Launch: grid=(1, 1, 1), block=(BLK_DIM=256, 1, 1).
* Shared memory: 4 * BLK_DIM floats:
* sh_win_count — sum of PS_KELLY_WIN_COUNT across envs
* sh_loss_count — sum of PS_KELLY_LOSS_COUNT across envs
* sh_sum_wins — sum of PS_KELLY_SUM_WINS across envs
* sh_sum_losses — sum of PS_KELLY_SUM_LOSSES across envs
* ══════════════════════════════════════════════════════════════════════════ */
#include <cuda_runtime.h>
#include "state_layout.cuh" /* PS_KELLY_WIN_COUNT, PS_KELLY_LOSS_COUNT,
* PS_KELLY_SUM_WINS, PS_KELLY_SUM_LOSSES */
#define BLK_DIM 256
#define EPS_F 1e-6f
extern "C" __global__
void kelly_bayesian_priors_update(
const float* __restrict__ portfolio_state,
int n_envs,
int ps_stride,
float* __restrict__ isv,
int wins_idx,
int losses_idx,
int sum_wins_idx,
int sum_losses_idx,
float sentinel_wins,
float sentinel_losses,
float sentinel_sum_wins,
float sentinel_sum_losses,
float count_min,
float count_max,
float sum_min,
float sum_max,
float alpha)
{
/* Single-block kernel — only block 0 does anything. */
if (blockIdx.x != 0) return;
extern __shared__ float shmem[];
float* sh_win_count = shmem; /* [BLK_DIM] */
float* sh_loss_count = shmem + BLK_DIM; /* [BLK_DIM] */
float* sh_sum_wins = shmem + 2 * BLK_DIM; /* [BLK_DIM] */
float* sh_sum_losses = shmem + 3 * BLK_DIM; /* [BLK_DIM] */
int tid = threadIdx.x;
/* Phase 1a: stride-based sweep — accumulate the four Kelly stats
* across all envs. PS_KELLY_* indices come from state_layout.cuh
* (the same constants kelly_cap_update_kernel uses). Counts are
* stored as floats in PS slots (matches Kelly stat accumulator
* convention in apply_kelly_cap). */
float local_wins = 0.0f;
float local_losses = 0.0f;
float local_sumw = 0.0f;
float local_suml = 0.0f;
for (int e = tid; e < n_envs; e += BLK_DIM) {
const float* ps = portfolio_state + (long long)e * ps_stride;
local_wins += ps[PS_KELLY_WIN_COUNT];
local_losses += ps[PS_KELLY_LOSS_COUNT];
local_sumw += ps[PS_KELLY_SUM_WINS];
local_suml += ps[PS_KELLY_SUM_LOSSES];
}
sh_win_count[tid] = local_wins;
sh_loss_count[tid] = local_losses;
sh_sum_wins[tid] = local_sumw;
sh_sum_losses[tid] = local_suml;
__syncthreads();
/* Phase 1b: block-tree-reduce (no atomicAdd) — halve every iter.
* All four streams fold via addition. */
for (int s = BLK_DIM >> 1; s > 0; s >>= 1) {
if (tid < s) {
sh_win_count[tid] += sh_win_count[tid + s];
sh_loss_count[tid] += sh_loss_count[tid + s];
sh_sum_wins[tid] += sh_sum_wins[tid + s];
sh_sum_losses[tid] += sh_sum_losses[tid + s];
}
__syncthreads();
}
/* Phase 2: thread 0 finalises EMA + Pearl-A + clamp + writes the four slots. */
if (tid != 0) return;
float total_wins = sh_win_count[0];
float total_losses = sh_loss_count[0];
float total_sumw = sh_sum_wins[0];
float total_suml = sh_sum_losses[0];
/* Guard: no realized trades observed across all envs — keep ISV slots
* unchanged (cold-start sentinel persists; consumers fall back to
* the static defaults via the NULL-tolerant guard pattern). */
if (total_wins + total_losses <= 0.0f) {
return;
}
/* Bilateral pre-EMA clamp on the targets so the blended value stays in
* the dimensional-safety bounds even before the post-EMA defensive
* clamp. Per `pearl_symmetric_clamp_audit.md`. */
float target_wins = fmaxf(count_min, fminf(total_wins, count_max));
float target_losses = fmaxf(count_min, fminf(total_losses, count_max));
float target_sumw = fmaxf(sum_min, fminf(total_sumw, sum_max));
float target_suml = fmaxf(sum_min, fminf(total_suml, sum_max));
/* Pearl-A first-observation bootstrap on each slot independently.
* The sentinel matches the pre-P1-Producer hardcoded constants
* (2.0/2.0/0.01/0.01) for bit-identical cold-start behavior. */
float cur_wins = isv[wins_idx];
float cur_losses = isv[losses_idx];
float cur_sumw = isv[sum_wins_idx];
float cur_suml = isv[sum_losses_idx];
float blended_wins = (fabsf(cur_wins - sentinel_wins) < EPS_F)
? target_wins
: ((1.0f - alpha) * cur_wins + alpha * target_wins);
float blended_losses = (fabsf(cur_losses - sentinel_losses) < EPS_F)
? target_losses
: ((1.0f - alpha) * cur_losses + alpha * target_losses);
float blended_sumw = (fabsf(cur_sumw - sentinel_sum_wins) < EPS_F)
? target_sumw
: ((1.0f - alpha) * cur_sumw + alpha * target_sumw);
float blended_suml = (fabsf(cur_suml - sentinel_sum_losses) < EPS_F)
? target_suml
: ((1.0f - alpha) * cur_suml + alpha * target_suml);
/* Defensive post-EMA bilateral clamp — handles malformed prior state
* outside the dimensional-safety range. */
blended_wins = fmaxf(count_min, fminf(blended_wins, count_max));
blended_losses = fmaxf(count_min, fminf(blended_losses, count_max));
blended_sumw = fmaxf(sum_min, fminf(blended_sumw, sum_max));
blended_suml = fmaxf(sum_min, fminf(blended_suml, sum_max));
isv[wins_idx] = blended_wins;
isv[losses_idx] = blended_losses;
isv[sum_wins_idx] = blended_sumw;
isv[sum_losses_idx] = blended_suml;
}

46
crates/ml/src/cuda_pipeline/kelly_cap_update_kernel.cu
View File

@@ -1,5 +1,7 @@
/* kelly_cap_update — effective Kelly cap, GPU-driven (Plan 1 Task 11).
* Reads: ISV[LEARNING_HEALTH_INDEX=12] (safety multiplier coupled to health)
* ISV[KELLY_PRIOR_*_INDEX=454..458) (Class A P1-Producer adaptive
* Bayesian priors — slow-EMA from realized fold trades)
* portfolio_state[n_envs, PS_STRIDE] — Kelly win/loss stats per env
* Writes: ISV[KELLY_CAP_EFF_INDEX=47] (shifted +3 from 44 by D.2 per-branch gamma, Plan 2 Task 3)
* Cold path (per-epoch boundary). Single-thread kernel.
@@ -9,9 +11,14 @@
* health-coupled safety multiplier. Result is a scalar cap in [0, 1].
* Downstream kernel readers use this cap to limit position sizing.
*
* Bayesian priors (matching trade_physics.cuh):
* prior_wins=2, prior_losses=2, prior_sum_wins=0.01, prior_sum_losses=0.01
* This ensures the cap never collapses to zero on cold start.
* Bayesian priors (Class A P1-Producer 2026-05-08): now read from
* ISV[454..458) where they are produced by `kelly_bayesian_priors_update_kernel`
* at the SAME per-epoch boundary, launched RIGHT BEFORE this kernel so the
* cold-path Kelly cap sees the freshly-blended priors. Sentinels match
* the pre-P1-Producer hardcoded constants
* (2.0 / 2.0 / 0.01 / 0.01) for bit-identical cold-start behavior. The
* fallback path below applies when an ISV slot is at sentinel (Pearl-A
* pre-bootstrap) or out of the dimensional-safety range — defensive only.
*/
#include "state_layout.cuh"
@@ -27,6 +34,21 @@
* Slot constant mirrors `sp5_isv_slots.rs::KELLY_WARMUP_FLOOR_INDEX = 330`. */
#define SP9_KELLY_WARMUP_FLOOR_INDEX 330
/* Class A P1-Producer cold-start fallback helper. Returns the ISV slot value
* if it is within the dimensional-safety range, else the static default
* (the pre-P1-Producer hardcoded value for bit-identical cold-start). */
__device__ __forceinline__ float kelly_prior_or_default(
float isv_val,
float fallback_default,
float lo_bound,
float hi_bound
) {
if (isv_val >= lo_bound && isv_val <= hi_bound) {
return isv_val;
}
return fallback_default;
}
extern "C" __global__ void kelly_cap_update(
const float* __restrict__ isv,
float* __restrict__ isv_out,
@@ -36,10 +58,20 @@ extern "C" __global__ void kelly_cap_update(
) {
if (threadIdx.x != 0 || blockIdx.x != 0) return;
const float prior_wins = 2.0f;
const float prior_losses = 2.0f;
const float prior_sum_wins = 0.01f;
const float prior_sum_losses = 0.01f;
/* Class A P1-Producer (2026-05-08): adaptive Bayesian priors from
* ISV[454..458). Producer kernel `kelly_bayesian_priors_update_kernel`
* runs at the same per-epoch boundary RIGHT BEFORE this kernel.
* Cold-start path: ISV slot at sentinel (Pearl-A pre-bootstrap) or
* out-of-range → fall back to the static default (matches pre-P1-Producer
* hardcoded values for bit-identical cold-start behavior). */
const float prior_wins = kelly_prior_or_default(
isv[KELLY_PRIOR_WINS_INDEX], KELLY_PRIOR_WINS_DEFAULT, 0.5f, 100.0f);
const float prior_losses = kelly_prior_or_default(
isv[KELLY_PRIOR_LOSSES_INDEX], KELLY_PRIOR_LOSSES_DEFAULT, 0.5f, 100.0f);
const float prior_sum_wins = kelly_prior_or_default(
isv[KELLY_PRIOR_SUM_WINS_INDEX], KELLY_PRIOR_SUM_WINS_DEFAULT, 0.001f, 1.0f);
const float prior_sum_losses = kelly_prior_or_default(
isv[KELLY_PRIOR_SUM_LOSSES_INDEX], KELLY_PRIOR_SUM_LOSSES_DEFAULT, 0.001f, 1.0f);
float health = fminf(fmaxf(isv[12], 0.0f), 1.0f);
/* Safety multiplier: 0.5 + 0.5*health — collapses cap during degradation */

97
crates/ml/src/cuda_pipeline/sp14_isv_slots.rs
View File

@@ -164,6 +164,65 @@ pub const REWARD_CAP_EMA_ALPHA: f32 = 0.01;
pub const SP14_P0A_SLOT_BASE: usize = 452;
pub const SP14_P0A_SLOT_END: usize = 454;
// ── Class A P1-Producer (2026-05-08): adaptive Bayesian Kelly priors ─────
// Replaces the hardcoded
// `prior_wins=2.0, prior_losses=2.0, prior_sum_wins=0.01, prior_sum_losses=0.01`
// in `kelly_cap_update_kernel.cu:39-42` AND
// `trade_physics.cuh::kelly_position_cap:304-307` with ISV-driven slow-EMA
// values produced by `kelly_bayesian_priors_update_kernel.cu` at per-fold
// boundary. The producer aggregates realized Kelly stats from
// `portfolio_state[n_envs, PS_STRIDE]` (the same buffer
// `kelly_cap_update_kernel` already reads from `PS_KELLY_*` indices) and
// EMA-blends into the slot. The static defaults remain as Pearl-A
// sentinels so cold-start (pre-first-observation) is bit-identical to the
// pre-P1-Producer behavior.
//
// Why slow EMA at fold cadence (α=0.005): the Bayesian prior is a *prior
// belief* about the Kelly distribution before observing the new fold's
// trades. It should change slowly across folds — a fast-moving "prior"
// would just be a noisy Kelly estimate. Per
// `pearl_controller_anchors_isv_driven.md` and
// `feedback_isv_for_adaptive_bounds.md` (the prior is a controller
// anchor, not a tuning constant).
//
// Producer kernel: `kelly_bayesian_priors_update_kernel.cu`.
// Consumers (atomic migration in same commit):
// - `kelly_cap_update_kernel.cu:39-42` → ISV[454..458)
// - `trade_physics.cuh::kelly_position_cap:304-307` → threaded
// `isv_signals_ptr` argument with NULL-tolerant cold-start fallback.
pub const KELLY_PRIOR_WINS_INDEX: usize = 454; // adaptive prior_wins (replaces 2.0f)
pub const KELLY_PRIOR_LOSSES_INDEX: usize = 455; // adaptive prior_losses (replaces 2.0f)
pub const KELLY_PRIOR_SUM_WINS_INDEX: usize = 456; // adaptive prior_sum_wins (replaces 0.01f)
pub const KELLY_PRIOR_SUM_LOSSES_INDEX: usize = 457; // adaptive prior_sum_losses (replaces 0.01f)
// Sentinels — match pre-P1-Producer hardcoded values for bit-identical
// cold-start behavior before the first valid observation lands.
pub const SENTINEL_KELLY_PRIOR_WINS: f32 = 2.0;
pub const SENTINEL_KELLY_PRIOR_LOSSES: f32 = 2.0;
pub const SENTINEL_KELLY_PRIOR_SUM_WINS: f32 = 0.01;
pub const SENTINEL_KELLY_PRIOR_SUM_LOSSES: f32 = 0.01;
// Bounds — Category-1 dimensional safety floors per
// `feedback_isv_for_adaptive_bounds.md`. Counts in [0.5, 100]: below 0.5
// the prior is effectively absent (kelly_f swings on first real trade);
// above 100 the prior dominates 10+ real trades (the maturity threshold
// for warmup_floor blend in kelly_position_cap). Sums in [0.001, 1.0]:
// below 1e-3 hits the same numerical floor as `fmaxf(avg_loss, 0.0001f)`
// in the Kelly formula; above 1.0 implies a per-trade win/loss > 100% of
// equity which would be risk-management-failure territory anyway.
pub const KELLY_PRIOR_COUNT_MIN: f32 = 0.5;
pub const KELLY_PRIOR_COUNT_MAX: f32 = 100.0;
pub const KELLY_PRIOR_SUM_MIN: f32 = 0.001;
pub const KELLY_PRIOR_SUM_MAX: f32 = 1.0;
// EMA blend rate — per-fold cadence, slower than reward_cap's α=0.01.
// The Bayesian prior is a prior belief; it should change slowly across
// folds (a fast-moving "prior" would just be a noisy Kelly estimate).
pub const KELLY_PRIOR_EMA_ALPHA: f32 = 0.005;
pub const SP14_P1_SLOT_BASE: usize = 454;
pub const SP14_P1_SLOT_END: usize = 458;
#[cfg(test)]
mod tests {
use super::*;
@@ -273,4 +332,42 @@ mod tests {
SP14_P0A_SLOT_END, ISV_TOTAL_DIM,
);
}
/// Lock Class A P1-Producer (2026-05-08) adaptive Bayesian-prior slot
/// layout. Four contiguous slots [454..458) replacing the hardcoded
/// `prior_wins=2.0 / prior_losses=2.0 / prior_sum_wins=0.01 /
/// prior_sum_losses=0.01` constants in `kelly_cap_update_kernel.cu` and
/// `trade_physics.cuh::kelly_position_cap`. Sentinels match the
/// pre-P1-Producer hardcoded values for bit-identical cold-start.
#[test]
fn sp14_p1_kelly_prior_slot_layout_locked() {
assert_eq!(SP14_P1_SLOT_BASE, 454);
assert_eq!(SP14_P1_SLOT_END, 458);
assert_eq!(KELLY_PRIOR_WINS_INDEX, 454);
assert_eq!(KELLY_PRIOR_LOSSES_INDEX, 455);
assert_eq!(KELLY_PRIOR_SUM_WINS_INDEX, 456);
assert_eq!(KELLY_PRIOR_SUM_LOSSES_INDEX, 457);
// Sentinels match pre-P1-Producer hardcoded constants for
// bit-identical cold-start behavior before first valid observation.
assert_eq!(SENTINEL_KELLY_PRIOR_WINS, 2.0);
assert_eq!(SENTINEL_KELLY_PRIOR_LOSSES, 2.0);
assert_eq!(SENTINEL_KELLY_PRIOR_SUM_WINS, 0.01);
assert_eq!(SENTINEL_KELLY_PRIOR_SUM_LOSSES, 0.01);
// Category-1 dimensional safety: counts [0.5, 100], sums [0.001, 1.0].
assert_eq!(KELLY_PRIOR_COUNT_MIN, 0.5);
assert_eq!(KELLY_PRIOR_COUNT_MAX, 100.0);
assert_eq!(KELLY_PRIOR_SUM_MIN, 0.001);
assert_eq!(KELLY_PRIOR_SUM_MAX, 1.0);
}
#[test]
fn all_sp14_p1_slots_fit_within_isv_total_dim() {
use crate::cuda_pipeline::gpu_dqn_trainer::ISV_TOTAL_DIM;
assert!(
SP14_P1_SLOT_END <= ISV_TOTAL_DIM,
"SP14_P1_SLOT_END={} exceeds ISV_TOTAL_DIM={} — bus too small for Class A P1-Producer Bayesian-prior slots; \
bump ISV_TOTAL_DIM in gpu_dqn_trainer.rs (and update layout_fingerprint_seed()).",
SP14_P1_SLOT_END, ISV_TOTAL_DIM,
);
}
}

37
crates/ml/src/cuda_pipeline/state_layout.cuh
View File

@@ -201,6 +201,13 @@
#define ISV_HOLD_RATE_TARGET_IDX 381 // SP13 v3 target Hold-pick rate (static, default 0.20)
#define ISV_HOLD_RATE_OBSERVED_EMA_IDX 382 // SP13 v3 observed Hold-pick rate EMA (per-fold, sentinel 0.0)
// NOTE: this macro is kept at 383 for the historic SP13-era boundary marker
// and is referenced only in commentary inside .cu/.cuh files. The actual
// ISV bus dimension is the Rust `ISV_TOTAL_DIM` constant in
// `gpu_dqn_trainer.rs` (currently 458, post Class A P1-Producer
// 2026-05-08). Kernels take `const float*` pointers and use raw indices —
// no .cu/.cuh file dimensions an array with this macro. The Rust constant
// is the single source of truth for the bus size.
#define ISV_TOTAL_DIM 383
// ────────────────────────────────────────────────────────────────────────────
@@ -299,6 +306,36 @@
#define REWARD_POS_CAP_MIN_BOUND 1.0f
#define REWARD_POS_CAP_MAX_BOUND 50.0f
// Class A P1-Producer (2026-05-08) — adaptive Bayesian Kelly priors.
// Replaces hardcoded `prior_wins=2.0f / prior_losses=2.0f /
// prior_sum_wins=0.01f / prior_sum_losses=0.01f` in kelly_cap_update_kernel
// and trade_physics.cuh::kelly_position_cap with ISV-driven slow-EMA values
// fed by the `kelly_bayesian_priors_update_kernel` producer at per-fold
// boundary (per-epoch boundary in current scheduling — same cadence as
// kelly_cap_update). Sentinels match the pre-P1-Producer hardcoded values
// so cold-start (pre-first-observation) is bit-identical to the prior
// behavior. Mirrors `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs`
// constants of the same names (locked by
// `sp14_p1_kelly_prior_slot_layout_locked` test).
#define KELLY_PRIOR_WINS_INDEX 454
#define KELLY_PRIOR_LOSSES_INDEX 455
#define KELLY_PRIOR_SUM_WINS_INDEX 456
#define KELLY_PRIOR_SUM_LOSSES_INDEX 457
#define SENTINEL_KELLY_PRIOR_WINS 2.0f
#define SENTINEL_KELLY_PRIOR_LOSSES 2.0f
#define SENTINEL_KELLY_PRIOR_SUM_WINS 0.01f
#define SENTINEL_KELLY_PRIOR_SUM_LOSSES 0.01f
// Cold-start fallback for the consumer-side ISV-NULL guard (mirrors the
// existing `kelly_warmup_floor_sp9` and `kelly_f_smooth` patterns in
// trade_physics.cuh). When `isv_signals_ptr == NULL` (validation envs that
// don't pass the bus through) OR the slot value is at sentinel, the
// consumer falls back to the static defaults below. Bit-identical
// pre-P1-Producer behavior.
#define KELLY_PRIOR_WINS_DEFAULT 2.0f
#define KELLY_PRIOR_LOSSES_DEFAULT 2.0f
#define KELLY_PRIOR_SUM_WINS_DEFAULT 0.01f
#define KELLY_PRIOR_SUM_LOSSES_DEFAULT 0.01f
// ── Compile-time checks ──
static_assert(SL_PADDING_START + SL_PADDING_DIM == SL_STATE_DIM,
"State layout dimensions must sum to SL_STATE_DIM");

72
crates/ml/src/cuda_pipeline/trade_physics.cuh
View File

@@ -263,10 +263,21 @@ __device__ __forceinline__ float compute_drawdown(
* the environment refuses to let the agent over-lever, rather than
* scoring the agent for matching a formula.
*
* Priors (prior_wins=2, prior_losses=2, prior_sum_wins=0.01,
* prior_sum_losses=0.01) prevent the cap from being degenerate on cold-start
* runs where no trades have completed yet — without priors Kelly would
* be undefined for the first ~10 trades.
* Priors (Class A P1-Producer 2026-05-08): ISV-driven adaptive Bayesian
* priors at ISV[454..458) replace the previously-hardcoded
* `prior_wins=2.0f, prior_losses=2.0f, prior_sum_wins=0.01f,
* prior_sum_losses=0.01f` constants. Producer kernel
* `kelly_bayesian_priors_update_kernel` aggregates realized Kelly stats
* across envs and slow-EMA-blends (α=0.005) into the four slots at the
* same per-epoch boundary as `kelly_cap_update_kernel`. Sentinels match
* the pre-P1-Producer hardcoded values for bit-identical cold-start.
* NULL-tolerant (`isv_signals_ptr == NULL`) and out-of-range guards
* fall back to the static defaults — the helper is safe to call from
* validation envs that don't pass the bus through (matches the existing
* `kelly_f_smooth` / `kelly_warmup_floor_sp9` pattern in
* `unified_env_step_core`). The priors prevent the cap from being
* degenerate on cold-start runs where no trades have completed yet —
* without priors Kelly would be undefined for the first ~10 trades.
*
* The safety_multiplier is composed by the caller from two orthogonal
* adaptive signals:
@@ -299,12 +310,36 @@ __device__ __forceinline__ float kelly_position_cap(
float max_position,
float safety_multiplier,
float conviction,
float health_floor
float health_floor,
/* ── Class A P1-Producer (2026-05-08): adaptive Bayesian priors ──
* NULL-tolerant: when NULL OR slot at sentinel/out-of-range, falls
* back to the static defaults (matches pre-P1-Producer hardcoded
* values for bit-identical cold-start). The fallback path mirrors
* the existing `kelly_f_smooth` / `kelly_warmup_floor_sp9` pattern
* already in use elsewhere in this file. */
const float* __restrict__ isv_signals_ptr /* [ISV_TOTAL_DIM]; NULL = static defaults */
) {
const float prior_wins = 2.0f;
const float prior_losses = 2.0f;
const float prior_sum_wins = 0.01f;
const float prior_sum_losses = 0.01f;
/* Class A P1-Producer (2026-05-08): adaptive Bayesian priors. NULL
* guard handles validation envs that don't pass the bus through;
* range guard handles cold-start (sentinel) and malformed prior
* state. Both fall back to the pre-P1-Producer static defaults
* for bit-identical cold-start behavior. */
float prior_wins, prior_losses, prior_sum_wins, prior_sum_losses;
if (isv_signals_ptr != NULL) {
float pw = isv_signals_ptr[KELLY_PRIOR_WINS_INDEX];
float pl = isv_signals_ptr[KELLY_PRIOR_LOSSES_INDEX];
float psw = isv_signals_ptr[KELLY_PRIOR_SUM_WINS_INDEX];
float psl = isv_signals_ptr[KELLY_PRIOR_SUM_LOSSES_INDEX];
prior_wins = (pw >= 0.5f && pw <= 100.0f) ? pw : KELLY_PRIOR_WINS_DEFAULT;
prior_losses = (pl >= 0.5f && pl <= 100.0f) ? pl : KELLY_PRIOR_LOSSES_DEFAULT;
prior_sum_wins = (psw >= 0.001f && psw <= 1.0f) ? psw : KELLY_PRIOR_SUM_WINS_DEFAULT;
prior_sum_losses = (psl >= 0.001f && psl <= 1.0f) ? psl : KELLY_PRIOR_SUM_LOSSES_DEFAULT;
} else {
prior_wins = KELLY_PRIOR_WINS_DEFAULT;
prior_losses = KELLY_PRIOR_LOSSES_DEFAULT;
prior_sum_wins = KELLY_PRIOR_SUM_WINS_DEFAULT;
prior_sum_losses = KELLY_PRIOR_SUM_LOSSES_DEFAULT;
}
float eff_wins = win_count + prior_wins;
float eff_losses = loss_count + prior_losses;
@@ -389,7 +424,11 @@ __device__ __forceinline__ float kelly_position_cap(
/* Apply Kelly cap: clamp a signed target_position to [-cap, +cap].
* Convenience for call sites that want a single clamp call.
* `conviction` ∈ [0, 1] threads through to kelly_position_cap's adaptive
* warmup_floor — see docstring there. */
* warmup_floor — see docstring there.
* `isv_signals_ptr` threads through to kelly_position_cap's adaptive
* Bayesian priors (Class A P1-Producer 2026-05-08). NULL-tolerant —
* validation envs that don't pass the bus through fall back to the
* pre-P1-Producer static defaults for bit-identical cold-start. */
__device__ __forceinline__ float apply_kelly_cap(
float target_position,
float win_count,
@@ -399,11 +438,12 @@ __device__ __forceinline__ float apply_kelly_cap(
float max_position,
float safety_multiplier,
float conviction,
float health_floor
float health_floor,
const float* __restrict__ isv_signals_ptr /* [ISV_TOTAL_DIM]; NULL = static defaults */
) {
float cap = kelly_position_cap(win_count, loss_count, sum_wins, sum_losses,
max_position, safety_multiplier, conviction,
health_floor);
health_floor, isv_signals_ptr);
return fmaxf(-cap, fminf(cap, target_position));
}
@@ -861,7 +901,13 @@ __device__ __forceinline__ void unified_env_step_core(
*sum_wins, *sum_losses,
max_position_physics, safety,
policy_conviction,
health_safety_sp9
health_safety_sp9,
/* Class A P1-Producer (2026-05-08): adaptive Bayesian priors.
* `isv_signals_ptr` is the existing parameter to
* `unified_env_step_core` (already used for SP5 Pearl 6
* `kelly_f_smooth` and SP9 Fix 37 `kelly_warmup_floor_sp9`
* just above). NULL-tolerant on the consumer side. */
isv_signals_ptr
);
}

42
crates/ml/src/trainers/dqn/state_reset_registry.rs
View File

@@ -1046,6 +1046,48 @@ impl StateResetRegistry {
category: ResetCategory::FoldReset,
description: "ISV[REWARD_NEG_CAP_ADAPTIVE_INDEX=453] — Class A P0-A adaptive negative reward cap (replaces hardcoded REWARD_NEG_CAP=-10.0f from state_layout.cuh:267). Produced by `reward_cap_update_kernel` as NEG = −REWARD_NEG_TO_POS_RATIO=2.0 × POS_CAP. The 2:1 ratio (Kahneman/Tversky meta-analysis ~2.0-2.25) is preserved at producer-time inside the kernel (single source of truth — no consumer applies the multiplier itself), per `pearl_audit_unboundedness_for_implicit_asymmetry`. FoldReset sentinel SENTINEL_REWARD_NEG_CAP=-10.0 (matches pre-P0-A hardcoded value for bit-identical cold-start). Resulting bounds [-100, -2] derived from POS bounds × ratio. Consumed alongside POS slot at the same 3 sites.",
},
// ── Class A P1-Producer (2026-05-08): adaptive Bayesian Kelly priors ──
// Four slots [454..458) replacing the hardcoded
// `prior_wins=2.0f / prior_losses=2.0f / prior_sum_wins=0.01f /
// prior_sum_losses=0.01f` constants in
// `kelly_cap_update_kernel.cu:39-42` and
// `trade_physics.cuh::kelly_position_cap:304-307`. Producer
// kernel `kelly_bayesian_priors_update_kernel` aggregates
// realized PS_KELLY_{WIN_COUNT, LOSS_COUNT, SUM_WINS,
// SUM_LOSSES} across envs from the same `portfolio_state` buffer
// `kelly_cap_update_kernel` reads from at the same per-epoch
// boundary, then slow-EMA-blends into the four slots. Pearl-A
// bootstrap sentinels (2.0/2.0/0.01/0.01) match the
// pre-P1-Producer hardcoded values for bit-identical cold-start
// behavior — consumers fall back to the static defaults via the
// NULL-tolerant + range guard pattern (mirrors the existing
// `kelly_f_smooth` / `kelly_warmup_floor_sp9` patterns in
// `unified_env_step_core`). Once the first valid observation
// lands the adaptive path takes over via slow EMA (α=0.005;
// per-fold cadence — the Bayesian prior is a *prior belief*
// and should change slowly across folds). Per
// `feedback_isv_for_adaptive_bounds.md` and
// `pearl_controller_anchors_isv_driven.md`.
RegistryEntry {
name: "sp14_p1_kelly_prior_wins",
category: ResetCategory::FoldReset,
description: "ISV[KELLY_PRIOR_WINS_INDEX=454] — Class A P1-Producer adaptive Bayesian prior_wins (replaces hardcoded prior_wins=2.0f from kelly_cap_update_kernel.cu:39 + trade_physics.cuh::kelly_position_cap:304). Produced by `kelly_bayesian_priors_update_kernel` from the realized PS_KELLY_WIN_COUNT field aggregated across envs in the per-epoch `portfolio_state` buffer. FoldReset sentinel SENTINEL_KELLY_PRIOR_WINS=2.0 — Pearl-A first-observation bootstrap (matches pre-P1-Producer hardcoded value for bit-identical cold-start). α=0.005 slow EMA thereafter. Bounds [0.5, 100] — Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`. Consumed by kelly_cap_update_kernel (cold-path epoch-boundary effective Kelly cap) and trade_physics.cuh::kelly_position_cap (per-step env Kelly cap, NULL-tolerant via threaded `isv_signals_ptr`).",
},
RegistryEntry {
name: "sp14_p1_kelly_prior_losses",
category: ResetCategory::FoldReset,
description: "ISV[KELLY_PRIOR_LOSSES_INDEX=455] — Class A P1-Producer adaptive Bayesian prior_losses (replaces hardcoded prior_losses=2.0f from kelly_cap_update_kernel.cu:40 + trade_physics.cuh::kelly_position_cap:305). Produced by `kelly_bayesian_priors_update_kernel` from the realized PS_KELLY_LOSS_COUNT field aggregated across envs in the per-epoch `portfolio_state` buffer. FoldReset sentinel SENTINEL_KELLY_PRIOR_LOSSES=2.0 — Pearl-A first-observation bootstrap (matches pre-P1-Producer hardcoded value for bit-identical cold-start). α=0.005 slow EMA thereafter. Bounds [0.5, 100] same as prior_wins. Consumed at the same 2 sites as prior_wins.",
},
RegistryEntry {
name: "sp14_p1_kelly_prior_sum_wins",
category: ResetCategory::FoldReset,
description: "ISV[KELLY_PRIOR_SUM_WINS_INDEX=456] — Class A P1-Producer adaptive Bayesian prior_sum_wins (replaces hardcoded prior_sum_wins=0.01f from kelly_cap_update_kernel.cu:41 + trade_physics.cuh::kelly_position_cap:306). Produced by `kelly_bayesian_priors_update_kernel` from the realized PS_KELLY_SUM_WINS field aggregated across envs in the per-epoch `portfolio_state` buffer. FoldReset sentinel SENTINEL_KELLY_PRIOR_SUM_WINS=0.01 — Pearl-A first-observation bootstrap (matches pre-P1-Producer hardcoded value for bit-identical cold-start). α=0.005 slow EMA thereafter. Bounds [0.001, 1.0] — sum below 1e-3 hits the `fmaxf(avg_loss, 0.0001f)` numerical floor; sum above 1.0 implies > 100% per-trade equity which is risk-management-failure territory. Consumed at the same 2 sites as prior_wins.",
},
RegistryEntry {
name: "sp14_p1_kelly_prior_sum_losses",
category: ResetCategory::FoldReset,
description: "ISV[KELLY_PRIOR_SUM_LOSSES_INDEX=457] — Class A P1-Producer adaptive Bayesian prior_sum_losses (replaces hardcoded prior_sum_losses=0.01f from kelly_cap_update_kernel.cu:42 + trade_physics.cuh::kelly_position_cap:307). Produced by `kelly_bayesian_priors_update_kernel` from the realized PS_KELLY_SUM_LOSSES field aggregated across envs in the per-epoch `portfolio_state` buffer. FoldReset sentinel SENTINEL_KELLY_PRIOR_SUM_LOSSES=0.01 — Pearl-A first-observation bootstrap (matches pre-P1-Producer hardcoded value for bit-identical cold-start). α=0.005 slow EMA thereafter. Bounds [0.001, 1.0] same as prior_sum_wins. Consumed at the same 2 sites as prior_wins.",
},
// ── 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

63
crates/ml/src/trainers/dqn/trainer/training_loop.rs
View File

@@ -383,6 +383,18 @@ impl DQNTrainer {
{
let ps_dev_ptr = collector.portfolio_states_dev_ptr();
let n_envs = collector.alloc_episodes() as i32;
// Class A P1-Producer (2026-05-08): adaptive Bayesian Kelly
// priors update. MUST run BEFORE `launch_kelly_cap_update`
// so that kernel sees the freshly-blended priors. Sweeps
// the same `portfolio_state` buffer with the same n_envs
// and writes ISV[454..458). Pearl-A bootstrap + α=0.005
// slow EMA. Replaces hardcoded
// prior_wins/prior_losses/prior_sum_wins/prior_sum_losses
// constants in kelly_cap_update_kernel.cu and
// trade_physics.cuh::kelly_position_cap.
if let Err(e) = fused.trainer().launch_kelly_bayesian_priors_update(ps_dev_ptr, n_envs) {
tracing::warn!(epoch, "launch_kelly_bayesian_priors_update failed (non-fatal): {e}");
}
if let Err(e) = fused.launch_kelly_cap_update(ps_dev_ptr, n_envs) {
tracing::warn!(epoch, "launch_kelly_cap_update failed (non-fatal): {e}");
}
@@ -8183,6 +8195,57 @@ impl DQNTrainer {
);
}
}
// Class A P1-Producer (2026-05-08): adaptive Bayesian Kelly
// priors. Four FoldReset slots [454..458) — sentinels match
// pre-P1-Producer hardcoded values (2.0/2.0/0.01/0.01) so
// the kernel's "first observation" check fires cleanly on
// the new fold's first launch (avoids cross-fold EMA
// contamination). Cold-start is bit-identical to
// pre-P1-Producer until the first valid observation lands.
"sp14_p1_kelly_prior_wins" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::{
KELLY_PRIOR_WINS_INDEX, SENTINEL_KELLY_PRIOR_WINS,
};
fused.trainer().write_isv_signal_at(
KELLY_PRIOR_WINS_INDEX,
SENTINEL_KELLY_PRIOR_WINS,
);
}
}
"sp14_p1_kelly_prior_losses" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::{
KELLY_PRIOR_LOSSES_INDEX, SENTINEL_KELLY_PRIOR_LOSSES,
};
fused.trainer().write_isv_signal_at(
KELLY_PRIOR_LOSSES_INDEX,
SENTINEL_KELLY_PRIOR_LOSSES,
);
}
}
"sp14_p1_kelly_prior_sum_wins" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::{
KELLY_PRIOR_SUM_WINS_INDEX, SENTINEL_KELLY_PRIOR_SUM_WINS,
};
fused.trainer().write_isv_signal_at(
KELLY_PRIOR_SUM_WINS_INDEX,
SENTINEL_KELLY_PRIOR_SUM_WINS,
);
}
}
"sp14_p1_kelly_prior_sum_losses" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::{
KELLY_PRIOR_SUM_LOSSES_INDEX, SENTINEL_KELLY_PRIOR_SUM_LOSSES,
};
fused.trainer().write_isv_signal_at(
KELLY_PRIOR_SUM_LOSSES_INDEX,
SENTINEL_KELLY_PRIOR_SUM_LOSSES,
);
}
}
// 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

390
crates/ml/tests/sp14_oracle_tests.rs
View File

@@ -787,6 +787,396 @@ mod sp14_p0a_reward_cap_gpu {
}
}
// ═══════════════════════════════════════════════════════════════════════════
// Class A P1-Producer (2026-05-08) — adaptive Bayesian Kelly priors producer
// tests.
//
// Verifies the `kelly_bayesian_priors_update_kernel` producer:
// 1. Pearl-A first-observation bootstrap: ISV at sentinel → REPLACES with
// the realized aggregated stats (no blend).
// 2. Slow EMA (α=0.005) blend after bootstrap.
// 3. Bounds enforced: counts in [0.5, 100], sums in [0.001, 1.0].
// 4. No realized trades (zero PS_KELLY_* fields) → ISV slots preserved
// bit-exactly.
//
// 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_p1_kelly_priors_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::{
KELLY_PRIOR_COUNT_MAX, KELLY_PRIOR_COUNT_MIN, KELLY_PRIOR_EMA_ALPHA,
KELLY_PRIOR_LOSSES_INDEX, KELLY_PRIOR_SUM_LOSSES_INDEX, KELLY_PRIOR_SUM_MAX,
KELLY_PRIOR_SUM_MIN, KELLY_PRIOR_SUM_WINS_INDEX, KELLY_PRIOR_WINS_INDEX,
SENTINEL_KELLY_PRIOR_LOSSES, SENTINEL_KELLY_PRIOR_SUM_LOSSES,
SENTINEL_KELLY_PRIOR_SUM_WINS, SENTINEL_KELLY_PRIOR_WINS,
};
// PS_STRIDE constant — matches state_layout.cuh::PS_STRIDE and the
// value used by `launch_kelly_cap_update`. Grown 41→43 by Plan 3 D.4c.
const PS_STRIDE: usize = 43;
// PS_KELLY_* slot offsets within a single env's portfolio_state row.
const PS_KELLY_WIN_COUNT: usize = 14;
const PS_KELLY_LOSS_COUNT: usize = 15;
const PS_KELLY_SUM_WINS: usize = 16;
const PS_KELLY_SUM_LOSSES: usize = 17;
const KELLY_BAYESIAN_PRIORS_UPDATE_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/kelly_bayesian_priors_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(KELLY_BAYESIAN_PRIORS_UPDATE_CUBIN.to_vec())
.expect("load kelly_bayesian_priors_update_kernel cubin");
module
.load_function("kelly_bayesian_priors_update")
.expect("load kelly_bayesian_priors_update function")
}
/// 4 arrays × 256 × 4 bytes = 4096 bytes shmem.
const BLK_DIM: u32 = 256;
const SMEM_BYTES: u32 = 4 * BLK_DIM * std::mem::size_of::<f32>() as u32;
/// Helper: launch with a fixed config.
#[allow(clippy::too_many_arguments)]
fn launch_kelly_priors(
stream: &Arc<CudaStream>,
kernel: &CudaFunction,
portfolio_state_ptr: u64,
n_envs: i32,
ps_stride: i32,
isv_ptr: u64,
wins_idx: i32,
losses_idx: i32,
sum_wins_idx: i32,
sum_losses_idx: i32,
sentinel_wins: f32,
sentinel_losses: f32,
sentinel_sum_wins: f32,
sentinel_sum_losses: f32,
count_min: f32,
count_max: f32,
sum_min: f32,
sum_max: f32,
alpha: f32,
) {
unsafe {
stream
.launch_builder(kernel)
.arg(&portfolio_state_ptr)
.arg(&n_envs)
.arg(&ps_stride)
.arg(&isv_ptr)
.arg(&wins_idx)
.arg(&losses_idx)
.arg(&sum_wins_idx)
.arg(&sum_losses_idx)
.arg(&sentinel_wins)
.arg(&sentinel_losses)
.arg(&sentinel_sum_wins)
.arg(&sentinel_sum_losses)
.arg(&count_min)
.arg(&count_max)
.arg(&sum_min)
.arg(&sum_max)
.arg(&alpha)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (BLK_DIM, 1, 1),
shared_mem_bytes: SMEM_BYTES,
})
.expect("launch kelly_bayesian_priors_update");
}
stream.synchronize().expect("sync after kelly_bayesian_priors_update");
}
/// Build a portfolio_state buffer of `n_envs * PS_STRIDE` entries with
/// the four PS_KELLY_* fields populated per env from the parallel slices.
/// All other PS_* fields are zero.
fn build_portfolio_state(
wins_per_env: &[f32],
losses_per_env: &[f32],
sum_wins_per_env: &[f32],
sum_losses_per_env: &[f32],
) -> Vec<f32> {
let n = wins_per_env.len();
assert_eq!(losses_per_env.len(), n);
assert_eq!(sum_wins_per_env.len(), n);
assert_eq!(sum_losses_per_env.len(), n);
let mut ps = vec![0.0_f32; n * PS_STRIDE];
for e in 0..n {
let row_off = e * PS_STRIDE;
ps[row_off + PS_KELLY_WIN_COUNT] = wins_per_env[e];
ps[row_off + PS_KELLY_LOSS_COUNT] = losses_per_env[e];
ps[row_off + PS_KELLY_SUM_WINS] = sum_wins_per_env[e];
ps[row_off + PS_KELLY_SUM_LOSSES] = sum_losses_per_env[e];
}
ps
}
/// Test 1 — Pearl-A first-observation bootstrap.
///
/// 4 envs with realized stats. Total wins = 1+2+3+4 = 10, losses = 0+1+0+1 = 2,
/// sum_wins = 0.05+0.10+0.15+0.20 = 0.50, sum_losses = 0+0.03+0+0.05 = 0.08.
/// Pre-EMA clamp: counts to [0.5, 100], sums to [0.001, 1.0] —
/// 10/2/0.5/0.08 all in range, pass through unchanged. Cold-start
/// (sentinels 2.0/2.0/0.01/0.01): Pearl-A REPLACES → ISV =
/// (10, 2, 0.5, 0.08).
#[test]
#[ignore = "requires GPU"]
fn kelly_priors_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];
// Cold-start: sentinels (matches pre-P1-Producer constants).
isv[KELLY_PRIOR_WINS_INDEX] = SENTINEL_KELLY_PRIOR_WINS;
isv[KELLY_PRIOR_LOSSES_INDEX] = SENTINEL_KELLY_PRIOR_LOSSES;
isv[KELLY_PRIOR_SUM_WINS_INDEX] = SENTINEL_KELLY_PRIOR_SUM_WINS;
isv[KELLY_PRIOR_SUM_LOSSES_INDEX] = SENTINEL_KELLY_PRIOR_SUM_LOSSES;
let wins_per_env = [1.0_f32, 2.0, 3.0, 4.0];
let losses_per_env = [0.0_f32, 1.0, 0.0, 1.0];
let sum_wins_per_env = [0.05_f32, 0.10, 0.15, 0.20];
let sum_losses_per_env = [0.0_f32, 0.03, 0.0, 0.05];
let n_envs = wins_per_env.len();
let ps = build_portfolio_state(&wins_per_env, &losses_per_env,
&sum_wins_per_env, &sum_losses_per_env);
let ps_buf = unsafe { MappedF32Buffer::new(ps.len()) }.expect("alloc ps");
ps_buf.write_from_slice(&ps);
let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
isv_buf.write_from_slice(&isv);
launch_kelly_priors(
&stream, &kernel,
ps_buf.dev_ptr, n_envs as i32, PS_STRIDE as i32, isv_buf.dev_ptr,
KELLY_PRIOR_WINS_INDEX as i32,
KELLY_PRIOR_LOSSES_INDEX as i32,
KELLY_PRIOR_SUM_WINS_INDEX as i32,
KELLY_PRIOR_SUM_LOSSES_INDEX as i32,
SENTINEL_KELLY_PRIOR_WINS, SENTINEL_KELLY_PRIOR_LOSSES,
SENTINEL_KELLY_PRIOR_SUM_WINS, SENTINEL_KELLY_PRIOR_SUM_LOSSES,
KELLY_PRIOR_COUNT_MIN, KELLY_PRIOR_COUNT_MAX,
KELLY_PRIOR_SUM_MIN, KELLY_PRIOR_SUM_MAX,
KELLY_PRIOR_EMA_ALPHA,
);
let result = isv_buf.read_all();
let pw = result[KELLY_PRIOR_WINS_INDEX];
let pl = result[KELLY_PRIOR_LOSSES_INDEX];
let psw = result[KELLY_PRIOR_SUM_WINS_INDEX];
let psl = result[KELLY_PRIOR_SUM_LOSSES_INDEX];
// Pearl-A: sentinel → REPLACE with target. Targets are the
// aggregated totals clamped to dimensional-safety bounds (all
// within range here, no clamping applied).
assert!((pw - 10.0).abs() < 1e-4,
"Pearl-A bootstrap: prior_wins expected 10.0, got {pw}");
assert!((pl - 2.0).abs() < 1e-4,
"Pearl-A bootstrap: prior_losses expected 2.0, got {pl}");
assert!((psw - 0.50).abs() < 1e-4,
"Pearl-A bootstrap: prior_sum_wins expected 0.50, got {psw}");
assert!((psl - 0.08).abs() < 1e-4,
"Pearl-A bootstrap: prior_sum_losses expected 0.08, got {psl}");
}
/// Test 2 — No realized trades → ISV slots preserved bit-exactly.
///
/// All envs have zero PS_KELLY_* fields → total_wins + total_losses == 0.
/// Kernel guard skips the EMA update; both slots stay at the seeded
/// values (NOT sentinels — to verify the guard preserves arbitrary
/// pre-existing state).
#[test]
#[ignore = "requires GPU"]
fn kelly_priors_no_realized_trades_preserves_isv() {
let stream = make_stream();
let kernel = load_kernel(&stream);
const ISV_DIM: usize = 1024;
// Seed with non-sentinel values to verify they survive the guard.
const SEED_WINS: f32 = 7.0;
const SEED_LOSSES: f32 = 5.0;
const SEED_SUM_WINS: f32 = 0.25;
const SEED_SUM_LOSSES: f32 = 0.15;
let mut isv = vec![0.0_f32; ISV_DIM];
isv[KELLY_PRIOR_WINS_INDEX] = SEED_WINS;
isv[KELLY_PRIOR_LOSSES_INDEX] = SEED_LOSSES;
isv[KELLY_PRIOR_SUM_WINS_INDEX] = SEED_SUM_WINS;
isv[KELLY_PRIOR_SUM_LOSSES_INDEX] = SEED_SUM_LOSSES;
// 3 envs with no realized trades.
let n_envs: usize = 3;
let ps = vec![0.0_f32; n_envs * PS_STRIDE];
let ps_buf = unsafe { MappedF32Buffer::new(ps.len()) }.expect("alloc ps");
ps_buf.write_from_slice(&ps);
let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
isv_buf.write_from_slice(&isv);
launch_kelly_priors(
&stream, &kernel,
ps_buf.dev_ptr, n_envs as i32, PS_STRIDE as i32, isv_buf.dev_ptr,
KELLY_PRIOR_WINS_INDEX as i32,
KELLY_PRIOR_LOSSES_INDEX as i32,
KELLY_PRIOR_SUM_WINS_INDEX as i32,
KELLY_PRIOR_SUM_LOSSES_INDEX as i32,
SENTINEL_KELLY_PRIOR_WINS, SENTINEL_KELLY_PRIOR_LOSSES,
SENTINEL_KELLY_PRIOR_SUM_WINS, SENTINEL_KELLY_PRIOR_SUM_LOSSES,
KELLY_PRIOR_COUNT_MIN, KELLY_PRIOR_COUNT_MAX,
KELLY_PRIOR_SUM_MIN, KELLY_PRIOR_SUM_MAX,
KELLY_PRIOR_EMA_ALPHA,
);
let result = isv_buf.read_all();
assert_eq!(result[KELLY_PRIOR_WINS_INDEX], SEED_WINS,
"prior_wins must be preserved bit-exactly when no realized trades");
assert_eq!(result[KELLY_PRIOR_LOSSES_INDEX], SEED_LOSSES,
"prior_losses must be preserved bit-exactly when no realized trades");
assert_eq!(result[KELLY_PRIOR_SUM_WINS_INDEX], SEED_SUM_WINS,
"prior_sum_wins must be preserved bit-exactly when no realized trades");
assert_eq!(result[KELLY_PRIOR_SUM_LOSSES_INDEX], SEED_SUM_LOSSES,
"prior_sum_losses must be preserved bit-exactly when no realized trades");
}
/// Test 3 — Bounds enforcement: extreme aggregate clamps to upper/lower bounds.
///
/// One env with huge counts (1e6 wins, 1e6 losses) and tiny sums (1e-9).
/// Pre-EMA clamp: wins/losses → 100 (count_max), sums → 0.001 (sum_min).
/// Cold-start (sentinels): Pearl-A REPLACES with clamped targets.
#[test]
#[ignore = "requires GPU"]
fn kelly_priors_bounds_clamp_extreme() {
let stream = make_stream();
let kernel = load_kernel(&stream);
const ISV_DIM: usize = 1024;
let mut isv = vec![0.0_f32; ISV_DIM];
isv[KELLY_PRIOR_WINS_INDEX] = SENTINEL_KELLY_PRIOR_WINS;
isv[KELLY_PRIOR_LOSSES_INDEX] = SENTINEL_KELLY_PRIOR_LOSSES;
isv[KELLY_PRIOR_SUM_WINS_INDEX] = SENTINEL_KELLY_PRIOR_SUM_WINS;
isv[KELLY_PRIOR_SUM_LOSSES_INDEX] = SENTINEL_KELLY_PRIOR_SUM_LOSSES;
// Single env with extreme aggregates.
let wins_per_env = [1.0e6_f32];
let losses_per_env = [1.0e6_f32];
let sum_wins_per_env = [1.0e-9_f32];
let sum_losses_per_env = [1.0e-9_f32];
let ps = build_portfolio_state(&wins_per_env, &losses_per_env,
&sum_wins_per_env, &sum_losses_per_env);
let ps_buf = unsafe { MappedF32Buffer::new(ps.len()) }.expect("alloc ps");
ps_buf.write_from_slice(&ps);
let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
isv_buf.write_from_slice(&isv);
launch_kelly_priors(
&stream, &kernel,
ps_buf.dev_ptr, 1, PS_STRIDE as i32, isv_buf.dev_ptr,
KELLY_PRIOR_WINS_INDEX as i32,
KELLY_PRIOR_LOSSES_INDEX as i32,
KELLY_PRIOR_SUM_WINS_INDEX as i32,
KELLY_PRIOR_SUM_LOSSES_INDEX as i32,
SENTINEL_KELLY_PRIOR_WINS, SENTINEL_KELLY_PRIOR_LOSSES,
SENTINEL_KELLY_PRIOR_SUM_WINS, SENTINEL_KELLY_PRIOR_SUM_LOSSES,
KELLY_PRIOR_COUNT_MIN, KELLY_PRIOR_COUNT_MAX,
KELLY_PRIOR_SUM_MIN, KELLY_PRIOR_SUM_MAX,
KELLY_PRIOR_EMA_ALPHA,
);
let result = isv_buf.read_all();
let pw = result[KELLY_PRIOR_WINS_INDEX];
let pl = result[KELLY_PRIOR_LOSSES_INDEX];
let psw = result[KELLY_PRIOR_SUM_WINS_INDEX];
let psl = result[KELLY_PRIOR_SUM_LOSSES_INDEX];
// 1e6 wins → clamped to count_max=100.
assert!((pw - KELLY_PRIOR_COUNT_MAX).abs() < 1e-4,
"prior_wins must clamp to count_max=100 on extreme; got {pw}");
assert!((pl - KELLY_PRIOR_COUNT_MAX).abs() < 1e-4,
"prior_losses must clamp to count_max=100 on extreme; got {pl}");
// 1e-9 sums → clamped to sum_min=0.001.
assert!((psw - KELLY_PRIOR_SUM_MIN).abs() < 1e-6,
"prior_sum_wins must clamp to sum_min=0.001 on tiny; got {psw}");
assert!((psl - KELLY_PRIOR_SUM_MIN).abs() < 1e-6,
"prior_sum_losses must clamp to sum_min=0.001 on tiny; got {psl}");
}
/// Test 4 — Slow EMA blend after bootstrap (α=0.005).
///
/// Pre-seed prior_wins=10.0 (NOT sentinel — bootstrap path NOT taken).
/// Single env with 5 wins. Target = 5 (in range). EMA blend:
/// (1 - 0.005) × 10 + 0.005 × 5 = 0.995 × 10 + 0.025 = 9.975.
#[test]
#[ignore = "requires GPU"]
fn kelly_priors_slow_ema_after_bootstrap() {
let stream = make_stream();
let kernel = load_kernel(&stream);
const ISV_DIM: usize = 1024;
const SEEDED_WINS: f32 = 10.0;
// Other slots stay at sentinel — only test EMA on prior_wins.
let mut isv = vec![0.0_f32; ISV_DIM];
isv[KELLY_PRIOR_WINS_INDEX] = SEEDED_WINS;
isv[KELLY_PRIOR_LOSSES_INDEX] = SENTINEL_KELLY_PRIOR_LOSSES;
isv[KELLY_PRIOR_SUM_WINS_INDEX] = SENTINEL_KELLY_PRIOR_SUM_WINS;
isv[KELLY_PRIOR_SUM_LOSSES_INDEX] = SENTINEL_KELLY_PRIOR_SUM_LOSSES;
// Single env with 5 wins.
let wins_per_env = [5.0_f32];
let losses_per_env = [1.0_f32];
let sum_wins_per_env = [0.10_f32];
let sum_losses_per_env = [0.02_f32];
let ps = build_portfolio_state(&wins_per_env, &losses_per_env,
&sum_wins_per_env, &sum_losses_per_env);
let ps_buf = unsafe { MappedF32Buffer::new(ps.len()) }.expect("alloc ps");
ps_buf.write_from_slice(&ps);
let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
isv_buf.write_from_slice(&isv);
launch_kelly_priors(
&stream, &kernel,
ps_buf.dev_ptr, 1, PS_STRIDE as i32, isv_buf.dev_ptr,
KELLY_PRIOR_WINS_INDEX as i32,
KELLY_PRIOR_LOSSES_INDEX as i32,
KELLY_PRIOR_SUM_WINS_INDEX as i32,
KELLY_PRIOR_SUM_LOSSES_INDEX as i32,
SENTINEL_KELLY_PRIOR_WINS, SENTINEL_KELLY_PRIOR_LOSSES,
SENTINEL_KELLY_PRIOR_SUM_WINS, SENTINEL_KELLY_PRIOR_SUM_LOSSES,
KELLY_PRIOR_COUNT_MIN, KELLY_PRIOR_COUNT_MAX,
KELLY_PRIOR_SUM_MIN, KELLY_PRIOR_SUM_MAX,
KELLY_PRIOR_EMA_ALPHA,
);
let result = isv_buf.read_all();
let pw = result[KELLY_PRIOR_WINS_INDEX];
// EMA: 0.995 × 10 + 0.005 × 5 = 9.975.
let expected = (1.0_f32 - KELLY_PRIOR_EMA_ALPHA) * SEEDED_WINS
+ KELLY_PRIOR_EMA_ALPHA * 5.0_f32;
assert!(
(pw - expected).abs() < 1e-4,
"Slow EMA blend (α={KELLY_PRIOR_EMA_ALPHA}): expected {expected}, got {pw}"
);
}
}
// ═══════════════════════════════════════════════════════════════════════════
// Test B.8: Layout-fingerprint regression for direction Q-head input bump.
//

147
docs/dqn-wire-up-audit.md
View File

@@ -7992,3 +7992,150 @@ P0-A made `REWARD_POS_CAP` adaptive via `ISV[REWARD_POS_CAP_ADAPTIVE_INDEX=452]`
1. **The legacy `compute_drawdown_penalty` is now partially adaptive but the SP15 `sp15_dd_penalty` quadratic path is fully ISV-driven via slots 420-421**. The two paths still coexist with different shapes (legacy: linear ramp scaled by POS_CAP; SP15: quadratic gated by λ_dd). The P1 batch already flagged this as future work. P0-A-downstream does not retire the legacy path — that's a separate refactor.
2. **`MIN_HOLD_PENALTY_RATIO = 0.6f` is a magic number at the call site** rather than a `#define` in `state_layout.cuh`. Kept local because it's strictly a relationship between two existing constants (3.0 / 5.0 = 0.6), not a new tunable signal. If a future batch lifts the ratio itself to ISV, the constant should move to `state_layout.cuh` first.
## Class A P1-Producer batch — adaptive Bayesian Kelly priors (2026-05-08)
The Class A P1 batch (commit `c4b6d6ef2`) wired Item 4 (var_floor q_gap-only adaptive) and explicitly deferred Items 1-3 to "a producer batch" because they required new ISV slot allocations. This entry lands one of those deferred items: the adaptive Bayesian Kelly priors. Item 2 (MIN_HOLD_TEMPERATURE EMA) was deferred separately — the audit-spec was wrong about the consumer site (see Concerns below).
### Background
`kelly_cap_update_kernel.cu:39-42` and `trade_physics.cuh::kelly_position_cap:304-307` both ran with hardcoded Bayesian priors:
```c
const float prior_wins = 2.0f;
const float prior_losses = 2.0f;
const float prior_sum_wins = 0.01f;
const float prior_sum_losses = 0.01f;
```
These are static "I don't know" priors frozen at sprint-1 defaults. Conceptually, Bayesian priors encode *prior belief* about the Kelly distribution before observing real trades. A static prior IS literally the prior — but it's also constant across folds, ignoring everything prior folds learned about the realized win/loss distribution shape.
The fix is an ISV-driven slow-EMA prior fed from the realized Kelly stats already aggregated in `portfolio_state[n_envs, PS_STRIDE]` (the same buffer `kelly_cap_update_kernel` reads from at the same per-epoch boundary). On the first fold, the sentinels match the pre-P1-Producer hardcoded values for bit-identical cold-start; once the first valid observation lands, the slow EMA blends in learned beliefs from realized distribution shape; on subsequent folds, the cold-start prior is no longer "I don't know" but "what we learned from the prior fold's distribution shape, slow-discounted."
### Slots allocated
| Slot | Constant | Sentinel | Bounds | Replaces |
|-----:|----------|---------:|--------|----------|
| 454 | `KELLY_PRIOR_WINS_INDEX` | 2.0 | [0.5, 100] | `prior_wins=2.0f` |
| 455 | `KELLY_PRIOR_LOSSES_INDEX` | 2.0 | [0.5, 100] | `prior_losses=2.0f` |
| 456 | `KELLY_PRIOR_SUM_WINS_INDEX` | 0.01 | [0.001, 1.0] | `prior_sum_wins=0.01f` |
| 457 | `KELLY_PRIOR_SUM_LOSSES_INDEX` | 0.01 | [0.001, 1.0] | `prior_sum_losses=0.01f` |
`ISV_TOTAL_DIM` bumped 454 → 458. Layout fingerprint regenerated.
Bounds rationale (Category-1 dimensional safety per `feedback_isv_for_adaptive_bounds.md`):
- Counts in [0.5, 100]: below 0.5 the prior is effectively absent (kelly_f swings on first real trade); above 100 the prior dominates 10+ real trades (the maturity threshold for warmup_floor blend in `kelly_position_cap`).
- Sums in [0.001, 1.0]: below 1e-3 hits the same numerical floor as `fmaxf(avg_loss, 0.0001f)` in the Kelly formula; above 1.0 implies a per-trade win/loss > 100% of equity which would be risk-management-failure territory.
### Producer
`crates/ml/src/cuda_pipeline/kelly_bayesian_priors_update_kernel.cu` (NEW, 250 LOC):
- Single-block 256-thread kernel; block-tree-reduce in shmem (no `atomicAdd` per `feedback_no_atomicadd.md`).
- Stride-based sweep accumulates four sums across all envs from the four `PS_KELLY_*` fields per env's portfolio_state row (PS_STRIDE=43 per state_layout.cuh).
- Phase 2 (thread 0): pre-EMA bilateral clamp on each target slot, then per-slot Pearl-A first-observation bootstrap (sentinel match within 1e-6 → REPLACE; otherwise EMA blend), then defensive post-EMA bilateral clamp.
- α=0.005 slow EMA (per-fold cadence; the Bayesian prior is a *prior belief* and should change slowly across folds — a fast-moving "prior" would just be a noisy Kelly estimate). Slower than `reward_cap_update`'s α=0.01.
- Cold-start fallback: `total_wins + total_losses == 0` → kernel guard skips the EMA update and ISV slots remain unchanged (sentinel persists; consumers fall back to the `state_layout.cuh::KELLY_PRIOR_*_DEFAULT` macros via the NULL-tolerant + range guard pattern).
### Consumers migrated atomically
**Site A** — `kelly_cap_update_kernel.cu:39-42`:
```c
// Before:
const float prior_wins = 2.0f; /* + 3 more */
// After:
const float prior_wins = kelly_prior_or_default(
isv[KELLY_PRIOR_WINS_INDEX], KELLY_PRIOR_WINS_DEFAULT, 0.5f, 100.0f);
/* + 3 more — same pattern */
```
The `kelly_prior_or_default` device helper returns the ISV slot value if it's within the dimensional-safety range, else the `state_layout.cuh::KELLY_PRIOR_*_DEFAULT` macro (which equals the pre-P1-Producer hardcoded value). Cold-start (sentinel) is in-range so the macro path only fires for malformed prior state.
**Site B** — `trade_physics.cuh::kelly_position_cap:304-307`:
`kelly_position_cap` gained an `isv_signals_ptr` parameter (NULL-tolerant). The function reads the four ISV slots and falls back to `KELLY_PRIOR_*_DEFAULT` macros under either `isv_signals_ptr == NULL` (validation envs that don't pass the bus through) OR out-of-range (cold-start sentinel + malformed state). `apply_kelly_cap` (the only caller of `kelly_position_cap`) gained the same parameter and passes it through. The single call site in `unified_env_step_core` (line 898) already had `isv_signals_ptr` — it's the same parameter already used for SP5 Pearl 6 `kelly_f_smooth` and SP9 Fix 37 `kelly_warmup_floor_sp9` just above. No new parameter threading needed.
### Producer launch site
Wired in `training_loop.rs:393` RIGHT BEFORE `launch_kelly_cap_update`:
```rust
if let Err(e) = fused.trainer().launch_kelly_bayesian_priors_update(ps_dev_ptr, n_envs) {
tracing::warn!(epoch, "launch_kelly_bayesian_priors_update failed (non-fatal): {e}");
}
if let Err(e) = fused.launch_kelly_cap_update(ps_dev_ptr, n_envs) {
/* ... */
}
```
Order matters: the priors producer MUST run before the cap update so the cap kernel sees the freshly-blended priors.
### State reset registry + dispatch arms (C.10 lesson)
Four FoldReset entries (`sp14_p1_kelly_prior_*`) added to `state_reset_registry.rs`, plus four matching dispatch arms in `training_loop.rs::reset_named_state` that write the sentinels (2.0/2.0/0.01/0.01) at fold boundary. Without these dispatch arms, the FoldReset entry exists but no actual reset fires — the C.10 lesson where slot drifts across folds and the layout-fingerprint smoke test eventually catches it as a runtime crash.
The `every_fold_and_soft_reset_entry_has_dispatch_arm` regression test confirms all four new entries have arms.
### Sites modified
| File | LOC delta | Change |
|------|-----------|--------|
| `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` | +97 / -0 | 4 new slot constants + sentinels + bounds + EMA α + 2 new tests |
| `crates/ml/src/cuda_pipeline/state_layout.cuh` | +37 / -0 | C #define mirrors for slots + sentinels + defaults |
| `crates/ml/src/cuda_pipeline/kelly_bayesian_priors_update_kernel.cu` | +250 / -0 | NEW producer kernel |
| `crates/ml/src/cuda_pipeline/kelly_cap_update_kernel.cu` | +37 / -8 | Consumer migration: ISV reads + helper for cold-start fallback |
| `crates/ml/src/cuda_pipeline/trade_physics.cuh` | +63 / -9 | `kelly_position_cap` + `apply_kelly_cap` gain `isv_signals_ptr` parameter; consumer migration |
| `crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs` | +120 / -1 | `KellyBayesianPriorsUpdateOps` struct + `new()` + `launch()` |
| `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` | +90 / -0 | Cubin static + struct field + constructor + `launch_kelly_bayesian_priors_update` wrapper |
| `crates/ml/build.rs` | +21 / -0 | Manifest entry for new cubin |
| `crates/ml/src/trainers/dqn/state_reset_registry.rs` | +35 / -0 | 4 FoldReset entries |
| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | +63 / -0 | 4 dispatch arms + producer launch wiring |
| `crates/ml/tests/sp14_oracle_tests.rs` | +390 / -0 | 4 GPU-gated oracle tests |
| `docs/dqn-wire-up-audit.md` | +90 | This entry |
### Architectural decisions
1. **No new buffers**: producer reads from existing `portfolio_state[n_envs, PS_STRIDE]`, the same buffer `kelly_cap_update_kernel` already aggregates Kelly stats from at the same boundary. Symmetry preserved; no buffer-allocation surface area added.
2. **Per-fold-end cadence (= per-epoch boundary in current scheduling)**: matches `kelly_cap_update`. The producer runs RIGHT BEFORE `kelly_cap_update` so that consumer sees the freshly-blended priors. Same ordering pattern as `aux_horizon_update_chain` runs before downstream consumers.
3. **α=0.005 (slower than reward_cap's α=0.01)**: Bayesian prior is a *prior belief* and should change slowly across folds. A fast-moving "prior" would just be a noisy Kelly estimate, defeating the purpose.
4. **NULL-tolerant + range guard on consumer side**: mirrors the existing `kelly_f_smooth` / `kelly_warmup_floor_sp9` patterns in `unified_env_step_core`. Validation envs that don't pass the bus through fall back to static defaults; in-range sentinels also fall back (within tolerance of the pre-P1-Producer hardcoded values). Single source of truth for the adaptive priors; consumers never re-derive bounds.
5. **Atomic per-batch commit (Item 1 of audit) — Item 2 deferred**: the audit asked for two items combined; Item 1 is independently completable, Item 2's audit-spec was wrong (see Concerns). Per `feedback_no_partial_refactor`, Item 1 is committed atomically with all consumers migrated.
### Verification
- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --tests --all-targets` — clean (warnings unchanged from prior baseline).
- `cargo build -p ml --release` — clean.
- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --lib --release sp14` — 10/10 pass (including new `sp14_p1_kelly_prior_slot_layout_locked` + `all_sp14_p1_slots_fit_within_isv_total_dim`).
- `cargo test -p ml --lib --release dispatch` — `every_fold_and_soft_reset_entry_has_dispatch_arm` passes (C.10 regression closed for the 4 new entries).
- `cargo test -p ml --test sp14_oracle_tests --release layout_fingerprint` — `layout_fingerprint_bumps_after_sp14_wire` passes.
- 4 new oracle tests in `sp14_oracle_tests.rs::sp14_p1_kelly_priors_gpu` (#[ignore = "requires GPU"]): Pearl-A bootstrap, no-realized-trades preserves ISV, bounds clamp, slow EMA after bootstrap. GPU smoke runs at deploy time.
### Cumulative WR-plateau fix series (continued)
- Class C bug 1 + P0-B (`8f218cab2`): replay buffer intent→realized + Kelly warmup floor wiring.
- P0-C (`316db416b`): MIN_HOLD_TARGET adaptive from AVG_WIN_HOLD_TIME.
- P0-A (`394de7d43`): REWARD_POS/NEG_CAP adaptive producer (slots 452-453).
- P1 wiring (`c4b6d6ef2`): var_floor q_gap-only adaptive (1 of 4 wireable, 3 deferred for slot allocation).
- P0-A-downstream (`657972a4b`): DD penalty + MIN_HOLD_PENALTY_MAX scale to POS_CAP_ADAPTIVE.
- **P1-Producer (this commit): adaptive Bayesian Kelly priors (slots 454-457).**
### Concerns
1. **Item 2 (MIN_HOLD_TEMPERATURE EMA) DEFERRED — audit-spec error**: The prompt described "MIN_HOLD_TEMPERATURE = 0.5f hardcoded somewhere" but the actual code has `MIN_HOLD_TEMPERATURE_{START=50.0f, END=5.0f, DECAY=20.0f}` defined in `state_layout.cuh:270-272` and consumed via the **per-epoch annealing schedule** `min_hold_temperature_for_epoch(epoch)` in `training_loop.rs:68-73`:
```rust
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()
}
```
This is already an epoch-driven adaptive schedule (50→5 over training) per the SP12 v3 design. The kernel takes T as a runtime scalar specifically to enable Phase 2 ISV-driven lift "without recompiling cubin" per the design comment in the same function — Phase 2 is anticipated, but the *signal* that should drive it is a separate spec decision. Per `feedback_no_quickfixes.md` (every issue gets a proper fix per established patterns), guessing at the right Phase 2 signal would not be principled — the right move is to defer until a follow-up spec defines the signal (likely a feedback control on observed hold-length distribution vs target band, but that's a design decision, not a wiring fix). Reporting back per the prompt's hard rule #8.
2. **Layout-fingerprint bump invalidates all pre-P1-Producer DQN checkpoints**: `LAYOUT_FINGERPRINT_CURRENT` regenerated from the new seed string with `KELLY_PRIOR_*=454..457` and `ISV_TOTAL_DIM=458`. Old checkpoints fail-fast on load (intentional, per spec §4.A.2 — no migration path). New training runs only.
3. **Validation envs receive the bus through `isv_signals_ptr`**: the migration extends `apply_kelly_cap` / `kelly_position_cap` signatures with `isv_signals_ptr`. NULL-tolerant for safety, but production code should pass the bus through so val and train see the same priors. `unified_env_step_core` (line 898) already does — same parameter as `kelly_f_smooth` / `kelly_warmup_floor_sp9` paths. No `backtest_env_kernel.cu` direct callers verified — the only caller in the codebase is `unified_env_step_core` (verified by `grep -nE "apply_kelly_cap\\("`).