diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index d227fda07..1d008daad 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -1250,22 +1250,32 @@ extern "C" __global__ void experience_action_select( * Q-gap = Q(best_direction) - Q(Flat). Higher = more conviction. * Used by env_step to scale target_position. * - * F6: Compute Q-gap with the same sign convention used by the action-selection - * softmax. When contrarian_active, this means the "best" action is the argmin - * of raw Q (equivalent to argmax of -Q). The conviction magnitude then reflects - * how confident the contrarian direction is — useful for position sizing - * consistently with the action actually being taken. */ + * Plan C Phase 2 Task 2 follow-up: q_gap reads `e_dir` (joint + * E[C51] + E[IQN]) — the same source the conviction block below + * consumes and the same source eval-mode direction argmax uses. + * Both signals feed env_step's Kelly-adjacent position sizing, so + * sourcing them from a shared scale keeps the downstream consumer + * coherent (per feedback_no_partial_refactor: shared-contract + * consumers migrate together). + * + * Contrarian sign flip dropped on this read. The Thompson direction + * branch above does NOT apply `q_sign` to `e_dir`; it picks + * `argmax(e_dir)` (eval) or argmax of joint Thompson sample + * (training) unconditionally. Symmetry with that selection logic — + * and with the conviction block — requires the gap math to follow + * suit. The legacy `q_sign_local` flip applied only when q_gap was + * sourced from `q_b0` and the direction-Boltzmann path still + * negated raw Q; both are unreachable on the Thompson path. */ if (out_q_gaps != NULL) { int flat_idx = DIR_FLAT; /* Flat direction's Q-value index within the 4-way direction branch */ - float q_sign_local = (contrarian_active != 0) ? -1.0f : 1.0f; - /* Find best action under the (possibly negated) Q. */ + /* Find best direction under the joint E[Q] posterior. */ int best_idx = 0; - float best_q = q_sign_local * q_b0[0]; + float best_q = e_dir[0]; for (int a = 1; a < b0_size; a++) { - float qa = q_sign_local * q_b0[a]; + float qa = e_dir[a]; if (qa > best_q) { best_q = qa; best_idx = a; } } - float gap = best_q - (q_sign_local * q_b0[flat_idx]); + float gap = best_q - e_dir[flat_idx]; out_q_gaps[i] = (gap > 0.0f) ? gap : 0.0f; } diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index c31747a67..ce519b245 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -1003,8 +1003,10 @@ need to retune if a systematic mismatch surfaces. The cold-start Magnitude (Branch 1), order (Branch 2), and urgency (Branch 3) selection paths are untouched — they still read `q_b0`/`q_b1`/`q_b2`/ `q_b3` from the cuBLAS Q-head and apply the existing Boltzmann-with- -adaptive-tau path. The `q_b0`-based `out_q_gaps` writer below is also -unchanged (still computes gap from cuBLAS direction Q, not `e_dir`). +adaptive-tau path. Partial-refactor follow-up (same date): `out_q_gaps` +also migrated to read `e_dir` (joint E[C51]+E[IQN]) — both Kelly-adjacent +signals feeding env_step now share scale; contrarian sign-flip dropped +symmetric with the unflipped Thompson direction selection. Until Tasks 3 (`gpu_dqn_trainer.rs`) and 4 (`gpu_backtest_evaluator .rs`) land, the kernel signature is intentionally non-callable end-