diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 73a151da1..302c056d3 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -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) diff --git a/crates/ml/src/cuda_pipeline/apply_fixed_alpha_ema_kernel.cu b/crates/ml/src/cuda_pipeline/apply_fixed_alpha_ema_kernel.cu new file mode 100644 index 000000000..5895d2447 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/apply_fixed_alpha_ema_kernel.cu @@ -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 + +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(); +} diff --git a/crates/ml/src/cuda_pipeline/aux_dir_acc_reduce_kernel.cu b/crates/ml/src/cuda_pipeline/aux_dir_acc_reduce_kernel.cu new file mode 100644 index 000000000..f21bf1bda --- /dev/null +++ b/crates/ml/src/cuda_pipeline/aux_dir_acc_reduce_kernel.cu @@ -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 + +#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(); + } +} diff --git a/crates/ml/src/cuda_pipeline/aux_pred_to_isv_tanh_kernel.cu b/crates/ml/src/cuda_pipeline/aux_pred_to_isv_tanh_kernel.cu new file mode 100644 index 000000000..5bb2762dd --- /dev/null +++ b/crates/ml/src/cuda_pipeline/aux_pred_to_isv_tanh_kernel.cu @@ -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 + +#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(); + } +} diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 4c65cbe94..71d98bcce 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -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; } diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 12aec6739..aca348d4d 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -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::() 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::() 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 diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 0c4d27484..d8d511b74 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -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::() 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. diff --git a/crates/ml/src/cuda_pipeline/hold_rate_observer_kernel.cu b/crates/ml/src/cuda_pipeline/hold_rate_observer_kernel.cu new file mode 100644 index 000000000..4ab5a1aac --- /dev/null +++ b/crates/ml/src/cuda_pipeline/hold_rate_observer_kernel.cu @@ -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 + +#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(); + } +} diff --git a/crates/ml/src/cuda_pipeline/mod.rs b/crates/ml/src/cuda_pipeline/mod.rs index 414f51a3d..ab724841b 100644 --- a/crates/ml/src/cuda_pipeline/mod.rs +++ b/crates/ml/src/cuda_pipeline/mod.rs @@ -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, diff --git a/crates/ml/src/cuda_pipeline/sp13_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp13_isv_slots.rs new file mode 100644 index 000000000..955a811ec --- /dev/null +++ b/crates/ml/src/cuda_pipeline/sp13_isv_slots.rs @@ -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; diff --git a/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs index ce1765ccf..5e302cf2f 100644 --- a/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs @@ -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 = (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 = (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); } } diff --git a/crates/ml/src/cuda_pipeline/state_layout.cuh b/crates/ml/src/cuda_pipeline/state_layout.cuh index 517d217b7..67986ba4e 100644 --- a/crates/ml/src/cuda_pipeline/state_layout.cuh +++ b/crates/ml/src/cuda_pipeline/state_layout.cuh @@ -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), diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 53013bd97..1caa27b9e 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -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 diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 1000b692b..2ea6f882f 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -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 diff --git a/crates/ml/tests/sp13_phase0_oracle_tests.rs b/crates/ml/tests/sp13_phase0_oracle_tests.rs new file mode 100644 index 000000000..7073dfa59 --- /dev/null +++ b/crates/ml/tests/sp13_phase0_oracle_tests.rs @@ -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 { + 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) -> 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, + 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::() 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 = vec![]; + let label: Vec = 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) -> 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, + 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::() 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 = 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 = 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 = 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) -> 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, + 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) -> 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, + 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::() 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 + ); +} diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index d97b79cb5..3f0614b6d 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -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.