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

feat(sp22): H6 Phase 2 — recenter state[121] to [-1, +1] (atomic)

Browse Source 
Phase 1 post-mortem traced an actual `pearl_first_observation_bootstrap`
violation in my own H6 implementation: state slot 121 wrote
`aux_softmax[env, 1] = p_up ∈ [0, 1]` with sentinel 0.5, but every
OTHER state slot uses 0 as the "no signal" baseline (zero-padding,
feature_mask, ofi-missing, mtf-missing). The encoder had to learn TWO
things about slot 121 (directional mapping + non-zero bias offset)
instead of one. Phase 2 fixes the encoding to match the project
convention BEFORE declaring H6 fully falsified.

Mechanism change
────────────────
- `aux_softmax_to_per_env_kernel.cu` writes `2*p_up - 1 ∈ [-1, +1]`
  instead of `p_up`. Still structurally bounded (softmax components
  in [0, 1] sum to 1).
- Cold-start + FoldReset sentinel: 0.5 → 0.0 via the same pure-GPU
  `fill_f32` path. No HtoD per
  `feedback_no_htod_htoh_only_mapped_pinned`.
- NULL-fallback in 3 state-gather kernels (training +
  backtest-per-step + backtest-chunk): 0.5f → 0.0f.
- Constant + device-function comment updates to document the
  recentered encoding.

Atomic per `feedback_no_partial_refactor`: the encoding contract
spans 5 source files; partial migration produces inconsistent slot
semantics between training and eval.

Verification gates (all clean)
──────────────────────────────
- cargo check -p ml --features cuda: 0 errors, 21 pre-existing warnings
  (parity with Phase 1 baseline)
- gpu_backtest_validation: 4/4 expected-passing tests still pass; 2
  pre-existing PnL-assertion failures bit-identical to Phase 1
  (confirms recentering does not perturb scripted-policy paths)
- compute-sanitizer --tool=memcheck: ERROR SUMMARY: 0 errors

Smoke dispatch deferred pending an orthogonal investigation into the
2 pre-existing gpu_backtest_validation failures (stale action
constants in the tests; addressed in a follow-up commit, NOT a
Phase 2 regression).

Verdict criteria (per spec, evaluated after smoke)
──────────────────────────────────────────────────
- WR > 50.5% within 3 epochs → recentering binding, H6 + Phase 2
  sufficient → justify A2.
- a_var for mag/ord/urg > 1e-3 → sub-branches gradient-coupled under
  recentered signal.
- WR pinned at 50.1–50.2% → Phase 2 falsified, pivot to amplitude
  scaling or deeper hypothesis.

Refs
────
- docs/plans/2026-05-12-sp22-h6-phase2-recenter.md (spec)
- docs/plans/2026-05-12-sp22-h6-phase2-recenter-runbook.md (this plan)
- pearl_first_observation_bootstrap (sentinel = 0)
- feedback_no_partial_refactor (5-file atomic)
- feedback_no_htod_htoh_only_mapped_pinned (fill_f32, not HtoD)

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-12 22:50:55 +02:00
parent cffc95152c
commit f9192f70a5
6 changed files with 148 additions and 56 deletions
Download Patch File Download Diff File Expand all files Collapse all files

15
crates/ml-core/src/state_layout.rs
View File

@@ -55,11 +55,16 @@ pub const PORTFOLIO_START: usize = MTF_START + MTF_DIM; // 106
pub const PLAN_ISV_START: usize = PORTFOLIO_START + PORTFOLIO_BASE_DIM; // 114
pub const PADDING_START: usize = PLAN_ISV_START + PORTFOLIO_PLAN_DIM; // 121
// SP22 H6 (2026-05-12) — aux directional probability from previous step.
// Lives in the first padding slot. PADDING_DIM stays 7 (no STATE_DIM change);
// `assemble_state` writes aux_dir_prob to [AUX_DIR_PROB_INDEX] and zeros the
// remaining 6 padding slots [PADDING_START+1..STATE_DIM).
// Sentinel = 0.5 (neutral; p_up = 50%) at cold-start and FoldReset.
// SP22 H6 Phase 2 (2026-05-12) — aux directional probability from previous
// step, RECENTERED. Lives in the first padding slot. PADDING_DIM stays 7
// (no STATE_DIM change); `assemble_state` writes aux_dir_prob to
// [AUX_DIR_PROB_INDEX] and zeros the remaining 6 padding slots
// [PADDING_START+1..STATE_DIM).
// Encoding: `2*p_up - 1` ∈ [-1, +1] (structurally bounded — softmax
// components in [0, 1] sum to 1). +1 = up with full conviction;
// -1 = down with full conviction; 0 = neutral.
// Sentinel = 0.0 at cold-start and FoldReset (matches every other state
// slot's "no signal = 0" baseline per `pearl_first_observation_bootstrap`).
pub const AUX_DIR_PROB_INDEX: usize = PADDING_START; // 121
const _: () = assert!(PADDING_START + PADDING_DIM == STATE_DIM, "layout must sum to STATE_DIM");

25
crates/ml/src/cuda_pipeline/aux_softmax_to_per_env_kernel.cu
View File

@@ -29,11 +29,15 @@
//
// Sentinel / cold-start / FoldReset
// ─────────────────────────────────
// The buffer is filled with 0.5 (neutral; p_up = 50%) at construction and at
// every FoldReset by the GPU `fill_f32` kernel (epsilon_greedy_kernel.cu).
// This kernel ALWAYS overwrites — no sentinel branch needed here. The first
// rollout step's state assembly reads 0.5 (no aux signal yet), then this
// kernel writes the real value for step 2's state, and so on.
// SP22 H6 Phase 2 (2026-05-12): buffer is filled with 0.0 (neutral
// under the recentered `2*p - 1` encoding; aux contributes nothing) at
// construction and at every FoldReset by the GPU `fill_f32` kernel
// (epsilon_greedy_kernel.cu). Pre-Phase-2 the sentinel was 0.5, which
// violated `pearl_first_observation_bootstrap` (every other state slot
// uses 0 as the "no signal" baseline). This kernel ALWAYS overwrites
// — no sentinel branch needed here. The first rollout step's state
// assembly reads 0.0 (no aux signal yet), then this kernel writes the
// real recentered value for step 2's state, and so on.
//
// Discipline
// ──────────
@@ -74,6 +78,13 @@ extern "C" __global__ void aux_softmax_to_per_env_kernel(
int env = blockIdx.x * blockDim.x + threadIdx.x;
if (env >= n_envs) return;
/* Write p_up = softmax[1]. Used by next step's state assembly. */
prev_aux_dir_prob[env] = aux_softmax[(size_t)env * (size_t)K + 1];
/* SP22 H6 Phase 2 (2026-05-12): write RECENTERED p_up = `2*p - 1`
* so slot 121 uses the same "no signal = 0" baseline as every
* other state slot, per `pearl_first_observation_bootstrap`.
* Range: [-1, +1] (softmax components non-negative and sum to 1,
* so 2*p - 1 is structurally bounded). Sentinel 0 = neutral
* (aux contributes nothing); +1 = up with full conviction;
* -1 = down with full conviction. */
prev_aux_dir_prob[env] =
2.0f * aux_softmax[(size_t)env * (size_t)K + 1] - 1.0f;
}

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

@@ -652,11 +652,14 @@ extern "C" __global__ void experience_state_gather(
const float* __restrict__ isv_signals_ptr, /* [12] pinned device-mapped ISV signals. NULL = static. */
const float* __restrict__ ofi_features, /* [total_bars, ofi_dim] OFI features. NULL = no OFI. */
int ofi_dim, /* OFI feature dimension (20 with deltas+book+dur). 0 = no OFI. */
/* SP22 H6 (2026-05-12): per-env aux directional probability from previous step.
* Length [N]; populated by `aux_softmax_to_per_env_kernel` after each aux forward.
* NULL = no aux signal yet (cold-start or producer not wired) → kernel uses
* 0.5 sentinel (neutral). See `state_layout.cuh::assemble_state` slot
* `SL_PADDING_START+0`. */
/* SP22 H6 Phase 2 (2026-05-12): per-env recentered aux directional
* probability from previous step. Encoding `2*p_up - 1 ∈ [-1, +1]`.
* Length [N]; populated by `aux_softmax_to_per_env_kernel` after each
* aux forward. NULL = no aux signal yet (cold-start or producer not
* wired) → kernel uses 0.0f sentinel (neutral; matches every other
* state slot's "no signal = 0" baseline per
* `pearl_first_observation_bootstrap`). See
* `state_layout.cuh::assemble_state` slot `SL_PADDING_START+0`. */
const float* __restrict__ aux_dir_prob_per_env
) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
@@ -974,13 +977,16 @@ extern "C" __global__ void experience_state_gather(
ofi[k] = ofi_row[k];
}
/* ── SP22 H6 aux→state bridge: previous-step p_up for this env ──
* NULL = no aux producer wired yet (cold-start, eval A3 fallback) → 0.5
* sentinel (neutral, no signal). Lag of 1 bar is intentional — H=200
* label horizon makes the signal slow-moving, so one-step lag is fine. */
/* ── SP22 H6 aux→state bridge: previous-step recentered p_up for this env ──
* Phase 2 (2026-05-12) encoding: `2*p_up - 1 ∈ [-1, +1]`; sentinel 0
* = neutral (matches every other state slot's "no signal = 0"
* baseline per `pearl_first_observation_bootstrap`). NULL = no aux
* producer wired yet (cold-start, eval A3 fallback) → 0.0f. Lag of
* 1 bar is intentional — H=200 label horizon makes the signal
* slow-moving, so one-step lag is fine. */
float aux_dir_prob = (aux_dir_prob_per_env != NULL)
? aux_dir_prob_per_env[i]
: 0.5f;
: 0.0f;
/* ── Final assembly: canonical layout via state_layout.cuh ── */
assemble_state(out, market, ofi, mtf, portfolio, plan_isv, aux_dir_prob);
@@ -8648,9 +8654,12 @@ extern "C" __global__ void backtest_state_gather(
int feat_dim, /* = market_dim + ofi_dim = 62 */
int current_step,
int padded_sd,
/* SP22 H6 (2026-05-12): per-window aux directional probability from
* previous step. Length [n_windows]. NULL = eval-time A3 fallback
* (Phase 1 — eval has no aux infrastructure yet) → 0.5 sentinel. */
/* SP22 H6 Phase 2 (2026-05-12): per-window recentered aux directional
* probability from previous step. Encoding `2*p_up - 1 ∈ [-1, +1]`.
* Length [n_windows]. NULL = eval-time A3 fallback (eval has no aux
* infrastructure yet) → 0.0f sentinel (neutral; matches every other
* state slot's "no signal = 0" baseline per
* `pearl_first_observation_bootstrap`). */
const float* __restrict__ aux_dir_prob_per_env
) {
int w = blockIdx.x * blockDim.x + threadIdx.x;
@@ -8727,10 +8736,12 @@ extern "C" __global__ void backtest_state_gather(
for (int k = 0; k < SL_PORTFOLIO_PLAN_DIM; k++)
pisv[k] = plan_isv[w * SL_PORTFOLIO_PLAN_DIM + k];
/* SP22 H6 — per-window aux p_up (NULL → 0.5 sentinel for Phase 1 A3 eval). */
/* SP22 H6 Phase 2 — per-window recentered aux p_up `2*p - 1 ∈ [-1, +1]`
* (NULL → 0.0f neutral sentinel for A3 eval; matches other state
* slots' "no signal = 0" baseline). */
float aux_dir_prob = (aux_dir_prob_per_env != NULL)
? aux_dir_prob_per_env[w]
: 0.5f;
: 0.0f;
/* Single shared assembly call — produces identical layout to training */
assemble_state(out, market, ofi, mtf, pf, pisv, aux_dir_prob);
@@ -8776,10 +8787,11 @@ extern "C" __global__ void backtest_state_gather_chunk(
int start_step,
int chunk_len,
int padded_sd,
/* SP22 H6 (2026-05-12): per-window aux directional probability.
* Constant within a chunk (matches portfolio/plan_isv constancy
* semantics — env updates aux only at chunk boundary, same as
* portfolio). Length [n_windows]. NULL = A3 fallback → 0.5. */
/* SP22 H6 Phase 2 (2026-05-12): per-window recentered aux directional
* probability. Encoding `2*p_up - 1 ∈ [-1, +1]`. Constant within a
* chunk (matches portfolio/plan_isv constancy semantics — env updates
* aux only at chunk boundary, same as portfolio). Length [n_windows].
* NULL = A3 fallback → 0.0f neutral sentinel. */
const float* __restrict__ aux_dir_prob_per_env
) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
@@ -8869,10 +8881,11 @@ extern "C" __global__ void backtest_state_gather_chunk(
for (int k = 0; k < SL_PORTFOLIO_PLAN_DIM; k++)
pisv[k] = plan_isv[w * SL_PORTFOLIO_PLAN_DIM + k];
/* SP22 H6 — per-window aux p_up (NULL → 0.5 for Phase 1 A3 eval). */
/* SP22 H6 Phase 2 — per-window recentered aux p_up `2*p - 1 ∈ [-1, +1]`
* (NULL → 0.0f neutral sentinel for A3 eval). */
float aux_dir_prob = (aux_dir_prob_per_env != NULL)
? aux_dir_prob_per_env[w]
: 0.5f;
: 0.0f;
/* Single shared assembly call — produces identical layout to training */
assemble_state(out, market, ofi, mtf, pf, pisv, aux_dir_prob);

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

@@ -1189,14 +1189,18 @@ pub struct GpuExperienceCollector {
/// softmax diff into ISV[375]), `exp_sp14_q_disagreement_update_kernel`
/// (vs Q-argmax, into ISV[383/384/389]).
exp_aux_nb_softmax_buf: cudarc::driver::CudaSlice<f32>,
/// SP22 H6 (2026-05-12) — per-env aux directional probability cache.
/// Sized `[alloc_episodes]` f32 (one slot per env). Written end-of-step
/// by `exp_aux_softmax_to_per_env_kernel` (gathers `aux_softmax[env, 1]`
/// = p_up). Read start-of-next-step by `experience_state_gather` for
/// SP22 H6 Phase 2 (2026-05-12) — per-env aux directional probability
/// cache, RECENTERED encoding. Sized `[alloc_episodes]` f32 (one slot
/// per env). Written end-of-step by `exp_aux_softmax_to_per_env_kernel`
/// (gathers `2*aux_softmax[env, 1] - 1` ∈ [-1, +1]). Read start-of-next
/// step by `experience_state_gather` for
/// `state[AUX_DIR_PROB_INDEX = SL_PADDING_START + 0 = 121]`. Cold-start
/// and FoldReset both initialize to 0.5 (neutral; p_up = 50%) via the
/// `fill_f32` GPU kernel — no HtoD per
/// `feedback_no_htod_htoh_only_mapped_pinned.md`.
/// and FoldReset both initialize to 0.0 (neutral; aux contributes
/// nothing) via the `fill_f32` GPU kernel — no HtoD per
/// `feedback_no_htod_htoh_only_mapped_pinned.md`. Pre-Phase-2 the
/// encoding was raw `p_up ∈ [0, 1]` with sentinel 0.5; recentered for
/// gradient-init parity with every other state slot per
/// `pearl_first_observation_bootstrap`.
prev_aux_dir_prob: cudarc::driver::CudaSlice<f32>,
/// SP22 H6 (2026-05-12) — copy kernel handle for
/// `aux_softmax_to_per_env_kernel`. Loaded from
@@ -2609,11 +2613,15 @@ impl GpuExperienceCollector {
alloc_episodes
)))?;
{
// Cold-start fill: write 0.5 sentinel to every slot. Pure GPU
// compute via `fill_f32` (no HtoD). Matches the launch shape used
// by `GpuActionSelector::run_branching` for the epsilon buffer.
// SP22 H6 Phase 2 (2026-05-12) cold-start fill: write 0.0
// sentinel (neutral under the recentered `2*p - 1` encoding;
// matches the "no signal = 0" baseline used by every other
// state slot per `pearl_first_observation_bootstrap`). Pure
// GPU compute via `fill_f32` (no HtoD). Matches the launch
// shape used by `GpuActionSelector::run_branching` for the
// epsilon buffer.
let n_envs_i32 = alloc_episodes as i32;
let sentinel: f32 = 0.5;
let sentinel: f32 = 0.0;
let blocks = ((alloc_episodes as u32).div_ceil(256)).max(1);
let prev_aux_dir_prob_ptr = prev_aux_dir_prob.raw_ptr();
unsafe {
@@ -5292,9 +5300,10 @@ impl GpuExperienceCollector {
let isv_ptr = self.isv_signals_dev_ptr;
let ofi_ptr = self.ofi_gpu.dev_ptr;
let ofi_dim_i32 = self.ofi_dim as i32;
// SP22 H6 (2026-05-12): per-env aux p_up cache (filled by
// SP22 H6 Phase 2 (2026-05-12): per-env recentered aux p_up
// cache (range [-1, +1]; filled by
// `exp_aux_softmax_to_per_env_kernel` at the END of the prior
// step; sentinel 0.5 at cold-start / FoldReset via fill_f32).
// step; sentinel 0.0 at cold-start / FoldReset via fill_f32).
let prev_aux_dir_prob_ptr = self.prev_aux_dir_prob.raw_ptr();
self.stream
@@ -7108,15 +7117,16 @@ impl GpuExperienceCollector {
upload_host_to_cuda_f32(
&self.stream, &portfolio_init, &mut self.portfolio_states, "portfolio_states reset")?;
// SP22 H6 (2026-05-12): FoldReset of the per-env aux directional
// probability cache. Re-seed every slot to the 0.5 neutral sentinel
// so the first state-gather call of the new fold reads "no aux
// signal yet" instead of the previous fold's last p_up. Pure GPU
// compute via fill_f32 — no HtoD per
// SP22 H6 Phase 2 (2026-05-12): FoldReset of the per-env aux
// directional probability cache. Re-seed every slot to the 0.0
// neutral sentinel (under the recentered `2*p - 1` encoding) so
// the first state-gather call of the new fold reads "no aux
// signal yet" with the same zero baseline used by every other
// state slot. Pure GPU compute via fill_f32 — no HtoD per
// `feedback_no_htod_htoh_only_mapped_pinned.md`.
{
let n_envs_i32 = self.alloc_episodes as i32;
let sentinel: f32 = 0.5;
let sentinel: f32 = 0.0;
let blocks = ((self.alloc_episodes as u32).div_ceil(256)).max(1);
let prev_aux_dir_prob_ptr = self.prev_aux_dir_prob.raw_ptr();
unsafe {

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

@@ -655,10 +655,12 @@ __device__ __forceinline__ void assemble_state(
for (int k = 0; k < SL_PORTFOLIO_PLAN_DIM; k++)
out[SL_PLAN_ISV_START + k] = plan_isv[k];
// Padding [121..128) — SP22 H6 (2026-05-12): slot 0 = aux directional
// probability (p_up from previous step's aux softmax); slots 1..7 zero
// for 8-alignment. Caller passes 0.5 sentinel at cold-start / FoldReset
// and for eval-time NULL fallback (A3).
// Padding [121..128) — SP22 H6 Phase 2 (2026-05-12): slot 0 = aux
// directional probability from previous step's aux softmax, RECENTERED
// to `2*p_up - 1 ∈ [-1, +1]` (matches every other state slot's
// "no signal = 0" baseline per `pearl_first_observation_bootstrap`).
// Slots 1..7 zero for 8-alignment. Caller passes 0.0 sentinel at
// cold-start / FoldReset and for eval-time NULL fallback (A3).
out[SL_PADDING_START + 0] = aux_dir_prob;
for (int k = 1; k < SL_PADDING_DIM; k++)
out[SL_PADDING_START + k] = 0.0f;

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

@@ -16523,3 +16523,54 @@ with healthy rewards.
If H3 lands and the WR moves but action-space stays narrow, V/A is
binding and gets fixed next. If H3 doesn't move WR either, V/A goes
first and we re-test in combination.
### Phase 2 (2026-05-12) — recenter state[121] to [-1, +1]
Post-mortem of the Phase 1 smoke verdict surfaced an actual pearl
violation in the H6 implementation itself, BEFORE pivoting to H3.
Per `pearl_first_observation_bootstrap`: "sentinel = 0; first
observation replaces directly." The Phase 1 design wrote
`state[121] = aux_softmax[env, 1] = p_up ∈ [0, 1]` with sentinel 0.5.
**Every other state slot uses 0 as its "no signal" baseline**
(zero-padding, feature_mask, ofi-missing, mtf-missing). Slot 121 alone
was off-pattern, forcing the encoder to learn TWO things — directional
mapping AND the non-zero bias offset — instead of one.
Phase 2 fix (atomic commit, 5 source files):
- `aux_softmax_to_per_env_kernel.cu` writes `2*p_up - 1 ∈ [-1, +1]`
(still structurally bounded since softmax components sum to 1)
- `gpu_experience_collector.rs` cold-start + FoldReset sentinel
`0.5 → 0.0`
- `experience_kernels.cu` NULL-fallback in three state-gather kernels
`0.5f → 0.0f`
- `state_layout.rs` + `state_layout.cuh` comment updates documenting
the recentered encoding and zero sentinel
Verification gates (all clean before commit):
| Gate | Result |
|---|---|
| `cargo check -p ml --features cuda` | 0 errors, 21 pre-existing warnings (parity with Phase 1 baseline) |
| `gpu_backtest_validation` | 4/4 expected-passing pass; 2 pre-existing PnL-assertion failures **bit-identical to Phase 1** (0.00023627281, 0.0001718998 vs 0.00021278858 — confirms recentering does not perturb scripted-policy paths) |
| `compute-sanitizer --tool=memcheck` | `ERROR SUMMARY: 0 errors` |
Smoke dispatch: deferred pending an orthogonal investigation into the
2 pre-existing `gpu_backtest_validation` PnL failures (separate
section below — those are stale action constants in the tests, not a
Phase 2 regression).
Verdict criteria (per spec, to be evaluated post-smoke):
- **WR > 50.5%** within 3 epochs → Phase 1 + Phase 2 sufficient;
recentering was the binding constraint; justify A2 (eval-side aux
integration) for production parity.
- **`a_var` for mag/ord/urg moves off 0** (> 1e-3) → secondary
success signal that sub-branches are gradient-coupled under
the recentered signal.
- **WR pinned at 50.1–50.2%** → Phase 2 falsified. The encoder
successfully consumes a recentered signal but still can't extract
directional alpha from one state slot in 3 epochs. Pivot to
amplitude scaling (multiply state[121] write by a scalar > 1) or
deeper hypothesis.