feat(alpha): stacker-threshold + Kelly-attenuation controller kernel

Phase E.2 Task 16. Engagement-rate self-correcting controller per
pearl_engagement_rate_self_correction. Single-block, single-thread
kernel; runs once per rollout-end boundary.

ISV slots driven:
  543  STACKER_THRESHOLD_INDEX        clamp [0, 0.5]  P-controller on rate
  545  TRADE_RATE_OBSERVED_EMA_INDEX  Pearl A+D floored Wiener-α
  546  STACKER_KELLY_ATTENUATION_INX  clamp [0.1, 1.0] P-controller on Sharpe

Reads ISV[544] TRADE_RATE_TARGET_INDEX (TrainingPersist anchor, set once
at training start).

Control law:
  observed = trade_count / max(decisions, 1)
  ISV[545] ← Pearl_A+D_floored(observed, prev, x_lag)
              [α* floor = 0.4 per pearl_wiener_alpha_floor_for_nonstationary;
               controller co-adapts with policy → need responsive EMA]
  err_rate = ISV[545] - ISV[544]
  ISV[543] ← clamp(0, 0.5, ISV[543] + k_threshold · err_rate)

  err_sharpe = rollout_sharpe - target_sharpe
  ISV[546] ← clamp(0.1, 1.0, prev_atten + k_atten · err_sharpe)
             where prev_atten = 1.0 if ISV[546] == 0.0 (sentinel-start)
             else ISV[546]

Wiener-α is INLINE (not via canonical apply_pearls_ad_kernel chain)
because the EMA is part of the control loop, not a separate diagnostic
slot. Lower latency, fewer kernels per step.

Floor at 0.1 on Kelly attenuation per
pearl_blend_formulas_must_have_permanent_floor — can't be 0, would
zero out position sizing permanently.

Pub launcher `launch_stacker_threshold_controller` in alpha_kernels.rs
with full safety asserts. Slot indices passed as i32 args (decouple
slot numbering from kernel).

GPU smoke test `stacker_threshold_controller_smoke_matches_hand_
computation` verifies 2-iteration sequence:
  Iter 1 (Pearl A):  observed=0.30 → ISV[545]=0.30; ISV[543]: 0.05 → 0.052
  Iter 2 (Pearl D):  observed=0.05 → ISV[545]=0.175 (α* hit floor 0.5)
                                     ISV[543]: 0.052 → 0.05275
Both within 1e-5 tolerance. Anchor slot 544 unchanged.

`cargo test -p ml --lib alpha_kernels`: 6 pass (compile witness + 5 GPU
smokes including this one) on RTX 3050 Ti in 2.18s.

Audit doc docs/isv-slots.md updated per Invariant 7.
This commit is contained in:
jgrusewski
2026-05-15 17:28:07 +02:00
parent 5c0bcb1fdb
commit 85792ed28a
4 changed files with 431 additions and 0 deletions

View File

@@ -1608,6 +1608,13 @@ fn main() {
// can't learn the kill-criteria gate, no architecture upgrade
// will save it.
"alpha_linear_q.cu",
// Phase E.2 Task 16 (2026-05-15): stacker-threshold controller +
// Kelly-attenuation controller. Engagement-rate self-correction —
// P-controllers on trade-rate error (drives slot 543) and Sharpe
// error (drives slot 546). Pearl A + Pearl D Wiener-α (floored
// at 0.4 per pearl_wiener_alpha_floor_for_nonstationary) on the
// observed-rate EMA at slot 545.
"stacker_threshold_controller.cu",
];
// ALL kernels get common header (BF16 types + wrappers)

View File

@@ -235,6 +235,72 @@ pub unsafe fn launch_alpha_linear_q_sgd_step(
Ok(())
}
/// Launch `stacker_threshold_controller_update` on `stream`.
///
/// Phase E.2 Task 16. Single-block, single-thread controller that runs
/// once per rollout-end boundary. Reads per-rollout scalars (trade
/// count, total decisions, realized Sharpe), updates ISV slots 543
/// (threshold), 545 (observed-rate EMA), 546 (Kelly attenuation) via
/// engagement-rate self-correction (P-controller).
///
/// # Safety
///
/// `isv_dev` and `wiener_state_dev` MUST be valid device pointers.
/// `wiener_state_dev` MUST point at a `[3]`-float buffer
/// (`[sample_var, diff_var, x_lag]` for slot 545's Pearl D state).
pub unsafe fn launch_stacker_threshold_controller(
stream: &cudarc::driver::CudaStream,
kernel: &cudarc::driver::CudaFunction,
rollout_trade_count: f32,
rollout_total_decisions: f32,
rollout_realized_sharpe: f32,
target_sharpe: f32,
k_threshold: f32,
k_atten: f32,
wiener_alpha_floor: f32,
alpha_meta: f32,
stacker_threshold_idx: i32,
trade_rate_target_idx: i32,
trade_rate_observed_idx: i32,
stacker_kelly_atten_idx: i32,
isv_dev: u64,
wiener_state_dev: u64,
) -> Result<(), MLError> {
use cudarc::driver::{LaunchConfig, PushKernelArg};
debug_assert!(isv_dev != 0, "isv_dev must be a valid device pointer");
debug_assert!(wiener_state_dev != 0, "wiener_state_dev must be a valid device pointer");
debug_assert!(wiener_alpha_floor >= 0.0 && wiener_alpha_floor < 1.0,
"wiener_alpha_floor must be in [0, 1)");
debug_assert!(alpha_meta > 0.0 && alpha_meta < 1.0,
"alpha_meta must be in (0, 1)");
let cfg = LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (1, 1, 1),
shared_mem_bytes: 0,
};
stream
.launch_builder(kernel)
.arg(&rollout_trade_count)
.arg(&rollout_total_decisions)
.arg(&rollout_realized_sharpe)
.arg(&target_sharpe)
.arg(&k_threshold)
.arg(&k_atten)
.arg(&wiener_alpha_floor)
.arg(&alpha_meta)
.arg(&stacker_threshold_idx)
.arg(&trade_rate_target_idx)
.arg(&trade_rate_observed_idx)
.arg(&stacker_kelly_atten_idx)
.arg(&isv_dev)
.arg(&wiener_state_dev)
.launch(cfg)
.map_err(|e| MLError::ModelError(format!("stacker_threshold_controller launch: {e}")))?;
Ok(())
}
/// Launch `alpha_clip_inplace_kernel` on `stream`. Element-wise
/// clamp(x, -bound, +bound) for gradient clipping or activation
/// stabilization.
@@ -365,6 +431,12 @@ static APPLY_PEARLS_CUBIN: &[u8] =
pub static ALPHA_LINEAR_Q_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/alpha_linear_q.cubin"));
/// Precompiled stacker-threshold + Kelly-attenuation controller cubin.
/// Phase E.2 Task 16. Loaded once at trainer startup; launched at each
/// rollout-end boundary.
pub static STACKER_THRESHOLD_CONTROLLER_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/stacker_threshold_controller.cubin"));
#[cfg(test)]
mod tests {
use super::*;
@@ -1026,4 +1098,174 @@ mod tests {
);
}
}
/// End-to-end GPU smoke for the stacker-threshold + Kelly-attenuation
/// controller (Phase E.2 Task 16). Two rollouts with synthetic stats
/// and hand-computed expected ISV transitions.
///
/// Setup:
/// - ISV[543] (threshold) = 0.05 initial
/// - ISV[544] (target rate) = 0.10 anchor
/// - ISV[545] (observed EMA) = 0.0 sentinel
/// - ISV[546] (Kelly atten) = 0.0 sentinel
/// - target_sharpe = 0.5, k_threshold = 0.01, k_atten = 0.005,
/// wiener_alpha_floor = 0.4, alpha_meta = 0.1
///
/// **Iter 1 (Pearl A bootstrap):**
/// - rollout: 30 trades / 100 decisions → observed = 0.30
/// - rollout_realized_sharpe = 0.7
/// - Pearl A: ISV[545] ← 0.30, wiener_state[2] ← 0.30
/// - err_rate = 0.30 - 0.10 = 0.20
/// - ISV[543] ← clamp(0, 0.5, 0.05 + 0.01·0.20) = 0.052
/// - err_sharpe = 0.7 - 0.5 = 0.2
/// - sentinel ISV[546]=0 → start = 1.0; clamp(0.1, 1.0, 1.0 + 0.005·0.2) = 1.0
///
/// **Iter 2 (Pearl D update):**
/// - rollout: 5 trades / 100 decisions → observed = 0.05
/// - rollout_realized_sharpe = 0.3
/// - sample_diff = 0.05 - 0.30 = -0.25
/// - x_lag_diff = 0.05 - 0.30 = -0.25
/// - new_sample_var = (1-0.1)·0 + 0.1·0.0625 = 0.00625
/// - new_diff_var = (1-0.1)·0 + 0.1·0.0625 = 0.00625
/// - alpha_opt_raw = 0.00625 / (0.0125 + 1e-8) ≈ 0.5
/// - alpha_opt = max(0.5, 0.4) = 0.5
/// - ISV[545] ← (1-0.5)·0.30 + 0.5·0.05 = 0.175
/// - err_rate = 0.175 - 0.10 = 0.075
/// - ISV[543] ← clamp(0, 0.5, 0.052 + 0.01·0.075) = 0.05275
/// - err_sharpe = 0.3 - 0.5 = -0.2
/// - ISV[546] ← clamp(0.1, 1.0, 1.0 + 0.005·-0.2) = 0.999
#[test]
fn stacker_threshold_controller_smoke_matches_hand_computation() {
use crate::cuda_pipeline::alpha_isv_slots::{
STACKER_KELLY_ATTENUATION_INDEX, STACKER_THRESHOLD_INDEX,
TRADE_RATE_OBSERVED_EMA_INDEX, TRADE_RATE_TARGET_INDEX,
};
let ctx = match CudaContext::new(0) {
Ok(c) => c,
Err(e) => {
eprintln!("Skipping GPU smoke — no CUDA device 0: {e}");
return;
}
};
let stream = ctx.default_stream();
let module = ctx
.load_cubin(STACKER_THRESHOLD_CONTROLLER_CUBIN.to_vec())
.expect("controller cubin load");
let kernel = module
.load_function("stacker_threshold_controller_update")
.expect("controller kernel load");
// ISV buffer sized to cover slot 546 + slack.
let mut isv_host = vec![0.0_f32; 552];
isv_host[STACKER_THRESHOLD_INDEX] = 0.05;
isv_host[TRADE_RATE_TARGET_INDEX] = 0.10;
// Slot 545 + 546 start at sentinel 0.
let mut isv_dev = stream.clone_htod(&isv_host).expect("upload isv");
let mut wiener_dev = stream
.alloc_zeros::<f32>(3)
.expect("alloc wiener state");
let target_sharpe = 0.5_f32;
let k_threshold = 0.01_f32;
let k_atten = 0.005_f32;
let wiener_alpha_floor = 0.4_f32;
let alpha_meta = 0.1_f32;
// --- Iter 1 ---
{
let (isv_ptr, _g0) = isv_dev.device_ptr_mut(&stream);
let (w_ptr, _g1) = wiener_dev.device_ptr_mut(&stream);
unsafe {
launch_stacker_threshold_controller(
&stream,
&kernel,
30.0, // trade_count
100.0, // total_decisions
0.7, // realized_sharpe
target_sharpe,
k_threshold,
k_atten,
wiener_alpha_floor,
alpha_meta,
STACKER_THRESHOLD_INDEX as i32,
TRADE_RATE_TARGET_INDEX as i32,
TRADE_RATE_OBSERVED_EMA_INDEX as i32,
STACKER_KELLY_ATTENUATION_INDEX as i32,
isv_ptr,
w_ptr,
)
.expect("controller iter 1");
}
}
stream.synchronize().expect("sync iter 1");
let isv1 = stream.clone_dtoh(&isv_dev).expect("dtoh iter 1");
assert!(
(isv1[TRADE_RATE_OBSERVED_EMA_INDEX] - 0.30).abs() < 1e-5,
"iter 1 Pearl A: observed_ema = {} (expected 0.30)",
isv1[TRADE_RATE_OBSERVED_EMA_INDEX]
);
assert!(
(isv1[STACKER_THRESHOLD_INDEX] - 0.052).abs() < 1e-5,
"iter 1 threshold = {} (expected 0.052)",
isv1[STACKER_THRESHOLD_INDEX]
);
assert!(
(isv1[STACKER_KELLY_ATTENUATION_INDEX] - 1.0).abs() < 1e-5,
"iter 1 atten = {} (expected 1.0 — clamped to ceiling)",
isv1[STACKER_KELLY_ATTENUATION_INDEX]
);
// --- Iter 2 (Pearl D) ---
{
let (isv_ptr, _g0) = isv_dev.device_ptr_mut(&stream);
let (w_ptr, _g1) = wiener_dev.device_ptr_mut(&stream);
unsafe {
launch_stacker_threshold_controller(
&stream,
&kernel,
5.0, // trade_count
100.0, // total_decisions
0.3, // realized_sharpe
target_sharpe,
k_threshold,
k_atten,
wiener_alpha_floor,
alpha_meta,
STACKER_THRESHOLD_INDEX as i32,
TRADE_RATE_TARGET_INDEX as i32,
TRADE_RATE_OBSERVED_EMA_INDEX as i32,
STACKER_KELLY_ATTENUATION_INDEX as i32,
isv_ptr,
w_ptr,
)
.expect("controller iter 2");
}
}
stream.synchronize().expect("sync iter 2");
let isv2 = stream.clone_dtoh(&isv_dev).expect("dtoh iter 2");
assert!(
(isv2[TRADE_RATE_OBSERVED_EMA_INDEX] - 0.175).abs() < 1e-4,
"iter 2 Pearl D: observed_ema = {} (expected 0.175)",
isv2[TRADE_RATE_OBSERVED_EMA_INDEX]
);
assert!(
(isv2[STACKER_THRESHOLD_INDEX] - 0.05275).abs() < 1e-5,
"iter 2 threshold = {} (expected 0.05275)",
isv2[STACKER_THRESHOLD_INDEX]
);
assert!(
(isv2[STACKER_KELLY_ATTENUATION_INDEX] - 0.999).abs() < 1e-5,
"iter 2 atten = {} (expected 0.999)",
isv2[STACKER_KELLY_ATTENUATION_INDEX]
);
// Target slot 544 must be unchanged (TrainingPersist anchor).
assert_eq!(
isv2[TRADE_RATE_TARGET_INDEX], 0.10,
"target slot 544 must not be modified by the controller"
);
}
}

View File

@@ -0,0 +1,146 @@
// crates/ml/src/cuda_pipeline/stacker_threshold_controller.cu
//
// Phase E.2 Task 16 (2026-05-15) — Stacker-threshold controller +
// Kelly-attenuation controller. Engagement-rate self-correction per
// `pearl_engagement_rate_self_correction`. Single block, single thread;
// runs once per rollout-end.
//
// ISV slots driven (see `alpha_isv_slots.rs`):
//
// 543 STACKER_THRESHOLD_INDEX — controller output [0, 0.5]
// 545 TRADE_RATE_OBSERVED_EMA_INDEX — diagnostic (Pearl A+D EMA)
// 546 STACKER_KELLY_ATTENUATION_INX — controller output [0.1, 1.0]
//
// ISV slots read:
//
// 544 TRADE_RATE_TARGET_INDEX — anchor (host-init, never reset)
//
// **Control law:**
//
// Trade rate (P-controller on the rate error):
// observed = trade_count / max(decisions, 1)
// ISV[545]_new ← Pearl A + Pearl D Wiener-α with floor=0.4 on observed
// (the controller co-adapts with the policy, so the floor
// keeps the EMA responsive per
// `pearl_wiener_alpha_floor_for_nonstationary`)
// err_rate = ISV[545]_new ISV[544] (signed)
// ISV[543] ← clamp(ISV[543] + k_threshold · err_rate, 0, 0.5)
//
// Kelly attenuation (P-controller on Sharpe error, floored):
// err_sharpe = rollout_sharpe target_sharpe
// ISV[546] ← clamp(ISV[546] + k_atten · err_sharpe, 0.1, 1.0)
//
// `k_threshold = 0.01`, `k_atten = 0.005` per plan defaults. Floored at
// the lower clamp bound per `pearl_blend_formulas_must_have_permanent_floor`
// (the Kelly attenuation can't be 0 — that would zero out position sizing
// permanently; floor at 0.1 keeps the policy minimally active).
//
// **Wiener-α inline** (Pearl A + Pearl D with floor): same math as
// `apply_pearls_ad_kernel`, with `alpha_meta` overridden to a higher
// constant for the variance updates because the rate observation has
// O(0.1) range vs SP4's O(1) — needs faster variance tracking.
#include <cuda_runtime.h>
extern "C" __global__ void stacker_threshold_controller_update(
/* Per-rollout raw inputs (mapped-pinned host scalars). */
float rollout_trade_count,
float rollout_total_decisions,
float rollout_realized_sharpe,
/* Kernel-arg controller parameters. */
float target_sharpe,
float k_threshold, // typ. 0.01
float k_atten, // typ. 0.005
float wiener_alpha_floor, // 0.4 per pearl_wiener_alpha_floor_for_nonstationary
float alpha_meta, // meta-EMA rate for Wiener variance updates
/* ISV slot indices passed as args (decouple slot numbering from kernel). */
int stacker_threshold_idx,
int trade_rate_target_idx,
int trade_rate_observed_idx,
int stacker_kelly_atten_idx,
/* Device buffers. */
float* __restrict__ isv,
float* __restrict__ wiener_state /* [sample_var, diff_var, x_lag] for slot 545 */
) {
if (blockIdx.x != 0 || threadIdx.x != 0) return;
constexpr float EPS_DIV = 1.0e-8f;
constexpr float THRESH_LO = 0.0f;
constexpr float THRESH_HI = 0.5f;
constexpr float ATTEN_LO = 0.1f;
constexpr float ATTEN_HI = 1.0f;
// ---- (1) Observed trade rate ----
const float observed_rate =
rollout_trade_count / fmaxf(rollout_total_decisions, 1.0f);
// ---- (2) Pearl A + Pearl D update on ISV[545] (observed-rate EMA) ----
//
// Pearl A: sentinel-bootstrap. If both prev_x_mean == 0 and x_lag == 0
// (untouched slot or fold-boundary reset), replace directly with the
// observation. Wiener state initialised to (sample_var=0, diff_var=0,
// x_lag=observation).
//
// Pearl D: variance-EMA at alpha_meta; α* = diff_var / (diff_var +
// sample_var + EPS_DIV); EMA blend `(1-α)·prev + α·obs`. Then APPLY
// FLOOR: `α* = max(α*, wiener_alpha_floor)` per
// `pearl_wiener_alpha_floor_for_nonstationary` — the controller
// co-adapts with the policy so we can't let α* drift to ~0 on a
// briefly-stationary segment.
const float prev_x_mean = isv[trade_rate_observed_idx];
const float sample_var = wiener_state[0];
const float diff_var = wiener_state[1];
const float x_lag = wiener_state[2];
float new_x_mean;
float new_sample_var;
float new_diff_var;
float new_x_lag;
if (prev_x_mean == 0.0f && x_lag == 0.0f) {
// Pearl A bootstrap.
new_x_mean = observed_rate;
new_sample_var = 0.0f;
new_diff_var = 0.0f;
new_x_lag = observed_rate;
} else {
const float sample_diff = observed_rate - prev_x_mean;
const float x_lag_diff = observed_rate - x_lag;
new_sample_var = (1.0f - alpha_meta) * sample_var
+ alpha_meta * sample_diff * sample_diff;
new_diff_var = (1.0f - alpha_meta) * diff_var
+ alpha_meta * x_lag_diff * x_lag_diff;
const float alpha_opt_raw =
new_diff_var / (new_diff_var + new_sample_var + EPS_DIV);
const float alpha_opt = fmaxf(alpha_opt_raw, wiener_alpha_floor);
new_x_mean = (1.0f - alpha_opt) * prev_x_mean + alpha_opt * observed_rate;
new_x_lag = observed_rate;
}
isv[trade_rate_observed_idx] = new_x_mean;
wiener_state[0] = new_sample_var;
wiener_state[1] = new_diff_var;
wiener_state[2] = new_x_lag;
// ---- (3) Update STACKER_THRESHOLD (slot 543) from rate error ----
const float target_rate = isv[trade_rate_target_idx];
const float err_rate = new_x_mean - target_rate;
const float old_threshold = isv[stacker_threshold_idx];
const float new_threshold =
fmaxf(THRESH_LO, fminf(THRESH_HI, old_threshold + k_threshold * err_rate));
isv[stacker_threshold_idx] = new_threshold;
// ---- (4) Update KELLY_ATTENUATION (slot 546) from Sharpe error ----
//
// Pearl A bootstrap: if slot is 0 (sentinel — untouched / fold-reset),
// start at the ceiling (1.0) so the policy is fully aggressive from the
// first rollout and the controller dials it down from there. NOT
// floored at sentinel — sentinel is the START state, the floor is the
// operational floor.
const float prev_atten_raw = isv[stacker_kelly_atten_idx];
const float prev_atten = prev_atten_raw == 0.0f ? ATTEN_HI : prev_atten_raw;
const float err_sharpe = rollout_realized_sharpe - target_sharpe;
const float new_atten =
fmaxf(ATTEN_LO, fminf(ATTEN_HI, prev_atten + k_atten * err_sharpe));
isv[stacker_kelly_atten_idx] = new_atten;
__threadfence_system();
}