feat(dqn): add Thompson + argmax-of-E[Q] device-inline functions

Mirrors the Phase 0 standalone test kernel's math into the production
kernel file. Used by Phase 2 direction-branch action selection.

crates/ml/src/cuda_pipeline/experience_kernels.cu

@@ -149,6 +149,68 @@ __device__ __forceinline__ int argmax_n(const float* arr, int n) {
 /* exposure_idx_to_fraction DELETED — replaced by compute_target_position()
  * from trade_physics.cuh which uses dynamic step = 2/(b0_size-1). */
 
+/* ================================================================== */
+/* Thompson sampling math (Plan C Phase 2)                            */
+/* ------------------------------------------------------------------ */
+/* Mirrors the math in thompson_test_kernel.cu so the production      */
+/* kernel can sample from C51/IQN posteriors and compute argmax-of-   */
+/* E[Q] for the direction branch. Single source of truth: keep these  */
+/* in lock-step with thompson_test_kernel.cu — Test 2.B verifies      */
+/* bit-for-bit equivalence between the two kernel files.              */
+/* No callers in this commit — Task 2 wires the direction-branch      */
+/* Thompson selector that consumes them.                              */
+/* ================================================================== */
+
+#ifndef N_IQN_QUANTILES
+#define N_IQN_QUANTILES 5
+#endif
+
+__device__ __forceinline__ float sample_c51_inverse_cdf(
+    const float* __restrict__ p,
+    const float* __restrict__ atom_values,
+    int n_atoms,
+    float u
+) {
+    float cum = 0.0f;
+    for (int a = 0; a < n_atoms; a++) {
+        cum += p[a];
+        if (u < cum) return atom_values[a];
+    }
+    return atom_values[n_atoms - 1];
+}
+
+__device__ __forceinline__ float sample_iqn_quantile_interp(
+    const float* __restrict__ q,
+    float u
+) {
+    const float taus[N_IQN_QUANTILES] = {0.05f, 0.25f, 0.50f, 0.75f, 0.95f};
+    if (u <= taus[0]) return q[0];
+    if (u >= taus[N_IQN_QUANTILES - 1]) return q[N_IQN_QUANTILES - 1];
+    for (int i = 0; i < N_IQN_QUANTILES - 1; i++) {
+        if (u <= taus[i + 1]) {
+            float w = (u - taus[i]) / (taus[i + 1] - taus[i]);
+            return q[i] * (1.0f - w) + q[i + 1] * w;
+        }
+    }
+    return q[N_IQN_QUANTILES - 1];
+}
+
+__device__ __forceinline__ float compute_e_c51_inline(
+    const float* __restrict__ p,
+    const float* __restrict__ atom_values,
+    int n_atoms
+) {
+    float e = 0.0f;
+    for (int a = 0; a < n_atoms; a++) e += p[a] * atom_values[a];
+    return e;
+}
+
+__device__ __forceinline__ float compute_e_iqn_inline(const float* __restrict__ q) {
+    float s = 0.0f;
+    for (int i = 0; i < N_IQN_QUANTILES; i++) s += q[i];
+    return s / (float)N_IQN_QUANTILES;
+}
+
 /* ================================================================== */
 /* Kernel 0a: domain_rand_episode_starts (GPU-native #1)              */
 /* ================================================================== */

docs/dqn-wire-up-audit.md

@@ -952,6 +952,24 @@ deleted (orphan after refactor; `feedback_wire_everything_up.md`). Tests 0.A
 and 0.B now run in ~0.15 s each on local RTX 3050 Ti (down from ~1.86 s and
 ~3.8 s respectively). Assertions unchanged. No production callers.
 
+Plan C Phase 2 Task 1 Thompson math device-inline funcs (2026-04-29):
+`cuda_pipeline/experience_kernels.cu` — copies the four `__device__
+__forceinline__` math primitives from the Phase 0 reference
+`thompson_test_kernel.cu` (`sample_c51_inverse_cdf`,
+`sample_iqn_quantile_interp`, `compute_e_c51_inline`,
+`compute_e_iqn_inline`) plus the `N_IQN_QUANTILES=5` `#ifndef`
+guard into the production kernel translation unit. Placed in the
+"Inline helpers" region between `argmax_n` and the first kernel
+(`domain_rand_episode_starts`). Bit-for-bit equivalent to the
+reference; Phase 2 Test 2.B verifies the production kernel matches
+`thompson_test_kernel.cu` exactly so the Phase 0 verification tests
+remain a valid oracle. The two `_inline` helpers re-name the reference's
+`compute_e_c51` / `compute_e_iqn` to make their intended usage explicit
+and reserve the un-suffixed names for any future non-inline callers.
+No callers in this commit — Task 2 wires the direction-branch Thompson
+selector inside `experience_action_select` that consumes them. Plan:
+`docs/superpowers/plans/2026-04-27-distributional-rl-thompson-plan-C-phase-2.md`.
+
 Persistent picked_action_history scatter (2026-04-26): `gpu_backtest_evaluator.rs`
 gains a window-major `picked_action_history_buf` (parallel layout to
 `actions_history_buf`) populated by an additional `scatter_intent_chunk`