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

feat(sp13): P0a atomic — Hold-pricing + dir_acc instrumentation (additive)

Browse Source 
Tests user's hypothesis (Hold being FREE is the bug, not Hold itself) by
pricing Hold via ISV-driven adaptive controller targeting 20% Hold-rate.

11 new SP13 ISV slots [372..383). 5 new GPU kernels:
- aux_dir_acc_reduce_kernel.cu (correct/pos_pred/pos_label/valid → 3 scalars)
- hold_rate_observer_kernel.cu (packed batch_actions decode, count(Hold)/B)
- apply_fixed_alpha_ema_kernel.cu (preserves short/long timescale split that
  Wiener-optimal apply_pearls_ad_kernel would collapse)
- aux_pred_to_isv_tanh_kernel.cu (mean(tanh(aux_pred)) → ISV[375])
- 3 reward-composition sites in experience_kernels.cu subtract isv[HOLD_COST]
  on Hold actions (segment_complete pre-asymmetric-cap, positioned-non-event
  per-bar, flat per-bar)

Host-side controller in training_loop.rs:
  excess = max(0, observed - target)
  hold_cost = HOLD_COST_BASE × (1 + 5 × excess), clamped [0.5×, 5.0×base]

Per-step observer + EMA chain in gpu_experience_collector.rs after
experience_action_select. Per-epoch HEALTH_DIAG emit:
  aux_dir_acc target/short/long/pred_tanh
  hold_pricing observed_rate/target/cost

4-way action space stays (ExposureLevel::Hold preserved). Replay buffer /
fxcache compatibility preserved. SP11 (11/11) + SP12 (14/14) tests no
regression. SP13 P0a oracle tests: 14/14 on RTX 3050 Ti.

Spec/plan: docs/superpowers/{specs,plans}/2026-05-04-sp13-redefine-success-for-predictive-skill.md (v3)
Audit: docs/dqn-wire-up-audit.md (SP13 P0a section appended)

v2 → v3 reframe: P0a.T3 v2 implementer's audit found DirectionAction enum
doesn't exist (codebase uses 8-variant fused ExposureLevel cascading through
77 files). v3 reframes from "eliminate Hold" (250 LOC + 32-test cascade) to
"price Hold" (additive, no contract change, no cross-crate cascade).

Tension with pearl_event_driven_reward_density_alignment acknowledged in spec
— per-bar Hold cost is exposure-NEGATIVE (pulls policy AWAY from Hold-default,
inverse of the pearl's failure mode), models real economic carry, ISV-bounded
by controller. Faithful reward modeling, not artificial shaping.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-05 00:52:54 +02:00
parent ba83fcd1f5
commit f934ea1719
16 changed files with 2385 additions and 36 deletions
Download Patch File Download Diff File Expand all files Collapse all files

57
crates/ml/build.rs
View File

@@ -663,6 +663,63 @@ fn main() {
// EPS_DIV) mirror sp4_wiener_ema.rs's host-reference for the
// GPU-vs-host oracle test.
"apply_pearls_kernel.cu",
// SP13 Phase 0a (2026-05-04): aux-head directional-accuracy
// reducer. Single-block tree-reduce kernel scoring the aux next-
// bar regression head's prediction sign against the next-bar
// return label sign. Four shared-memory int arrays
// (correct/pos_pred/pos_label/valid) reduce in lockstep — no
// atomicAdd per `feedback_no_atomicadd.md`. Outputs three
// batch-level fractions (dir_acc, pos_pred_frac, pos_label_frac)
// into a 3-element mapped-pinned buffer. P0a reads regression-
// mode aux scalar; Layer B swaps the source pointer to
// `softmax[1] - softmax[0]` after the head becomes 2-class
// softmax. Consumed by `launch_aux_dir_acc_reduce` in
// `gpu_dqn_trainer.rs`.
"aux_dir_acc_reduce_kernel.cu",
// SP13 v3 Phase 0a P0a.T3 (2026-05-04): Hold-rate observer.
// Single-block tree-reduce kernel counting per-bar Hold-action
// picks (decoded from packed `batch_actions[i]` via
// `dir = action_idx / (NUM_MAGNITUDES * NUM_ORD * NUM_URG)` —
// P0a.T4 ABI alignment, 2026-05-04) and writing the fraction
// `count(Hold) / batch_size` to a 1-element mapped-pinned buffer.
// One shared-memory int array reduces in lockstep — no atomicAdd
// per `feedback_no_atomicadd.md`. The fixed-α EMA applicator
// smooths the observation into ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382];
// the host-side controller reads ISV[382] + ISV[381] and writes the
// priced Hold cost back to ISV[HOLD_COST_INDEX=380]. The reward
// composition site in `experience_kernels.cu` then subtracts the
// cost on every Hold-action bar — pricing the action so the
// policy uses Hold deliberately rather than as a free CQL-bias-
// anchored default. Replaces v2's atomic Hold elimination per
// the v3 reframe (preserves the 4-way action space — no cross-
// crate cascade). Consumed by `launch_hold_rate_observer` in
// `gpu_dqn_trainer.rs`.
"hold_rate_observer_kernel.cu",
// SP13 Phase 0a P0a.T4 (2026-05-04): fixed-α EMA applicator.
// Sibling of `apply_pearls_kernel` — needed because Pearls A+D
// collapse two EMAs of the same signal to identical values
// (variance ratio is α-independent), which would deadlock the
// SP13 stagnation detector that compares slots 373 (α=0.3) vs
// 374 (α=0.05). Per-thread first-observation-sentinel bootstrap
// (caller passes the slot's registry sentinel — 0.5 for dir-acc,
// 0.0 for hold-rate) plus a fixed-α blend. No atomicAdd per
// `feedback_no_atomicadd.md`; pure GPU compute per
// `feedback_no_cpu_compute_strict.md`. Consumed by
// `launch_apply_fixed_alpha_ema` in `gpu_dqn_trainer.rs`.
"apply_fixed_alpha_ema_kernel.cu",
// SP13 Phase 0a P0a.T4 (2026-05-04): aux-head per-bar prediction
// → ISV[AUX_DIR_PREDICTION_INDEX=375] tanh-bounded scalar
// producer. Single-block tree-reduce reads the captured-graph
// `aux_nb_pred_buf [B]` tile and writes
// `mean(tanh(aux_pred[i]))` ∈ [-1, +1] to the SHARED ISV slot 375
// (per-step state, not an EMA — slot is overwritten each launch).
// tanh squash bounds the scalar so the downstream Q-head
// consumer sees a well-conditioned signal regardless of
// `label_scale` magnitude. No atomicAdd per
// `feedback_no_atomicadd.md`; pure GPU compute per
// `feedback_no_cpu_compute_strict.md`. Consumed by
// `launch_aux_pred_to_isv_tanh` in `gpu_dqn_trainer.rs`.
"aux_pred_to_isv_tanh_kernel.cu",
];
// ALL kernels get common header (BF16 types + wrappers)

93
crates/ml/src/cuda_pipeline/apply_fixed_alpha_ema_kernel.cu Normal file
View File

@@ -0,0 +1,93 @@
// crates/ml/src/cuda_pipeline/apply_fixed_alpha_ema_kernel.cu
//
// SP13 Phase 0a P0a.T4 (2026-05-04): fixed-alpha EMA applicator.
//
// Sibling of `apply_pearls_ad_kernel.cu` (the SP4 Wiener-optimal Pearls
// A+D applicator) for cases where the EMA needs DISTINCT TIMESCALES
// instead of a Wiener-optimal blend. SP13 Phase 0a needs both a SHORT
// (α=0.3) and a SLOW (α=0.05) EMA of the aux-head directional accuracy:
// the stagnation detector compares the two to decide whether `aux_w`
// should keep ramping or relax back toward the base. Pearl D would
// collapse both EMAs to the same value because the underlying signal is
// the same and the variance ratio doesn't depend on α — so the slot
// pair would report `short - long ≈ 0` permanently and the controller
// could never read "improving" or "stalled". Fixed-α EMAs preserve the
// timescale separation by construction.
//
// First-observation bootstrap per `pearl_first_observation_bootstrap`:
// when `prev_state == sentinel` the kernel REPLACES `prev_state` with
// the new sample directly (rather than blending `(1-α)·sentinel + α·obs`
// which would carry the sentinel into the EMA forever for the
// fold's first non-zero observation). The sentinel is passed as a
// kernel arg because each consumer slot has its own (slot 373 / 374
// use 0.5 — the random-guessing baseline for binary directional
// accuracy; slot 382 uses 0.0 — the empty-batch fallback for the
// hold-rate observer). The kernel has no compile-time knowledge of
// the slot's reset semantics; the launcher passes the same sentinel
// the registry entry uses on fold reset.
//
// One launch updates N consecutive ISV slots, each with its own
// (sample[i], prev_isv[isv_offset+i]) pair. For SP13 P0a: N=1 per
// EMA (one scalar dir-acc fed into both slots 373 and 374; two
// launches in lockstep). The kernel parallelises across slots —
// each thread updates one slot — so larger N (e.g. per-batch tile
// reductions) just scales the grid; no inter-thread coordination
// needed.
//
// Single-block, BLOCK_SIZE up to 256 (sized at launch from N).
// Per-thread bootstrap test + blend; no shared memory, no atomicAdd
// (per `feedback_no_atomicadd.md`). Pure GPU compute per
// `feedback_no_cpu_compute_strict.md`. Stream-ordered with the
// producer that wrote `sample[]`; the launcher MUST issue this
// kernel on the SAME stream so the producer's `__threadfence_system()`
// orders against the read here.
#include <cuda_runtime.h>
extern "C" __global__ void apply_fixed_alpha_ema_kernel(
/* New per-slot observations [N]. The producer kernel that wrote
* this buffer MUST have issued `__threadfence_system()` before
* returning, so this thread's read is well-ordered against the
* write. */
const float* __restrict__ sample,
/* Number of slots to update (the launcher's grid covers `N`
* threads). N=1 is the common SP13 case (a single dir-acc scalar
* fanned out into two ISV slots via two launches with different α). */
int n,
/* Offset into the ISV array where this kernel begins writing. The
* caller's slot indices are `[isv_offset, isv_offset + N)`. */
int isv_offset,
/* Fixed blend rate. SP13: 0.3 for the short EMA (slot 373), 0.05
* for the slow EMA (slot 374) and the hold-rate EMA (slot 382). */
float alpha,
/* Sentinel value for first-observation bootstrap. When
* `isv[isv_offset + i] == sentinel` (per `pearl_first_observation_bootstrap`),
* the kernel replaces directly with `sample[i]` instead of blending.
* 0.5 for the SP13 dir-acc EMAs (slots 373/374), 0.0 for the
* hold-rate EMA (slot 382). */
float sentinel,
/* ISV array — both read (prev state) and write (new state) at
* `[isv_offset, isv_offset + N)`. */
float* __restrict__ isv)
{
const int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i >= n) return;
const int slot = isv_offset + i;
const float prev = isv[slot];
const float new_obs = sample[i];
/* Pearl A first-observation bootstrap: replace directly when the
* stored state still matches the registry sentinel. For an EMA
* whose sentinel is the new fold's "no-observation-yet" baseline
* (0.5 for dir-acc, 0.0 for hold-rate), blending α·obs with the
* sentinel would keep the sentinel's bias in the EMA across the
* fold boundary. Direct replacement on the first real observation
* gives the new fold a clean start. */
const float new_state = (prev == sentinel)
? new_obs
: (1.0f - alpha) * prev + alpha * new_obs;
isv[slot] = new_state;
__threadfence_system();
}

144
crates/ml/src/cuda_pipeline/aux_dir_acc_reduce_kernel.cu Normal file
View File

@@ -0,0 +1,144 @@
// crates/ml/src/cuda_pipeline/aux_dir_acc_reduce_kernel.cu
//
// SP13 Phase 0a (2026-05-04): aux-head directional-accuracy reducer.
//
// Single-block tree-reduce kernel that scores the auxiliary next-bar
// regression head's per-bar prediction sign against the next-bar return
// label sign. Produces three batch-level fractions:
// out_3[0] = dir_acc — fraction of valid bars where pred_sign == label_sign
// out_3[1] = pos_pred_frac — fraction of valid bars where pred > 0
// out_3[2] = pos_label_frac — fraction of valid bars where label > 0
//
// "Valid" bars are those whose `next_bar_label != 0.0f` — a zero label is
// treated as "no signal to score against" (the producer of the label
// emits 0 when next-bar return rounds to flat) and excluded from the
// denominator entirely. When the entire batch is invalid (denom == 0),
// all three outputs are the random-baseline sentinel 0.5 (matches
// `DIR_ACC_EMA_SENTINEL` in `sp13_isv_slots.rs` per
// `pearl_first_observation_bootstrap` — the EMA's first observation
// replaces the sentinel directly).
//
// Sign convention for the prediction: `(p > 0.0f) ? 1 : 0`. Zero (and
// negative-zero) predict class 0 (negative). Tested by the
// `dir_acc_zero_pred_is_negative` GPU oracle.
//
// Reads regression-mode aux scalar in P0a (the existing aux_heads
// next-bar regression head writes one f32 per bar). Layer B refactors the
// aux head to a 2-class softmax classifier and updates this kernel to
// read `(softmax[1] > softmax[0])` instead — the producer/consumer
// contract on the rest of the pipeline (3-element output, sentinel
// semantics, valid-bar masking) is unchanged.
//
// Single block, BLOCK_SIZE=256. Four shared-memory int arrays
// (correct, pos_pred, pos_label, valid) all participate in one fused
// log2(BLOCK_SIZE) tree reduce — no atomicAdd per
// `feedback_no_atomicadd.md`. Pure GPU compute per
// `feedback_no_cpu_compute_strict.md`. Shared memory footprint:
// 4 × 256 × sizeof(int) = 4096 bytes.
#include <cuda_runtime.h>
#define BLOCK_SIZE 256
extern "C" __global__ void aux_dir_acc_reduce_kernel(
/* Per-bar aux-head prediction. P0a regression mode: one f32 per bar
* (the next-bar return regression scalar). Layer B switches the head
* to 2-class softmax and the launcher swaps the source pointer to
* `softmax[1] - softmax[0]` — same buffer shape, same kernel. */
const float* __restrict__ aux_pred,
/* Per-bar next-bar return label, sign-encoded. +1 / -1 for valid
* directional bars; 0 for "no signal" bars which are excluded from
* the denominator (see header comment). */
const float* __restrict__ next_bar_label,
/* Number of bars in the batch. The kernel strides each thread over
* `[0, batch_size)` so any batch_size up to ~2^31 is valid; in
* practice batches are O(1k-10k). */
int batch_size,
/* 3-element output buffer. Mapped-pinned in production (see the
* launcher in `gpu_dqn_trainer.rs`). The kernel writes
* out_3[0] = dir_acc
* out_3[1] = pos_pred_frac
* out_3[2] = pos_label_frac
* in thread 0 after the tree-reduce completes. */
float* __restrict__ out_3)
{
/* Single-block reducer — guard against accidental multi-block launch. */
if (blockIdx.x != 0) return;
const int tid = threadIdx.x;
const int bdim = blockDim.x;
/* Four int arrays packed back-to-back in dynamic shared memory.
* `extern __shared__ int shared[]` is sized at launch via
* `shared_mem_bytes = 4 × bdim × sizeof(int)` — the launcher MUST
* pass that exact value (the test scaffold and the production
* launcher both do). */
extern __shared__ int shared[];
int* sh_correct = &shared[0 * bdim];
int* sh_pos_pred = &shared[1 * bdim];
int* sh_pos_label = &shared[2 * bdim];
int* sh_valid = &shared[3 * bdim];
/* Per-thread strided accumulation. Each thread folds its slice of
* `[0, batch_size)` into four private int counters, then writes them
* into shared memory for the tree reduce. Bars whose label is zero
* (no directional signal to score against) are skipped entirely —
* `valid` does NOT increment, and `correct/pos_pred/pos_label` are
* not touched for those bars. */
int local_correct = 0;
int local_pos_pred = 0;
int local_pos_label = 0;
int local_valid = 0;
for (int i = tid; i < batch_size; i += bdim) {
const float p = aux_pred[i];
const float l = next_bar_label[i];
if (l == 0.0f) continue;
local_valid += 1;
const int pred_pos = (p > 0.0f) ? 1 : 0;
const int label_pos = (l > 0.0f) ? 1 : 0;
if (pred_pos == label_pos) local_correct += 1;
local_pos_pred += pred_pos;
local_pos_label += label_pos;
}
sh_correct[tid] = local_correct;
sh_pos_pred[tid] = local_pos_pred;
sh_pos_label[tid] = local_pos_label;
sh_valid[tid] = local_valid;
__syncthreads();
/* Standard log2(BLOCK_SIZE) tree reduction over all four arrays in
* lockstep. One `__syncthreads()` per pass keeps every thread's view
* of all four arrays consistent — splitting them into separate
* loops would cost 4× the syncs without any compute saving. */
for (int s = bdim / 2; s > 0; s >>= 1) {
if (tid < s) {
sh_correct[tid] += sh_correct[tid + s];
sh_pos_pred[tid] += sh_pos_pred[tid + s];
sh_pos_label[tid] += sh_pos_label[tid + s];
sh_valid[tid] += sh_valid[tid + s];
}
__syncthreads();
}
/* Thread 0 finalises. When the batch has no valid bars (empty batch
* or all-zero labels) the denom is 0; all three outputs fall back to
* 0.5 — the random-baseline sentinel matching DIR_ACC_EMA_SENTINEL
* in `sp13_isv_slots.rs`. The downstream EMA's first-observation
* replacement (Pearl A) consumes the sentinel naturally. */
if (tid == 0) {
const float denom = (float)sh_valid[0];
const float sentinel = 0.5f;
if (denom > 0.0f) {
out_3[0] = (float)sh_correct[0] / denom;
out_3[1] = (float)sh_pos_pred[0] / denom;
out_3[2] = (float)sh_pos_label[0] / denom;
} else {
out_3[0] = sentinel;
out_3[1] = sentinel;
out_3[2] = sentinel;
}
__threadfence_system();
}
}

111
crates/ml/src/cuda_pipeline/aux_pred_to_isv_tanh_kernel.cu Normal file
View File

@@ -0,0 +1,111 @@
// crates/ml/src/cuda_pipeline/aux_pred_to_isv_tanh_kernel.cu
//
// SP13 Phase 0a P0a.T4 (2026-05-04): aux-head per-bar prediction →
// ISV[AUX_DIR_PREDICTION_INDEX=375] tanh-bounded scalar producer.
//
// Reads the aux next-bar regression head's `aux_pred [B]` tile (the
// per-bar return-prediction scalar produced inside the captured forward
// graph) and writes `mean(tanh(aux_pred[i]))` ∈ [-1, +1] to the SHARED
// ISV scalar at slot 375. The tanh squash bounds the scalar so the
// downstream consumer (the direction Q-head input layer mirroring the
// SP13 spec's "feed aux prediction back into the policy" wiring) sees
// a well-conditioned signal rather than a raw regression scalar that
// could grow with `label_scale`.
//
// Why a batch mean instead of per-bar broadcast: ISV is a single
// `[ISV_TOTAL_DIM]` shared array (see `gpu_dqn_trainer.rs::isv_signals_pinned`
// allocation) — there is NO per-batch tile. Every per-step ISV producer
// reduces a per-bar tile to a single scalar before writing (cf.
// `h_s2_rms_ema_kernel.cu` reducing `save_h_s2 [B, SH2]` to a single
// `producer_step_scratch_buf` slot, or `aux_label_scale_ema_update`
// reducing `aux_nb_label_buf [B]` to one EMA slot). The mean of
// per-bar predictions is the natural batch-aggregate signal for a
// shared scalar — it sits at zero when the head has no consensus
// direction, drifts toward +1/-1 as the head develops conviction.
//
// Sentinel-friendly: per `pearl_first_observation_bootstrap`,
// `(prev == 0.0)` is the cold-start sentinel for slot 375 (constructor
// zero-init; no fold-reset entry — the slot is per-step overwritten so
// any non-zero produces a clean replacement on the first launch via
// the existing per-step write semantics). No additional bootstrap
// branch is needed here because the kernel ALWAYS overwrites the
// slot with the current step's batch mean — it's a per-step state,
// not an EMA, and the registry entry skips slot 375 explicitly per
// the SP13 P0a registry comments.
//
// Single block, BLOCK_SIZE=256, one shared-memory float array — no
// atomicAdd per `feedback_no_atomicadd.md`. Pure GPU compute per
// `feedback_no_cpu_compute_strict.md`. Stream-ordered with the
// producer that wrote `aux_pred`; same-stream barrier suffices
// (cf. `aux_heads_loss_ema_update`'s same-stream contract with
// `aux_next_bar_loss_reduce`).
//
// Cost per launch: 256 threads × ceil(B/256) strided tanh + sum
// passes, plus log2(256)=8 reduce steps + 1 global write. B is
// O(1k-10k) in production training; total ~1-2µs per step, well
// under the per-step ISV-producer budget.
#include <cuda_runtime.h>
#define BLOCK_SIZE 256
extern "C" __global__ void aux_pred_to_isv_tanh_kernel(
/* Per-bar aux next-bar regression scalar [B] — produced inside
* the captured forward graph by `aux_next_bar_forward`, valid
* after the launcher's same-stream sync against the producer. */
const float* __restrict__ aux_pred,
/* Number of bars in the batch. The strided loop covers any
* `batch_size > 0`; empty batch falls through to the sentinel
* write (see thread-0 finaliser). */
int batch_size,
/* ISV slot index for the SHARED scalar — `AUX_DIR_PREDICTION_INDEX = 375`
* in the SP13 layout. The kernel writes
* isv[isv_slot_offset] = mean(tanh(aux_pred[i]))
* after the tree-reduce, with `__threadfence_system()` to make the
* write visible to subsequent same-stream consumers. */
int isv_slot_offset,
/* ISV array (whole bus). The kernel only touches slot
* `isv_slot_offset`; no other slots are read or written. */
float* __restrict__ isv)
{
/* Single-block reducer — guard against accidental multi-block launch. */
if (blockIdx.x != 0) return;
extern __shared__ float sh_sum[];
const int tid = threadIdx.x;
const int bdim = blockDim.x;
/* Per-thread strided accumulation: each thread folds tanh(aux_pred[i])
* over its slice of `[0, batch_size)`. tanh squashes any unbounded
* regression scalar into [-1, +1] before the reduce so the batch
* mean is bounded by construction (the pre-squash mean could grow
* unbounded with `label_scale` and contaminate downstream
* consumers). */
float local_sum = 0.0f;
for (int i = tid; i < batch_size; i += bdim) {
local_sum += tanhf(aux_pred[i]);
}
sh_sum[tid] = local_sum;
__syncthreads();
/* Standard log2(BLOCK_SIZE) tree reduction. */
for (int s = bdim / 2; s > 0; s >>= 1) {
if (tid < s) {
sh_sum[tid] += sh_sum[tid + s];
}
__syncthreads();
}
/* Thread 0 writes the batch mean. Empty-batch path writes 0.0
* (matches the constructor-zero sentinel for slot 375 — the
* downstream Q-head consumer treats 0.0 as "no aux signal yet"
* via the same logic that handles the cold-start step before
* the first forward pass). */
if (tid == 0) {
const float mean = (batch_size > 0)
? (sh_sum[0] / (float)batch_size)
: 0.0f;
isv[isv_slot_offset] = mean;
__threadfence_system();
}
}

36
crates/ml/src/cuda_pipeline/experience_kernels.cu
View File

@@ -2768,6 +2768,18 @@ extern "C" __global__ void experience_env_step(
* `exiting_trade=1` upstream (line ~2371), so this branch fires
* for both — the segregation matters only for component attribution. */
float base_reward = 2.0f * vol_normalized_return;
/* SP13 v3 P0a.T3 (2026-05-04): per-bar Hold cost subtraction at the
* voluntary segment_complete branch. Goes BEFORE the SP12 asymmetric
* cap so the cost participates in the bounded clamp — keeps total
* reward in the same [-10, +5] range and avoids a long Hold run
* silently saturating the lower cap. Hold action rarely coincides
* with segment_complete (Hold keeps current position, doesn't close
* trades) but trail-fire can force-exit while the policy picks Hold;
* this branch covers that edge. Per-bar (non-segment_complete)
* Hold-cost subtraction is in the per-bar branches below. */
if (dir_idx == DIR_HOLD && isv_signals_ptr != NULL) {
base_reward -= isv_signals_ptr[ISV_HOLD_COST_IDX];
}
/* SP12 v3 fix (2026-05-04): asymmetric bounded cap (loss aversion).
* SP11 (commit 35db31089) made the cap symmetric ±10 to fix slot-63
* PopArt EMA inflation. That fix preserved stability but erased the
@@ -3148,6 +3160,20 @@ extern "C" __global__ void experience_env_step(
* (which is a separate concern; spec carve-out: "ONLY micro_reward
* and opp_cost are the per-bar shaping anti-patterns"). */
r_micro = 0.0f;
/* SP13 v3 P0a.T3 (2026-05-04): per-bar Hold cost subtraction.
* Hold action keeps the current position with zero PnL signal —
* pre-SP13 this was a free no-op the policy could pick indefinitely
* (CQL bias anchors toward minimum-variance Hold ~45% of bars across
* SP1-SP12). Hold-pricing makes Hold uneconomic when overused: the
* controller in `training_loop.rs` raises ISV[HOLD_COST_INDEX] when
* observed Hold-rate > target (default 0.20). The cost is bounded
* structurally (HOLD_COST_BASE × CEIL_RATIO = 0.005 max per bar),
* well below the SP12 asymmetric cap [-10, +5] — no per-bar clamp
* required. Spec §"Change 1: Price Hold (replaces v2's Eliminate
* Hold)". */
if (dir_idx == DIR_HOLD && isv_signals_ptr != NULL) {
r_micro -= isv_signals_ptr[ISV_HOLD_COST_IDX];
}
reward = r_micro;
micro_reward_per_sample[out_off] = r_micro;
reward_components_per_sample[out_off * 6 + 3] = r_micro;
@@ -3202,6 +3228,16 @@ extern "C" __global__ void experience_env_step(
* preserved: only the Flat opp_cost EMITTER is zeroed; the
* conviction Welford observer is independent. */
r_opp_cost = 0.0f;
/* SP13 v3 P0a.T3 (2026-05-04): per-bar Hold cost subtraction on
* the flat-non-event branch. Mirrors the positioned-branch
* subtraction above so a policy picking Hold while flat (i.e.
* "stay flat") pays the same per-bar cost as a policy picking
* Hold while positioned (i.e. "stay long/short"). The cost is
* structurally bounded by HOLD_COST_BASE × CEIL_RATIO = 0.005
* per bar, well below the SP12 asymmetric cap [-10, +5]. */
if (dir_idx == DIR_HOLD && isv_signals_ptr != NULL) {
r_opp_cost -= isv_signals_ptr[ISV_HOLD_COST_IDX];
}
reward = r_opp_cost;
reward_components_per_sample[out_off * 6 + 4] = r_opp_cost;
}

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

@@ -588,6 +588,70 @@ static SP11_POPART_COMPONENT_EMA_CUBIN: &[u8] =
static SP11_NOVELTY_SIMHASH_PROJ_INIT_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/novelty_simhash_proj_init_kernel.cubin"));
/// SP13 v3 Phase 0a P0a.T3 (2026-05-04): Hold-rate observer reducer.
///
/// Single-block tree-reduce kernel that decodes per-bar packed factored
/// `batch_actions[i]` (`dir = action_idx / (NUM_MAGNITUDES * NUM_ORD * NUM_URG)`)
/// and writes the resulting fraction `count(Hold) / batch_size ∈ [0, 1]`
/// to a 1-element mapped-pinned scratch. The chained
/// `apply_fixed_alpha_ema_kernel` (α=0.05, sentinel 0.0) blends the
/// per-step observation into `ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382]`;
/// the host-side Hold-cost controller in `training_loop.rs` reads the
/// EMA + the static target at `ISV[HOLD_RATE_TARGET_INDEX=381]` and
/// writes the priced Hold cost back to `ISV[HOLD_COST_INDEX=380]`. The
/// reward composition site in `experience_kernels.cu` then subtracts
/// the cost on every Hold-action bar — pricing the action so the policy
/// uses Hold deliberately rather than as a free CQL-bias-anchored
/// default.
///
/// Loaded by `GpuExperienceCollector` (per-step on the action-select
/// stream) — the per-step training-time observation is collector-owned
/// because `batch_actions` lives there. The trainer carries no copy
/// (avoids a dead `CudaFunction` field per `feedback_no_stubs.md`); the
/// `pub(crate)` visibility lets the collector's constructor reach the
/// cubin without re-declaring it. See `hold_rate_observer_kernel.cu`
/// for kernel contract details.
pub(crate) static SP13_HOLD_RATE_OBSERVER_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/hold_rate_observer_kernel.cubin"));
/// SP13 Phase 0a (2026-05-04): aux-head directional-accuracy reducer.
/// Single-block tree-reduce kernel that scores the aux next-bar
/// regression head's per-bar prediction sign against the next-bar return
/// label sign and writes (dir_acc, pos_pred_frac, pos_label_frac) into a
/// 3-element mapped-pinned buffer. Consumed by
/// `GpuDqnTrainer::launch_aux_dir_acc_reduce` (the per-step launcher
/// wired by P0a.T4 — this static include lands the cubin handle for the
/// constructor in P0a.T2). Produces directional-accuracy signal feeding
/// the SP13 EMAs at ISV[373..375) (short/long EMAs, sentinel 0.5 per
/// `pearl_first_observation_bootstrap`). See
/// `aux_dir_acc_reduce_kernel.cu` for kernel contract details.
static SP13_AUX_DIR_ACC_REDUCE_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/aux_dir_acc_reduce_kernel.cubin"));
/// SP13 Phase 0a P0a.T4 (2026-05-04): fixed-α EMA applicator. Sibling
/// of `apply_pearls_kernel` for cases where Pearls A+D's Wiener-optimal
/// blend would collapse two EMAs of the same signal to identical
/// values (defeating the SP13 stagnation detector that compares
/// slot 373's α=0.3 EMA against slot 374's α=0.05 EMA). Per-thread
/// first-observation-sentinel bootstrap + fixed-α blend; caller passes
/// the slot's registry sentinel (0.5 for dir-acc, 0.0 for hold-rate).
/// Consumed by `GpuDqnTrainer::launch_apply_fixed_alpha_ema`. See
/// `apply_fixed_alpha_ema_kernel.cu` for kernel contract details.
pub(crate) static SP13_APPLY_FIXED_ALPHA_EMA_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/apply_fixed_alpha_ema_kernel.cubin"));
/// SP13 Phase 0a P0a.T4 (2026-05-04): aux-head per-bar prediction →
/// ISV[AUX_DIR_PREDICTION_INDEX=375] tanh-bounded scalar producer.
/// Single-block tree-reduce reads the captured-graph `aux_nb_pred_buf [B]`
/// tile and writes `mean(tanh(aux_pred[i]))` ∈ [-1, +1] to the SHARED
/// ISV slot 375 (per-step overwrite — not an EMA). tanh squash bounds
/// the scalar so the downstream Q-head consumer sees a well-conditioned
/// signal regardless of `label_scale` magnitude. Consumed by
/// `GpuDqnTrainer::launch_aux_pred_to_isv_tanh`. See
/// `aux_pred_to_isv_tanh_kernel.cu` for kernel contract details.
static SP13_AUX_PRED_TO_ISV_TANH_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/aux_pred_to_isv_tanh_kernel.cubin"));
/// SP11 Fix 39 (2026-05-04, Task A2): SimHash novelty signal — lookup +
/// update kernels sharing one cubin. Lookup reads
/// `1/sqrt(1+count)` ∈ [0, 1] for each (state, action) bucket; update
@@ -983,7 +1047,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 = 367; // SP5 + Layer D D1+D2+D3 + SP7 + SP8 + SP9 + SP10 + SP11: 173 + 194 (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])
pub(crate) const ISV_TOTAL_DIM: usize = 383; // SP5 + Layer D D1+D2+D3 + SP7 + SP8 + SP9 + SP10 + SP11 + SP13 + SP13 v3: 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]; 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).
@@ -2040,7 +2104,11 @@ const fn layout_fingerprint_seed() -> &'static [u8] {
SABOTEUR_ENGAGEMENT_RATE=358;PNL_REWARD_MAGNITUDE_EMA=359;\
POPART_COMPONENT_MAG_EMA=360;\
REWARD_COMPONENT_VAR_EMA_BASE=361;\
ISV_TOTAL_DIM=367;\
TARGET_DIR_ACC=372;AUX_DIR_ACC_SHORT_EMA=373;AUX_DIR_ACC_LONG_EMA=374;\
AUX_DIR_PREDICTION=375;DIR_SKILL_BONUS_ALPHA=376;DIR_SKILL_BONUS_BETA=377;\
LUCK_WIN_DISCOUNT=378;SKILL_BONUS_CAP_RATIO=379;\
HOLD_COST=380;HOLD_RATE_TARGET=381;HOLD_RATE_OBSERVED_EMA=382;\
ISV_TOTAL_DIM=383;\
PARAM_W_A_H_S1=0;PARAM_B_A_H_S1=1;PARAM_W_B_H_S1=2;PARAM_B_B_H_S1=3;\
PARAM_W_RESIDUAL_H_S1=4;PARAM_GAMMA_H_S1=5;PARAM_BETA_H_S1=6;\
PARAM_W_A_H_S2=7;PARAM_B_A_H_S2=8;PARAM_W_B_H_S2=9;PARAM_B_B_H_S2=10;\
@@ -4926,6 +4994,47 @@ pub struct GpuDqnTrainer {
/// unwired).
pub(crate) sp11_popart_component_dev_ptr: u64,
pub(crate) sp11_popart_component_len: usize,
/// SP13 Phase 0a (2026-05-04): aux-head directional-accuracy
/// reducer kernel. Single-block tree-reduce loaded from
/// `aux_dir_acc_reduce_kernel.cubin`. Reads the aux next-bar
/// regression-head prediction buffer + a per-bar sign-encoded
/// `next_bar_label` buffer and writes
/// `(dir_acc, pos_pred_frac, pos_label_frac)` to
/// `aux_dir_acc_buf` (mapped-pinned, 3 floats). Producer is wired
/// by P0a.T4; the kernel handle here lands the constructor edge
/// per `feedback_no_partial_refactor.md` (P0a is one atomic
/// commit). See the kernel source for the per-bar prediction sign
/// convention (`(p > 0.0f) ? 1 : 0`) and the all-zero-label
/// sentinel (0.5, matching `DIR_ACC_EMA_SENTINEL`).
aux_dir_acc_reduce: CudaFunction,
/// SP13 Phase 0a (2026-05-04): mapped-pinned 3-element output
/// buffer for `aux_dir_acc_reduce`. Layout `[dir_acc,
/// pos_pred_frac, pos_label_frac]` per the kernel contract. Per
/// `feedback_no_htod_htoh_only_mapped_pinned.md`: every CPU↔GPU
/// scalar read-back goes through `MappedF32Buffer` (host writes
/// the launcher reads via `read_all()`'s volatile read after stream
/// sync). Constructor-zero-initialised; the kernel overwrites all
/// three slots on every launch (or all three to the sentinel 0.5
/// when the batch is empty / all-zero-label).
pub(crate) aux_dir_acc_buf: super::mapped_pinned::MappedF32Buffer,
/// SP13 Phase 0a P0a.T4 (2026-05-04): fixed-α EMA applicator
/// kernel handle. Loaded from `apply_fixed_alpha_ema_kernel.cubin`.
/// Consumed by `launch_apply_fixed_alpha_ema` for the SP13 dir-acc
/// EMAs (slots 373/374 with α=0.3 / 0.05 and sentinel 0.5) and the
/// hold-rate EMA (slot 382 with α=0.05 and sentinel 0.0). Sibling
/// of `apply_pearls_ad_kernel`; see field declaration there for the
/// Wiener-optimal counterpart. Fixed-α form needed because Pearls
/// A+D collapse two EMAs of the same signal to identical values,
/// defeating the stagnation detector that compares slot 373 vs 374.
apply_fixed_alpha_ema_kernel: CudaFunction,
/// SP13 Phase 0a P0a.T4 (2026-05-04): aux-head per-bar prediction
/// → ISV[AUX_DIR_PREDICTION_INDEX=375] tanh-bounded scalar producer
/// kernel handle. Loaded from `aux_pred_to_isv_tanh_kernel.cubin`.
/// Single-block tree-reduce reads `aux_nb_pred_buf [B]` and writes
/// `mean(tanh(aux_pred[i]))` ∈ [-1, +1] to ISV[375]. Per-step state
/// (not an EMA — slot is overwritten each launch); no FoldReset
/// registry entry. Consumed by `launch_aux_pred_to_isv_tanh`.
aux_pred_to_isv_tanh_kernel: CudaFunction,
/// SP11 Fix 39 (2026-05-04, Task A2): reward-subsystem controller
/// kernel. Single-block, 10-thread producer reading 5 canary ISV slots
/// [350..360) and writing 10 floats to scratch[SCRATCH_SP11_CONTROLLER_BASE..+10).
@@ -13044,6 +13153,248 @@ impl GpuDqnTrainer {
Ok(())
}
/// SP13 Phase 0a (2026-05-04): launch the aux-head directional-
/// accuracy reducer.
///
/// Reads `aux_pred [B]` (the per-bar aux next-bar regression
/// prediction tile, P0a regression mode) and `next_bar_label [B]`
/// (per-bar sign-encoded label: +1 / 0 / -1) and writes
/// `(dir_acc, pos_pred_frac, pos_label_frac)` into the 3-element
/// output buffer at `out_3_dev`. Both predictions and labels live
/// on the GPU; the 3-element output is the trainer's mapped-pinned
/// `aux_dir_acc_buf` in production (the launcher is invoked with
/// `aux_dir_acc_buf.dev_ptr` from `training_loop.rs` after the
/// captured forward graph populates `aux_nb_pred_buf` /
/// `aux_nb_label_buf`). The launcher takes raw `u64` device
/// pointers per Decision D / `feedback_no_htod_htoh_only_mapped_pinned.md`
/// because the production source for `out_3_dev` is a
/// `MappedF32Buffer` (which exposes only `dev_ptr: u64`).
///
/// Single-block tree-reduce (BLOCK_SIZE=256) — four shared-memory
/// int arrays (correct/pos_pred/pos_label/valid) reduce in lockstep
/// per `feedback_no_atomicadd.md`. The launcher computes
/// `shared_mem_bytes = 4 × bdim × sizeof(i32)` to match the kernel's
/// `extern __shared__ int shared[]` declaration; passing a smaller
/// value would make the kernel read past the dynamic shared-memory
/// region and corrupt other data.
///
/// Per `feedback_cudarc_f64_f32_abi.md`: `batch_size` is `i32`
/// (matches the kernel signature exactly — no implicit cast).
pub(crate) fn launch_aux_dir_acc_reduce(
&self,
aux_pred_dev: u64,
next_bar_label_dev: u64,
batch_size: i32,
out_3_dev: u64,
) -> Result<(), MLError> {
let bdim: u32 = 256;
let cfg = LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (bdim, 1, 1),
shared_mem_bytes: 4 * bdim * std::mem::size_of::<i32>() as u32,
};
unsafe {
self.stream
.launch_builder(&self.aux_dir_acc_reduce)
.arg(&aux_pred_dev)
.arg(&next_bar_label_dev)
.arg(&batch_size)
.arg(&out_3_dev)
.launch(cfg)
.map_err(|e| MLError::ModelError(format!("sp13 aux_dir_acc_reduce launch: {e}")))?;
}
Ok(())
}
/// SP13 Phase 0a P0a.T4 (2026-05-04): launch the fixed-α EMA
/// applicator.
///
/// One launch updates `n` consecutive ISV slots starting at
/// `isv_offset`, blending each with the matching `sample[i]` at
/// fixed `alpha`. Caller passes the registry sentinel for the
/// destination slots (`pearl_first_observation_bootstrap`):
/// - 0.5 for the SP13 dir-acc EMAs at slots 373 / 374 (random-
/// guessing baseline; sentinel means "no observation yet")
/// - 0.0 for the hold-rate EMA at slot 382 (empty-batch fallback)
///
/// Sibling of `launch_apply_pearls` — needed because the Wiener-
/// optimal blend collapses two EMAs of the same signal to identical
/// values, defeating the SP13 stagnation detector at slot 374
/// (compares against slot 373's faster EMA). The fixed-α form
/// preserves the timescale separation that makes "improving vs
/// stalled" diagnosable.
///
/// Stream-ordered with the producer that wrote `sample_dev`. The
/// producer kernel MUST issue `__threadfence_system()` before
/// returning so this kernel's per-thread `sample[i]` read is
/// visible.
pub(crate) fn launch_apply_fixed_alpha_ema(
&self,
sample_dev: u64,
n: i32,
isv_offset: i32,
alpha: f32,
sentinel: f32,
) -> Result<(), MLError> {
debug_assert!(n >= 0,
"launch_apply_fixed_alpha_ema: n must be non-negative, got {n}");
if n == 0 {
return Ok(());
}
debug_assert!(self.isv_signals_dev_ptr != 0,
"launch_apply_fixed_alpha_ema: isv_signals_dev_ptr must be allocated by constructor");
// 1D thread grid over the slot count. SP13 P0a always launches
// with n=1 (single dir-acc scalar fed into both EMAs in
// lockstep, hold-rate scalar into slot 382), but the kernel
// generalises to n-slot blocks for future use without an
// ABI change.
let bdim: u32 = (n as u32).min(256).max(1);
let grid: u32 = ((n as u32) + bdim - 1) / bdim;
let cfg = LaunchConfig {
grid_dim: (grid, 1, 1),
block_dim: (bdim, 1, 1),
shared_mem_bytes: 0,
};
let isv_dev = self.isv_signals_dev_ptr;
unsafe {
self.stream
.launch_builder(&self.apply_fixed_alpha_ema_kernel)
.arg(&sample_dev)
.arg(&n)
.arg(&isv_offset)
.arg(&alpha)
.arg(&sentinel)
.arg(&isv_dev)
.launch(cfg)
.map_err(|e| MLError::ModelError(format!("sp13 apply_fixed_alpha_ema launch: {e}")))?;
}
Ok(())
}
/// SP13 Phase 0a P0a.T4 (2026-05-04): launch the aux-head
/// per-bar prediction → ISV[AUX_DIR_PREDICTION_INDEX=375] tanh-
/// bounded scalar producer.
///
/// Reads `aux_pred [B]` (production source: `aux_nb_pred_buf`,
/// populated by the captured forward graph's
/// `aux_next_bar_forward`) and writes `mean(tanh(aux_pred[i]))`
/// ∈ [-1, +1] to the SHARED ISV slot at `isv_slot_offset`. Per-step
/// state, not an EMA — slot is overwritten each launch (no
/// FoldReset registry entry, per the SP13 P0a registry comments).
///
/// Single-block tree-reduce (BLOCK_SIZE=256) — one shared-memory
/// float array reduces in lockstep per `feedback_no_atomicadd.md`.
/// `shared_mem_bytes = bdim × sizeof(f32)`.
///
/// Stream-ordered with the producer that wrote `aux_pred_dev`;
/// same-stream contract mirrors `launch_aux_heads_loss_ema`.
pub(crate) fn launch_aux_pred_to_isv_tanh(
&self,
aux_pred_dev: u64,
batch_size: i32,
isv_slot_offset: i32,
) -> Result<(), MLError> {
debug_assert!(self.isv_signals_dev_ptr != 0,
"launch_aux_pred_to_isv_tanh: isv_signals_dev_ptr must be allocated by constructor");
let bdim: u32 = 256;
let cfg = LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (bdim, 1, 1),
shared_mem_bytes: bdim * std::mem::size_of::<f32>() as u32,
};
let isv_dev = self.isv_signals_dev_ptr;
unsafe {
self.stream
.launch_builder(&self.aux_pred_to_isv_tanh_kernel)
.arg(&aux_pred_dev)
.arg(&batch_size)
.arg(&isv_slot_offset)
.arg(&isv_dev)
.launch(cfg)
.map_err(|e| MLError::ModelError(format!("sp13 aux_pred_to_isv_tanh launch: {e}")))?;
}
Ok(())
}
/// SP13 Phase 0a P0a.T4 (2026-05-04): orchestrator for the per-step
/// aux directional-accuracy producer chain.
///
/// Single-call API for `training_loop.rs` to fire the SP13 P0a
/// per-step metrics:
/// 1. `launch_aux_dir_acc_reduce` — reduce aux_nb_pred_buf +
/// aux_nb_label_buf into `aux_dir_acc_buf [3]`
/// (dir_acc, pos_pred_frac, pos_label_frac).
/// 2. `launch_apply_fixed_alpha_ema` — short EMA (α=0.3) into
/// ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373], sentinel 0.5.
/// 3. `launch_apply_fixed_alpha_ema` — slow EMA (α=0.05) into
/// ISV[AUX_DIR_ACC_LONG_EMA_INDEX=374], sentinel 0.5.
/// 4. `launch_aux_pred_to_isv_tanh` — mean(tanh(aux_pred)) →
/// ISV[AUX_DIR_PREDICTION_INDEX=375].
///
/// All launches are stream-ordered; the producer's
/// `__threadfence_system()` orders against the consumer reads on
/// the same stream so no host sync is needed between steps.
/// Caller invokes once per training step BEFORE the HEALTH_DIAG
/// snapshot reads slot 373/374/375 (matches the post-cascade-fix
/// invariant from commit `a5f23b28f`).
///
/// Pre-condition: the captured forward graph has run this step,
/// so `aux_nb_pred_buf` and `aux_nb_label_buf` are populated.
/// Producer-only — no consumer kernel reads slots 373/374/375 in
/// this commit; Phase 0b wires the controller and the Q-head
/// input layer that reads slot 375.
pub fn launch_sp13_aux_dir_metrics(
&self,
) -> Result<(), MLError> {
use crate::cuda_pipeline::sp13_isv_slots::{
AUX_DIR_ACC_SHORT_EMA_INDEX, AUX_DIR_ACC_LONG_EMA_INDEX,
AUX_DIR_PREDICTION_INDEX,
DIR_ACC_EMA_SENTINEL,
};
let aux_pred_dev = self.aux_nb_pred_buf.raw_ptr();
let aux_label_dev = self.aux_nb_label_buf.raw_ptr();
let out_3_dev = self.aux_dir_acc_buf.dev_ptr;
let batch_size_i32 = self.config.batch_size as i32;
// 1. Reduce → mapped-pinned aux_dir_acc_buf [3]
self.launch_aux_dir_acc_reduce(
aux_pred_dev, aux_label_dev, batch_size_i32, out_3_dev,
)?;
// 2. Short EMA (α=0.3) — fast tracker for aux-w controller deficit term.
// n=1: only out_3[0] = dir_acc feeds the EMA; pos_pred/pos_label slots
// are diagnostic-only and don't have ISV destinations.
self.launch_apply_fixed_alpha_ema(
out_3_dev,
1,
AUX_DIR_ACC_SHORT_EMA_INDEX as i32,
0.3,
DIR_ACC_EMA_SENTINEL,
)?;
// 3. Slow EMA (α=0.05) — stagnation detector comparator.
self.launch_apply_fixed_alpha_ema(
out_3_dev,
1,
AUX_DIR_ACC_LONG_EMA_INDEX as i32,
0.05,
DIR_ACC_EMA_SENTINEL,
)?;
// 4. mean(tanh(aux_pred)) → ISV[375]. SHARED scalar (per-step
// overwrite, no EMA — see kernel header for layout rationale).
self.launch_aux_pred_to_isv_tanh(
aux_pred_dev,
batch_size_i32,
AUX_DIR_PREDICTION_INDEX as i32,
)?;
Ok(())
}
/// SP11 Fix 39 (2026-05-04, Task A2): launch the reward-subsystem
/// controller producer + chained Pearls A+D output smoothing.
///
@@ -16487,6 +16838,67 @@ impl GpuDqnTrainer {
module.load_function("popart_component_ema_kernel")
.map_err(|e| MLError::ModelError(format!("sp11 popart_component_ema load: {e}")))?
};
// SP13 Phase 0a (2026-05-04): aux-head directional-accuracy
// reducer kernel + 3-element mapped-pinned readback buffer.
// Producer reads the aux next-bar regression-head prediction
// tile + the per-bar sign-encoded next_bar_label tile and
// writes (dir_acc, pos_pred_frac, pos_label_frac) to the
// 3-float mapped-pinned buffer. P0a.T4 wires the per-step
// launch + ISV[373..375) EMA producer; this constructor edge
// is the cubin/kernel handle + buffer alloc — both atomic with
// the rest of the P0a chain per `feedback_no_partial_refactor.md`.
let aux_dir_acc_reduce = {
let module = stream.context()
.load_cubin(SP13_AUX_DIR_ACC_REDUCE_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!("sp13 aux_dir_acc_reduce cubin load: {e}")))?;
module.load_function("aux_dir_acc_reduce_kernel")
.map_err(|e| MLError::ModelError(format!("sp13 aux_dir_acc_reduce load: {e}")))?
};
let aux_dir_acc_buf = unsafe {
super::mapped_pinned::MappedF32Buffer::new(3)
}.map_err(|e| MLError::ModelError(
format!("SP13 aux_dir_acc_buf alloc (3 f32): {e}")
))?;
// SP13 v3 P0a.T3 (2026-05-04): the per-step Hold-rate observer
// chain lives on the `GpuExperienceCollector` stream — the
// observer reads `batch_actions [N]` (collector-owned) and the
// chained fixed-α EMA writes ISV[382]. The collector loads its
// own copy of the cubins from `SP13_HOLD_RATE_OBSERVER_CUBIN` and
// `SP13_APPLY_FIXED_ALPHA_EMA_CUBIN`, so the trainer carries no
// hold-rate-specific kernel handle or scratch buffer (avoids a
// dead `CudaFunction` field per `feedback_no_stubs.md`). The
// trainer's host-side controller in `training_loop.rs:3604`
// reads the slot-382 EMA + slot-381 target via
// `read_isv_signal_at` at the per-epoch metrics block and
// writes the priced cost to ISV[HOLD_COST_INDEX=380]; the
// reward-composition site in `experience_kernels.cu` consumes
// ISV[380] from the same shared `isv_signals_dev_ptr`. The
// ISV[380] / ISV[381] / ISV[382] sentinels are seeded by
// `state_reset_registry`'s SP13 entries and `seed_static_signals`
// (Invariant-1 anchors at ISV[380, 381]).
// SP13 Phase 0a P0a.T4 (2026-05-04): fixed-α EMA applicator
// kernel handle + aux-pred → ISV[375] tanh producer handle.
// Both load from their cubins and stay live for the trainer's
// lifetime. P0a.T4 wires the per-step launches in
// `launch_sp13_aux_dir_metrics` and the hold-rate observer
// chain in `gpu_experience_collector.rs`; this constructor edge
// is the cubin/kernel handle, all-atomic with the rest of the
// P0a chain per `feedback_no_partial_refactor.md`.
let apply_fixed_alpha_ema_kernel = {
let module = stream.context()
.load_cubin(SP13_APPLY_FIXED_ALPHA_EMA_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!("sp13 apply_fixed_alpha_ema cubin load: {e}")))?;
module.load_function("apply_fixed_alpha_ema_kernel")
.map_err(|e| MLError::ModelError(format!("sp13 apply_fixed_alpha_ema_kernel load: {e}")))?
};
let aux_pred_to_isv_tanh_kernel = {
let module = stream.context()
.load_cubin(SP13_AUX_PRED_TO_ISV_TANH_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!("sp13 aux_pred_to_isv_tanh cubin load: {e}")))?;
module.load_function("aux_pred_to_isv_tanh_kernel")
.map_err(|e| MLError::ModelError(format!("sp13 aux_pred_to_isv_tanh_kernel load: {e}")))?
};
// SP11 Fix 39 B1b fix-up (2026-05-04): allocate trainer-side
// popart-component-per-sample mapped-pinned placeholder buffer.
// The canonical per-bar buffer is owned by the experience
@@ -18766,6 +19178,39 @@ impl GpuDqnTrainer {
*sig_ptr.add(KELLY_DIVERGENCE_TARGET_INDEX) = 2.0_f32;
*sig_ptr.add(KELLY_TEMPORAL_TARGET_INDEX) = 5.0_f32;
// SP13 directional-skill Invariant-1 anchors. Constants, not
// stateful EMAs — must never reach 0 sentinel between folds
// (FoldReset rewrites them in `reset_named_state`).
use crate::cuda_pipeline::sp5_isv_slots::{
TARGET_DIR_ACC_INDEX, TARGET_DIR_ACC_DEFAULT,
DIR_SKILL_BONUS_ALPHA_INDEX, DIR_SKILL_BONUS_ALPHA_DEFAULT,
DIR_SKILL_BONUS_BETA_INDEX, DIR_SKILL_BONUS_BETA_DEFAULT,
LUCK_WIN_DISCOUNT_INDEX, LUCK_WIN_DISCOUNT_DEFAULT,
SKILL_BONUS_CAP_RATIO_INDEX, SKILL_BONUS_CAP_RATIO_DEFAULT,
};
*sig_ptr.add(TARGET_DIR_ACC_INDEX) = TARGET_DIR_ACC_DEFAULT;
*sig_ptr.add(DIR_SKILL_BONUS_ALPHA_INDEX) = DIR_SKILL_BONUS_ALPHA_DEFAULT;
*sig_ptr.add(DIR_SKILL_BONUS_BETA_INDEX) = DIR_SKILL_BONUS_BETA_DEFAULT;
*sig_ptr.add(LUCK_WIN_DISCOUNT_INDEX) = LUCK_WIN_DISCOUNT_DEFAULT;
*sig_ptr.add(SKILL_BONUS_CAP_RATIO_INDEX) = SKILL_BONUS_CAP_RATIO_DEFAULT;
// SP13 v3 P0a.T3 (2026-05-04): Hold-pricing controller anchors.
// HOLD_COST_INDEX is rewritten every step by the host-side
// controller in `training_loop.rs`; the constructor write
// seeds it with the baseline cost so the first
// `experience_env_step` (before the controller fires) sees
// a meaningful value rather than 0 (which would silently
// disable Hold-pricing on the first batch). HOLD_RATE_TARGET
// is a static MFT default 0.20 — never reaches sentinel 0
// between folds. The observed-rate EMA at slot 382 is
// FoldReset (sentinel 0.0) and registered separately.
use crate::cuda_pipeline::sp5_isv_slots::{
HOLD_COST_INDEX, HOLD_COST_BASE,
HOLD_RATE_TARGET_INDEX, HOLD_RATE_TARGET_DEFAULT,
};
*sig_ptr.add(HOLD_COST_INDEX) = HOLD_COST_BASE;
*sig_ptr.add(HOLD_RATE_TARGET_INDEX) = HOLD_RATE_TARGET_DEFAULT;
// Layout fingerprint (ISV[58..60)). Compile-time structural hash of
// the slot layout; checkpoint load fails-fast on mismatch.
// Stored as a u64 split across two f32 lanes using raw bit-cast so
@@ -19913,6 +20358,20 @@ impl GpuDqnTrainer {
sp11_popart_component_ema_kernel,
sp11_popart_component_dev_ptr: 0,
sp11_popart_component_len: 0,
// SP13 Phase 0a (2026-05-04): aux-head directional-accuracy
// reducer kernel + 3-element mapped-pinned readback buffer.
// Both fields atomic with the SP13 ISV slot constants and
// SP13 state-reset registry entries staged by P0a.T1; the
// per-step launcher is wired by P0a.T4.
aux_dir_acc_reduce,
aux_dir_acc_buf,
// SP13 Phase 0a P0a.T4 (2026-05-04): fixed-α EMA applicator +
// aux-pred → ISV[375] tanh producer kernel handles. Wired into
// `launch_sp13_aux_dir_metrics` (called per step from
// `training_loop.rs`) and the hold-rate observer chain in
// `gpu_experience_collector.rs` per-step action-select site.
apply_fixed_alpha_ema_kernel,
aux_pred_to_isv_tanh_kernel,
// SP11 Fix 39 (Task A2): controller + SimHash novelty kernels
// and their backing buffers. Projection matrix was populated
// on-device above by `launch_novelty_simhash_proj_init`; hash

141
crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
View File

@@ -833,6 +833,34 @@ pub struct GpuExperienceCollector {
/// `feedback_no_cpu_compute_strict.md`.
apply_pearls_ad_kernel: CudaFunction,
/// SP13 v3 Phase 0a P0a.T4 (2026-05-04): Hold-rate observer kernel
/// handle on the collector's stream. Reads the per-bar PACKED
/// factored `batch_actions [B]` tile (written by
/// `experience_action_select` to `out_actions[i] = dir*M*O*U + ...`)
/// and writes the Hold-pick fraction into `hold_rate_buf [1]`.
/// Same cubin as the trainer's handle (`SP13_HOLD_RATE_OBSERVER_CUBIN`)
/// — loaded on this stream so the per-step launch chains directly
/// after `experience_action_select` without a cross-stream sync.
sp13_hold_rate_observer_kernel: CudaFunction,
/// SP13 Phase 0a P0a.T4 (2026-05-04): fixed-α EMA applicator kernel
/// handle on the collector's stream. Reads `hold_rate_buf [1]` and
/// blends into ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382] with α=0.05
/// and sentinel 0.0 (matching the registry entry
/// `sp13_hold_rate_observed_ema`). Same cubin as the trainer's
/// `apply_fixed_alpha_ema_kernel` field — Pearls A+D would be the
/// wrong tool here because the spec needs the slow-EMA timescale
/// (α=0.05) for the host-side controller's gain math. Loaded on
/// this stream so the per-step EMA fires directly after the
/// observer launch with no cross-stream sync.
sp13_apply_fixed_alpha_ema_kernel: CudaFunction,
/// SP13 v3 Phase 0a P0a.T4 (2026-05-04): mapped-pinned 1-element
/// scratch for the Hold-rate observer's per-step output. Per
/// `feedback_no_htod_htoh_only_mapped_pinned.md`: kernel writes via
/// `dev_ptr` (with `__threadfence_system()`); the chained fixed-α
/// EMA reads via the same `dev_ptr`. Constructor-zero-initialised;
/// the kernel overwrites it on every launch.
sp13_hold_rate_buf: MappedF32Buffer,
/// B.2 Plan 3 Task 3: trade_attempt_rate_ema_update GPU kernel.
/// Single-block (1 thread) reduction + adaptive EMA of Flat→Positioned
/// transition rate into ISV[71]. Launched alongside reward_component_ema
@@ -1516,6 +1544,39 @@ impl GpuExperienceCollector {
.map_err(|e| MLError::ModelError(format!("apply_pearls_ad_kernel load: {e}")))?
};
// SP13 v3 Phase 0a P0a.T4 (2026-05-04): Hold-rate observer +
// fixed-α EMA applicator on the collector's stream. The observer
// chains directly after `experience_action_select` (which writes
// `batch_actions [B]` packed factored action_idx) so the kernels
// MUST live on the same stream — cross-stream sync would defeat
// the no-host-sync per-step contract. Same cubins as the trainer's
// handles; loaded twice (once per stream) per the cudarc Function/
// Module pattern. The 1-element mapped-pinned `sp13_hold_rate_buf`
// is owned by the collector because it's a per-step transient
// scratch and the EMA writeback is into the trainer's ISV (read
// via `isv_signals_dev_ptr` set after construction by
// `set_isv_signals_ptr`).
let sp13_hold_rate_observer_kernel = {
use super::gpu_dqn_trainer::SP13_HOLD_RATE_OBSERVER_CUBIN;
let m = stream.context()
.load_cubin(SP13_HOLD_RATE_OBSERVER_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!("sp13 hold_rate_observer cubin (collector): {e}")))?;
m.load_function("hold_rate_observer_kernel")
.map_err(|e| MLError::ModelError(format!("sp13 hold_rate_observer_kernel load (collector): {e}")))?
};
let sp13_apply_fixed_alpha_ema_kernel = {
use super::gpu_dqn_trainer::SP13_APPLY_FIXED_ALPHA_EMA_CUBIN;
let m = stream.context()
.load_cubin(SP13_APPLY_FIXED_ALPHA_EMA_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!("sp13 apply_fixed_alpha_ema cubin (collector): {e}")))?;
m.load_function("apply_fixed_alpha_ema_kernel")
.map_err(|e| MLError::ModelError(format!("sp13 apply_fixed_alpha_ema_kernel load (collector): {e}")))?
};
let sp13_hold_rate_buf = unsafe { MappedF32Buffer::new(1) }
.map_err(|e| MLError::ModelError(
format!("SP13 v3 sp13_hold_rate_buf alloc (1 f32, collector): {e}")
))?;
// B.2 Plan 3 Task 3: load trade_attempt_rate_ema_update kernel.
let trade_attempt_rate_ema_kernel = {
use super::gpu_dqn_trainer::TRADE_RATE_EMA_CUBIN;
@@ -1798,6 +1859,13 @@ impl GpuExperienceCollector {
bn_tanh_concat_fn,
reward_component_ema_kernel,
apply_pearls_ad_kernel,
// SP13 v3 P0a.T4 (2026-05-04): Hold-rate observer + fixed-α
// EMA applicator + 1-elem mapped-pinned scratch — chained
// after `experience_action_select` per step in
// `collect_experiences_gpu`.
sp13_hold_rate_observer_kernel,
sp13_apply_fixed_alpha_ema_kernel,
sp13_hold_rate_buf,
trade_attempt_rate_ema_kernel,
plan_threshold_update_kernel,
state_kl_kernel,
@@ -3992,6 +4060,79 @@ impl GpuExperienceCollector {
}
}
// ── 4c. SP13 v3 Phase 0a P0a.T4 (2026-05-04): Hold-rate
// observer + fixed-α EMA into ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382].
//
// Per-step on the SAME stream as `experience_action_select` (and
// the optional expert override above) so the observer reads
// `batch_actions [N]` AFTER both producers have written. Single-
// block tree-reduce → 1-elem mapped-pinned `sp13_hold_rate_buf`,
// then a 1-thread fixed-α (α=0.05, sentinel 0.0) EMA blend into
// the trainer's ISV via `isv_signals_dev_ptr` (set after
// construction by `set_isv_signals_ptr`). Skip both launches
// when ISV is not yet wired (test scaffold path) — the EMA
// would have nowhere to write.
//
// The trainer's host-side controller in `training_loop.rs:3604`
// reads slot 382 + slot 381 at the per-epoch metrics block and
// writes the priced cost to ISV[HOLD_COST_INDEX=380]; the
// reward-composition site in `experience_kernels.cu` then
// subtracts the cost on every Hold-action bar — pricing the
// action so the policy uses Hold deliberately rather than as a
// free CQL-bias-anchored default.
if self.isv_signals_dev_ptr != 0 && !self.seed_phase_active_cache {
use crate::cuda_pipeline::sp13_isv_slots::HOLD_RATE_OBSERVED_EMA_INDEX;
let actions_dev = self.batch_actions.raw_ptr();
let out_dev = self.sp13_hold_rate_buf.dev_ptr;
let bdim: u32 = 256;
// Step 1: count(decoded_dir == DIR_HOLD) / N → sp13_hold_rate_buf[0].
unsafe {
self.stream
.launch_builder(&self.sp13_hold_rate_observer_kernel)
.arg(&actions_dev)
.arg(&n_i32)
.arg(&out_dev)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (bdim, 1, 1),
shared_mem_bytes: bdim * std::mem::size_of::<i32>() as u32,
})
.map_err(|e| MLError::ModelError(format!(
"sp13 v3 hold_rate_observer t={t}: {e}"
)))?;
}
// Step 2: fixed-α EMA into ISV[382]. n=1 (single scalar);
// sentinel 0.0 matches the registry entry for slot 382;
// α=0.05 matches the spec's slow-EMA timescale for the
// host-side Hold-cost controller.
let n_slots: i32 = 1;
let isv_off: i32 = HOLD_RATE_OBSERVED_EMA_INDEX as i32;
let alpha: f32 = 0.05;
let sentinel: f32 = 0.0;
let isv_dev = self.isv_signals_dev_ptr;
unsafe {
self.stream
.launch_builder(&self.sp13_apply_fixed_alpha_ema_kernel)
.arg(&out_dev)
.arg(&n_slots)
.arg(&isv_off)
.arg(&alpha)
.arg(&sentinel)
.arg(&isv_dev)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (1, 1, 1),
shared_mem_bytes: 0,
})
.map_err(|e| MLError::ModelError(format!(
"sp13 v3 hold_rate apply_fixed_alpha_ema t={t}: {e}"
)))?;
}
}
// ── 5. Environment step (reward v5: trade-aware hybrid) ──────
// max_pos already defined above (action_select block)
// D4/N4: adaptive counterfactual ratio — 0.5 at healthy, 0.8 at collapse.

116
crates/ml/src/cuda_pipeline/hold_rate_observer_kernel.cu Normal file
View File

@@ -0,0 +1,116 @@
// crates/ml/src/cuda_pipeline/hold_rate_observer_kernel.cu
//
// SP13 Phase 0a v3 (2026-05-04): Hold-rate observer.
//
// Single-block tree-reduce kernel that counts how many bars in a batch
// picked the Hold direction action and writes the resulting fraction
// `count(Hold) / batch_size ∈ [0, 1]` to `out_1[0]`. Pearls A+D then
// smooth this per-step observation into
// `ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382]`; the host-side controller in
// `training_loop.rs` reads the EMA + the static target at
// `ISV[HOLD_RATE_TARGET_INDEX=381]` and writes the priced Hold cost back
// to `ISV[HOLD_COST_INDEX=380]`. The reward-composition site in
// `experience_kernels.cu` then subtracts the cost on every Hold-action
// bar — pricing the action so the policy uses Hold deliberately rather
// than as a free CQL-bias-anchored default.
//
// P0a.T4 (2026-05-04): the kernel reads the production PACKED factored
// `batch_actions [B]` tile (`int* action_idx`) and decodes direction
// inline as `dir = action_idx / (NUM_MAGNITUDES * NUM_ORD * NUM_URG)`.
// The earlier T3-staged signature took an unpacked `actions_dir [B]`
// tile that production never materialises; `experience_action_select`
// (the upstream producer) writes only the packed factored form to
// `out_actions[i] = dir*b1*b2*b3 + mag*b2*b3 + ord*b3 + urg`. Per
// `feedback_no_partial_refactor.md`, every consumer of the packed
// contract migrates atomically — this kernel + the T3 oracle tests are
// updated in lockstep with the launcher. Bucket sizes
// (`NUM_MAGNITUDES`, `NUM_ORD`, `NUM_URG`) live in `state_layout.cuh`
// so the kernel doesn't need an extra launcher arg.
//
// Sign convention:
// - decoded direction `dir == DIR_HOLD` (=1) → contributes 1 to count
// - any other direction (Short=0 / Long=2 / Flat=3) → contributes 0
// - empty batch (`batch_size <= 0`) → out = 0.0 (Pearl A's first-
// observation replacement consumes 0.0 cleanly because the
// companion sentinel for slot 382 is also 0.0; see registry entry
// `sp13_hold_rate_observed_ema`).
//
// Single block, BLOCK_SIZE=256, one shared-memory int array — no
// atomicAdd per `feedback_no_atomicadd.md`. Pure GPU compute per
// `feedback_no_cpu_compute_strict.md`. Mirrors the `aux_dir_acc_reduce`
// kernel's reducer shape (T2, 2026-05-04) at smaller surface — this
// kernel only needs ONE counter (Hold-count), the dir-acc kernel needed
// four (correct/pos_pred/pos_label/valid).
//
// `state_layout.cuh` provides the `DIR_HOLD` constant (==1) shared by
// every kernel that includes the header — the same constant
// `experience_action_select` writes when the policy picks Hold and the
// reward-composition site reads to gate the per-bar Hold-cost subtraction.
#include <cuda_runtime.h>
#include "state_layout.cuh"
#define BLOCK_SIZE 256
extern "C" __global__ void hold_rate_observer_kernel(
/* Per-bar PACKED factored action tile [B], values in
* `[0, NUM_DIRECTIONS * NUM_MAGNITUDES * NUM_ORD * NUM_URG)` =
* `[0, 4 * 3 * 3 * 3) = [0, 108)` for the production branch sizes.
* Encoding: `action_idx = dir*M*O*U + mag*O*U + ord*U + urg` —
* matches `experience_action_select`'s `out_actions[i]` write at
* `experience_kernels.cu:1453`. The launcher passes the post-
* action-selection `batch_actions` buffer; the kernel strides each
* thread over `[0, batch_size)` so any batch_size up to ~2^31 is
* valid. */
const int* __restrict__ batch_actions,
/* Number of bars in the batch. */
int batch_size,
/* 1-element output buffer. Mapped-pinned in production (see the
* launcher in `gpu_dqn_trainer.rs`). The kernel writes
* out_1[0] = count(decoded_dir == DIR_HOLD) / batch_size
* in thread 0 after the tree-reduce completes. */
float* __restrict__ out_1)
{
/* Single-block reducer — guard against accidental multi-block launch. */
if (blockIdx.x != 0) return;
const int tid = threadIdx.x;
const int bdim = blockDim.x;
/* Dynamic shared memory: one int array of `bdim` slots. The launcher
* passes `shared_mem_bytes = bdim × sizeof(int)`. */
extern __shared__ int sh_hold[];
/* Per-thread strided accumulation over the batch. Decode the packed
* factored action_idx and check direction; the divisor matches the
* authoritative encoding in `experience_action_select`. */
const int dir_divisor = NUM_MAGNITUDES * NUM_ORD * NUM_URG;
int local_hold = 0;
for (int i = tid; i < batch_size; i += bdim) {
const int action_idx = batch_actions[i];
const int dir = action_idx / dir_divisor;
if (dir == DIR_HOLD) local_hold += 1;
}
sh_hold[tid] = local_hold;
__syncthreads();
/* Standard log2(BLOCK_SIZE) tree reduction. One `__syncthreads()` per
* pass keeps every thread's view consistent. */
for (int s = bdim / 2; s > 0; s >>= 1) {
if (tid < s) {
sh_hold[tid] += sh_hold[tid + s];
}
__syncthreads();
}
/* Thread 0 finalises. Empty batch → 0.0 sentinel matches the slot-382
* fold-reset sentinel (see `sp13_hold_rate_observed_ema` registry
* entry); Pearl A's first-observation replacement consumes 0.0 cleanly. */
if (tid == 0) {
out_1[0] = (batch_size > 0)
? ((float)sh_hold[0] / (float)batch_size)
: 0.0f;
__threadfence_system();
}
}

1
crates/ml/src/cuda_pipeline/mod.rs
View File

@@ -61,6 +61,7 @@ pub mod meta_q_network;
pub mod sp4_isv_slots;
pub mod sp5_isv_slots;
pub mod sp11_isv_slots;
pub mod sp13_isv_slots;
pub use sp4_isv_slots::{
TARGET_Q_BOUND_INDEX, ATOM_POS_BOUND_BASE, WEIGHT_BOUND_BASE,
ADAM_M_BOUND_BASE, ADAM_V_BOUND_BASE, WD_RATE_BASE,

76
crates/ml/src/cuda_pipeline/sp13_isv_slots.rs Normal file
View File

@@ -0,0 +1,76 @@
//! SP13 — Redefine success for predictive skill ISV slot constants.
//!
//! Slots [372..383) populated by the SP13 chain:
//! - 372: directional accuracy target (static, default 0.55)
//! - 373: aux dir-acc fast EMA (per-fold, sentinel 0.5)
//! - 374: aux dir-acc slow EMA (per-fold, sentinel 0.5) — stagnation detector
//! - 375: aux head per-bar prediction (per-bar overwrite, [-1, +1])
//! - 376: dir-skill bonus α (per-fold, default 1.0) — magnitude on correct calls
//! - 377: dir-skill bonus β (per-fold, default 1.0) — penalty on wrong calls
//! - 378: lucky-win discount factor (per-fold, default 0.3)
//! - 379: skill-bonus cap ratio (per-fold, default 0.3) — bonus ≤ ratio × |α|
//! - 380: per-bar Hold cost (controller output, sentinel = HOLD_COST_BASE)
//! - 381: target Hold-pick rate (static, default 0.20 — MFT-ish)
//! - 382: observed Hold-pick rate EMA (per-fold, sentinel 0.0)
//!
//! Slots 367..372 are an intentional gap separating the SP11 controller block
//! (ends at 367) from this SP13 block. The gap matches the codebase pattern of
//! reserving range markers for future SP-block boundaries.
//!
//! v3 Hold-pricing (slots 380..383, 2026-05-04): replaces v2's atomic Hold
//! elimination after the T3 implementer's audit found `DirectionAction` doesn't
//! exist in the codebase (8-variant fused `ExposureLevel` with cross-crate
//! consumers). Instead of removing Hold, the v3 chain prices it via an ISV-
//! driven adaptive cost controller targeting a Hold-rate (~20% — MFT-ish
//! deliberate use). When the model abuses Hold (rate > target), cost rises
//! until Hold becomes uneconomic; when rate ≤ target, cost relaxes to baseline
//! infrastructure carry. Per `feedback_isv_for_adaptive_bounds`: base cost is
//! the only constant; rate target, observed rate, and cost output are ISV-
//! driven. Spec §"Change 1: Price Hold (replaces v2's Eliminate Hold)".
//!
//! Spec: docs/superpowers/specs/2026-05-04-sp13-redefine-success-for-predictive-skill.md
pub const TARGET_DIR_ACC_INDEX: usize = 372;
pub const AUX_DIR_ACC_SHORT_EMA_INDEX: usize = 373;
pub const AUX_DIR_ACC_LONG_EMA_INDEX: usize = 374;
pub const AUX_DIR_PREDICTION_INDEX: usize = 375;
pub const DIR_SKILL_BONUS_ALPHA_INDEX: usize = 376;
pub const DIR_SKILL_BONUS_BETA_INDEX: usize = 377;
pub const LUCK_WIN_DISCOUNT_INDEX: usize = 378;
pub const SKILL_BONUS_CAP_RATIO_INDEX: usize = 379;
// ── v3 Hold-pricing controller (slots 380..383, 2026-05-04) ──────────
pub const HOLD_COST_INDEX: usize = 380;
pub const HOLD_RATE_TARGET_INDEX: usize = 381;
pub const HOLD_RATE_OBSERVED_EMA_INDEX: usize = 382;
/// Static-init defaults written ONCE at trainer construction (and re-written
/// at fold boundary by the constructor path — these are Invariant-1 numerical
/// anchors, not stateful EMAs).
pub const TARGET_DIR_ACC_DEFAULT: f32 = 0.55;
pub const DIR_SKILL_BONUS_ALPHA_DEFAULT: f32 = 1.0;
pub const DIR_SKILL_BONUS_BETA_DEFAULT: f32 = 1.0;
pub const LUCK_WIN_DISCOUNT_DEFAULT: f32 = 0.3;
pub const SKILL_BONUS_CAP_RATIO_DEFAULT: f32 = 0.3;
/// v3 Hold-pricing controller anchors (Invariant-1 numerical constants per
/// `feedback_isv_for_adaptive_bounds.md`). The controller scales `HOLD_COST_BASE`
/// by `1 + GAIN × max(0, observed − target)` clamped to
/// `[FLOOR_RATIO × BASE, CEIL_RATIO × BASE]`. Base cost is small per-bar
/// (~10 ticks of price for ES.FUT at ~0.25 tick value) but cumulative across
/// long Hold runs becomes meaningful relative to the SP12 capped trade reward
/// (±5..10).
pub const HOLD_COST_BASE: f32 = 0.001;
pub const HOLD_RATE_TARGET_DEFAULT: f32 = 0.20;
pub const HOLD_COST_CONTROLLER_GAIN: f32 = 5.0;
pub const HOLD_COST_FLOOR_RATIO: f32 = 0.5; // never below 0.5 × base
pub const HOLD_COST_CEIL_RATIO: f32 = 5.0; // never above 5.0 × base
/// Sentinel for both dir-acc EMAs at fold boundary: 0.5 is the random-guessing
/// baseline for binary directional accuracy. First observation replaces this
/// directly per `pearl_first_observation_bootstrap` so the EMA tracks the new
/// fold's signal without bias from the previous fold's terminal accuracy.
pub const DIR_ACC_EMA_SENTINEL: f32 = 0.5;
pub const SP13_SLOT_BASE: usize = 372;
pub const SP13_SLOT_END: usize = 383;

120
crates/ml/src/cuda_pipeline/sp5_isv_slots.rs
View File

@@ -21,14 +21,18 @@
//! 340..367 SP11 (Fix 39) reward-subsystem controller (27 slots, fold-reset) — see sp11_isv_slots
//! [includes B1b fix-up slot 360 = POPART_COMPONENT_MAG_EMA_INDEX
//! and B1b smoke-recovery 6 var EMAs at 361..367 for z-score normalization]
//! 367..372 Intentional 5-slot boundary gap separating SP11 controller block from SP13
//! 372..380 SP13 (2026-05-04) directional-skill instrumentation (8 slots) — see sp13_isv_slots
//! 380..383 SP13 v3 (2026-05-04) Hold-pricing controller (3 slots) — see sp13_isv_slots
//!
//! Total: 191 new SP5 ISV slots (52 + 24 + 20 + 4 + 4 + 6 + 4 + 4 + 3 + 16 + 8 + 9 + 9 + 1 + 27).
//! Total: 202 new SP5 ISV slots (52 + 24 + 20 + 4 + 4 + 6 + 4 + 4 + 3 + 16 + 8 + 9 + 9 + 1 + 27 + 8 + 3).
// Re-export SP11 slot constants so existing consumers can import every
// Re-export SP11 + SP13 slot constants so existing consumers can import every
// slot constant from a single location (`crate::cuda_pipeline::sp5_isv_slots::*`).
// SP11 owns its file but participates in the SP5 linear-span window
// (SP5_SLOT_BASE..SP5_SLOT_END covers all SP5/SP7/SP8/SP9/SP10/SP11 slots).
// SP11 / SP13 own their files but participate in the SP5 linear-span window
// (SP5_SLOT_BASE..SP5_SLOT_END covers all SP5/SP7/SP8/SP9/SP10/SP11/SP13 slots).
pub use crate::cuda_pipeline::sp11_isv_slots::*;
pub use crate::cuda_pipeline::sp13_isv_slots::*;
pub const SP5_SLOT_BASE: usize = 174;
@@ -304,7 +308,27 @@ pub const EVAL_THOMPSON_TEMP_INDEX: usize = 339; // [1] eval Thompson se
// handled by the existing bulk memset of `wiener_state_buf` covered by the
// `sp4_wiener_state` registry entry's dispatch arm.
pub const SP5_SLOT_END: usize = 367;
// ── SP13 (2026-05-04): redefine success for predictive skill ──────────
//
// 8 new SP13 ISV slots at ISV[372..380) for the directional-accuracy
// instrumentation + reward-composition controllers. Slots 367..372 are an
// intentional 5-slot gap separating the SP11 controller block (ends at 367)
// from this SP13 block — a structural boundary marker, not allocated. Per
// `pearl_controller_anchors_isv_driven.md` + `feedback_isv_for_adaptive_bounds`:
// every adaptive bound lives on ISV; the static defaults at slots 372 and
// 376..379 are constructor-written Invariant-1 anchors (not stateful EMAs)
// and will be lifted to controllers in Layer C if Phase 0a confirms the
// hypothesis. All 8 slot constants live in `crate::cuda_pipeline::sp13_isv_slots`
// (re-exported above).
//
// SP13 v3 (P0a.T3, 2026-05-04): 3 additional slots at ISV[380..383) for the
// Hold-pricing adaptive controller — `HOLD_COST_INDEX=380` (controller output,
// per-bar Hold cost), `HOLD_RATE_TARGET_INDEX=381` (static MFT target 0.20),
// `HOLD_RATE_OBSERVED_EMA_INDEX=382` (per-fold-reset Hold-pick rate EMA from
// `hold_rate_observer_kernel`). Replaces v2's atomic Hold elimination after
// the T3 implementer audited and found `DirectionAction` doesn't exist — the
// 4-way action space stays; Hold becomes priced rather than removed.
pub const SP5_SLOT_END: usize = 383;
/// Wiener-buffer producer-count constant. Sizes `wiener_state_buf` via
/// `(SP4_PRODUCER_COUNT + SP5_PRODUCER_COUNT) * SP4_WIENER_FLOATS_PER_SLOT`.
@@ -402,9 +426,9 @@ pub const SP5_SLOT_END: usize = 367;
/// FoldReset (sentinel 0; Pearl A bootstrap on first observation).
/// Constants for the SP11 block live in `crate::cuda_pipeline::sp11_isv_slots`
/// and are re-exported above.
pub const SP5_PRODUCER_COUNT: usize = 193;
// linear span = SP5_SLOT_END - SP5_SLOT_BASE = 367 - 174 = 193 wiener triples
// unique-slot count = 191 (52 per-branch + 24 Adam + 20 IQN τ + 4 trail
pub const SP5_PRODUCER_COUNT: usize = 209;
// linear span = SP5_SLOT_END - SP5_SLOT_BASE = 383 - 174 = 209 wiener triples
// unique-slot count = 202 (52 per-branch + 24 Adam + 20 IQN τ + 4 trail
// + 4 num_atoms + 6 Kelly + 4 Layer D D1 PnL aggregation
// + 4 Layer D D2 health composition
// + 3 Layer D D3 training metrics EMA
@@ -415,7 +439,11 @@ pub const SP5_PRODUCER_COUNT: usize = 193;
// + 9 SP9 Kelly cold-start warmup floor + targets + eval_dist
// + 1 SP10 eval Thompson temperature
// + 27 SP11 reward-subsystem controller [includes B1b fix-up
// slot 360 + B1b smoke-recovery 6 per-component var EMAs 361..367])
// slot 360 + B1b smoke-recovery 6 per-component var EMAs 361..367]
// + 8 SP13 directional-skill instrumentation [372..380); slots
// 367..372 are an intentional 5-slot gap, reserved-unused
// in both ISV and the wiener buffer
// + 3 SP13 v3 Hold-pricing controller [380..383))
// ── Convenience accessors ────────────────────────────────────────────
#[inline] pub const fn atom_v_center(b: usize) -> usize { ATOM_V_CENTER_BASE + b }
@@ -480,7 +508,11 @@ pub const SP5_LAYOUT_FINGERPRINT_FRAGMENT: &str =
SABOTEUR_ENGAGEMENT_RATE=358;PNL_REWARD_MAGNITUDE_EMA=359;\
POPART_COMPONENT_MAG_EMA=360;\
REWARD_COMPONENT_VAR_EMA_BASE=361;\
ISV_TOTAL_DIM=367";
TARGET_DIR_ACC=372;AUX_DIR_ACC_SHORT_EMA=373;AUX_DIR_ACC_LONG_EMA=374;\
AUX_DIR_PREDICTION=375;DIR_SKILL_BONUS_ALPHA=376;DIR_SKILL_BONUS_BETA=377;\
LUCK_WIN_DISCOUNT=378;SKILL_BONUS_CAP_RATIO=379;\
HOLD_COST=380;HOLD_RATE_TARGET=381;HOLD_RATE_OBSERVED_EMA=382;\
ISV_TOTAL_DIM=383";
#[cfg(test)]
mod tests {
@@ -593,33 +625,53 @@ mod tests {
slots.insert(s);
}
// 1. Exactly 191 unique slots (164 pre-SP11 + 27 SP11).
assert_eq!(slots.len(), 191, "expected 191 unique slots, got {}", slots.len());
// SP13 (2026-05-04): include the 8 directional-skill slots + the
// 3 v3 Hold-pricing controller slots — `SP13_SLOT_BASE..SP13_SLOT_END`
// covers both blocks contiguously since v3 extended the upper end
// from 380 to 383.
for s in SP13_SLOT_BASE..SP13_SLOT_END {
slots.insert(s);
}
// 1. Exactly 202 unique slots (164 pre-SP11 + 27 SP11 + 8 SP13 + 3 SP13 v3).
assert_eq!(slots.len(), 202, "expected 202 unique slots, got {}", slots.len());
// 2. Min slot is SP5_SLOT_BASE = 174.
assert_eq!(*slots.iter().min().unwrap(), 174);
// 3. Max slot is SP5_SLOT_END - 1 = 366.
assert_eq!(*slots.iter().max().unwrap(), 366);
// 3. Max slot is SP5_SLOT_END - 1 = 382.
assert_eq!(*slots.iter().max().unwrap(), 382);
// 4. Intentional carve-out gap (278, 279) is absent.
assert!(!slots.contains(&278), "slot 278 must be absent (carve-out gap)");
assert!(!slots.contains(&279), "slot 279 must be absent (carve-out gap)");
// 5. Set equals {174..278} ∪ {280..367} — no holes (other than the
// carve-out gap 278..280), no overlaps. Layer D D1 extended the
// upper end 286 → 290; D2 extended it 290 → 294; D3 extends it
// 294 → 297; SP7 T1 extended it 297 → 313; SP7 activation-flag
// fix extends it 313 → 321; SP8 (Fix 36) extends it 321 → 330;
// SP9 (Fix 37) extends it 330 → 339; SP10 (Fix 38) extends it
// 339 → 340; SP11 (Fix 39) extends it 340 → 360; SP11 B1b fix-up
// 5. SP13 boundary gap (367..372) is absent — separates SP11's controller
// block from SP13's directional-skill block.
for s in 367..372 {
assert!(!slots.contains(&s),
"slot {} must be absent (SP11→SP13 boundary gap)", s);
}
// 6. Set equals {174..278} ∪ {280..367} ∪ {372..383} — no holes (other
// than the carve-out gap 278..280 and the SP11→SP13 boundary gap
// 367..372), no overlaps. Layer D D1 extended the upper end
// 286 → 290; D2 extended it 290 → 294; D3 extends it 294 → 297;
// SP7 T1 extended it 297 → 313; SP7 activation-flag fix extends
// it 313 → 321; SP8 (Fix 36) extends it 321 → 330; SP9 (Fix 37)
// extends it 330 → 339; SP10 (Fix 38) extends it 339 → 340;
// SP11 (Fix 39) extends it 340 → 360; SP11 B1b fix-up
// (2026-05-04) extends it 360 → 361 with POPART_COMPONENT_MAG_EMA_INDEX;
// SP11 B1b smoke-recovery (2026-05-04) extends it 361 → 367 with
// `REWARD_COMPONENT_VAR_EMA_BASE..+6` for the z-score normalised
// mag-ratio canary.
let expected: HashSet<usize> = (174..278).chain(280..367).collect();
// mag-ratio canary; SP13 (2026-05-04) appends 372 → 380 with
// 8 directional-skill slots after the boundary gap 367..372;
// SP13 v3 P0a.T3 (2026-05-04) appends 380 → 383 with the 3 Hold-
// pricing controller slots (HOLD_COST=380, HOLD_RATE_TARGET=381,
// HOLD_RATE_OBSERVED_EMA=382).
let expected: HashSet<usize> = (174..278).chain(280..367).chain(372..383).collect();
assert_eq!(slots, expected,
"slot set does not match expected {{174..278}} ∪ {{280..367}}");
"slot set does not match expected {{174..278}} ∪ {{280..367}} ∪ {{372..383}}");
}
#[test]
@@ -646,11 +698,11 @@ mod tests {
// is cross-fold-persistent — the contracts must remain disjoint).
assert!(PNL_TOTAL_INDEX > LOSS_RATE_SMOOTH_INDEX);
// SP5_SLOT_END must reflect the post-SP8 end-of-block.
assert_eq!(SP5_SLOT_END, 367);
assert_eq!(SP5_SLOT_END, 383);
// SP5_PRODUCER_COUNT is the wiener-buffer linear span (slot-range
// width including the 2-slot carve-out gap), NOT the unique-slot
// count. See SP5_PRODUCER_COUNT docstring for the rationale.
assert_eq!(SP5_PRODUCER_COUNT, 193);
assert_eq!(SP5_PRODUCER_COUNT, 209);
assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE);
}
@@ -669,7 +721,7 @@ mod tests {
// 4-slot block is internally contiguous.
assert_eq!(GRAD_NORM_NORM_INDEX - HEALTH_SCORE_INDEX, 3);
// SP5_SLOT_END must reflect the post-SP8 end-of-block.
assert_eq!(SP5_SLOT_END, 367);
assert_eq!(SP5_SLOT_END, 383);
// SP5_PRODUCER_COUNT linear-span check matches the new end.
assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE);
}
@@ -690,12 +742,12 @@ mod tests {
assert_eq!(MAX_DD_EMA_INDEX - TRAINING_SHARPE_EMA_INDEX, 1);
assert_eq!(LOW_DD_RATIO_INDEX - MAX_DD_EMA_INDEX, 1);
// SP5_SLOT_END must reflect the post-SP8 end-of-block.
assert_eq!(SP5_SLOT_END, 367);
assert_eq!(SP5_SLOT_END, 383);
// SP5_PRODUCER_COUNT linear-span check matches the new end. The
// wiener buffer must cover the entire linear span — including the
// 2-slot carve-out gap (278..280) and the Pearl 6 reserved-but-
// unused 6-float block (slots 280..286 don't call apply_pearls).
assert_eq!(SP5_PRODUCER_COUNT, 193);
assert_eq!(SP5_PRODUCER_COUNT, 209);
assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE);
}
@@ -715,7 +767,7 @@ mod tests {
assert_eq!(lb_c51_active(b), LB_C51_ACTIVE_BASE + b);
}
// SP5_SLOT_END / SP5_PRODUCER_COUNT cover the new range.
assert_eq!(SP5_SLOT_END, 367);
assert_eq!(SP5_SLOT_END, 383);
assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE);
}
@@ -735,7 +787,7 @@ mod tests {
assert_eq!(lb_max_budget_c51(b), LB_MAX_BUDGET_C51_BASE + b);
}
// Layout: 1 + 4 + 4 = 9 contiguous slots (321..330).
assert_eq!(SP5_SLOT_END, 367);
assert_eq!(SP5_SLOT_END, 383);
assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE);
}
@@ -768,7 +820,7 @@ mod tests {
// Strictly above the SP8 MAX_BUDGET block.
assert!(KELLY_WARMUP_FLOOR_INDEX > LB_MAX_BUDGET_C51_BASE + 3);
// Block end at ISV[340) post-SP10 (Fix 38).
assert_eq!(SP5_SLOT_END, 367);
assert_eq!(SP5_SLOT_END, 383);
// Linear span matches producer count (no gaps in SP9 block).
assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE);
}
@@ -782,9 +834,9 @@ mod tests {
// Strictly above the SP9 eval_dist block (last entry at 338).
assert!(EVAL_THOMPSON_TEMP_INDEX > EVAL_DIST_F_INDEX);
// SP5_SLOT_END reflects the post-SP10 end-of-block.
assert_eq!(SP5_SLOT_END, 367);
assert_eq!(SP5_SLOT_END, 383);
// Linear span matches producer count (single new slot at top).
assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE);
assert_eq!(SP5_PRODUCER_COUNT, 193);
assert_eq!(SP5_PRODUCER_COUNT, 209);
}
}

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

@@ -135,6 +135,22 @@
#define MAG_FULL 2 // 1.00× max_position
#define NUM_MAGNITUDES 3
// ────────────────────────────────────────────────────────────────────────────
// Order-type and urgency branch sizes. Production default in
// `ml-dqn::dqn::DQNConfig` is 3 each (`num_order_types = 3`,
// `num_urgency_levels = 3`); the runtime config has never been tuned away
// from these and the kernels that consume packed factored actions assume
// the same value tile-wide. SP13 P0a.T4 (2026-05-04) lifts these into
// `state_layout.cuh` so kernels that need to decode the packed
// `action_idx = dir*M*O*U + mag*O*U + ord*U + urg` (e.g.
// `hold_rate_observer`) can compute the divisor without taking new
// launcher args. If runtime config ever changes, every consumer of the
// packed encoding must migrate atomically per
// `feedback_no_partial_refactor.md`.
// ────────────────────────────────────────────────────────────────────────────
#define NUM_ORD 3
#define NUM_URG 3
// ────────────────────────────────────────────────────────────────────────────
// ISV bus slot indices referenced by kernels that receive the ISV pointer.
// Kept in state_layout.cuh because both experience_kernels.cu and
@@ -152,6 +168,30 @@
#define ISV_SEED_FRAC_EMA_IDX 84 // == SEED_FRAC_EMA_INDEX — adaptive EMA of (1 - done/target) ∈ [0, 1] (Plan 3 Task 8 B.3; consumed by Task 9 CQL ramp)
#define ISV_EVAL_THOMPSON_TEMP_IDX 339 // == EVAL_THOMPSON_TEMP_INDEX — eval Thompson selector temperature (SP10 / Fix 38 2026-05-03; ISV-driven temperature blend on direction-branch Thompson sample)
// ────────────────────────────────────────────────────────────────────────────
// === SP13 SLOT INDICES === (2026-05-04, redefine success for predictive skill)
// Mirror of crates/ml/src/cuda_pipeline/sp13_isv_slots.rs. Slots [372..383)
// at the top of the SP5 ISV linear span; [367..372) is an intentional 5-slot
// boundary gap separating the SP11 controller block from this SP13 block.
// Slots 380..383 are the v3 P0a.T3 (2026-05-04) Hold-pricing controller
// (replaces v2's atomic Hold elimination — preserves the 4-way action space,
// prices Hold instead of removing it).
// Spec: docs/superpowers/specs/2026-05-04-sp13-redefine-success-for-predictive-skill.md
// ────────────────────────────────────────────────────────────────────────────
#define ISV_TARGET_DIR_ACC_IDX 372 // directional accuracy target (static, default 0.55)
#define ISV_AUX_DIR_ACC_SHORT_EMA_IDX 373 // aux dir-acc fast EMA (per-fold, sentinel 0.5)
#define ISV_AUX_DIR_ACC_LONG_EMA_IDX 374 // aux dir-acc slow EMA (per-fold, sentinel 0.5) — stagnation detector
#define ISV_AUX_DIR_PREDICTION_IDX 375 // aux head per-bar prediction (per-bar overwrite, [-1, +1])
#define ISV_DIR_SKILL_BONUS_ALPHA_IDX 376 // dir-skill bonus α (per-fold, default 1.0) — magnitude on correct calls
#define ISV_DIR_SKILL_BONUS_BETA_IDX 377 // dir-skill bonus β (per-fold, default 1.0) — penalty on wrong calls
#define ISV_LUCK_WIN_DISCOUNT_IDX 378 // lucky-win discount factor (per-fold, default 0.3)
#define ISV_SKILL_BONUS_CAP_RATIO_IDX 379 // skill-bonus cap ratio (per-fold, default 0.3) — bonus ≤ ratio × |α|
#define ISV_HOLD_COST_IDX 380 // SP13 v3 per-bar Hold cost (controller output, sentinel = HOLD_COST_BASE=0.001)
#define ISV_HOLD_RATE_TARGET_IDX 381 // SP13 v3 target Hold-pick rate (static, default 0.20)
#define ISV_HOLD_RATE_OBSERVED_EMA_IDX 382 // SP13 v3 observed Hold-pick rate EMA (per-fold, sentinel 0.0)
#define ISV_TOTAL_DIM 383
// ────────────────────────────────────────────────────────────────────────────
// Feature-group ranges — Plan 4 Task 1A (E.1 VSN prerequisite).
// Referenced by: VSN feature selection (E.1), attention diagnostics (E.5),

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

@@ -905,6 +905,51 @@ impl StateResetRegistry {
category: ResetCategory::FoldReset,
description: "ISV[REWARD_COMPONENT_VAR_EMA_BASE=361..367) — SP11 Fix 39 B1b smoke-recovery (2026-05-04, spec §4 amendment 'Why z-score' lines 564-619) per-reward-component variance EMAs (6 contiguous slots: popart at 361, then cf 362, trail 363, micro 364, opp_cost 365, bonus 366). Slot 361 is produced alongside the popart-component magnitude EMA by the extended `popart_component_ema_kernel` via single-pass two-tree-reduce Welford (mean → var); slots 362..367 are produced alongside the non-popart magnitudes by the extended `reward_component_ema_kernel` via per-thread Welford (component 0/popart variance is split off because slot 63 is the pre-SP11 total-reward magnitude EMA, not popart-component magnitude — see slot-360 fix-up rationale). All 6 variances are smoothed by Pearls A+D via chained `apply_pearls_ad_kernel` (n_slots=1 for popart variance at slot 361, n_slots=5 for the contiguous block at slots 362..367). Consumed by `reward_component_mag_ratio_compute_kernel` for z-score normalisation: `z[c] = mag[c] / max(sqrt(var[c]), EPS_DIV)` then `ratio[c] = z[c] / sum_z`. Replaces the linear `winner_weight = mag_ratio` formula that amplified popart's intrinsic O(100) magnitude over the other components' O(0.1-2) — w_pop saturated to MAX_WEIGHT, curiosity_b exploded, sharpe collapsed in B1b smoke `smoke-test-6wd2c` on commit `61b2fa962`. The single registry entry covers the whole 6-slot block — the dispatch arm zeroes all 6 with one loop, matching the `sp11_reward_component_mag_ratios` block-reset pattern. FoldReset sentinel 0; Pearl A bootstraps the variance from the first observation alongside its paired magnitude.",
},
// ── SP13 (2026-05-04): redefine success for predictive skill ─────
//
// Phase 0a Task T1: register the two dir-acc EMA slots so the
// fold-boundary path resets them to the random-guessing sentinel
// 0.5 (NOT 0.0 — directional accuracy has a meaningful baseline at
// 50%, and Pearl A's first-observation replacement fires on the
// first epoch of the new fold to track the new fold's signal
// without contamination from the previous fold's terminal value).
//
// Slots 372 (TARGET_DIR_ACC), 376..380 (DIR_SKILL_BONUS_ALPHA/BETA,
// LUCK_WIN_DISCOUNT, SKILL_BONUS_CAP_RATIO) are constructor-
// initialised Invariant-1 anchors and need NO registry entry —
// they must never reach sentinel 0 between folds, so the
// constructor-write path keeps them populated with their defaults
// for the lifetime of the trainer. Slot 375 (AUX_DIR_PREDICTION)
// is per-bar overwritten by the aux head and needs no reset
// entry either. Spec: docs/superpowers/specs/2026-05-04-sp13-redefine-success-for-predictive-skill.md
RegistryEntry {
name: "sp13_aux_dir_acc_short_ema",
category: ResetCategory::FoldReset,
description: "ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373] — SP13 Phase 0a fast EMA of directional accuracy (α=0.3). Produced by `apply_fixed_alpha_ema_kernel` chained off `aux_dir_acc_reduce_kernel` (the fixed-α applicator preserves the short/long timescale split that Wiener-optimal `apply_pearls_ad_kernel` would collapse). Consumed by the Phase 0b aux-w controller as the deficit numerator (`deficit = max(0, target − short)`) and the stagnation comparator. FoldReset sentinel 0.5 (random-guessing baseline) so Pearl A bootstraps from the new fold's first observation per `pearl_first_observation_bootstrap.md` — the EMA tracks the new fold's signal without bias from the previous fold's terminal accuracy. Spec §implementation phases.",
},
RegistryEntry {
name: "sp13_aux_dir_acc_long_ema",
category: ResetCategory::FoldReset,
description: "ISV[AUX_DIR_ACC_LONG_EMA_INDEX=374] — SP13 Phase 0a slow EMA of directional accuracy (α=0.05) feeding the stagnation detector. Produced by `apply_fixed_alpha_ema_kernel` chained off `aux_dir_acc_reduce_kernel` (the fixed-α applicator preserves the short/long timescale split that Wiener-optimal `apply_pearls_ad_kernel` would collapse). Consumed by the Phase 0b aux-w controller's stagnation term (`stagnation = max(0, 1 − (short − long) / max(0.005, deficit))`) so a stalled short EMA decays aux_w back toward base — the deliberate 'we tried, it didn't help' signal that distinguishes Phase 0b's controller from SP11's inverted formula. FoldReset sentinel 0.5 (random-guessing baseline) so Pearl A bootstraps from the new fold's first observation per `pearl_first_observation_bootstrap.md`. Spec §implementation phases.",
},
// SP13 v3 P0a.T3 (2026-05-04): Hold-pricing observed-rate EMA.
// Slots 380 (HOLD_COST) and 381 (HOLD_RATE_TARGET) are
// constructor-initialised Invariant-1 anchors and need NO
// registry entry — slot 380 is rewritten every step by the
// host-side controller in `training_loop.rs`, and slot 381 is
// a static default (0.20) that must never reach sentinel 0
// between folds. Slot 382 is the per-fold observed-rate EMA
// produced by `apply_fixed_alpha_ema_kernel` chained off
// `hold_rate_observer_kernel`; sentinel 0.0 lets Pearl A
// bootstrap from the new fold's first observation per
// `pearl_first_observation_bootstrap.md` so the EMA tracks
// the new fold's Hold-rate without bias from the previous
// fold's terminal value.
RegistryEntry {
name: "sp13_hold_rate_observed_ema",
category: ResetCategory::FoldReset,
description: "ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382] — SP13 v3 P0a.T3 per-fold EMA of the per-step Hold-pick rate `count(action_dir == DIR_HOLD) / batch_size`. Produced by `apply_fixed_alpha_ema_kernel` chained off `hold_rate_observer_kernel` (added in P0a.T3). Consumed by the host-side Hold-cost controller in `training_loop.rs`: `excess = max(0, observed − target); hold_cost = HOLD_COST_BASE × (1 + 5 × excess)` clamped to `[0.5×base, 5.0×base]`, written to ISV[HOLD_COST_INDEX=380]. The reward-composition site in `experience_kernels.cu` then subtracts the cost on every Hold-action bar — pricing the action so the policy uses Hold deliberately rather than as a free CQL-bias-anchored default. FoldReset sentinel 0.0 so Pearl A bootstraps from the new fold's first observation; the controller relaxes back to baseline cost (HOLD_COST_BASE × 0.5 floor) when observed ≤ target. Spec §'Change 1: Price Hold (replaces v2's Eliminate Hold)'.",
},
// SP5 Task A1: Wiener-state companion reset. The wiener_state_buf
// covers SP4 (SP4_PRODUCER_COUNT=71 producers × 3 = 213 floats) +
// SP5 (SP5_PRODUCER_COUNT × 3 floats) = (71 + SP5_PRODUCER_COUNT) × 3

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

@@ -3563,6 +3563,63 @@ impl DQNTrainer {
"SP11 mag_ratio_compute launch failed");
}
// SP13 v3 P0a.T3 (2026-05-04): Hold-pricing
// controller. Reads the observed Hold-pick rate
// EMA at ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382]
// (populated by `hold_rate_observer_kernel` +
// chained Pearls A+D — observer launch wires
// alongside the dir_acc launch in P0a.T4) and
// the static Hold-rate target at
// ISV[HOLD_RATE_TARGET_INDEX=381] (constructor-
// initialised to 0.20). Computes the priced
// per-bar Hold cost
// excess = max(0, observed − target)
// cost = HOLD_COST_BASE × (1 + 5 × excess)
// cost = clamp(cost, base × 0.5, base × 5.0)
// and writes it back to ISV[HOLD_COST_INDEX=380].
// The reward composition site in
// `experience_kernels.cu` reads slot 380 each
// bar where action == DIR_HOLD and subtracts
// the cost — pricing the action so the policy
// uses Hold deliberately rather than as a free
// CQL-bias-anchored default.
//
// Cold-start (epoch 1, slot 382 still at fold-
// reset sentinel 0.0): observed=0, target=0.20,
// excess=0 → cost = base × 1.0 = HOLD_COST_BASE,
// clamped to [base × 0.5, base × 5.0] →
// baseline carry cost. Pearl A's first-
// observation replacement consumes the sentinel
// when the observer first fires.
//
// Per `feedback_isv_for_adaptive_bounds`: base
// cost is the only constant; rate target,
// observed rate, and cost output are ISV-driven.
// Per `feedback_no_cpu_compute_strict`: the
// formula here is host-side controller arithmetic
// on already-smoothed scalars — no per-sample
// GPU compute path is bypassed (the observer +
// Pearls A+D + per-bar reward subtraction all
// run on GPU).
{
use crate::cuda_pipeline::sp5_isv_slots::{
HOLD_COST_INDEX, HOLD_RATE_TARGET_INDEX,
HOLD_RATE_OBSERVED_EMA_INDEX,
HOLD_COST_BASE, HOLD_COST_CONTROLLER_GAIN,
HOLD_COST_FLOOR_RATIO, HOLD_COST_CEIL_RATIO,
};
let trainer = fused.trainer();
let observed = trainer.read_isv_signal_at(HOLD_RATE_OBSERVED_EMA_INDEX);
let target = trainer.read_isv_signal_at(HOLD_RATE_TARGET_INDEX);
let excess = (observed - target).max(0.0);
let cost = HOLD_COST_BASE * (1.0 + HOLD_COST_CONTROLLER_GAIN * excess);
let cost = cost.clamp(
HOLD_COST_BASE * HOLD_COST_FLOOR_RATIO,
HOLD_COST_BASE * HOLD_COST_CEIL_RATIO,
);
trainer.write_isv_signal_at(HOLD_COST_INDEX, cost);
}
// SP11 Fix 39 (2026-05-04, A1.2): per-bar
// saboteur engagement canary producer. Reads
// the per-bar saboteur Δreward signal written
@@ -3914,6 +3971,30 @@ impl DQNTrainer {
}
}
// SP13 Phase 0a P0a.T4 (2026-05-04): per-step aux
// directional-accuracy producer chain. The orchestrator
// reads the captured-forward-graph `aux_nb_pred_buf` +
// `aux_nb_label_buf` (just populated above by the same
// per-step block that fires `launch_aux_heads_loss_ema`),
// reduces them into the trainer's mapped-pinned
// `aux_dir_acc_buf [3]`, then chains:
// - fixed-α EMA (α=0.3, sentinel 0.5) → ISV[373]
// - fixed-α EMA (α=0.05, sentinel 0.5) → ISV[374]
// - mean(tanh(aux_pred)) → ISV[375]
// All on the same stream as the producer, so no host
// sync is required between captured graph and consumer.
// Producer-only — no consumer kernel reads slots 373/374/
// 375 in this commit; Phase 0b wires the controller and
// the Q-head input layer that reads slot 375. Same
// post-cascade-fix invariant as `launch_aux_heads_loss_ema`:
// ISV producer launches MUST run BEFORE the HEALTH_DIAG
// line emit further down.
if let Some(ref fused) = self.fused_ctx {
if let Err(e) = fused.trainer().launch_sp13_aux_dir_metrics() {
tracing::warn!("SP13 P0a.T4 aux_dir_metrics launch failed: {e}");
}
}
// Phase 3 T3.5: MoE expert-utilisation EMA + gate entropy EMA.
// Reads `moe_gate_softmax_buf [B, K]` (valid after the captured
// forward graph ran) and EMA-updates ISV[118..127).
@@ -4709,6 +4790,55 @@ impl DQNTrainer {
);
}
// SP13 Phase 0a P0a.T4 (2026-05-04): aux directional-accuracy +
// Hold-pricing HEALTH_DIAG. Reads ISV slots populated by the
// per-step producer chains above:
// ISV[373..375) — aux dir-acc fast/slow EMAs (α=0.3 / 0.05)
// ISV[375] — mean(tanh(aux_pred)) per-step state
// ISV[372] — directional accuracy target (static 0.55)
// ISV[380..383) — Hold-pricing controller output / target /
// observed-rate EMA (slot 380 written by the
// per-epoch host controller at training_loop:3604;
// slot 382 written by the per-step
// gpu_experience_collector chain)
// Diagnostic only — exposes the SP13 chain's signal flow before
// Phase 0b wires the consumers (aux-w controller, dir-skill
// bonus, Q-head slot-375 input).
{
use crate::cuda_pipeline::sp13_isv_slots::{
TARGET_DIR_ACC_INDEX,
AUX_DIR_ACC_SHORT_EMA_INDEX, AUX_DIR_ACC_LONG_EMA_INDEX,
AUX_DIR_PREDICTION_INDEX,
HOLD_COST_INDEX, HOLD_RATE_TARGET_INDEX,
HOLD_RATE_OBSERVED_EMA_INDEX,
};
let (
target_dir_acc, dir_short, dir_long, aux_pred_tanh,
hold_cost, hold_target, hold_observed,
) = if let Some(ref fused) = self.fused_ctx {
let trainer = fused.trainer();
(
trainer.read_isv_signal_at(TARGET_DIR_ACC_INDEX),
trainer.read_isv_signal_at(AUX_DIR_ACC_SHORT_EMA_INDEX),
trainer.read_isv_signal_at(AUX_DIR_ACC_LONG_EMA_INDEX),
trainer.read_isv_signal_at(AUX_DIR_PREDICTION_INDEX),
trainer.read_isv_signal_at(HOLD_COST_INDEX),
trainer.read_isv_signal_at(HOLD_RATE_TARGET_INDEX),
trainer.read_isv_signal_at(HOLD_RATE_OBSERVED_EMA_INDEX),
)
} else {
(0.55, 0.5, 0.5, 0.0, 0.001, 0.20, 0.0)
};
tracing::info!(
"HEALTH_DIAG[{}]: aux_dir_acc target={:.4} short={:.4} long={:.4} pred_tanh={:.4}",
epoch, target_dir_acc, dir_short, dir_long, aux_pred_tanh,
);
tracing::info!(
"HEALTH_DIAG[{}]: hold_pricing observed_rate={:.4} target={:.4} cost={:.6}",
epoch, hold_observed, hold_target, hold_cost,
);
}
// Phase 3 T3.5: MoE expert-utilisation + gate entropy HEALTH_DIAG.
// Extended 2026-04-27: also emits the adaptive λ_eff slot so the
// controller is observable per epoch — see `launch_moe_lambda_eff_update`.
@@ -7488,6 +7618,54 @@ impl DQNTrainer {
}
}
}
// SP13 (2026-05-04) Phase 0a: directional-accuracy EMAs.
// Sentinel 0.5 (NOT 0.0) is the random-guessing baseline for
// binary directional accuracy. Pearl A's first-observation
// replacement fires on the first epoch of the new fold so the
// EMA tracks the new fold's signal without bias from the
// previous fold's terminal value. Companion Wiener-state slots
// (offsets in `wiener_state_buf`) are reset by the existing bulk
// memset covered by the `sp5_wiener_state` registry entry's
// dispatch arm — both halves of the Pearl A sentinel contract
// reset together per `feedback_no_partial_refactor.md`.
"sp13_aux_dir_acc_short_ema" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp5_isv_slots::{
AUX_DIR_ACC_SHORT_EMA_INDEX, DIR_ACC_EMA_SENTINEL,
};
fused.trainer().write_isv_signal_at(
AUX_DIR_ACC_SHORT_EMA_INDEX,
DIR_ACC_EMA_SENTINEL,
);
}
}
"sp13_aux_dir_acc_long_ema" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp5_isv_slots::{
AUX_DIR_ACC_LONG_EMA_INDEX, DIR_ACC_EMA_SENTINEL,
};
fused.trainer().write_isv_signal_at(
AUX_DIR_ACC_LONG_EMA_INDEX,
DIR_ACC_EMA_SENTINEL,
);
}
}
// SP13 v3 P0a.T3 (2026-05-04): Hold-pricing observed-rate EMA.
// Sentinel 0.0 — no Hold-pick observations at fold start. The
// companion wiener-state slot (offset = SP4_PRODUCER_COUNT × 3
// + (382 − SP5_SLOT_BASE) × 3) resets in the existing bulk
// memset covered by the `sp5_wiener_state` registry entry's
// dispatch arm; both halves of the Pearl A sentinel contract
// reset together per `feedback_no_partial_refactor.md`.
"sp13_hold_rate_observed_ema" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp5_isv_slots::HOLD_RATE_OBSERVED_EMA_INDEX;
fused.trainer().write_isv_signal_at(
HOLD_RATE_OBSERVED_EMA_INDEX,
0.0,
);
}
}
// SP11 Task A2 (2026-05-04): novelty visit-count hash table
// reset arm — closes the A0 deferral per the registry entry's
// docstring. 1M f32 slots, zero-filled via mapped-pinned host

731
crates/ml/tests/sp13_phase0_oracle_tests.rs Normal file
View File

@@ -0,0 +1,731 @@
#![allow(unsafe_code)] // CUDA kernel launch + mapped-pinned memory.
//! SP13 Phase 0a directional-accuracy reducer GPU oracle tests.
//!
//! Validates the single-block tree-reduce kernel introduced by P0a.T2:
//! `aux_dir_acc_reduce_kernel.cu`
//!
//! Six oracle cases drive synthetic `(aux_pred, next_bar_label)` pairs
//! through the production kernel and verify the GPU-computed
//! `(dir_acc, pos_pred_frac, pos_label_frac)` triple against analytically-
//! known expected values. Per `feedback_no_cpu_test_fallbacks.md`: GPU
//! oracle only, no CPU reference impl. Per
//! `feedback_no_htod_htoh_only_mapped_pinned.md`: every CPU↔GPU buffer
//! is a `MappedF32Buffer` (cuMemHostAlloc with DEVICEMAP|PORTABLE);
//! zero `htod_copy`, zero `dtoh_sync_copy`. Tests are NOT exempt from
//! this rule.
//!
//! All tests are `#[ignore = "requires GPU"]`-gated to match every other
//! GPU oracle test in this crate (sp4/sp5/sp11/sp12). Run on a GPU host:
//!
//! SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 \
//! cargo test -p ml --test sp13_phase0_oracle_tests --features cuda \
//! -- --ignored --nocapture
//!
//! Sign convention under test (kernel side: `(p > 0.0f) ? 1 : 0`):
//! - `p > 0` → predicts class 1 (positive)
//! - `p <= 0` → predicts class 0 (negative) — incl. `0.0` and `-0.0`
//! - `label == 0.0` → no-signal bar, excluded from denominator entirely
//!
//! Sentinel:
//! - When the entire batch is invalid (empty / all-zero labels),
//! `(dir_acc, pos_pred_frac, pos_label_frac) = (0.5, 0.5, 0.5)`.
//! Matches `DIR_ACC_EMA_SENTINEL` in `sp13_isv_slots.rs` per
//! `pearl_first_observation_bootstrap.md`.
#![cfg(feature = "cuda")]
use std::sync::Arc;
use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg};
use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedI32Buffer};
/// Test-only cubin built by `crates/ml/build.rs`.
const SP13_AUX_DIR_ACC_REDUCE_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/aux_dir_acc_reduce_kernel.cubin"));
/// Resolve a CUDA stream against device 0. Mirrors `make_test_stream` in
/// every other oracle test in this crate (`sp4_producer_unit_tests.rs`,
/// `sp11_producer_unit_tests.rs`, `sp12_reward_math_tests.rs`).
fn make_test_stream() -> Arc<CudaStream> {
let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
ctx.default_stream()
}
/// Load the reducer kernel function from the embedded cubin.
fn load_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
let module = stream
.context()
.load_cubin(SP13_AUX_DIR_ACC_REDUCE_CUBIN.to_vec())
.expect("load aux_dir_acc_reduce cubin");
module
.load_function("aux_dir_acc_reduce_kernel")
.expect("load aux_dir_acc_reduce_kernel function")
}
/// Drive one launch of the production reducer kernel with the given
/// per-bar prediction + label slices and return
/// `[dir_acc, pos_pred_frac, pos_label_frac]`.
///
/// All buffers are mapped-pinned (`MappedF32Buffer`) per
/// `feedback_no_htod_htoh_only_mapped_pinned.md`. The kernel writes the
/// 3-element output through `dev_ptr` with `__threadfence_system()`; the
/// test reads through `read_all()`'s volatile read after stream sync.
///
/// Block dim 256, shared mem `4 × 256 × sizeof(int) = 4 KB` — matches
/// the production launcher's config exactly.
fn run_dir_acc(
stream: &Arc<CudaStream>,
f: &CudaFunction,
pred: &[f32],
label: &[f32],
) -> [f32; 3] {
assert_eq!(
pred.len(),
label.len(),
"pred and label must have matching length; got {} vs {}",
pred.len(),
label.len()
);
let n = pred.len();
// Empty-batch case: even if `n == 0` the kernel still needs valid
// device pointers for the input args (CUDA dereferences them only
// inside the strided loop, which doesn't execute, but the launch
// builder still passes the pointers). Allocate 1-element placeholders
// so the dev_ptr is non-zero.
let alloc_len = n.max(1);
// Safety: CUDA context active on this thread (resolved via the
// stream's context above). All CPU↔GPU buffers are mapped-pinned.
let pred_buf = unsafe { MappedF32Buffer::new(alloc_len) }
.expect("alloc pred buffer (mapped-pinned)");
let label_buf = unsafe { MappedF32Buffer::new(alloc_len) }
.expect("alloc label buffer (mapped-pinned)");
if n > 0 {
pred_buf.write_from_slice(pred);
label_buf.write_from_slice(label);
}
let out_buf = unsafe { MappedF32Buffer::new(3) }
.expect("alloc out buffer (3 f32 mapped-pinned)");
let pred_dev = pred_buf.dev_ptr;
let label_dev = label_buf.dev_ptr;
let out_dev = out_buf.dev_ptr;
let batch_size_i32: i32 = n as i32;
let bdim: u32 = 256;
unsafe {
stream
.launch_builder(f)
.arg(&pred_dev)
.arg(&label_dev)
.arg(&batch_size_i32)
.arg(&out_dev)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (bdim, 1, 1),
shared_mem_bytes: 4 * bdim * std::mem::size_of::<i32>() as u32,
})
.expect("launch aux_dir_acc_reduce_kernel");
}
stream.synchronize().expect("sync after aux_dir_acc_reduce launch");
let out = out_buf.read_all();
[out[0], out[1], out[2]]
}
// ─── Six P0a.T2 oracle cases ─────────────────────────────────────────────
/// All four bars predict the correct sign → dir_acc = 1.0.
#[test]
#[ignore = "requires GPU"]
fn dir_acc_all_correct_returns_one() {
let stream = make_test_stream();
let f = load_kernel(&stream);
let pred = vec![ 1.0_f32, 1.0, -1.0, -1.0];
let label = vec![ 1.0_f32, 1.0, -1.0, -1.0];
let out = run_dir_acc(&stream, &f, &pred, &label);
assert!(
(out[0] - 1.0).abs() < 1e-6,
"expected dir_acc = 1.0, got {}",
out[0]
);
}
/// All four bars predict the wrong sign → dir_acc = 0.0.
#[test]
#[ignore = "requires GPU"]
fn dir_acc_all_wrong_returns_zero() {
let stream = make_test_stream();
let f = load_kernel(&stream);
let pred = vec![ 1.0_f32, 1.0, -1.0, -1.0];
let label = vec![-1.0_f32, -1.0, 1.0, 1.0];
let out = run_dir_acc(&stream, &f, &pred, &label);
assert!(
out[0].abs() < 1e-6,
"expected dir_acc = 0.0, got {}",
out[0]
);
}
/// Half right + asymmetric prediction/label distributions.
///
/// pred = [+, +, +, +] → 4 positive predictions / 4 valid bars = 1.0
/// label = [+, +, -, -] → 2 positive labels / 4 valid bars = 0.5
/// dir_acc = 2/4 = 0.5 (bars 0+1 correct, bars 2+3 wrong).
#[test]
#[ignore = "requires GPU"]
fn dir_acc_half_correct() {
let stream = make_test_stream();
let f = load_kernel(&stream);
let pred = vec![ 1.0_f32, 1.0, 1.0, 1.0];
let label = vec![ 1.0_f32, 1.0, -1.0, -1.0];
let out = run_dir_acc(&stream, &f, &pred, &label);
assert!(
(out[0] - 0.5).abs() < 1e-6,
"expected dir_acc = 0.5, got {}",
out[0]
);
assert!(
(out[1] - 1.0).abs() < 1e-6,
"expected pos_pred_frac = 1.0 (4/4 positive predictions), got {}",
out[1]
);
assert!(
(out[2] - 0.5).abs() < 1e-6,
"expected pos_label_frac = 0.5 (2/4 positive labels), got {}",
out[2]
);
}
/// Bars whose label is exactly 0.0 are excluded from the denominator
/// entirely.
///
/// pred = [+, +, -]
/// label = [+, 0, -] ← bar 1 is no-signal and skipped
/// valid bars = 2 (indexes 0 and 2). Both correct → dir_acc = 2/2 = 1.0.
/// If the kernel mistakenly counted bar 1 as `label_pos = 0` (matching
/// `pred_pos = 1`?) it would be wrong; if it counted bar 1 as valid with
/// `label_pos = 0` and `pred_pos = 1`, dir_acc would drop to 2/3 ≈ 0.667.
/// Either failure mode is caught by the strict 1.0 expectation.
#[test]
#[ignore = "requires GPU"]
fn dir_acc_skips_zero_labels() {
let stream = make_test_stream();
let f = load_kernel(&stream);
let pred = vec![ 1.0_f32, 1.0, -1.0];
let label = vec![ 1.0_f32, 0.0, -1.0];
let out = run_dir_acc(&stream, &f, &pred, &label);
assert!(
(out[0] - 1.0).abs() < 1e-6,
"expected dir_acc = 1.0 (bar 1 with label=0 excluded), got {}",
out[0]
);
}
/// Sign convention: `(p > 0.0f) ? 1 : 0` treats both `+0.0` and `-0.0`
/// as the negative class.
///
/// pred = [0.0, -0.0] ← both predict class 0 (negative)
/// label = [-1.0, -1.0] ← both negative
/// dir_acc = 2/2 = 1.0 (both correct).
#[test]
#[ignore = "requires GPU"]
fn dir_acc_zero_pred_is_negative() {
let stream = make_test_stream();
let f = load_kernel(&stream);
let pred = vec![ 0.0_f32, -0.0];
let label = vec![-1.0_f32, -1.0];
let out = run_dir_acc(&stream, &f, &pred, &label);
assert!(
(out[0] - 1.0).abs() < 1e-6,
"expected dir_acc = 1.0 (zero pred → class 0 → matches negative label), got {}",
out[0]
);
}
/// Empty batch (or all-zero-label batch) returns the random-baseline
/// sentinel 0.5 across all three outputs. Matches
/// `DIR_ACC_EMA_SENTINEL` in `sp13_isv_slots.rs` per
/// `pearl_first_observation_bootstrap.md` — the downstream EMA's first-
/// observation replacement consumes the sentinel naturally.
#[test]
#[ignore = "requires GPU"]
fn dir_acc_empty_batch_returns_sentinel() {
let stream = make_test_stream();
let f = load_kernel(&stream);
let pred: Vec<f32> = vec![];
let label: Vec<f32> = vec![];
let out = run_dir_acc(&stream, &f, &pred, &label);
assert!(
(out[0] - 0.5).abs() < 1e-6,
"expected dir_acc sentinel = 0.5 on empty batch, got {}",
out[0]
);
assert!(
(out[1] - 0.5).abs() < 1e-6,
"expected pos_pred_frac sentinel = 0.5 on empty batch, got {}",
out[1]
);
assert!(
(out[2] - 0.5).abs() < 1e-6,
"expected pos_label_frac sentinel = 0.5 on empty batch, got {}",
out[2]
);
}
// ─── SP13 v3 P0a.T3 Hold-rate observer oracle cases ─────────────────────
//
// Validates the second SP13 Phase 0a single-block tree-reduce:
// `hold_rate_observer_kernel.cu`
//
// Drives synthetic per-bar PACKED factored-action tiles (`batch_actions [B]`
// of `i32`) through the production kernel and verifies the GPU-computed
// fraction `count(decoded_dir == DIR_HOLD) / batch_size` against
// analytically-known expected values. Same `MappedF32Buffer` discipline
// as the dir-acc oracle above; the actions tile is allocated through a
// dedicated `MappedI32Buffer` helper since the production input is `i32`
// rather than `f32`.
//
// Action encoding under test (mirrors `experience_action_select` at
// `experience_kernels.cu:1453`):
//
// action_idx = dir * (NUM_MAGNITUDES * NUM_ORD * NUM_URG)
// + mag * (NUM_ORD * NUM_URG)
// + ord * NUM_URG
// + urg
//
// Production branch sizes are NUM_MAGNITUDES = NUM_ORD = NUM_URG = 3
// (see `state_layout.cuh`); the per-direction stride is 27.
//
// Sign convention under test:
// - decoded `dir == DIR_HOLD` (=1) — bars whose packed action_idx
// decodes to direction Hold contribute 1 to the count
// - any other direction (`DIR_SHORT=0`, `DIR_LONG=2`, `DIR_FLAT=3`)
// contributes 0
// - empty batch → out = 0.0 (matches the Pearl A sentinel for slot
// 382 / `HOLD_RATE_OBSERVED_EMA_INDEX`)
const SP13_HOLD_RATE_OBSERVER_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/hold_rate_observer_kernel.cubin"));
/// Production branch sizes (mirror `state_layout.cuh`: `NUM_MAGNITUDES`,
/// `NUM_ORD`, `NUM_URG`). Authoritative source — every consumer of the
/// packed factored encoding must use the same values per
/// `feedback_no_partial_refactor.md`.
const TEST_NUM_MAGNITUDES: i32 = 3;
const TEST_NUM_ORD: i32 = 3;
const TEST_NUM_URG: i32 = 3;
/// Pack a 4-tuple `(dir, mag, ord, urg)` into a factored action_idx using
/// the same encoding as `experience_action_select`. Direction stride is
/// `M * O * U`; magnitude stride is `O * U`; order stride is `U`.
fn pack_action(dir: i32, mag: i32, ord: i32, urg: i32) -> i32 {
dir * (TEST_NUM_MAGNITUDES * TEST_NUM_ORD * TEST_NUM_URG)
+ mag * (TEST_NUM_ORD * TEST_NUM_URG)
+ ord * TEST_NUM_URG
+ urg
}
fn load_hold_rate_observer(stream: &Arc<CudaStream>) -> CudaFunction {
let module = stream
.context()
.load_cubin(SP13_HOLD_RATE_OBSERVER_CUBIN.to_vec())
.expect("load hold_rate_observer cubin");
module
.load_function("hold_rate_observer_kernel")
.expect("load hold_rate_observer_kernel function")
}
/// Drive one launch of the production Hold-rate observer reducer with
/// the given per-bar direction-action slice and return the Hold-pick
/// fraction.
///
/// Block dim 256, shared mem `256 × sizeof(i32) = 1 KB` — matches the
/// production launcher's config exactly. All CPU↔GPU buffers are
/// mapped-pinned (`MappedI32Buffer` for the i32 actions tile,
/// `MappedF32Buffer` for the f32 output) per
/// `feedback_no_htod_htoh_only_mapped_pinned.md`.
fn run_hold_rate_observer(
stream: &Arc<CudaStream>,
f: &CudaFunction,
actions: &[i32],
) -> f32 {
let n = actions.len();
// Empty-batch case: kernel still needs a non-null device pointer for
// the actions arg (strided loop is empty but launch builder still
// passes the pointer). Mirrors `run_dir_acc`.
let alloc_len = n.max(1);
// Safety: CUDA context active on this thread (resolved via the
// stream's context above). All CPU↔GPU buffers are mapped-pinned.
let actions_buf = unsafe { MappedI32Buffer::new(alloc_len) }
.expect("alloc actions buffer (mapped-pinned)");
if n > 0 {
actions_buf.write_from_slice(actions);
}
let out_buf = unsafe { MappedF32Buffer::new(1) }
.expect("alloc out buffer (1 f32 mapped-pinned)");
let actions_dev = actions_buf.dev_ptr;
let out_dev = out_buf.dev_ptr;
let batch_size_i32: i32 = n as i32;
let bdim: u32 = 256;
unsafe {
stream
.launch_builder(f)
.arg(&actions_dev)
.arg(&batch_size_i32)
.arg(&out_dev)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (bdim, 1, 1),
shared_mem_bytes: bdim * std::mem::size_of::<i32>() as u32,
})
.expect("launch hold_rate_observer_kernel");
}
stream.synchronize().expect("sync after hold_rate_observer launch");
out_buf.read_all()[0]
}
/// Every action in the batch decodes to Hold direction → out = 1.0.
/// Uses non-zero (mag, ord, urg) sub-indices to exercise the divisor —
/// would catch a kernel that mistakenly treated the raw action_idx as the
/// direction (it's only the direction *after* dividing by M*O*U).
#[test]
#[ignore = "requires GPU"]
fn hold_rate_observer_all_hold_returns_one() {
let stream = make_test_stream();
let f = load_hold_rate_observer(&stream);
// Mix Hold actions across different (mag, ord, urg) sub-indices to
// verify the kernel's divisor. action_idx for Hold ranges over
// [27, 54): every value in that interval decodes to dir=1.
let actions: Vec<i32> = vec![
pack_action(1, 0, 0, 0), // 27
pack_action(1, 1, 0, 0), // 36
pack_action(1, 0, 2, 0), // 33
pack_action(1, 2, 2, 2), // 53 (Hold high-corner)
pack_action(1, 0, 0, 1), // 28
pack_action(1, 1, 1, 1), // 40
pack_action(1, 2, 0, 0), // 45
pack_action(1, 0, 1, 2), // 32
pack_action(1, 1, 2, 1), // 43
pack_action(1, 0, 0, 2), // 29
];
let out = run_hold_rate_observer(&stream, &f, &actions);
assert!(
(out - 1.0).abs() < 1e-6,
"expected hold_rate = 1.0 on all-Hold batch (decoded), got {}",
out
);
}
/// No action in the batch decodes to Hold (mix of Short=0/Long=2/Flat=3
/// across all (mag, ord, urg)) → out = 0.0. Verifies that no off-by-one
/// in the divisor flips a non-Hold packed value into the Hold band.
#[test]
#[ignore = "requires GPU"]
fn hold_rate_observer_no_hold_returns_zero() {
let stream = make_test_stream();
let f = load_hold_rate_observer(&stream);
// Mix non-Hold direction values across the (mag, ord, urg) sub-space.
// Short band = [0, 27), Long band = [54, 81), Flat band = [81, 108).
// None of these decode to dir=1.
let actions: Vec<i32> = vec![
pack_action(0, 0, 0, 0), // 0 (Short low-corner)
pack_action(0, 2, 2, 2), // 26 (Short high-corner)
pack_action(2, 0, 0, 0), // 54 (Long low-corner)
pack_action(2, 2, 2, 2), // 80 (Long high-corner)
pack_action(3, 0, 0, 0), // 81 (Flat low-corner)
pack_action(3, 2, 2, 2), // 107 (Flat high-corner)
];
let out = run_hold_rate_observer(&stream, &f, &actions);
assert!(
out.abs() < 1e-6,
"expected hold_rate = 0.0 on no-Hold batch (decoded), got {}",
out
);
}
/// Half the batch decodes to Hold → out = 0.5. Verifies the tree-reduce
/// arithmetic produces the correct fraction (not just the boolean any-
/// Hold/no-Hold extremes the previous two tests cover) AND that every
/// non-Hold direction (Short/Long/Flat) decodes correctly to non-Hold —
/// would catch an off-by-one comparing against DIR_SHORT instead of
/// DIR_HOLD.
#[test]
#[ignore = "requires GPU"]
fn hold_rate_observer_half_hold_returns_half() {
let stream = make_test_stream();
let f = load_hold_rate_observer(&stream);
// 4 Hold + 4 non-Hold = 0.5 fraction. Each Hold has different
// (mag, ord, urg); the 4 non-Hold span all 3 alternative directions.
let actions: Vec<i32> = vec![
pack_action(1, 0, 0, 0), // Hold
pack_action(0, 1, 1, 1), // Short
pack_action(1, 1, 1, 1), // Hold
pack_action(2, 0, 0, 0), // Long
pack_action(1, 2, 2, 2), // Hold
pack_action(3, 1, 1, 1), // Flat
pack_action(1, 0, 1, 2), // Hold
pack_action(0, 2, 2, 2), // Short
];
let out = run_hold_rate_observer(&stream, &f, &actions);
assert!(
(out - 0.5).abs() < 1e-6,
"expected hold_rate = 0.5 on half-Hold batch (decoded), got {}",
out
);
}
// ─── SP13 P0a.T4 fixed-α EMA applicator oracle cases ────────────────────
//
// Validates the SP13 fixed-α EMA applicator kernel introduced by P0a.T4:
// `apply_fixed_alpha_ema_kernel.cu`
//
// Drives synthetic (sample, prev_state, alpha, sentinel) tuples through
// the production kernel and verifies the GPU-computed
// `(1-α) · prev + α · sample` blend (or sentinel-bootstrap direct
// replacement) against analytically-known expected values. Same
// `MappedF32Buffer` discipline as the other oracles; the kernel is N-slot
// generalisable but P0a always launches with N=1, so the tests stay at
// N=1 to cover the production launch shape.
const SP13_APPLY_FIXED_ALPHA_EMA_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/apply_fixed_alpha_ema_kernel.cubin"));
fn load_apply_fixed_alpha_ema(stream: &Arc<CudaStream>) -> CudaFunction {
let module = stream
.context()
.load_cubin(SP13_APPLY_FIXED_ALPHA_EMA_CUBIN.to_vec())
.expect("load apply_fixed_alpha_ema cubin");
module
.load_function("apply_fixed_alpha_ema_kernel")
.expect("load apply_fixed_alpha_ema_kernel function")
}
/// Drive one launch of the fixed-α EMA applicator on a 1-slot ISV
/// scratch buffer pre-seeded with `prev_state`. Returns the post-launch
/// ISV slot value.
fn run_fixed_alpha_ema(
stream: &Arc<CudaStream>,
f: &CudaFunction,
sample: f32,
prev_state: f32,
alpha: f32,
sentinel: f32,
) -> f32 {
// Allocate a 1-elem mapped-pinned ISV scratch + 1-elem sample buffer.
let isv_buf = unsafe { MappedF32Buffer::new(1) }
.expect("alloc isv buffer (mapped-pinned)");
isv_buf.write_from_slice(&[prev_state]);
let sample_buf = unsafe { MappedF32Buffer::new(1) }
.expect("alloc sample buffer (mapped-pinned)");
sample_buf.write_from_slice(&[sample]);
let sample_dev = sample_buf.dev_ptr;
let isv_dev = isv_buf.dev_ptr;
let n_slots: i32 = 1;
let isv_offset: i32 = 0;
unsafe {
stream
.launch_builder(f)
.arg(&sample_dev)
.arg(&n_slots)
.arg(&isv_offset)
.arg(&alpha)
.arg(&sentinel)
.arg(&isv_dev)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (1, 1, 1),
shared_mem_bytes: 0,
})
.expect("launch apply_fixed_alpha_ema_kernel");
}
stream.synchronize().expect("sync after apply_fixed_alpha_ema launch");
isv_buf.read_all()[0]
}
/// Standard blend: prev=0.5, sample=1.0, α=0.5 → 0.75.
/// Sentinel is intentionally distinct from `prev` so the bootstrap
/// branch does NOT fire — this test exercises the steady-state blend
/// formula `(1-α) · prev + α · sample`.
#[test]
#[ignore = "requires GPU"]
fn fixed_alpha_ema_steady_state_blend() {
let stream = make_test_stream();
let f = load_apply_fixed_alpha_ema(&stream);
let out = run_fixed_alpha_ema(
&stream, &f,
/* sample */ 1.0,
/* prev */ 0.5,
/* alpha */ 0.5,
/* sentinel */ -1.0, // far from prev — bootstrap MUST NOT fire
);
assert!(
(out - 0.75).abs() < 1e-6,
"expected (1-0.5)*0.5 + 0.5*1.0 = 0.75, got {}",
out
);
}
/// First-observation bootstrap: prev == sentinel → replace directly.
/// Mirrors `pearl_first_observation_bootstrap` — the SP13 dir-acc
/// EMAs use sentinel=0.5 and the hold-rate EMA uses sentinel=0.0;
/// after fold reset the slot holds the sentinel, and the very first
/// non-sentinel observation must replace it directly so the EMA tracks
/// the new fold's signal without the sentinel's bias.
#[test]
#[ignore = "requires GPU"]
fn fixed_alpha_ema_sentinel_bootstrap() {
let stream = make_test_stream();
let f = load_apply_fixed_alpha_ema(&stream);
// SP13 dir-acc sentinel = 0.5. First real observation = 0.62 (>
// baseline). The kernel must REPLACE prev directly, not blend —
// a blend at α=0.3 would produce 0.5*0.7 + 0.62*0.3 = 0.536 which
// would carry the sentinel bias into the EMA forever.
let out = run_fixed_alpha_ema(
&stream, &f,
/* sample */ 0.62,
/* prev */ 0.5,
/* alpha */ 0.3,
/* sentinel */ 0.5, // matches prev — bootstrap MUST fire
);
assert!(
(out - 0.62).abs() < 1e-6,
"expected sentinel bootstrap to replace directly with 0.62, got {}",
out
);
}
// ─── SP13 P0a.T4 aux-pred → ISV[375] tanh oracle cases ──────────────────
//
// Validates the SP13 aux-head per-bar prediction → ISV[AUX_DIR_PREDICTION_INDEX=375]
// tanh-bounded scalar producer kernel introduced by P0a.T4:
// `aux_pred_to_isv_tanh_kernel.cu`
//
// Drives synthetic per-bar `aux_pred [B]` tiles through the production
// kernel and verifies the GPU-computed `mean(tanh(aux_pred[i]))` ∈ [-1, +1]
// against analytically-known expected values.
const SP13_AUX_PRED_TO_ISV_TANH_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/aux_pred_to_isv_tanh_kernel.cubin"));
fn load_aux_pred_to_isv_tanh(stream: &Arc<CudaStream>) -> CudaFunction {
let module = stream
.context()
.load_cubin(SP13_AUX_PRED_TO_ISV_TANH_CUBIN.to_vec())
.expect("load aux_pred_to_isv_tanh cubin");
module
.load_function("aux_pred_to_isv_tanh_kernel")
.expect("load aux_pred_to_isv_tanh_kernel function")
}
/// Drive one launch of the aux-pred → ISV[375] tanh producer with the
/// given per-bar prediction slice. Returns ISV slot 0 post-launch.
fn run_aux_pred_to_isv_tanh(
stream: &Arc<CudaStream>,
f: &CudaFunction,
aux_pred: &[f32],
) -> f32 {
let n = aux_pred.len();
let alloc_len = n.max(1);
let pred_buf = unsafe { MappedF32Buffer::new(alloc_len) }
.expect("alloc pred buffer (mapped-pinned)");
if n > 0 {
pred_buf.write_from_slice(aux_pred);
}
// 1-elem ISV scratch — kernel only writes slot 0.
let isv_buf = unsafe { MappedF32Buffer::new(1) }
.expect("alloc isv buffer (mapped-pinned)");
let pred_dev = pred_buf.dev_ptr;
let isv_dev = isv_buf.dev_ptr;
let batch_i32: i32 = n as i32;
let isv_off_i32: i32 = 0;
let bdim: u32 = 256;
unsafe {
stream
.launch_builder(f)
.arg(&pred_dev)
.arg(&batch_i32)
.arg(&isv_off_i32)
.arg(&isv_dev)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (bdim, 1, 1),
shared_mem_bytes: bdim * std::mem::size_of::<f32>() as u32,
})
.expect("launch aux_pred_to_isv_tanh_kernel");
}
stream.synchronize().expect("sync after aux_pred_to_isv_tanh launch");
isv_buf.read_all()[0]
}
/// Symmetric batch around zero → mean(tanh) ≈ 0.0.
/// `tanh(±x)` is odd, so an exactly-symmetric input sums to zero.
#[test]
#[ignore = "requires GPU"]
fn aux_pred_tanh_symmetric_batch_returns_zero() {
let stream = make_test_stream();
let f = load_aux_pred_to_isv_tanh(&stream);
let pred = vec![1.0_f32, -1.0, 0.5, -0.5, 2.0, -2.0];
let out = run_aux_pred_to_isv_tanh(&stream, &f, &pred);
assert!(
out.abs() < 1e-6,
"expected mean(tanh) = 0.0 on symmetric batch, got {}",
out
);
}
/// All-zero batch → mean(tanh(0)) = 0.0.
#[test]
#[ignore = "requires GPU"]
fn aux_pred_tanh_all_zero_returns_zero() {
let stream = make_test_stream();
let f = load_aux_pred_to_isv_tanh(&stream);
let pred = vec![0.0_f32; 16];
let out = run_aux_pred_to_isv_tanh(&stream, &f, &pred);
assert!(
out.abs() < 1e-6,
"expected mean(tanh(0)) = 0.0, got {}",
out
);
}
/// Bounded by [-1, +1] regardless of input magnitude. tanh saturates
/// at ±1 for large |x|, so mean(tanh(large)) ≈ ±1; large negative
/// inputs land near -1 (well within the [-1, +1] bound).
#[test]
#[ignore = "requires GPU"]
fn aux_pred_tanh_large_magnitude_saturates_in_bounds() {
let stream = make_test_stream();
let f = load_aux_pred_to_isv_tanh(&stream);
let pred = vec![10.0_f32; 32]; // tanh(10) ≈ 1.0 to ~5 nines
let out = run_aux_pred_to_isv_tanh(&stream, &f, &pred);
assert!(
(out - 1.0).abs() < 1e-4 && out <= 1.0,
"expected mean(tanh(10)) ≈ 1.0 (saturated), got {}",
out
);
let pred_neg = vec![-10.0_f32; 32];
let out_neg = run_aux_pred_to_isv_tanh(&stream, &f, &pred_neg);
assert!(
(out_neg + 1.0).abs() < 1e-4 && out_neg >= -1.0,
"expected mean(tanh(-10)) ≈ -1.0 (saturated), got {}",
out_neg
);
}

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

@@ -5732,3 +5732,72 @@ The B1b sharpe-degradation residual is **resolved**. The B1b
controller (z-score normalised mag-ratio canary) and the SP11
reward cap fix together close the within-fold degradation root
cause.
---
## SP13 P0a — Hold-pricing + dir_acc instrumentation (2026-05-04)
**Spec**: `docs/superpowers/specs/2026-05-04-sp13-redefine-success-for-predictive-skill.md` v3 (commit `ba83fcd1f`)
**Plan**: `docs/superpowers/plans/2026-05-04-sp13-redefine-success-for-predictive-skill.md` v3
**Hypothesis**: directional signal IS in the data; ~46% WR persists because Hold is FREE — CQL bias makes it the lazy default. Pricing Hold forces deliberate use; if WR climbs, hypothesis confirmed.
### Architectural changes
| Change | Mechanism |
|---|---|
| **Price Hold** (replaces v2's Eliminate Hold) | ISV-driven adaptive controller in `training_loop.rs`. Per-bar `reward -= isv[HOLD_COST_INDEX]` at 3 reward-composition sites in `experience_kernels.cu` when action == DIR_HOLD. 4-way action space stays — `ExposureLevel::Hold` preserved (audit revealed v2's `DirectionAction` enum doesn't exist; fused 8-variant `ExposureLevel` cascades through 77 files). |
| **Aux dir_acc instrumentation** | New `aux_dir_acc_reduce_kernel.cu` (single-block tree-reduce of correct/pos_pred/pos_label/valid → 3 scalars). Dual-EMA (short α=0.3, long α=0.05) feeds Phase 0b stagnation detector. |
| **Aux signal → Q-head input** | New `aux_pred_to_isv_tanh_kernel.cu` reduces aux scalar to ISV[375] for direction-Q-head consumption (Layer B will replace tanh(scalar) with softmax logit-diff). |
| **Hold-rate observer** | New `hold_rate_observer_kernel.cu` reads packed `batch_actions [B]`, decodes direction inline (`dir = action_idx / (NUM_MAGNITUDES * NUM_ORD * NUM_URG)`), reduces to `count(Hold) / B` scalar. Wired per-step in `gpu_experience_collector.rs` after `experience_action_select`. |
| **Fixed-α EMA applicator** | New `apply_fixed_alpha_ema_kernel.cu` — sibling of `apply_pearls_kernel.cu`. Wiener-optimal α* would collapse the short/long pair to identical values, defeating the stagnation detector; fixed-α preserves the timescale split. Sentinel-aware first-observation bootstrap per `pearl_first_observation_bootstrap`. |
### ISV slot allocation
11 new slots `[372..383)`:
| Slot | Name | Role | Reset |
|---|---|---|---|
| 372 | `TARGET_DIR_ACC_INDEX` | Default 0.55 — dir_acc target | static (constructor) |
| 373 | `AUX_DIR_ACC_SHORT_EMA_INDEX` | Fast EMA (α=0.3) | per-fold sentinel 0.5 |
| 374 | `AUX_DIR_ACC_LONG_EMA_INDEX` | Slow EMA (α=0.05) — stagnation detector | per-fold sentinel 0.5 |
| 375 | `AUX_DIR_PREDICTION_INDEX` | tanh(aux_pred) → direction Q-head input | per-bar overwrite |
| 376 | `DIR_SKILL_BONUS_ALPHA_INDEX` | Layer C bonus magnitude (correct dir) | static 1.0 |
| 377 | `DIR_SKILL_BONUS_BETA_INDEX` | Layer C penalty magnitude (wrong dir) | static 1.0 |
| 378 | `LUCK_WIN_DISCOUNT_INDEX` | Layer C lucky-win discount | static 0.3 |
| 379 | `SKILL_BONUS_CAP_RATIO_INDEX` | Layer C bonus cap as fraction of \|alpha\| | static 0.3 |
| 380 | `HOLD_COST_INDEX` | Per-bar Hold cost — controller output | static (constructor seeds with HOLD_COST_BASE=0.001) |
| 381 | `HOLD_RATE_TARGET_INDEX` | Default 0.20 — controller target | static |
| 382 | `HOLD_RATE_OBSERVED_EMA_INDEX` | Per-step EMA of Hold-pick rate | per-fold sentinel 0.0 |
`SP5_SLOT_END = 383`, `ISV_TOTAL_DIM = 383`. Layout fingerprint extended with all 11 slot names.
### Test coverage
`crates/ml/tests/sp13_phase0_oracle_tests.rs` — **14 GPU oracle tests on RTX 3050 Ti, all passing**:
- 6 dir_acc reduction tests (T2)
- 3 hold_rate observer tests (T3 v3 — packed `batch_actions` decode)
- 2 fixed-α EMA tests (T4 — steady-state blend + sentinel bootstrap)
- 3 aux_pred → tanh tests (T4 — symmetric/all-zero/saturating)
SP12 14/14 + SP11 11/11 oracle tests still green (no regression).
### What's deferred to subsequent tasks
- **Layer B**: aux head regression → 30-bar binary classification (replaces `tanh(scalar)` in slot 375 with softmax `logit_diff`). 5-epoch smoke after green Phase 0a.
- **Layer C**: alpha-vs-benchmark (DROPPED v2) → direction-skill bonus + luck-win discount with bounded calibration (slots 376-379 populated but unused until then).
- **Layer D**: 30-epoch L40S validation + close-out commit with 3 new pearls.
### Pearls (post-validation)
To be written if Phase 0a confirms hypothesis:
- `pearl_redefine_success_for_predictive_skill`
- `pearl_skill_bonus_must_be_alpha_bounded` (calibration lesson from spec v1→v2 review)
- `pearl_reward_quadrant_audit_required` (meta-pearl from v1 review — every reward change needs a 4-quadrant worked example)
### Tension acknowledged in spec
Per-bar Hold cost is per-bar reward shaping, which `pearl_event_driven_reward_density_alignment` warns against. Spec justifies as: (a) economically realistic carry cost, (b) ISV-bounded by controller, (c) inverse-direction from the pearl's failure mode (pulls policy AWAY from Hold-default, doesn't create exposure-positive bias). Faithful reward modeling, not artificial shaping.
### Files (atomic P0a commit)
15 files / 2316 lines added: 5 new kernels + 9 modifications + 1 mod.rs wire-up. Per `feedback_no_partial_refactor`: all consumers of the new ISV slots wired in this single commit.