refactor(F6/D7): contrarian Q-gap uses sign-flipped Q for consistent conviction

crates/ml/src/cuda_pipeline/experience_kernels.cu

@@ -1012,20 +1012,23 @@ 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.
      *
-     * D7/N7: During contrarian override, the action selected is argmin rather
-     * than argmax. The raw Q(best) - Q(Flat) gap would mislead the position
-     * sizer into full conviction for what's actually the "deliberately wrong"
-     * action. Zero out the gap so contrarian trades size minimal. */
+     * 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. */
     if (out_q_gaps != NULL) {
-        if (contrarian_active != 0) {
-            out_q_gaps[i] = 0.0f;
-        } else {
-            int flat_idx = 3; /* Flat = index 3 for direction(4): Short(0)/Hold(1)/Long(2)/Flat(3) */
-            float q_best = q_b0[argmax_n(q_b0, b0_size)];
-            float q_flat = q_b0[flat_idx];
-            float gap = q_best - q_flat;
-            out_q_gaps[i] = (gap > 0.0f) ? gap : 0.0f;
+        int flat_idx = 3; /* Flat = index 3 for direction(4): Short(0)/Hold(1)/Long(2)/Flat(3) */
+        float q_sign_local = (contrarian_active != 0) ? -1.0f : 1.0f;
+        /* Find best action under the (possibly negated) Q. */
+        int best_idx = 0;
+        float best_q = q_sign_local * q_b0[0];
+        for (int a = 1; a < b0_size; a++) {
+            float qa = q_sign_local * q_b0[a];
+            if (qa > best_q) { best_q = qa; best_idx = a; }
         }
+        float gap = best_q - (q_sign_local * q_b0[flat_idx]);
+        out_q_gaps[i] = (gap > 0.0f) ? gap : 0.0f;
     }
 }
 

crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs

@@ -7742,7 +7742,7 @@ impl GpuDqnTrainer {
 
         let n = ring.param_count();
         let num_bytes = n * std::mem::size_of::<f32>();
-        let mut new_weights = self.stream.alloc_zeros::<f32>(n)
+        let mut new_weights = unsafe { self.stream.alloc::<f32>(n) }
             .map_err(|e| MLError::ModelError(format!("snapshot alloc: {e}")))?;
         {
             let (src_ptr, _sg) = self.params_buf.device_ptr(&self.stream);

crates/ml/src/cuda_pipeline/q_snapshot.rs

@@ -88,7 +88,7 @@ impl SnapshotRing {
 
         let n = self.param_count;
         let num_bytes = n * std::mem::size_of::<f32>();
-        let mut new_weights = self.stream.alloc_zeros::<f32>(n)
+        let mut new_weights = unsafe { self.stream.alloc::<f32>(n) }
             .map_err(|e| MLError::ModelError(format!("snapshot alloc: {e}")))?;
         {
             let (src_ptr, _src_guard) = params_src.device_ptr(&self.stream);