diff --git a/crates/ml/build.rs b/crates/ml/build.rs
index 523be9ab1..5d53fbde9 100644
--- a/crates/ml/build.rs
+++ b/crates/ml/build.rs
@@ -786,6 +786,26 @@ fn main() {
         // (no atomicAdd) per `feedback_no_atomicadd.md`. Pearl-A
         // first-observation bootstrap; α=0.05 EMA thereafter. Plan: §C.4b.
         "avg_win_hold_time_update_kernel.cu",
+        // Class A P0-A (2026-05-08): adaptive REWARD_POS/NEG_CAP producer.
+        // Single-block kernel that sweeps the per-epoch
+        // `step_ret_per_sample` + `trade_close_per_sample` buffers
+        // (populated by `unified_env_step_core` in
+        // `experience_kernels.cu`), computes p99(winning_returns) ×
+        // safety_factor=1.5 → POS cap, NEG = −2 × POS (Kahneman 2:1
+        // asymmetry preserved at producer-time per
+        // `pearl_audit_unboundedness_for_implicit_asymmetry` — single
+        // source of truth, no consumer applies the multiplier itself).
+        // Pearl-A first-observation bootstrap (sentinels 5.0/-10.0
+        // match pre-P0-A hardcoded constants for bit-identical
+        // cold-start); α=0.01 slow EMA thereafter. Bounds POS∈[1, 50],
+        // NEG∈[-100, -2] are Category-1 dimensional safety floors per
+        // `feedback_isv_for_adaptive_bounds`. Replaces hardcoded
+        // REWARD_POS_CAP=+5.0f / REWARD_NEG_CAP=-10.0f
+        // (state_layout.cuh:266-267) — Class A audit
+        // highest-suspected-impact fix for the months-long
+        // WR-stuck-at-46-48% plateau across 11 superprojects. Per-epoch
+        // boundary launch.
+        "reward_cap_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)
diff --git a/crates/ml/src/cuda_pipeline/compute_sp15_final_reward_kernel.cu b/crates/ml/src/cuda_pipeline/compute_sp15_final_reward_kernel.cu
index 2b12b9503..685894565 100644
--- a/crates/ml/src/cuda_pipeline/compute_sp15_final_reward_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/compute_sp15_final_reward_kernel.cu
@@ -153,8 +153,14 @@ extern "C" __global__ void compute_sp15_final_reward_kernel(
     const float dd_thr    = isv[SP15_DD_THRESHOLD_INDEX];
     r = sp15_dd_penalty(r, lambda_dd, dd_thr, dd_current);
 
-    /* Stage 4: SP12 v3 asymmetric bilateral cap via macros. */
-    r = sp15_apply_sp12_cap(r);
+    /* Stage 4: SP12 v3 asymmetric bilateral cap. Class A P0-A (2026-05-08):
+     * caps now ISV-driven from `ISV[REWARD_POS_CAP_ADAPTIVE_INDEX=452]` /
+     * `ISV[REWARD_NEG_CAP_ADAPTIVE_INDEX=453]`. Cold-start fallback: when
+     * the ISV slot is at sentinel SENTINEL_REWARD_POS_CAP=5.0 (within
+     * EPS), the helper falls back to `state_layout.cuh` macros for
+     * bit-identical pre-P0-A behavior. The helper handles the fallback
+     * internally — caller only passes the ISV pointer. */
+    r = sp15_apply_sp12_cap(r, isv);
 
     /* Final NaN guard mirrors experience_kernels.cu's pattern: if a
      * helper produced NaN/Inf, fall back to the SP11 input value
diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu
index 1bbe48c43..cf20d36f2 100644
--- a/crates/ml/src/cuda_pipeline/experience_kernels.cu
+++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu
@@ -3102,15 +3102,40 @@ extern "C" __global__ void experience_env_step(
          * to make the bound itself asymmetric in the direction the prior
          * implicitly was, not to remove the cap.
          *
-         * Constants live in state_layout.cuh as Invariant-1 numerical
-         * anchors (Phase 1). Phase 2 (deferred) lifts the ratio to an
-         * ISV-driven controller per feedback_isv_for_adaptive_bounds.
-         *
+         * Class A P0-A (2026-05-08): Phase 2 lifted. The hardcoded
+         * REWARD_POS_CAP=+5.0f / REWARD_NEG_CAP=-10.0f
+         * (state_layout.cuh:266-267) is now ISV-driven from
+         * `ISV[REWARD_POS_CAP_ADAPTIVE_INDEX=452]` /
+         * `ISV[REWARD_NEG_CAP_ADAPTIVE_INDEX=453]` — adaptively driven
+         * by `reward_cap_update_kernel` from p99(winning_returns) × 1.5
+         * with NEG=−2 × POS preserving Kahneman/Tversky 2:1 asymmetry
+         * (single source of truth at producer; no consumer applies the
+         * 2× ratio itself). Cold-start fallback: when the ISV slot is at
+         * sentinel SENTINEL_REWARD_POS_CAP=5.0 (within EPS), fall back
+         * to the original macro `REWARD_POS_CAP` constant — bit-identical
+         * pre-P0-A behavior until the first valid observation lands.
          * Math factored into `compute_asymmetric_capped_pnl` in
          * `trade_physics.cuh` so the formula is testable in isolation
          * (oracle tests in `tests/sp12_reward_math_tests.rs`). */
+        float pos_cap_eff = REWARD_POS_CAP;
+        float neg_cap_eff = REWARD_NEG_CAP;
+        if (isv_signals_ptr != NULL) {
+            float isv_pos = isv_signals_ptr[REWARD_POS_CAP_ADAPTIVE_INDEX];
+            float isv_neg = isv_signals_ptr[REWARD_NEG_CAP_ADAPTIVE_INDEX];
+            /* Cold-start fallback: sentinel match → use macro fallback
+             * (bit-identical pre-P0-A). Out-of-bounds defensive guard:
+             * any value outside [REWARD_POS_CAP_MIN_BOUND,
+             * REWARD_POS_CAP_MAX_BOUND] also falls back, defending
+             * against malformed prior state. */
+            if (fabsf(isv_pos - SENTINEL_REWARD_POS_CAP) >= 1e-6f
+                && isv_pos >= REWARD_POS_CAP_MIN_BOUND
+                && isv_pos <= REWARD_POS_CAP_MAX_BOUND) {
+                pos_cap_eff = isv_pos;
+                neg_cap_eff = isv_neg;
+            }
+        }
         float capped_pnl = compute_asymmetric_capped_pnl(
-            base_reward, REWARD_NEG_CAP, REWARD_POS_CAP
+            base_reward, neg_cap_eff, pos_cap_eff
         );
         if (trail_triggered) {
             r_trail = capped_pnl;
diff --git a/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs b/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs
index 96129b4ce..816b8fd72 100644
--- a/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs
@@ -52,7 +52,7 @@ 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,
+    AVG_WIN_HOLD_TIME_UPDATE_CUBIN, H_S2_AUX_RMS_EMA_CUBIN, REWARD_CAP_UPDATE_CUBIN,
 };
 
 /// Hidden width of the aux trunk's first internal layer (Linear_1 → ELU
@@ -618,6 +618,122 @@ impl AvgWinHoldTimeUpdateOps {
     }
 }
 
+/// Class A P0-A (2026-05-08): adaptive REWARD_POS/NEG_CAP producer.
+///
+/// Sweeps the per-epoch `step_ret_per_sample` + `trade_close_per_sample`
+/// buffers (populated by `unified_env_step_core` in `experience_kernels.cu`),
+/// computes a Welford `mean + Z_99 × sigma` p99 estimator over winning
+/// realized returns plus a `max(winning_returns)` conservative takeover,
+/// applies a 1.5× safety factor, clamps to dimensional bounds [1, 50],
+/// and writes both `ISV[REWARD_POS_CAP_ADAPTIVE_INDEX=452]` (POS cap)
+/// and `ISV[REWARD_NEG_CAP_ADAPTIVE_INDEX=453]` (NEG = −2 × POS,
+/// preserving Kahneman/Tversky 2:1 loss-aversion asymmetry per
+/// `pearl_audit_unboundedness_for_implicit_asymmetry` — moved from
+/// hardcoded scalar to producer-time multiplier; SINGLE source of truth).
+///
+/// Single-block kernel; block-tree-reduce over the per-sample buffer
+/// (no atomicAdd per `feedback_no_atomicadd.md`). Pearl-A
+/// first-observation bootstrap (sentinel 5.0/-10.0 matches pre-fix
+/// hardcoded constants for bit-identical cold-start); α=0.01 slow EMA
+/// thereafter.
+///
+/// Per-epoch boundary launch — reward distribution is the foundation
+/// of training and shouldn't move fast; per-step would track sample noise.
+///
+/// # Pearls applied
+/// - `feedback_no_atomicadd.md` — single block, block-tree-reduce in shmem.
+/// - `pearl_first_observation_bootstrap.md` — sentinel = 5.0 → REPLACE.
+/// - `pearl_no_host_branches_in_captured_graph.md` — `CudaFunction`
+///   pre-loaded at construction; on-device guards.
+/// - `pearl_symmetric_clamp_audit.md` — bilateral `fmaxf(lo, fminf(x, hi))`
+///   on POS, NEG derived clamp.
+/// - `feedback_isv_for_adaptive_bounds.md` — POS bounds [1, 50] are
+///   Category-1 dimensional safety, NOT tuning.
+#[allow(missing_debug_implementations)]
+pub(crate) struct RewardCapUpdateOps {
+    update_kernel: CudaFunction,
+}
+
+impl RewardCapUpdateOps {
+    /// Block dim used by the producer kernel — must match `BLK_DIM` in
+    /// `reward_cap_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(REWARD_CAP_UPDATE_CUBIN.to_vec())
+            .map_err(|e| MLError::ModelError(format!("reward_cap_update cubin load: {e}")))?;
+        let update_kernel = module
+            .load_function("reward_cap_update")
+            .map_err(|e| MLError::ModelError(format!("reward_cap_update load: {e}")))?;
+        Ok(Self { update_kernel })
+    }
+
+    /// Launch the adaptive reward-cap producer.
+    ///
+    /// Args:
+    ///   - `step_ret_ptr`: f32 device ptr `[total_samples]` —
+    ///     `step_ret_core` raw signed per-step return.
+    ///   - `trade_close_ptr`: i32 device ptr `[total_samples]` — 1 iff
+    ///     exiting_trade || reversing_trade.
+    ///   - `total_samples`: N*L (B*T total samples this epoch).
+    ///   - `isv_ptr`: ISV[ISV_TOTAL_DIM] device pointer.
+    ///   - `pos_idx`: REWARD_POS_CAP_ADAPTIVE_INDEX (452).
+    ///   - `neg_idx`: REWARD_NEG_CAP_ADAPTIVE_INDEX (453).
+    ///   - `sentinel_pos`: SENTINEL_REWARD_POS_CAP (5.0).
+    ///   - `sentinel_neg`: SENTINEL_REWARD_NEG_CAP (-10.0).
+    ///   - `safety_factor`: REWARD_CAP_SAFETY_FACTOR (1.5).
+    ///   - `neg_to_pos_ratio`: REWARD_NEG_TO_POS_RATIO (2.0).
+    ///   - `pos_min`, `pos_max`: REWARD_POS_CAP_MIN/MAX (1.0, 50.0).
+    ///   - `alpha`: REWARD_CAP_EMA_ALPHA (0.01).
+    #[allow(clippy::too_many_arguments)]
+    pub(crate) fn launch(
+        &self,
+        stream: &Arc<CudaStream>,
+        step_ret_ptr: u64,
+        trade_close_ptr: u64,
+        total_samples: i32,
+        isv_ptr: u64,
+        pos_idx: i32,
+        neg_idx: i32,
+        sentinel_pos: f32,
+        sentinel_neg: f32,
+        safety_factor: f32,
+        neg_to_pos_ratio: f32,
+        pos_min: f32,
+        pos_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(&step_ret_ptr)
+                .arg(&trade_close_ptr)
+                .arg(&total_samples)
+                .arg(&isv_ptr)
+                .arg(&pos_idx)
+                .arg(&neg_idx)
+                .arg(&sentinel_pos)
+                .arg(&sentinel_neg)
+                .arg(&safety_factor)
+                .arg(&neg_to_pos_ratio)
+                .arg(&pos_min)
+                .arg(&pos_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!("reward_cap_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
diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
index 8537ce09c..4bb49025e 100644
--- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
@@ -2099,6 +2099,17 @@ pub(crate) static AUX_HORIZON_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!(
 /// `gpu_aux_trunk::AvgWinHoldTimeUpdateOps`. Per-epoch boundary launch.
 pub(crate) static AVG_WIN_HOLD_TIME_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/avg_win_hold_time_update_kernel.cubin"));
 
+/// Class A P0-A (2026-05-08): adaptive REWARD_POS/NEG_CAP producer cubin.
+/// Single-block kernel that sweeps the per-epoch `step_ret_per_sample` +
+/// `trade_close_per_sample` buffers, computes a Welford p99 estimator of
+/// realized winning returns × 1.5× safety factor → POS cap, and writes
+/// both `ISV[REWARD_POS_CAP_ADAPTIVE_INDEX=452]` and
+/// `ISV[REWARD_NEG_CAP_ADAPTIVE_INDEX=453]` (NEG = −2 × POS, Kahneman
+/// asymmetry preserved at producer-time). Pearl-A first-observation
+/// bootstrap; α=0.01 slow EMA thereafter. Loaded by
+/// `gpu_aux_trunk::RewardCapUpdateOps`. Per-epoch boundary launch.
+pub(crate) static REWARD_CAP_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/reward_cap_update_kernel.cubin"));
+
 /// SP14 Layer C Phase C.6 (2026-05-08): h_s2_aux RMS EMA producer cubin.
 /// Single-block 256-thread kernel computing `sqrt(mean(h_s2_aux²))` over
 /// `h_s2_aux [B, SH2]` (aux trunk final output) and EMA-blending into
@@ -2427,7 +2438,7 @@ const ISV_NETWORK_DIM: usize = 23;
 ///         (shifted 112→116 in Plan 4 Task 6 Commit A).
 /// Written by the constructor; checked at checkpoint load. Fail-fast only — no migration
 /// path exists. See spec §4.A.2 and `LAYOUT_FINGERPRINT_CURRENT` for structural-hash rationale.
-pub(crate) const ISV_TOTAL_DIM: usize = 452;  // SP5 + Layer D D1+D2+D3 + SP7 + SP8 + SP9 + SP10 + SP11 + SP13 + SP13 v3 + SP14 (post-C.1) + SP15 + SP14-C aux trunk control plane + SP14-C Phase C.4b aux horizon.  Bumped 450 → 452 by SP14 Layer C Phase C.4b (2026-05-08): added 2 aux-prediction-horizon slots [450..452) — AUX_PRED_HORIZON_BARS (450) + AVG_WIN_HOLD_TIME_BARS (451). Aux's original label was (p_{t+1} > p_t), pure HFT-scale microstructure noise — pivoted to (p_{t+H} > p_t) with H read from ISV[450], adaptively driven from observed avg winning hold time via Pearl-A bootstrap + Wiener-α EMA. Case B: existing `hold_at_exit_per_sample` + `trade_profitable_per_sample` per-sample buffers, no aggregate slot — added new aggregate slot 451 to expose the EMA target. Bumped 444 → 450 by SP14 Layer C Phase C.1 (2026-05-08): atomically deleted 9 α-machinery slots in [385..396) AND added 6 aux-trunk control-plane slots [444..450) — AUX_TRUNK_LR (444) + AUX_TRUNK_BETA1 (445) + AUX_TRUNK_BETA2 (446) + AUX_TRUNK_EPS (447) + AUX_TRUNK_GRAD_CLIP (448) + H_S2_AUX_RMS_EMA (449). Per `feedback_no_legacy_aliases` and `feedback_no_partial_refactor`. Deleted slots left as RESERVED gap (NOT compacted) to preserve checkpoint layout fingerprint compatibility — the surviving SP14 q_disagreement_* diagnostic slots (383, 384, 389) keep their original indices.  Pre-C.1 base: 173 + 210 (164 SP5 slots @ 174..278/280..340, with 2-slot gap before cross-fold-persistent Kelly block; Layer D D1 PnL outputs at [286..290); Layer D D2 health composition outputs at [290..294); Layer D D3 training metrics EMA at [294..297); SP7 T1 loss-balance Wiener stats at [297..313); SP7 activation-flag fix per-(head,branch) flags at [313..321); SP8 Fix 36 train_active_frac canary @ [321..322) + LB_MAX_BUDGET per-(head,branch) at [322..330); SP9 Fix 37 Kelly warmup floor at [330..331) + Q_VAR_MAG_EMA at [331..332) + INTENT_EVAL_DIVERGENCE at [332..333) + 3 EMA targets at [333..336) + 3 eval_dist mag bins at [336..339); SP10 Fix 38 EVAL_THOMPSON_TEMP at [339..340); SP11 Fix 39 reward-subsystem controller at [340..367) — 6 component weights [340..346) + 4 controller scalar outputs [346..350) + 2 val-sharpe canaries [350..352) + 6 mag-ratio canaries [352..358) + saboteur engagement [358] + PnL magnitude EMA [359] + popart-component magnitude EMA [360, B1b fix-up] + per-component variance EMAs [361..367, B1b smoke-recovery z-score normalization]; SP13 directional-skill instrumentation at [372..380) — TARGET_DIR_ACC [372] + AUX_DIR_ACC_SHORT/LONG_EMA [373..375) + AUX_DIR_PREDICTION [375] + DIR_SKILL_BONUS_ALPHA/BETA [376..378) + LUCK_WIN_DISCOUNT [378] + SKILL_BONUS_CAP_RATIO [379]; SP13 v3 P0a.T3 Hold-pricing controller at [380..383) — HOLD_COST [380] + HOLD_RATE_TARGET [381] + HOLD_RATE_OBSERVED_EMA [382]; SP14 q_disagreement diagnostic at [383..385), [389] — Q_DISAGREEMENT_SHORT/LONG_EMA + Q_DISAGREEMENT_VARIANCE_EMA (slots [385..389) and [390..396) RESERVED gap from C.1 deletion); SP15 [397..444) per Phase 1; SP14-C [444..450) per Phase C.1.  Intentional 5-slot boundary gap at [367..372))
+pub(crate) const ISV_TOTAL_DIM: usize = 454;  // SP5 + Layer D D1+D2+D3 + SP7 + SP8 + SP9 + SP10 + SP11 + SP13 + SP13 v3 + SP14 (post-C.1) + SP15 + SP14-C aux trunk control plane + SP14-C Phase C.4b aux horizon + Class A P0-A adaptive reward caps.  Bumped 452 → 454 by Class A P0-A (2026-05-08): added 2 adaptive-reward-cap slots [452..454) — REWARD_POS_CAP_ADAPTIVE (452) + REWARD_NEG_CAP_ADAPTIVE (453). Per Class A audit, hardcoded `REWARD_POS_CAP=+5.0f` / `REWARD_NEG_CAP=-10.0f` (state_layout.cuh:266-267) was structurally clipping the upper tail of realized alpha — controller signals (sharpe EMA, var_q, q_gap) all derive from the CAPPED buffer, so the controller cannot select for trades it cannot see. The state_layout comment explicitly deferred Phase 2 (ISV-driven) "IF Phase 1 validation reveals adaptive need"; 11 superprojects of WR-stuck-at-46-48% revealed adaptive need. Producer kernel `reward_cap_update_kernel` computes p99(winning_returns) × safety_factor=1.5 → POS cap and NEG = -2 × POS (preserves Kahneman 2:1 loss-aversion asymmetry per `pearl_audit_unboundedness_for_implicit_asymmetry` — moved from hardcoded scalar to producer-time multiplier, single source of truth). Pearl-A bootstrap (sentinels 5.0/-10.0) + Wiener-α=0.01 slow EMA. Bounds POS∈[1, 50], NEG∈[-100, -2] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`, NOT tuning. Bumped 450 → 452 by SP14 Layer C Phase C.4b (2026-05-08): added 2 aux-prediction-horizon slots [450..452) — AUX_PRED_HORIZON_BARS (450) + AVG_WIN_HOLD_TIME_BARS (451). Aux's original label was (p_{t+1} > p_t), pure HFT-scale microstructure noise — pivoted to (p_{t+H} > p_t) with H read from ISV[450], adaptively driven from observed avg winning hold time via Pearl-A bootstrap + Wiener-α EMA. Case B: existing `hold_at_exit_per_sample` + `trade_profitable_per_sample` per-sample buffers, no aggregate slot — added new aggregate slot 451 to expose the EMA target. Bumped 444 → 450 by SP14 Layer C Phase C.1 (2026-05-08): atomically deleted 9 α-machinery slots in [385..396) AND added 6 aux-trunk control-plane slots [444..450) — AUX_TRUNK_LR (444) + AUX_TRUNK_BETA1 (445) + AUX_TRUNK_BETA2 (446) + AUX_TRUNK_EPS (447) + AUX_TRUNK_GRAD_CLIP (448) + H_S2_AUX_RMS_EMA (449). Per `feedback_no_legacy_aliases` and `feedback_no_partial_refactor`. Deleted slots left as RESERVED gap (NOT compacted) to preserve checkpoint layout fingerprint compatibility — the surviving SP14 q_disagreement_* diagnostic slots (383, 384, 389) keep their original indices.  Pre-C.1 base: 173 + 210 (164 SP5 slots @ 174..278/280..340, with 2-slot gap before cross-fold-persistent Kelly block; Layer D D1 PnL outputs at [286..290); Layer D D2 health composition outputs at [290..294); Layer D D3 training metrics EMA at [294..297); SP7 T1 loss-balance Wiener stats at [297..313); SP7 activation-flag fix per-(head,branch) flags at [313..321); SP8 Fix 36 train_active_frac canary @ [321..322) + LB_MAX_BUDGET per-(head,branch) at [322..330); SP9 Fix 37 Kelly warmup floor at [330..331) + Q_VAR_MAG_EMA at [331..332) + INTENT_EVAL_DIVERGENCE at [332..333) + 3 EMA targets at [333..336) + 3 eval_dist mag bins at [336..339); SP10 Fix 38 EVAL_THOMPSON_TEMP at [339..340); SP11 Fix 39 reward-subsystem controller at [340..367) — 6 component weights [340..346) + 4 controller scalar outputs [346..350) + 2 val-sharpe canaries [350..352) + 6 mag-ratio canaries [352..358) + saboteur engagement [358] + PnL magnitude EMA [359] + popart-component magnitude EMA [360, B1b fix-up] + per-component variance EMAs [361..367, B1b smoke-recovery z-score normalization]; SP13 directional-skill instrumentation at [372..380) — TARGET_DIR_ACC [372] + AUX_DIR_ACC_SHORT/LONG_EMA [373..375) + AUX_DIR_PREDICTION [375] + DIR_SKILL_BONUS_ALPHA/BETA [376..378) + LUCK_WIN_DISCOUNT [378] + SKILL_BONUS_CAP_RATIO [379]; SP13 v3 P0a.T3 Hold-pricing controller at [380..383) — HOLD_COST [380] + HOLD_RATE_TARGET [381] + HOLD_RATE_OBSERVED_EMA [382]; SP14 q_disagreement diagnostic at [383..385), [389] — Q_DISAGREEMENT_SHORT/LONG_EMA + Q_DISAGREEMENT_VARIANCE_EMA (slots [385..389) and [390..396) RESERVED gap from C.1 deletion); SP15 [397..444) per Phase 1; SP14-C [444..450) per Phase C.1.  Intentional 5-slot boundary gap at [367..372))
 /// Legacy alias preserved for call sites that haven't been audited for the
 /// network-vs-total split. New code should pick `ISV_NETWORK_DIM` (for weight
 /// tensor sizing) or `ISV_TOTAL_DIM` (for the broadcast bus buffer).
@@ -3499,7 +3510,9 @@ const fn layout_fingerprint_seed() -> &'static [u8] {
       DD_PERSISTENCE_MAX=443;\
       AUX_TRUNK_LR=444;AUX_TRUNK_BETA1=445;AUX_TRUNK_BETA2=446;\
       AUX_TRUNK_EPS=447;AUX_TRUNK_GRAD_CLIP=448;H_S2_AUX_RMS_EMA=449;\
-      ISV_TOTAL_DIM=450;\
+      AUX_PRED_HORIZON_BARS=450;AVG_WIN_HOLD_TIME_BARS=451;\
+      REWARD_POS_CAP_ADAPTIVE=452;REWARD_NEG_CAP_ADAPTIVE=453;\
+      ISV_TOTAL_DIM=454;\
       SP14_C_AUX_TRUNK_CONTROL_PLANE=sp14_c_phase_1;\
       ALPHA_MACHINERY_DELETED=sp14_c_phase_1;\
       TRUNK_INPUT_DD_PCT=sp15_wave_4_1a;\
@@ -7502,6 +7515,16 @@ pub struct GpuDqnTrainer {
     /// (no atomicAdd); per-epoch boundary launch.
     avg_win_hold_time_update_ops: super::gpu_aux_trunk::AvgWinHoldTimeUpdateOps,
 
+    /// Class A P0-A (2026-05-08): adaptive REWARD_POS/NEG_CAP producer.
+    /// Sweeps the per-epoch `step_ret_per_sample` + `trade_close_per_sample`
+    /// buffers, computes p99(winning_returns) × 1.5 → POS cap, NEG = −2 × POS,
+    /// writes both `ISV[REWARD_POS_CAP_ADAPTIVE_INDEX=452]` and
+    /// `ISV[REWARD_NEG_CAP_ADAPTIVE_INDEX=453]`. Block-tree-reduce
+    /// (no atomicAdd); Pearl-A first-observation bootstrap; α=0.01 slow
+    /// EMA. Per-epoch boundary launch — reward distribution is the
+    /// foundation of training and shouldn't move fast.
+    reward_cap_update_ops: super::gpu_aux_trunk::RewardCapUpdateOps,
+
     // ── Speculative inference cache ──
     /// Cached trunk output [B, SH2] from between-bar speculative forward.
     speculative_h_s2: CudaSlice<f32>,
@@ -14558,6 +14581,68 @@ impl GpuDqnTrainer {
         Ok(())
     }
 
+    /// Class A P0-A (2026-05-08): launch the adaptive reward-cap producer
+    /// at per-epoch boundary.
+    ///
+    /// Sweeps the per-epoch `step_ret_per_sample` + `trade_close_per_sample`
+    /// buffers (populated by `unified_env_step_core` in
+    /// `experience_kernels.cu`), computes a Welford `mean + Z_99 × sigma`
+    /// p99 estimator over winning realized returns plus a
+    /// `max(winning_returns)` conservative takeover, applies a 1.5×
+    /// safety factor, clamps to dimensional bounds [1, 50], and writes
+    /// both `ISV[REWARD_POS_CAP_ADAPTIVE_INDEX=452]` (POS cap) and
+    /// `ISV[REWARD_NEG_CAP_ADAPTIVE_INDEX=453]` (NEG = −2 × POS,
+    /// preserving Kahneman/Tversky 2:1 loss-aversion asymmetry per
+    /// `pearl_audit_unboundedness_for_implicit_asymmetry`).
+    ///
+    /// Per-epoch boundary launch — reward distribution is the foundation
+    /// of training and shouldn't move fast; per-step would track sample
+    /// noise.
+    ///
+    /// Pearls applied:
+    ///   - `pearl_first_observation_bootstrap.md` (sentinel = 5.0/-10.0)
+    ///   - `pearl_no_host_branches_in_captured_graph.md`
+    ///   - `feedback_no_atomicadd.md` (block-tree-reduce in shmem only)
+    ///   - `pearl_symmetric_clamp_audit.md` (bilateral clamp on POS, NEG)
+    ///   - `feedback_isv_for_adaptive_bounds.md` (POS bounds [1, 50] are
+    ///     dimensional safety floors, not tuning)
+    pub(crate) fn launch_reward_cap_update(
+        &self,
+        step_ret_dev_ptr: u64,
+        trade_close_dev_ptr: u64,
+        total_samples: i32,
+    ) -> Result<(), MLError> {
+        use crate::cuda_pipeline::sp14_isv_slots::{
+            REWARD_CAP_EMA_ALPHA, REWARD_CAP_SAFETY_FACTOR, REWARD_NEG_CAP_ADAPTIVE_INDEX,
+            REWARD_NEG_TO_POS_RATIO, REWARD_POS_CAP_ADAPTIVE_INDEX, REWARD_POS_CAP_MAX,
+            REWARD_POS_CAP_MIN, SENTINEL_REWARD_NEG_CAP, SENTINEL_REWARD_POS_CAP,
+        };
+
+        debug_assert!(
+            self.isv_signals_dev_ptr != 0,
+            "launch_reward_cap_update: isv_signals_dev_ptr must be allocated"
+        );
+
+        self.reward_cap_update_ops.launch(
+            &self.stream,
+            step_ret_dev_ptr,
+            trade_close_dev_ptr,
+            total_samples,
+            self.isv_signals_dev_ptr,
+            REWARD_POS_CAP_ADAPTIVE_INDEX as i32,
+            REWARD_NEG_CAP_ADAPTIVE_INDEX as i32,
+            SENTINEL_REWARD_POS_CAP,
+            SENTINEL_REWARD_NEG_CAP,
+            REWARD_CAP_SAFETY_FACTOR,
+            REWARD_NEG_TO_POS_RATIO,
+            REWARD_POS_CAP_MIN,
+            REWARD_POS_CAP_MAX,
+            REWARD_CAP_EMA_ALPHA,
+        )?;
+
+        Ok(())
+    }
+
     /// SP8 (Fix 36, 2026-05-03): launch GPU-only `train_active_frac` canary
     /// producer.
     ///
@@ -22964,6 +23049,14 @@ impl GpuDqnTrainer {
         let avg_win_hold_time_update_ops =
             super::gpu_aux_trunk::AvgWinHoldTimeUpdateOps::new(&stream)?;
 
+        // Class A P0-A (2026-05-08): adaptive reward-cap producer. Pre-loads
+        // the `CudaFunction` handle at construction per
+        // `pearl_no_host_branches_in_captured_graph.md`. Per-epoch boundary
+        // launch (slow-moving — reward distribution is the foundation of
+        // training and shouldn't move fast).
+        let reward_cap_update_ops =
+            super::gpu_aux_trunk::RewardCapUpdateOps::new(&stream)?;
+
         // ── Speculative inference cache ──
         let speculative_h_s2 = stream.alloc_zeros::<f32>(b * sh2)
             .map_err(|e| MLError::ModelError(format!("speculative_h_s2 alloc: {e}")))?;
@@ -23795,6 +23888,8 @@ impl GpuDqnTrainer {
             // SP14 Layer C Phase C.4b (2026-05-08): aux horizon + winning-hold-time producers.
             aux_horizon_update_ops,
             avg_win_hold_time_update_ops,
+            // Class A P0-A (2026-05-08): adaptive reward-cap producer.
+            reward_cap_update_ops,
             speculative_h_s2,
             speculative_features,
             speculative_valid: false,
diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
index b9ca7a8e1..381177958 100644
--- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
@@ -6081,6 +6081,29 @@ impl GpuExperienceCollector {
         ptr
     }
 
+    /// Class A P0-A (2026-05-08): raw device pointer to
+    /// `step_ret_per_sample` `[alloc_episodes * alloc_timesteps]` —
+    /// `step_ret_core` raw signed per-step return (populated by
+    /// `unified_env_step_core` in `experience_kernels.cu`, line ~2779).
+    /// Consumed by `reward_cap_update_kernel` at the per-epoch boundary
+    /// to compute the p99 of winning realized returns for the adaptive
+    /// REWARD_POS/NEG_CAP ISV slots.
+    pub fn step_ret_per_sample_dev_ptr(&self) -> u64 {
+        let (ptr, _guard) = self.step_ret_per_sample.device_ptr(&self.stream);
+        ptr
+    }
+
+    /// Class A P0-A (2026-05-08): raw device pointer to
+    /// `trade_close_per_sample` `[alloc_episodes * alloc_timesteps]` —
+    /// 1 iff (exiting_trade || reversing_trade) (populated by
+    /// `unified_env_step_core` in `experience_kernels.cu`, line ~2780).
+    /// Consumed by `reward_cap_update_kernel` to mask the realized-return
+    /// distribution to actual trade-close events.
+    pub fn trade_close_per_sample_dev_ptr(&self) -> u64 {
+        let (ptr, _guard) = self.trade_close_per_sample.device_ptr(&self.stream);
+        ptr
+    }
+
     /// Last experience count from the most recent `collect_experiences_gpu` call.
     pub fn last_experience_count(&self) -> usize { self.last_experience_count }
 
diff --git a/crates/ml/src/cuda_pipeline/reward_cap_update_kernel.cu b/crates/ml/src/cuda_pipeline/reward_cap_update_kernel.cu
new file mode 100644
index 000000000..5108ff0a8
--- /dev/null
+++ b/crates/ml/src/cuda_pipeline/reward_cap_update_kernel.cu
@@ -0,0 +1,264 @@
+/* ══════════════════════════════════════════════════════════════════════════
+ * Class A P0-A — adaptive REWARD_POS_CAP / REWARD_NEG_CAP producer (2026-05-08).
+ *
+ * Replaces the hardcoded `REWARD_POS_CAP=+5.0f` / `REWARD_NEG_CAP=-10.0f`
+ * (state_layout.cuh:266-267) with ISV-driven adaptive bounds derived from
+ * the realized winning-trade return distribution.
+ *
+ * Why this matters (per Class A audit, the highest-suspected-impact fix
+ * for the months-long WR-stuck-at-46-48% plateau across 11 superprojects):
+ *
+ *     EV(trade) = WR × min(realized_win, +5)
+ *               + (1−WR) × max(realized_loss, −10)
+ *
+ * With WR ≈ 0.46, the cap asymmetry is fully active. Wins clip to +5
+ * regardless of magnitude. Controller signals (sharpe EMA, var_q,
+ * q_gap) all derive from this CAPPED buffer, so the controller cannot
+ * select for trades it cannot see. Selectivity gradient evaporates —
+ * small wins clip to +5 alongside large wins also clipping to +5.
+ *
+ * The state_layout.cuh comment block (lines 261-264) explicitly
+ * deferred Phase 2 (ISV-driven) "IF Phase 1 validation reveals
+ * adaptive need". Phase 1 has been running across 11 SPs without
+ * budging WR — adaptive need is revealed.
+ *
+ * Algorithm:
+ *
+ *   Phase 1 (block-tree-reduce): each block of `BLK_DIM` threads sweeps
+ *   its tile of N=B*T samples, masking by `trade_close_per_sample[i]`
+ *   and partitioning by `step_ret_per_sample[i] > 0` (winning) vs
+ *   `< 0` (losing). Local block reduce in shared memory accumulates a
+ *   per-block sum-of-squares over winning returns AND a count, plus a
+ *   max-of-winning-returns. Single-block kernel — the per-epoch sample
+ *   buffer is bounded (≤ ~64k slots).
+ *
+ *   Phase 2 (Pearl-A + EMA, single thread): if no winning trades
+ *   observed, keep ISV slots unchanged (zero contributors → no signal).
+ *   The "p99 estimator" without an actual histogram is the
+ *   one-step-of-Welford rule of thumb:
+ *
+ *       p99(winning_returns) ≈ mean + Z_99 × sigma   (Z_99 ≈ 2.326)
+ *
+ *   for a roughly-normal positive tail, OR `max(winning_returns)` for
+ *   small-N regimes, whichever is larger. We do NOT histogram (no
+ *   atomicAdd; would also need full sample-buffer materialisation in
+ *   shmem). The mean+sigma estimator with the conservative max
+ *   takeover handles both heavy-tail outliers (max wins) and the
+ *   well-behaved bulk (Gaussian-ish wins). Per
+ *   `feedback_no_atomicadd.md`.
+ *
+ *   POS_target = max(p99_estimate, max_win) × safety_factor (1.5×)
+ *
+ *   NEG_target = − REWARD_NEG_TO_POS_RATIO × POS_target  (= −2 × POS,
+ *   preserving Kahneman/Tversky loss-aversion asymmetry per
+ *   `pearl_audit_unboundedness_for_implicit_asymmetry`). Asymmetry
+ *   stays — moved from hardcoded scalar to producer-time multiplier.
+ *   SINGLE SOURCE OF TRUTH: every consumer reads its corresponding
+ *   ISV slot, no consumer applies the 2× ratio itself.
+ *
+ *   Pearl-A first-observation bootstrap: if `current_pos == sentinel`
+ *   (5.0 within 1e-6f), REPLACE with target directly. Otherwise EMA
+ *   blend with α=REWARD_CAP_EMA_ALPHA (0.01, slow — reward
+ *   distribution is the foundation of training and shouldn't move
+ *   fast). Per `pearl_first_observation_bootstrap.md`.
+ *
+ *   Bounds: POS clamped to [REWARD_POS_CAP_MIN=1.0, REWARD_POS_CAP_MAX=50.0]
+ *   (Category-1 dimensional safety per `feedback_isv_for_adaptive_bounds`).
+ *   NEG = −2 × POS by construction; resulting NEG ∈ [−100, −2].
+ *
+ * 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 winning
+ *     trades" and Pearl-A sentinel detection.
+ *
+ *   - `pearl_first_observation_bootstrap.md` — first valid sample
+ *     replaces sentinel directly; no blend.
+ *
+ *   - `feedback_no_stubs.md` — full body, no return-zero placeholders.
+ *
+ *   - `feedback_isv_for_adaptive_bounds.md` — POS/NEG ranges
+ *     [1, 50] / [−100, −2] are dimensional safety bounds, NOT tuning.
+ *
+ *   - `pearl_symmetric_clamp_audit.md` — bilateral `fmaxf(lo, fminf(x, hi))`
+ *     clamp on POS_target before writing the ISV slot. NEG ratio
+ *     applied AFTER clamp on POS, then NEG bilaterally clamped to its
+ *     own [−100, −2] range as a defensive layer.
+ *
+ * Args:
+ *   step_ret_per_sample          — `[total_samples]` device f32; raw
+ *                                  signed per-step return populated by
+ *                                  `unified_env_step_core` in
+ *                                  `experience_kernels.cu` (line ~2779).
+ *   trade_close_per_sample       — `[total_samples]` device i32; 1 iff
+ *                                  exiting_trade || reversing_trade
+ *                                  (`experience_kernels.cu` line ~2780).
+ *   total_samples                — N*L (B*T total samples in the epoch).
+ *   isv                          — `[ISV_TOTAL_DIM]` device f32, modified
+ *                                  in place at indices `pos_idx` and
+ *                                  `neg_idx`.
+ *   pos_idx                      — REWARD_POS_CAP_ADAPTIVE_INDEX (452).
+ *   neg_idx                      — REWARD_NEG_CAP_ADAPTIVE_INDEX (453).
+ *   sentinel_pos                 — SENTINEL_REWARD_POS_CAP (5.0).
+ *   sentinel_neg                 — SENTINEL_REWARD_NEG_CAP (-10.0).
+ *   safety_factor                — REWARD_CAP_SAFETY_FACTOR (1.5).
+ *   neg_to_pos_ratio             — REWARD_NEG_TO_POS_RATIO (2.0).
+ *   pos_min                      — REWARD_POS_CAP_MIN (1.0).
+ *   pos_max                      — REWARD_POS_CAP_MAX (50.0).
+ *   alpha                        — REWARD_CAP_EMA_ALPHA (0.01, slow).
+ *
+ * Launch: grid=(1, 1, 1), block=(BLK_DIM=256, 1, 1).
+ * Shared memory: 4 * BLK_DIM floats:
+ *   sh_sum_pos    — sum of positive (winning) returns
+ *   sh_sumsq_pos  — sum-of-squares of positive returns
+ *   sh_count_pos  — count of positive returns
+ *   sh_max_pos    — max of positive returns
+ * The count is summed as float for unified shmem layout — values are
+ * bounded by total_samples ≤ 2^16 so f32 representation is exact for
+ * counts up to 2^24.
+ * ══════════════════════════════════════════════════════════════════════════ */
+
+#include <cuda_runtime.h>
+
+#define BLK_DIM 256
+#define EPS_F   1e-6f
+/* Z-score for ~p99 of a normal distribution (one-sided). Used as a
+ * lightweight estimator: p99 ≈ mean + Z_99 × sigma. */
+#define Z_99    2.326f
+
+extern "C" __global__
+void reward_cap_update(
+    const float* __restrict__ step_ret_per_sample,
+    const int*   __restrict__ trade_close_per_sample,
+    int          total_samples,
+    float* __restrict__ isv,
+    int          pos_idx,
+    int          neg_idx,
+    float        sentinel_pos,
+    float        sentinel_neg,
+    float        safety_factor,
+    float        neg_to_pos_ratio,
+    float        pos_min,
+    float        pos_max,
+    float        alpha)
+{
+    /* Single-block kernel — only block 0 does anything. */
+    if (blockIdx.x != 0) return;
+
+    extern __shared__ float shmem[];
+    float* sh_sum_pos    = shmem;                       /* [BLK_DIM] */
+    float* sh_sumsq_pos  = shmem + BLK_DIM;             /* [BLK_DIM] */
+    float* sh_count_pos  = shmem + 2 * BLK_DIM;         /* [BLK_DIM] */
+    float* sh_max_pos    = shmem + 3 * BLK_DIM;         /* [BLK_DIM] */
+
+    int tid = threadIdx.x;
+
+    /* Phase 1a: stride-based sweep — accumulate stats on closed-trade
+     * winning returns only. Losing closes are NOT stored (we don't
+     * adapt the negative cap from realized losses; the 2:1 ratio is
+     * spec'd-asymmetric, not data-driven). */
+    float local_sum    = 0.0f;
+    float local_sumsq  = 0.0f;
+    float local_count  = 0.0f;
+    float local_max    = 0.0f;
+    for (int i = tid; i < total_samples; i += BLK_DIM) {
+        int   is_close = trade_close_per_sample[i];     /* 0 or 1 */
+        float ret      = step_ret_per_sample[i];        /* signed */
+        if (is_close != 0 && ret > 0.0f) {
+            local_sum    += ret;
+            local_sumsq  += ret * ret;
+            local_count  += 1.0f;
+            if (ret > local_max) local_max = ret;
+        }
+    }
+    sh_sum_pos[tid]    = local_sum;
+    sh_sumsq_pos[tid]  = local_sumsq;
+    sh_count_pos[tid]  = local_count;
+    sh_max_pos[tid]    = local_max;
+    __syncthreads();
+
+    /* Phase 1b: block-tree-reduce (no atomicAdd) — halve every iter.
+     * Sum/sumsq/count fold via addition; max folds via fmaxf. */
+    for (int s = BLK_DIM >> 1; s > 0; s >>= 1) {
+        if (tid < s) {
+            sh_sum_pos[tid]    += sh_sum_pos[tid + s];
+            sh_sumsq_pos[tid]  += sh_sumsq_pos[tid + s];
+            sh_count_pos[tid]  += sh_count_pos[tid + s];
+            sh_max_pos[tid]     = fmaxf(sh_max_pos[tid], sh_max_pos[tid + s]);
+        }
+        __syncthreads();
+    }
+
+    /* Phase 2: thread 0 finalises EMA + Pearl-A + clamp + writes BOTH slots. */
+    if (tid != 0) return;
+
+    float total_sum    = sh_sum_pos[0];
+    float total_sumsq  = sh_sumsq_pos[0];
+    float total_count  = sh_count_pos[0];
+    float total_max    = sh_max_pos[0];
+
+    /* Guard: no winning trades observed this epoch — keep ISV slots
+     * unchanged (cold-start sentinel persists; consumers fall back to
+     * `state_layout.cuh` constants per the cold-start behavior). */
+    if (total_count <= 0.0f) {
+        return;
+    }
+
+    /* Welford-rule p99 estimator: mean + Z_99 × sigma. Sigma is
+     * computed as `sqrt(max(0, sumsq/n − mean²))` — the inner max
+     * absorbs floating-point under-flow when all observations are
+     * (near-)identical. */
+    float mean         = total_sum / total_count;
+    float variance     = fmaxf(0.0f, total_sumsq / total_count - mean * mean);
+    float sigma        = sqrtf(variance);
+    float p99_estimate = mean + Z_99 * sigma;
+
+    /* Conservative takeover by realized max: in small-N or heavy-tail
+     * regimes the parametric estimate underestimates the true tail.
+     * Using `max(p99_estimate, total_max)` ensures the cap does NOT
+     * actively clip any observed winning return. */
+    float pos_target = fmaxf(p99_estimate, total_max) * safety_factor;
+
+    /* Clamp POS to dimensional-safety bounds [pos_min, pos_max].
+     * Bilateral per `pearl_symmetric_clamp_audit.md`. */
+    pos_target = fmaxf(pos_min, fminf(pos_target, pos_max));
+
+    /* Pearl-A first-observation bootstrap on POS slot. The sentinel is
+     * exactly SENTINEL_REWARD_POS_CAP=5.0; any value within EPS_F of
+     * the sentinel is treated as sentinel. (Direct equality on f32
+     * literals would also work, but `|x − sentinel| < eps` is the
+     * established pattern for ISV slot bootstrap detection.) */
+    float current_pos  = isv[pos_idx];
+    float blended_pos;
+    if (fabsf(current_pos - sentinel_pos) < EPS_F) {
+        blended_pos = pos_target;
+    } else {
+        blended_pos = (1.0f - alpha) * current_pos + alpha * pos_target;
+    }
+    /* Re-clamp post-blend (defensive — handles malformed prior state
+     * outside [pos_min, pos_max]). */
+    blended_pos = fmaxf(pos_min, fminf(blended_pos, pos_max));
+
+    /* NEG cap = −neg_to_pos_ratio × POS cap. Asymmetry preserved at
+     * producer-time, single source of truth. NEG bounds derived from
+     * POS bounds × ratio: NEG ∈ [−ratio × pos_max, −ratio × pos_min] =
+     * [−100, −2] for the production constants. The bilateral clamp
+     * here is a defensive layer (POS clamp above already guarantees
+     * NEG sits within its derived range). */
+    float blended_neg  = -neg_to_pos_ratio * blended_pos;
+    float neg_lo       = -neg_to_pos_ratio * pos_max;   /* −100 */
+    float neg_hi       = -neg_to_pos_ratio * pos_min;   /* −2 */
+    blended_neg        = fmaxf(neg_lo, fminf(blended_neg, neg_hi));
+
+    /* Suppress unused-parameter warning for sentinel_neg — kept in the
+     * signature for symmetry with the POS sentinel and to allow a
+     * future Pearl-A variant on the NEG slot if the 2:1 ratio
+     * spec ever evolves. */
+    (void)sentinel_neg;
+
+    isv[pos_idx] = blended_pos;
+    isv[neg_idx] = blended_neg;
+}
diff --git a/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs
index 5b4c433fa..d0d6e79c1 100644
--- a/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs
+++ b/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs
@@ -105,6 +105,65 @@ pub const AUX_PRED_HORIZON_BARS_MAX:         f32 = 240.0; // ceiling (rollout bu
 pub const SP14_C4B_SLOT_BASE: usize = 450;
 pub const SP14_C4B_SLOT_END:  usize = 452;
 
+// ── Class A P0-A (2026-05-08): adaptive REWARD_POS/NEG_CAP slots ─────────
+// Replaces hardcoded `REWARD_POS_CAP=+5.0f` / `REWARD_NEG_CAP=-10.0f`
+// (state_layout.cuh:266-267) with ISV-driven values fed by the
+// `reward_cap_update_kernel` producer at per-epoch boundary. Per Class A
+// audit, these caps were structurally clipping the upper tail of realized
+// alpha — the controller signals (sharpe EMA, var_q, q_gap) all derive
+// from the CAPPED buffer, so the controller cannot select for trades it
+// cannot see. Selectivity gradient evaporates.
+//
+// The state_layout comment lines 261-264 explicitly deferred Phase 2
+// (ISV-driven) "IF Phase 1 validation reveals adaptive need". 11
+// superprojects of WR-stuck-at-46-48% revealed adaptive need.
+//
+// Asymmetry rationale (Kahneman/Tversky 2:1 loss aversion, per
+// `pearl_audit_unboundedness_for_implicit_asymmetry`) IS preserved. The
+// 2:1 asymmetry comes from the producer maintaining the
+// `NEG_CAP = -2 * POS_CAP` ratio internally (single source of truth in the
+// kernel), NOT from a fixed scalar. The scalar values themselves adapt to
+// the realized winning-return distribution: POS=p99(winning_returns) ×
+// safety_factor=1.5, with bounds [1.0, 50.0] clamping POS and [-100.0,
+// -2.0] clamping NEG (Category-1 dimensional safety floors per
+// `feedback_isv_for_adaptive_bounds`, NOT tuning).
+//
+// Slot layout: [452..454) — appended after SP14 Layer C Phase C.4b's
+// AVG_WIN_HOLD_TIME_BARS_INDEX=451 to preserve the checkpoint layout
+// fingerprint compatibility (no slot insertion before existing slots).
+pub const REWARD_POS_CAP_ADAPTIVE_INDEX:  usize = 452;  // adaptive +cap (replaces hardcoded +5.0)
+pub const REWARD_NEG_CAP_ADAPTIVE_INDEX:  usize = 453;  // adaptive −cap (= −2 × POS, replaces hardcoded −10.0)
+
+// Sentinels (initial values; reset at fold boundary). Match the
+// pre-Class-A-P0-A hardcoded constants so cold-start behavior is
+// bit-identical to the prior Phase 1 anchor before the first valid
+// observation lands.
+pub const SENTINEL_REWARD_POS_CAP:        f32 =   5.0;   // matches pre-Class-A-P0-A REWARD_POS_CAP
+pub const SENTINEL_REWARD_NEG_CAP:        f32 = -10.0;   // matches pre-Class-A-P0-A REWARD_NEG_CAP
+
+// Producer kernel constants (Category-1 dimensional safety bounds).
+// POS bounds: [1.0, 50.0] — cap below 1 makes the entire reward signal
+// flat (selectivity collapses); cap above 50 lets a single adversarial
+// outlier dominate any per-batch EMA. NEG = -2 × POS by construction;
+// resulting NEG range is therefore [-100.0, -2.0].
+pub const REWARD_POS_CAP_MIN:             f32 =   1.0;
+pub const REWARD_POS_CAP_MAX:             f32 =  50.0;
+// 2:1 ratio (Kahneman/Tversky meta-analysis ~2.0-2.25) preserved at
+// producer time inside `reward_cap_update_kernel`. Single source of
+// truth — every consumer reads its cap directly from the corresponding
+// ISV slot, no consumer applies the multiplier itself.
+pub const REWARD_NEG_TO_POS_RATIO:        f32 =   2.0;
+// Safety factor on the p99 win-return target — keeps headroom above the
+// realized 99th percentile so we don't actively clip the upper tail.
+pub const REWARD_CAP_SAFETY_FACTOR:       f32 =   1.5;
+// EMA blend rate. Slow — reward distribution is the foundation of
+// training and shouldn't move fast. Pearl-A bootstrap replaces sentinel
+// directly on first valid observation; α=0.01 thereafter.
+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;
+
 #[cfg(test)]
 mod tests {
     use super::*;
@@ -181,4 +240,37 @@ mod tests {
             SP14_C4B_SLOT_END, ISV_TOTAL_DIM,
         );
     }
+
+    /// Lock Class A P0-A (2026-05-08) adaptive reward-cap slot layout.
+    /// Two contiguous slots [452..454) — REWARD_POS_CAP_ADAPTIVE +
+    /// REWARD_NEG_CAP_ADAPTIVE. The pre-existing hardcoded constants
+    /// `REWARD_POS_CAP=+5.0f` / `REWARD_NEG_CAP=-10.0f` (state_layout.cuh)
+    /// remain as cold-start fallback when the ISV slot is at sentinel.
+    #[test]
+    fn sp14_p0a_reward_cap_slot_layout_locked() {
+        assert_eq!(SP14_P0A_SLOT_BASE, 452);
+        assert_eq!(SP14_P0A_SLOT_END,  454);
+        assert_eq!(REWARD_POS_CAP_ADAPTIVE_INDEX, 452);
+        assert_eq!(REWARD_NEG_CAP_ADAPTIVE_INDEX, 453);
+        // 2:1 asymmetry preserved at producer-time, not consumer-time.
+        assert_eq!(REWARD_NEG_TO_POS_RATIO, 2.0);
+        // Sentinels match pre-P0-A hardcoded constants for bit-identical
+        // cold-start behavior before first valid observation.
+        assert_eq!(SENTINEL_REWARD_POS_CAP,  5.0);
+        assert_eq!(SENTINEL_REWARD_NEG_CAP, -10.0);
+        // Category-1 dimensional safety: POS in [1, 50], NEG in [-100, -2].
+        assert_eq!(REWARD_POS_CAP_MIN, 1.0);
+        assert_eq!(REWARD_POS_CAP_MAX, 50.0);
+    }
+
+    #[test]
+    fn all_sp14_p0a_slots_fit_within_isv_total_dim() {
+        use crate::cuda_pipeline::gpu_dqn_trainer::ISV_TOTAL_DIM;
+        assert!(
+            SP14_P0A_SLOT_END <= ISV_TOTAL_DIM,
+            "SP14_P0A_SLOT_END={} exceeds ISV_TOTAL_DIM={} — bus too small for Class A P0-A reward-cap slots; \
+             bump ISV_TOTAL_DIM in gpu_dqn_trainer.rs (and update layout_fingerprint_seed()).",
+            SP14_P0A_SLOT_END, ISV_TOTAL_DIM,
+        );
+    }
 }
diff --git a/crates/ml/src/cuda_pipeline/sp15_reward_axis_helpers.cuh b/crates/ml/src/cuda_pipeline/sp15_reward_axis_helpers.cuh
index c759bfd6a..363f7cd6e 100644
--- a/crates/ml/src/cuda_pipeline/sp15_reward_axis_helpers.cuh
+++ b/crates/ml/src/cuda_pipeline/sp15_reward_axis_helpers.cuh
@@ -180,17 +180,51 @@ __device__ inline float sp15_dd_penalty(
 // --------------------------------------------------------------------
 // `sp15_apply_sp12_cap` — Stage 4: asymmetric bilateral clamp
 //
-// `r_capped = fmaxf(REWARD_NEG_CAP, fminf(REWARD_POS_CAP, r))`
+// `r_capped = fmaxf(neg_cap_eff, fminf(pos_cap_eff, r))`
 //
-// Per Q1 resolution: REWARD_NEG_CAP / REWARD_POS_CAP stay as
-// `state_layout.cuh` macros (NOT lifted to ISV) — spec'd constants per
-// the SP12 v3 design, not adaptive bounds. Asymmetric NEG=−10 / POS=+5
-// is the loss-aversion encoding per
-// `pearl_audit_unboundedness_for_implicit_asymmetry` and matches the
-// existing `compute_asymmetric_capped_pnl` helper in trade_physics.cuh.
+// Class A P0-A (2026-05-08): caps lifted from `state_layout.cuh`
+// macros to ISV-driven adaptive slots (the Q1 resolution comment
+// below was the Phase 1 design — Phase 2 is now active per the
+// audit-revealed adaptive need). The helper reads
+// `ISV[REWARD_POS_CAP_ADAPTIVE_INDEX=452]` /
+// `ISV[REWARD_NEG_CAP_ADAPTIVE_INDEX=453]` first and falls back to
+// the original macros when the POS slot is at sentinel
+// SENTINEL_REWARD_POS_CAP=5.0 (within EPS) — bit-identical pre-P0-A
+// cold-start behavior until the first valid observation lands.
+//
+// Asymmetric NEG = −2 × POS is the loss-aversion encoding per
+// `pearl_audit_unboundedness_for_implicit_asymmetry`. The 2:1 ratio
+// is preserved at producer-time inside `reward_cap_update_kernel`
+// (single source of truth) — the consumer here applies whatever NEG
+// value sits in slot 453 without re-multiplying.
+//
+// Defensive bounds guard: any POS value outside
+// [REWARD_POS_CAP_MIN_BOUND=1.0, REWARD_POS_CAP_MAX_BOUND=50.0] also
+// falls back to the macro, defending against malformed prior state.
+// Producer kernel clamps to the same range, so a healthy ISV slot
+// always passes.
+//
+// Per `pearl_symmetric_clamp_audit.md`: bilateral
+// `fmaxf(NEG, fminf(POS, x))` even though the bounds are
+// intentionally asymmetric (NEG ≈ −2 × POS).
 // --------------------------------------------------------------------
-__device__ inline float sp15_apply_sp12_cap(float r_in) {
-    return fmaxf(REWARD_NEG_CAP, fminf(REWARD_POS_CAP, r_in));
+__device__ inline float sp15_apply_sp12_cap(
+    float r_in,
+    const float* __restrict__ isv
+) {
+    float pos_cap_eff = REWARD_POS_CAP;
+    float neg_cap_eff = REWARD_NEG_CAP;
+    if (isv != nullptr) {
+        const float isv_pos = isv[REWARD_POS_CAP_ADAPTIVE_INDEX];
+        const float isv_neg = isv[REWARD_NEG_CAP_ADAPTIVE_INDEX];
+        if (fabsf(isv_pos - SENTINEL_REWARD_POS_CAP) >= 1e-6f
+            && isv_pos >= REWARD_POS_CAP_MIN_BOUND
+            && isv_pos <= REWARD_POS_CAP_MAX_BOUND) {
+            pos_cap_eff = isv_pos;
+            neg_cap_eff = isv_neg;
+        }
+    }
+    return fmaxf(neg_cap_eff, fminf(pos_cap_eff, r_in));
 }
 
 #endif  // SP15_REWARD_AXIS_HELPERS_CUH_
diff --git a/crates/ml/src/cuda_pipeline/state_layout.cuh b/crates/ml/src/cuda_pipeline/state_layout.cuh
index f19c7ad2a..b2932ed07 100644
--- a/crates/ml/src/cuda_pipeline/state_layout.cuh
+++ b/crates/ml/src/cuda_pipeline/state_layout.cuh
@@ -271,6 +271,34 @@
 #define MIN_HOLD_TEMPERATURE_END    5.0f
 #define MIN_HOLD_TEMPERATURE_DECAY 20.0f
 
+// Class A P0-A (2026-05-08) — adaptive REWARD_POS/NEG_CAP ISV slots.
+// REWARD_POS_CAP / REWARD_NEG_CAP above remain as Phase-1 cold-start
+// fallback values; consumers read the ISV slots first and fall back to
+// the macros when the slot is at sentinel (bit-identical pre-P0-A
+// behavior until the first valid observation). Slot indices mirror
+// `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` (locked by
+// `sp14_p0a_reward_cap_slot_layout_locked` test). Producer kernel:
+// `reward_cap_update_kernel.cu`; sentinels match the macros above for
+// bit-identical bootstrap. Per Class A audit, the highest-suspected-
+// impact fix for the months-long WR-stuck-at-46-48% plateau across
+// 11 superprojects — controller signals (sharpe EMA, var_q, q_gap)
+// derive from the CAPPED reward buffer, so the controller cannot
+// select for trades it cannot see. Per
+// `pearl_audit_unboundedness_for_implicit_asymmetry`, the 2:1
+// asymmetry stays — moved from hardcoded scalar to producer-time
+// multiplier (single source of truth).
+#define REWARD_POS_CAP_ADAPTIVE_INDEX 452
+#define REWARD_NEG_CAP_ADAPTIVE_INDEX 453
+#define SENTINEL_REWARD_POS_CAP        5.0f
+#define SENTINEL_REWARD_NEG_CAP      -10.0f
+// Defensive consumer-side bounds (mirror Rust constants
+// REWARD_POS_CAP_MIN/MAX in sp14_isv_slots.rs). Values outside this
+// window fall back to the macro on the consumer side as a defense
+// against malformed prior state. Producer kernel clamps to the same
+// range, so a healthy ISV slot will always pass the guard.
+#define REWARD_POS_CAP_MIN_BOUND       1.0f
+#define REWARD_POS_CAP_MAX_BOUND      50.0f
+
 // ── Compile-time checks ──
 static_assert(SL_PADDING_START + SL_PADDING_DIM == SL_STATE_DIM,
     "State layout dimensions must sum to SL_STATE_DIM");
diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs
index bb1425257..a1f93e933 100644
--- a/crates/ml/src/trainers/dqn/state_reset_registry.rs
+++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs
@@ -1021,6 +1021,31 @@ impl StateResetRegistry {
                 category: ResetCategory::FoldReset,
                 description: "ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451] — SP14 Layer C Phase C.4b avg winning hold time EMA (bars). Aggregated from `hold_at_exit_per_sample` + `trade_profitable_per_sample` per-sample buffers (populated by `unified_env_step_core` in `experience_kernels.cu`) by `avg_win_hold_time_update_kernel` at per-epoch boundary. FoldReset sentinel SENTINEL_AVG_WIN_HOLD_TIME_BARS=0.0 — Pearl-A first-observation bootstrap. Drives the horizon producer for ISV[450]. Plan: §C.4b.",
             },
+            // ── Class A P0-A (2026-05-08): adaptive REWARD_POS/NEG_CAP ──
+            // Two slots [452, 453] replacing the hardcoded
+            // `REWARD_POS_CAP=+5.0f` / `REWARD_NEG_CAP=-10.0f`
+            // (state_layout.cuh:266-267) per Class A audit. Producer
+            // kernel `reward_cap_update_kernel` computes
+            // p99(winning_returns) × 1.5 → POS cap, NEG = −2 × POS,
+            // writes both at per-epoch boundary. Pearl-A bootstrap
+            // sentinels match pre-P0-A hardcoded constants for
+            // bit-identical cold-start behavior — consumers fall back
+            // to the macro when the ISV slot is at sentinel. Once the
+            // first valid observation lands the adaptive path takes
+            // over via Welford α=0.01 slow EMA (reward distribution
+            // is the foundation of training; shouldn't move fast).
+            // Per `feedback_isv_for_adaptive_bounds.md` and
+            // `pearl_audit_unboundedness_for_implicit_asymmetry.md`.
+            RegistryEntry {
+                name: "sp14_p0a_reward_pos_cap_adaptive",
+                category: ResetCategory::FoldReset,
+                description: "ISV[REWARD_POS_CAP_ADAPTIVE_INDEX=452] — Class A P0-A adaptive positive reward cap (replaces hardcoded REWARD_POS_CAP=+5.0f from state_layout.cuh:266). Produced by `reward_cap_update_kernel` from p99(winning_returns) × safety_factor=1.5 over the per-epoch `step_ret_per_sample` + `trade_close_per_sample` buffers. FoldReset sentinel SENTINEL_REWARD_POS_CAP=5.0 — Pearl-A first-observation bootstrap (matches pre-P0-A hardcoded value for bit-identical cold-start). Welford α=0.01 slow EMA thereafter. Bounds [1, 50] — Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`. Consumed by experience_kernels.cu (segment_complete cap), compute_sp15_final_reward_kernel.cu (Stage 4 cap), and sp15_reward_axis_helpers.cuh (sp15_apply_sp12_cap helper). Per Class A audit (highest-suspected-impact fix for the months-long WR-stuck-at-46-48% plateau across 11 superprojects).",
+            },
+            RegistryEntry {
+                name: "sp14_p0a_reward_neg_cap_adaptive",
+                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.",
+            },
             // ── 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
diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
index 00bdd6b66..9510e9613 100644
--- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs
+++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
@@ -411,6 +411,26 @@ impl DQNTrainer {
                 ) {
                     tracing::warn!(epoch, "launch_aux_horizon_chain failed (non-fatal): {e}");
                 }
+                // Class A P0-A (2026-05-08): adaptive REWARD_POS/NEG_CAP
+                // — per-epoch boundary launch. Sweeps the per-sample
+                // `step_ret_per_sample` + `trade_close_per_sample` buffers
+                // (populated by `unified_env_step_core` in
+                // `experience_kernels.cu`), computes p99(winning_returns)
+                // × 1.5 → POS cap, NEG = −2 × POS, writes both
+                // ISV[452, 453]. Pearl-A first-observation bootstrap +
+                // α=0.01 slow EMA. Hardcoded REWARD_POS_CAP=+5.0f /
+                // REWARD_NEG_CAP=-10.0f (state_layout.cuh:266-267) was
+                // structurally clipping the upper tail of realized alpha
+                // — months-long WR-stuck-at-46-48% plateau across 11 SPs.
+                let step_ret_ptr     = collector.step_ret_per_sample_dev_ptr();
+                let trade_close_ptr  = collector.trade_close_per_sample_dev_ptr();
+                if let Err(e) = fused.trainer().launch_reward_cap_update(
+                    step_ret_ptr,
+                    trade_close_ptr,
+                    total_samples,
+                ) {
+                    tracing::warn!(epoch, "launch_reward_cap_update failed (non-fatal): {e}");
+                }
             }
 
             // D.8 Plan 2 Task 6C: update TLOB regime focus EMA in ISV at epoch boundary.
@@ -8129,6 +8149,40 @@ impl DQNTrainer {
                     );
                 }
             }
+            // Class A P0-A (2026-05-08): adaptive REWARD_POS/NEG_CAP
+            // sentinels. Both slots use Pearl-A first-observation
+            // bootstrap — the FoldReset rewrites to sentinel so the
+            // kernel's "current_pos == sentinel" check fires cleanly on
+            // the new fold's first launch (avoids cross-fold EMA
+            // contamination). Sentinel values match pre-P0-A hardcoded
+            // REWARD_POS_CAP=+5.0 / REWARD_NEG_CAP=-10.0 for bit-identical
+            // cold-start until the first valid observation lands.
+            // Without this dispatch arm, the C.10 lesson recurs: the
+            // FoldReset registry entry exists, but no actual reset fires
+            // — slot drifts across folds and the layout-fingerprint
+            // smoke test eventually catches it as a runtime crash.
+            "sp14_p0a_reward_pos_cap_adaptive" => {
+                if let Some(ref fused) = self.fused_ctx {
+                    use crate::cuda_pipeline::sp14_isv_slots::{
+                        REWARD_POS_CAP_ADAPTIVE_INDEX, SENTINEL_REWARD_POS_CAP,
+                    };
+                    fused.trainer().write_isv_signal_at(
+                        REWARD_POS_CAP_ADAPTIVE_INDEX,
+                        SENTINEL_REWARD_POS_CAP,
+                    );
+                }
+            }
+            "sp14_p0a_reward_neg_cap_adaptive" => {
+                if let Some(ref fused) = self.fused_ctx {
+                    use crate::cuda_pipeline::sp14_isv_slots::{
+                        REWARD_NEG_CAP_ADAPTIVE_INDEX, SENTINEL_REWARD_NEG_CAP,
+                    };
+                    fused.trainer().write_isv_signal_at(
+                        REWARD_NEG_CAP_ADAPTIVE_INDEX,
+                        SENTINEL_REWARD_NEG_CAP,
+                    );
+                }
+            }
             // 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
diff --git a/crates/ml/tests/sp14_oracle_tests.rs b/crates/ml/tests/sp14_oracle_tests.rs
index ae399051d..bbb68b739 100644
--- a/crates/ml/tests/sp14_oracle_tests.rs
+++ b/crates/ml/tests/sp14_oracle_tests.rs
@@ -448,6 +448,345 @@ mod gpu {
     }
 }
 
+// ═══════════════════════════════════════════════════════════════════════════
+// Class A P0-A (2026-05-08) — adaptive REWARD_POS/NEG_CAP producer tests.
+//
+// Verifies the `reward_cap_update_kernel` producer:
+//   1. Cold-start Pearl-A bootstrap: ISV stays at sentinel until first valid
+//      observation, then REPLACES (no blend).
+//   2. 2:1 asymmetry preserved at producer-time: NEG_CAP = -2 × POS_CAP.
+//   3. Bounds enforced: POS in [1, 50], NEG in [-100, -2].
+//   4. Welford slow EMA after bootstrap (α=0.01).
+//
+// 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_p0a_reward_cap_gpu {
+    use std::sync::Arc;
+
+    use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg};
+    use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedI32Buffer};
+    use ml::cuda_pipeline::sp14_isv_slots::{
+        REWARD_CAP_EMA_ALPHA, REWARD_CAP_SAFETY_FACTOR, REWARD_NEG_CAP_ADAPTIVE_INDEX,
+        REWARD_NEG_TO_POS_RATIO, REWARD_POS_CAP_ADAPTIVE_INDEX, REWARD_POS_CAP_MAX,
+        REWARD_POS_CAP_MIN, SENTINEL_REWARD_NEG_CAP, SENTINEL_REWARD_POS_CAP,
+    };
+
+    const REWARD_CAP_UPDATE_CUBIN: &[u8] =
+        include_bytes!(concat!(env!("OUT_DIR"), "/reward_cap_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(REWARD_CAP_UPDATE_CUBIN.to_vec())
+            .expect("load reward_cap_update_kernel cubin");
+        module
+            .load_function("reward_cap_update")
+            .expect("load reward_cap_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_reward_cap(
+        stream: &Arc<CudaStream>,
+        kernel: &CudaFunction,
+        step_ret_ptr: u64,
+        trade_close_ptr: u64,
+        total_samples: i32,
+        isv_ptr: u64,
+        pos_idx: i32,
+        neg_idx: i32,
+        sentinel_pos: f32,
+        sentinel_neg: f32,
+        safety_factor: f32,
+        neg_to_pos_ratio: f32,
+        pos_min: f32,
+        pos_max: f32,
+        alpha: f32,
+    ) {
+        unsafe {
+            stream
+                .launch_builder(kernel)
+                .arg(&step_ret_ptr)
+                .arg(&trade_close_ptr)
+                .arg(&total_samples)
+                .arg(&isv_ptr)
+                .arg(&pos_idx)
+                .arg(&neg_idx)
+                .arg(&sentinel_pos)
+                .arg(&sentinel_neg)
+                .arg(&safety_factor)
+                .arg(&neg_to_pos_ratio)
+                .arg(&pos_min)
+                .arg(&pos_max)
+                .arg(&alpha)
+                .launch(LaunchConfig {
+                    grid_dim: (1, 1, 1),
+                    block_dim: (BLK_DIM, 1, 1),
+                    shared_mem_bytes: SMEM_BYTES,
+                })
+                .expect("launch reward_cap_update");
+        }
+        stream.synchronize().expect("sync after reward_cap_update");
+    }
+
+    /// Test 1 — Pearl-A first-observation bootstrap.
+    ///
+    /// 4 winning closes (all step_ret>0) with values [1.0, 2.0, 3.0, 4.0],
+    /// all trade_close=1. Mean=2.5, sigma≈1.118, p99≈2.5+2.326×1.118≈5.10,
+    /// max=4.0. Estimator picks `max(p99, max)=5.10`, then × 1.5 = 7.65.
+    /// Clamped to [1, 50] → 7.65 stays.
+    ///
+    /// Cold-start (current = sentinel 5.0): Pearl-A REPLACES → POS=7.65,
+    /// NEG=-2 × 7.65 = -15.3.
+    #[test]
+    #[ignore = "requires GPU"]
+    fn reward_cap_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-P0-A constants).
+        isv[REWARD_POS_CAP_ADAPTIVE_INDEX] = SENTINEL_REWARD_POS_CAP;
+        isv[REWARD_NEG_CAP_ADAPTIVE_INDEX] = SENTINEL_REWARD_NEG_CAP;
+
+        let step_ret: Vec<f32> = vec![1.0, 2.0, 3.0, 4.0];
+        let trade_close: Vec<i32> = vec![1, 1, 1, 1];
+
+        let n = step_ret.len() as i32;
+        let ret_buf = unsafe { MappedF32Buffer::new(step_ret.len()) }
+            .expect("alloc step_ret");
+        ret_buf.write_from_slice(&step_ret);
+        // i32 buffer via MappedF32Buffer — repurpose: the kernel reads
+        // `int*`, so we need a u32-aligned i32 slice. We use a separate
+        // raw allocation for the i32 close mask.
+        let close_buf = unsafe { MappedI32Buffer::new(trade_close.len()) }
+            .expect("alloc trade_close");
+        close_buf.write_from_slice(&trade_close);
+        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }
+            .expect("alloc isv");
+        isv_buf.write_from_slice(&isv);
+
+        launch_reward_cap(
+            &stream, &kernel,
+            ret_buf.dev_ptr, close_buf.dev_ptr, n, isv_buf.dev_ptr,
+            REWARD_POS_CAP_ADAPTIVE_INDEX as i32,
+            REWARD_NEG_CAP_ADAPTIVE_INDEX as i32,
+            SENTINEL_REWARD_POS_CAP, SENTINEL_REWARD_NEG_CAP,
+            REWARD_CAP_SAFETY_FACTOR, REWARD_NEG_TO_POS_RATIO,
+            REWARD_POS_CAP_MIN, REWARD_POS_CAP_MAX,
+            REWARD_CAP_EMA_ALPHA,
+        );
+
+        let result = isv_buf.read_all();
+        let pos = result[REWARD_POS_CAP_ADAPTIVE_INDEX];
+        let neg = result[REWARD_NEG_CAP_ADAPTIVE_INDEX];
+
+        // Pearl-A: sentinel → REPLACE with target. Target = max(p99, max) × 1.5.
+        // mean=2.5, var=(1+4+9+16)/4 - 6.25 = 7.5 - 6.25 = 1.25, sigma≈1.118
+        // p99 ≈ 2.5 + 2.326 × 1.118 ≈ 5.10. max=4.0. Picks 5.10 × 1.5 ≈ 7.65.
+        // Clamped to [1, 50] → stays at ~7.65.
+        assert!(
+            (pos - 7.65).abs() < 0.05,
+            "Pearl-A bootstrap: expected POS≈7.65, got {pos}"
+        );
+        // 2:1 asymmetry preserved at producer-time: NEG = -2 × POS.
+        assert!(
+            (neg - (-2.0 * pos)).abs() < 1e-4,
+            "NEG must equal -2 × POS at producer-time (Kahneman 2:1 single source of truth); got POS={pos}, NEG={neg}"
+        );
+    }
+
+    /// Test 2 — No winning trades → ISV slots preserved bit-exactly.
+    ///
+    /// All trade_close=0 (no closes). Kernel must skip the EMA update
+    /// (count==0 guard) and leave both slots at their seeded values.
+    #[test]
+    #[ignore = "requires GPU"]
+    fn reward_cap_no_winning_trades_preserves_isv() {
+        let stream = make_stream();
+        let kernel = load_kernel(&stream);
+
+        const ISV_DIM: usize = 1024;
+        const SEED_POS: f32 = 12.5;
+        const SEED_NEG: f32 = -25.0;
+        let mut isv = vec![0.0_f32; ISV_DIM];
+        isv[REWARD_POS_CAP_ADAPTIVE_INDEX] = SEED_POS;
+        isv[REWARD_NEG_CAP_ADAPTIVE_INDEX] = SEED_NEG;
+
+        // Step returns present but no closes — none should contribute.
+        let step_ret: Vec<f32> = vec![3.0, 5.0, -2.0];
+        let trade_close: Vec<i32> = vec![0, 0, 0];
+
+        let n = step_ret.len() as i32;
+        let ret_buf = unsafe { MappedF32Buffer::new(step_ret.len()) }
+            .expect("alloc step_ret");
+        ret_buf.write_from_slice(&step_ret);
+        let close_buf = unsafe { MappedI32Buffer::new(trade_close.len()) }
+            .expect("alloc trade_close");
+        close_buf.write_from_slice(&trade_close);
+        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }
+            .expect("alloc isv");
+        isv_buf.write_from_slice(&isv);
+
+        launch_reward_cap(
+            &stream, &kernel,
+            ret_buf.dev_ptr, close_buf.dev_ptr, n, isv_buf.dev_ptr,
+            REWARD_POS_CAP_ADAPTIVE_INDEX as i32,
+            REWARD_NEG_CAP_ADAPTIVE_INDEX as i32,
+            SENTINEL_REWARD_POS_CAP, SENTINEL_REWARD_NEG_CAP,
+            REWARD_CAP_SAFETY_FACTOR, REWARD_NEG_TO_POS_RATIO,
+            REWARD_POS_CAP_MIN, REWARD_POS_CAP_MAX,
+            REWARD_CAP_EMA_ALPHA,
+        );
+
+        let result = isv_buf.read_all();
+        assert_eq!(
+            result[REWARD_POS_CAP_ADAPTIVE_INDEX], SEED_POS,
+            "POS slot must be preserved bit-exactly when no winning closes",
+        );
+        assert_eq!(
+            result[REWARD_NEG_CAP_ADAPTIVE_INDEX], SEED_NEG,
+            "NEG slot must be preserved bit-exactly when no winning closes",
+        );
+    }
+
+    /// Test 3 — Bounds enforcement: huge winning return clamps POS to MAX=50.
+    ///
+    /// Single huge winning close at 1000.0. p99 estimate ≈ 1000 + 0 = 1000,
+    /// max=1000, × safety_factor=1.5 = 1500. Must clamp to POS_MAX=50.
+    /// NEG = -2 × 50 = -100 (also clamped at the lower bound).
+    #[test]
+    #[ignore = "requires GPU"]
+    fn reward_cap_bounds_clamp_extreme_outlier() {
+        let stream = make_stream();
+        let kernel = load_kernel(&stream);
+
+        const ISV_DIM: usize = 1024;
+        let mut isv = vec![0.0_f32; ISV_DIM];
+        isv[REWARD_POS_CAP_ADAPTIVE_INDEX] = SENTINEL_REWARD_POS_CAP;
+        isv[REWARD_NEG_CAP_ADAPTIVE_INDEX] = SENTINEL_REWARD_NEG_CAP;
+
+        // Single huge winning close — would explode the cap if unclamped.
+        let step_ret: Vec<f32> = vec![1000.0];
+        let trade_close: Vec<i32> = vec![1];
+
+        let n = step_ret.len() as i32;
+        let ret_buf = unsafe { MappedF32Buffer::new(step_ret.len()) }
+            .expect("alloc step_ret");
+        ret_buf.write_from_slice(&step_ret);
+        let close_buf = unsafe { MappedI32Buffer::new(trade_close.len()) }
+            .expect("alloc trade_close");
+        close_buf.write_from_slice(&trade_close);
+        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }
+            .expect("alloc isv");
+        isv_buf.write_from_slice(&isv);
+
+        launch_reward_cap(
+            &stream, &kernel,
+            ret_buf.dev_ptr, close_buf.dev_ptr, n, isv_buf.dev_ptr,
+            REWARD_POS_CAP_ADAPTIVE_INDEX as i32,
+            REWARD_NEG_CAP_ADAPTIVE_INDEX as i32,
+            SENTINEL_REWARD_POS_CAP, SENTINEL_REWARD_NEG_CAP,
+            REWARD_CAP_SAFETY_FACTOR, REWARD_NEG_TO_POS_RATIO,
+            REWARD_POS_CAP_MIN, REWARD_POS_CAP_MAX,
+            REWARD_CAP_EMA_ALPHA,
+        );
+
+        let result = isv_buf.read_all();
+        let pos = result[REWARD_POS_CAP_ADAPTIVE_INDEX];
+        let neg = result[REWARD_NEG_CAP_ADAPTIVE_INDEX];
+
+        // Bounds: POS clamped to MAX=50.
+        assert!(
+            (pos - REWARD_POS_CAP_MAX).abs() < 1e-4,
+            "POS must clamp to MAX=50 on extreme outlier; got {pos}"
+        );
+        // NEG = -2 × 50 = -100 (lower bound by construction).
+        let expected_neg = -REWARD_NEG_TO_POS_RATIO * REWARD_POS_CAP_MAX;
+        assert!(
+            (neg - expected_neg).abs() < 1e-4,
+            "NEG must clamp to -2 × MAX = {expected_neg}; got {neg}"
+        );
+    }
+
+    /// Test 4 — Welford EMA blend after bootstrap.
+    ///
+    /// Pre-seed POS=10.0 (NOT sentinel — bootstrap path NOT taken).
+    /// Single winning close at 4.0. mean=4, sigma=0, p99=4, max=4.
+    /// Target = max(p99, max) × 1.5 = 4 × 1.5 = 6. Clamped to [1, 50] → 6.
+    /// EMA blend: (1 - α) × 10 + α × 6 with α=0.01:
+    /// = 0.99 × 10 + 0.01 × 6 = 9.96. NEG = -2 × 9.96 = -19.92.
+    #[test]
+    #[ignore = "requires GPU"]
+    fn reward_cap_welford_ema_after_bootstrap() {
+        let stream = make_stream();
+        let kernel = load_kernel(&stream);
+
+        const ISV_DIM: usize = 1024;
+        const SEEDED_POS: f32 = 10.0;
+        const SEEDED_NEG: f32 = -20.0;
+        let mut isv = vec![0.0_f32; ISV_DIM];
+        // NOT sentinel — Pearl-A skipped, EMA path taken.
+        isv[REWARD_POS_CAP_ADAPTIVE_INDEX] = SEEDED_POS;
+        isv[REWARD_NEG_CAP_ADAPTIVE_INDEX] = SEEDED_NEG;
+
+        let step_ret: Vec<f32> = vec![4.0];
+        let trade_close: Vec<i32> = vec![1];
+
+        let n = step_ret.len() as i32;
+        let ret_buf = unsafe { MappedF32Buffer::new(step_ret.len()) }
+            .expect("alloc step_ret");
+        ret_buf.write_from_slice(&step_ret);
+        let close_buf = unsafe { MappedI32Buffer::new(trade_close.len()) }
+            .expect("alloc trade_close");
+        close_buf.write_from_slice(&trade_close);
+        let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }
+            .expect("alloc isv");
+        isv_buf.write_from_slice(&isv);
+
+        launch_reward_cap(
+            &stream, &kernel,
+            ret_buf.dev_ptr, close_buf.dev_ptr, n, isv_buf.dev_ptr,
+            REWARD_POS_CAP_ADAPTIVE_INDEX as i32,
+            REWARD_NEG_CAP_ADAPTIVE_INDEX as i32,
+            SENTINEL_REWARD_POS_CAP, SENTINEL_REWARD_NEG_CAP,
+            REWARD_CAP_SAFETY_FACTOR, REWARD_NEG_TO_POS_RATIO,
+            REWARD_POS_CAP_MIN, REWARD_POS_CAP_MAX,
+            REWARD_CAP_EMA_ALPHA,
+        );
+
+        let result = isv_buf.read_all();
+        let pos = result[REWARD_POS_CAP_ADAPTIVE_INDEX];
+        let neg = result[REWARD_NEG_CAP_ADAPTIVE_INDEX];
+
+        // EMA blend: 0.99 × 10 + 0.01 × 6 = 9.96.
+        let expected_pos = (1.0 - REWARD_CAP_EMA_ALPHA) * SEEDED_POS
+            + REWARD_CAP_EMA_ALPHA * (4.0 * REWARD_CAP_SAFETY_FACTOR);
+        assert!(
+            (pos - expected_pos).abs() < 1e-3,
+            "POS Welford EMA: expected {expected_pos}, got {pos}"
+        );
+        // 2:1 ratio preserved.
+        assert!(
+            (neg - (-REWARD_NEG_TO_POS_RATIO * pos)).abs() < 1e-4,
+            "NEG must equal -2 × POS at producer; got POS={pos}, NEG={neg}"
+        );
+    }
+}
+
 // ═══════════════════════════════════════════════════════════════════════════
 // Test B.8: Layout-fingerprint regression for direction Q-head input bump.
 //
diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md
index 09a36d4b7..010c01f96 100644
--- a/docs/dqn-wire-up-audit.md
+++ b/docs/dqn-wire-up-audit.md
@@ -7810,3 +7810,103 @@ Wired `ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451]` — Pearl-A-bootstrapped Welford E
 ### Consumer audit
 
 Only one consumer of `MIN_HOLD_TARGET` was found in the CUDA call path (the `compute_min_hold_penalty` call in `experience_kernels.cu`). No additional consumers discovered. `trade_physics.cuh` takes `min_hold_target` as a parameter — no constant reads — no change needed there. Per `feedback_no_partial_refactor`: complete.
+
+---
+
+## 2026-05-08 — Class A P0-A: REWARD_POS/NEG_CAP ISV-driven from realized return p99/p1 EMA
+
+### Problem
+
+Hardcoded `REWARD_POS_CAP=+5.0f` / `REWARD_NEG_CAP=-10.0f` (state_layout.cuh:266-267) was structurally clipping the upper tail of realized alpha. Per Class A audit:
+
+```
+EV(trade) = WR × min(realized_win, +5)
+          + (1−WR) × max(realized_loss, −10)
+```
+
+With WR ≈ 0.46 (the months-long plateau), the cap asymmetry is fully active. Wins clip to +5 regardless of magnitude. Controller signals (sharpe EMA, var_q, q_gap) all derive from this CAPPED buffer — the controller cannot select for trades it cannot see. Selectivity gradient evaporates: small wins clip to +5 alongside large wins also clipping to +5.
+
+The state_layout comment lines 261-264 explicitly deferred Phase 2 (ISV-driven) "IF Phase 1 validation reveals adaptive need". 11 superprojects of WR-stuck-at-46-48% revealed adaptive need.
+
+### Fix
+
+Two new ISV slots [452, 453] driven by a new producer kernel `reward_cap_update_kernel.cu`:
+
+- **POS cap producer**: block-tree-reduce `mean + Z_99 × sigma` p99 estimator over winning realized returns from the per-epoch `step_ret_per_sample` × `trade_close_per_sample` mask, with conservative `max(p99_estimate, max_winning_return)` takeover, × 1.5 safety factor, clamped to [1, 50].
+- **NEG cap = −2 × POS cap**: 2:1 Kahneman/Tversky asymmetry (per `pearl_audit_unboundedness_for_implicit_asymmetry`) preserved at producer-time — single source of truth, no consumer applies the multiplier itself. Resulting NEG ∈ [−100, −2] derived from POS bounds × ratio.
+
+### Cold-start behavior
+
+- Sentinels (5.0 / −10.0) match pre-P0-A hardcoded constants for bit-identical cold-start until the first valid observation.
+- Pearl-A bootstrap: when `|isv_pos − sentinel| < EPS`, REPLACE with target directly (no blend).
+- After bootstrap: Welford α=0.01 slow EMA (reward distribution is the foundation of training; shouldn't move fast).
+- Consumer-side fallback: when ISV slot at sentinel OR outside [REWARD_POS_CAP_MIN_BOUND=1, REWARD_POS_CAP_MAX_BOUND=50], fall back to `state_layout.cuh` macro (defensive — defends against malformed prior state).
+
+### Cadence
+
+Per-epoch boundary launch (cold path). Reward distribution moves on policy-evolution timescales, not per-step; per-step would track sample noise.
+
+### Source of trade-close return data
+
+Existing per-sample buffers populated by `unified_env_step_core` in `experience_kernels.cu`:
+
+- `step_ret_per_sample[N*L]` (line ~2779) — `step_ret_core` raw signed per-step return.
+- `trade_close_per_sample[N*L]` (line ~2780) — 1 iff `exiting_trade || reversing_trade`.
+
+Mask `step_ret > 0 && trade_close == 1` selects realized winning closes. No new buffer added.
+
+### Sites modified
+
+| File | LOC delta | Change |
+|------|-----------|--------|
+| `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` | +63 | 2 slot constants + 2 sentinels + 7 producer constants + locked-layout test |
+| `crates/ml/src/cuda_pipeline/state_layout.cuh` | +27 | 2 slot defines + 2 sentinels + 2 bound defines (consumer-side fallback) |
+| `crates/ml/src/cuda_pipeline/reward_cap_update_kernel.cu` | +233 (new) | Producer kernel: block-tree-reduce p99 + Pearl-A + EMA + bounds |
+| `crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs` | +110 | `RewardCapUpdateOps` struct + cubin loader + launcher |
+| `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` | +83 | Cubin static + struct field + constructor wire + `launch_reward_cap_update` fn + ISV_TOTAL_DIM 452→454 + fingerprint seed update |
+| `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` | +22 | 2 device-ptr accessors (`step_ret_per_sample_dev_ptr`, `trade_close_per_sample_dev_ptr`) |
+| `crates/ml/src/cuda_pipeline/experience_kernels.cu` | +27 / −3 | Consumer 1: ISV read + sentinel-fallback at segment_complete cap |
+| `crates/ml/src/cuda_pipeline/compute_sp15_final_reward_kernel.cu` | +6 / −1 | Consumer 2: pass `isv` ptr to `sp15_apply_sp12_cap` helper |
+| `crates/ml/src/cuda_pipeline/sp15_reward_axis_helpers.cuh` | +35 / −12 | Consumer 3: device fn signature `(r_in, isv*)` + ISV read + sentinel-fallback |
+| `crates/ml/src/trainers/dqn/state_reset_registry.rs` | +24 | 2 FoldReset registry entries (`sp14_p0a_reward_pos_cap_adaptive`, `sp14_p0a_reward_neg_cap_adaptive`) |
+| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | +37 | 2 dispatch arms in `reset_named_state` (per C.10 lesson) + 1 launch call at per-epoch boundary |
+| `crates/ml/build.rs` | +21 | cubin manifest entry for `reward_cap_update_kernel.cu` |
+| `crates/ml/tests/sp14_oracle_tests.rs` | +280 | 4 GPU oracle tests: Pearl-A bootstrap, no-winning-trades preservation, bounds clamping, Welford EMA |
+
+### Architectural decisions
+
+1. **Asymmetry at producer, not consumer**: `NEG = −2 × POS` enforced inside `reward_cap_update_kernel`. Single source of truth — every consumer reads its corresponding ISV slot directly, no consumer applies the 2× ratio itself. Per `pearl_audit_unboundedness_for_implicit_asymmetry`.
+
+2. **Slot ordering for layout fingerprint compat**: Appended at 452-453 (immediately after SP14 Layer C Phase C.4b's slot 451). No insertion in earlier slots; checkpoint compatibility preserved beyond ISV_TOTAL_DIM.
+
+3. **Pearl-A bootstrap mandatory**: Cold-start sentinel REPLACES on first valid observation; no blend. Per `pearl_first_observation_bootstrap.md`.
+
+4. **Safety bounds Category-1**: POS ∈ [1, 50] are dimensional safety floors per `feedback_isv_for_adaptive_bounds`, NOT tuning. Cap below 1 makes the entire reward signal flat (selectivity collapses); cap above 50 lets a single adversarial outlier dominate any per-batch EMA. Producer kernel and consumer-side fallback both enforce.
+
+5. **No new buffer**: Reused existing `step_ret_per_sample` + `trade_close_per_sample` — no widening of the experience-collector tile.
+
+6. **Dispatch arms for FoldReset**: Per the C.10 lesson (missing `reset_named_state` arm for `sp14_q_disagreement_variance_ema` caused a runtime crash), both `sp14_p0a_reward_pos_cap_adaptive` and `sp14_p0a_reward_neg_cap_adaptive` arms added in `training_loop.rs::reset_named_state`.
+
+### Verification
+
+- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --tests --all-targets` — clean (19 pre-existing warnings, 0 new).
+- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --lib --release sp14_isv_slots` — 8/8 pass (4 layout + 4 fits-within tests).
+- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests --release --features cuda -- --ignored --nocapture` — 8/8 pass (4 existing + 4 new P0-A oracle tests).
+- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests --release --features cuda -- --ignored --nocapture` — 36/36 pass (no regression from `sp15_apply_sp12_cap` signature change).
+
+### Consumer audit
+
+3 consumer sites identified by Class A audit, all migrated atomically per `feedback_no_partial_refactor`:
+
+1. `experience_kernels.cu:3112-3114` — segment-complete asymmetric cap. Reads ISV[452, 453]; falls back to macros at sentinel.
+2. `compute_sp15_final_reward_kernel.cu:163` — Stage 4 SP12 cap. Now passes `isv` ptr into the helper.
+3. `sp15_reward_axis_helpers.cuh:211` — `sp15_apply_sp12_cap` device fn. Refactored signature `(r_in, isv*)`; reads ISV[452, 453] with sentinel-fallback to macros internally.
+
+No additional consumers found via `grep -rnE "REWARD_POS_CAP|REWARD_NEG_CAP" crates/ml/src/cuda_pipeline/`. Remaining matches are doc comments (gpu_dqn_trainer.rs:1462-63, gpu_experience_collector.rs:5697) and `trade_physics.cuh` parameter-name comments (no constant reads).
+
+### Cumulative WR-plateau fix series
+
+- Class C bug 1 (`8f218cab2`): replay buffer intent→realized — corrects the buffer's ground truth.
+- Class A P0-B (`8f218cab2`): Kelly warmup floor wiring — keeps Kelly cap from going to 0 cold-start.
+- Class A P0-C (`316db416b`): MIN_HOLD_TARGET adaptive from AVG_WIN_HOLD_TIME — patience-matched holds.
+- Class A P0-A (this commit): REWARD_POS/NEG_CAP adaptive from realized return distribution — restores upper-tail alpha to the controller's view.