From 85792ed28a934e17d4a4ab3c04b2172ccd357d0f Mon Sep 17 00:00:00 2001
From: jgrusewski <jgrusewski@localhost>
Date: Fri, 15 May 2026 17:28:07 +0200
Subject: [PATCH] feat(alpha): stacker-threshold + Kelly-attenuation controller
 kernel
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

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.
---
 crates/ml/build.rs                            |   7 +
 crates/ml/src/cuda_pipeline/alpha_kernels.rs  | 242 ++++++++++++++++++
 .../stacker_threshold_controller.cu           | 146 +++++++++++
 docs/isv-slots.md                             |  36 +++
 4 files changed, 431 insertions(+)
 create mode 100644 crates/ml/src/cuda_pipeline/stacker_threshold_controller.cu

diff --git a/crates/ml/build.rs b/crates/ml/build.rs
index fd81ebb3f..5604e04fc 100644
--- a/crates/ml/build.rs
+++ b/crates/ml/build.rs
@@ -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)
diff --git a/crates/ml/src/cuda_pipeline/alpha_kernels.rs b/crates/ml/src/cuda_pipeline/alpha_kernels.rs
index 0ba8d700c..2b7e43ead 100644
--- a/crates/ml/src/cuda_pipeline/alpha_kernels.rs
+++ b/crates/ml/src/cuda_pipeline/alpha_kernels.rs
@@ -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"
+        );
+    }
 }
diff --git a/crates/ml/src/cuda_pipeline/stacker_threshold_controller.cu b/crates/ml/src/cuda_pipeline/stacker_threshold_controller.cu
new file mode 100644
index 000000000..88002d1d9
--- /dev/null
+++ b/crates/ml/src/cuda_pipeline/stacker_threshold_controller.cu
@@ -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();
+}
diff --git a/docs/isv-slots.md b/docs/isv-slots.md
index 26bb78f7d..6e53b1fc3 100644
--- a/docs/isv-slots.md
+++ b/docs/isv-slots.md
@@ -720,3 +720,39 @@ Also fixed the diagnostic: `weight_norm` was direction-insensitive (orthogonal r
 `early_mvmt` grew monotonically (0.005 → 0.021) — direction-sensitive diagnostic confirms genuine policy learning across episodes.
 
 Next: tune for the real H=600 / 1000-episode run; if PASS holds, proceed to Task 13 (H=6000 scale-up).
+
+## Phase E.2 Task 16 — stacker-threshold controller kernel (2026-05-15)
+
+`crates/ml/src/cuda_pipeline/stacker_threshold_controller.cu` is the engagement-rate self-correcting controller for the alpha trading system. Single-block, single-thread; runs once per rollout-end boundary.
+
+Drives:
+- ISV[543] `STACKER_THRESHOLD_INDEX` (clamped [0, 0.5]) — P-controller on rate error
+- ISV[545] `TRADE_RATE_OBSERVED_EMA_INDEX` — Pearl A bootstrap + Pearl D Wiener-α (floored at 0.4 per `pearl_wiener_alpha_floor_for_nonstationary`)
+- ISV[546] `STACKER_KELLY_ATTENUATION_INDEX` (clamped [0.1, 1.0]) — P-controller on Sharpe error
+
+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_mean, x_lag)
+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, ISV[546] + k_atten · err_sharpe)  (sentinel 0 → start at 1.0)
+```
+
+Wiener-α is computed INLINE (not via the canonical `apply_pearls_ad_kernel` chain) because the controller's job IS to drive the slot — the EMA is part of the control loop, not a separate diagnostic concern. Lower latency, fewer kernels per step.
+
+**Test:** `stacker_threshold_controller_smoke_matches_hand_computation` in `alpha_kernels.rs::tests` verifies two iterations:
+- Iter 1 (Pearl A): observed = 0.30 → ISV[545] = 0.30 raw; threshold 0.05 → 0.052
+- Iter 2 (Pearl D): observed = 0.05 → ISV[545] = 0.175 (α* hit floor 0.5); threshold 0.052 → 0.05275
+
+Both within 1e-5 tolerance. Anchor slot 544 unchanged after both iterations.
+
+**Cubin:** `target/release/build/ml-*/out/stacker_threshold_controller.cubin`.
+
+**Wiring:** Task 17 — invoke at each rollout-end in `alpha_dqn_h600_smoke.rs`; initialize slot 544 = 0.08 (8% target trade rate) at training start.