THE K=3 HEAD NOW REACHES THE POLICY. Phase C-2 is the consumer-side
flip — slot 121's semantic changes from K=2's recentered p_up to K=3's
p_Profit, and slots [122, 123] gain new meaning as (p_Stop, p_Timeout).
Architectural constraint: aux_dir_prob_per_env (K=2 buffer) is ALSO
consumed by experience_env_step's β reward at lines 2335 + 3724-3725.
Cannot repurpose that pointer to point at the K=3 [N, 3] buffer —
env_step β consumer would read wrong-stride memory. This commit adds
SEPARATE arg aux_outcome_probs_per_env [N, 3] to state_gather kernels
with NULL fallback.
Branching semantic at state_gather read site:
- aux_outcome_probs_per_env != NULL → K=3 active path: writes slots
[121..124) via new assemble_state_outcome_k3 helper
- aux_outcome_probs_per_env == NULL → K=2 fallback: writes slot 121
via existing assemble_state from the recentered scalar
Changes:
- state_layout.cuh: NEW __device__ helper assemble_state_outcome_k3
— mirrors assemble_state except padding slots [121..124) get
p_Profit/p_Stop/p_Timeout (raw softmax probs [0, 1]), slots
[124..128) zero for 8-alignment.
- experience_kernels.cu: training-side experience_state_gather +
eval-side backtest_state_gather both get new trailing arg
aux_outcome_probs_per_env (NULL-tolerant). Read site branches:
K=3 reads 3 floats / env → assemble_state_outcome_k3; K=2 fallback
preserves legacy assemble_state call.
- gpu_experience_collector.rs: training launcher passes
self.prev_aux_outcome_probs.raw_ptr() → K=3 active in training.
- gpu_backtest_evaluator.rs: eval launcher passes NULL → K=2
fallback in eval (eval has no aux producer infra yet).
K=2 head still alive:
- prev_aux_dir_prob still populated by aux_softmax_to_per_env_kernel
- experience_env_step still reads it for β reward (independent
consumer untouched)
- EGF chain still reads exp_aux_nb_softmax_buf
- Only K=2's slot 121 contribution to policy state is suppressed
End-to-end K=3 chain now active:
Label producer (A2) → per-(env, t) ring (B4b-1) → replay buffer
scatter (B4b-2) → PER direct gather → trainer aux_to_label_buf
→ loss_reduce (B4) sparse CE on real labels → backward (B4)
per-sample partials → Adam SAXPY (B1+B4) updates W1/b1/W2/b2 at
[163..167)
PARALLEL:
Collector rollout K=3 forward (B3) → softmax tile → C-1 producer
→ prev_aux_outcome_probs [N, 3] → C-2 state gather → state[121..124]
→ policy reads in next step
The K=3 head closes the loop: learns from real labels via replay,
AND predictions reach policy via state assembly. Trainable +
observable.
Verification:
- cargo check -p ml clean.
- cargo test -p ml --lib → 1016/0 green.
Remaining vNext work:
- Phase D: 12-weight W atom-shift (4 actions × 3 outcomes)
- Phase E: dW backward + Adam for W[4, 3]
- Phase F: Validation smoke at structural prior
- B5b (deferred): plan_params input concat
Audit: docs/dqn-wire-up-audit.md Phase C-2 section.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
53462a28d9