fix: revert accumulator to device memory — mapped memory breaks read-modify-write

GPU read-modify-write (acc_buf[0] += loss) requires device memory.
Mapped host memory doesn't support atomic-free += from GPU kernels
on all architectures — H100 showed avg_grad=0.000000.

Reverted to CudaSlice<f32> with DtoH at epoch boundary (gated by
FOXHUNT_GPU_SYNC_DIAG). Removed __threadfence_system from accumulator
writes (not needed for device memory).

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

crates/ml/src/cuda_pipeline/gpu_training_guard.rs

@@ -188,9 +188,9 @@ pub struct GpuTrainingGuard {
     qdiv_has_prev: bool,
 
     /// Accumulator buffer (3 floats: loss_sum, grad_norm_sum, step_count).
-    /// Pinned device-mapped memory — GPU writes via dev_ptr, CPU reads via
-    /// read_volatile. Zero memcpy_dtoh sync at epoch boundary.
-    acc_mapped: MappedBuffer,
+    /// Device memory — GPU does read-modify-write (+=) which requires device
+    /// memory, not mapped host memory. DtoH at epoch boundary only.
+    acc_buf: CudaSlice<f32>,
 
     /// Welford running-mean accumulator for Q-value estimation (CPU-side, no GPU tensor).
     q_count: usize,
@@ -222,8 +222,10 @@ impl GpuTrainingGuard {
             .load_function("qvalue_divergence_check")
             .map_err(|e| MLError::ModelError(format!("qvalue_divergence_check load: {e}")))?;
 
-        // Allocate pinned device-mapped accumulator (zero DtoH sync at epoch boundary)
-        let acc_mapped = unsafe { MappedBuffer::new(3)? };
+        // Device memory accumulator — GPU does +=, DtoH at epoch boundary.
+        let acc_buf = stream.alloc_zeros::<f32>(3).map_err(|e| {
+            MLError::ModelError(format!("alloc acc_buf: {e}"))
+        })?;
 
         // Allocate mapped pinned double-buffers
         // Safety: CUDA context is active (we just loaded the module on this device).
@@ -244,7 +246,7 @@ impl GpuTrainingGuard {
             qdiv_mapped,
             qdiv_buf_idx: 0,
             qdiv_has_prev: false,
-            acc_mapped,
+            acc_buf,
             q_count: 0,
             q_mean_val: 0.0,
             stream,
@@ -316,7 +318,7 @@ impl GpuTrainingGuard {
         };
 
         // Pass raw device pointers to bypass cudarc event tracking on graph-captured buffers.
-        let acc_ptr = self.acc_mapped.dev_ptr;
+        let acc_ptr = self.acc_buf.raw_ptr();
         unsafe {
             self.stream
                 .launch_builder(&self.fused_check_accum_func)
@@ -340,19 +342,16 @@ impl GpuTrainingGuard {
 
     /// Read the epoch-boundary loss and grad_norm averages from the accumulator.
     ///
-    /// Zero-copy: reads from pinned device-mapped host memory via read_volatile.
-    /// Stream sync ensures all GPU writes are visible before CPU reads.
+    /// Read the epoch-boundary loss and grad_norm averages from the accumulator.
+    /// DtoH readback — device memory requires memcpy.
     pub fn read_accumulators(&mut self) -> Result<(f64, f64), MLError> {
-        // Epoch boundary: all training steps have completed by the time
-        // the training loop calls this. __threadfence_system() in the kernel
-        // ensures mapped memory is CPU-visible. Only sync in debug mode.
-        if std::env::var("FOXHUNT_GPU_SYNC_DIAG").map_or(true, |v| v != "0") {
-            unsafe { cudarc::driver::sys::cuStreamSynchronize(self.stream.cu_stream()); }
-        }
+        let mut host = [0.0_f32; 3];
+        self.stream.memcpy_dtoh(&self.acc_buf, &mut host)
+            .map_err(|e| MLError::ModelError(format!("acc_buf DtoH: {e}")))?;
 
-        let loss_sum = self.acc_mapped.read(0) as f64;
-        let grad_sum = self.acc_mapped.read(1) as f64;
-        let steps = self.acc_mapped.read(2) as f64;
+        let loss_sum = host[0] as f64;
+        let grad_sum = host[1] as f64;
+        let steps = host[2] as f64;
 
         if steps <= 0.0 {
             return Ok((0.0, 0.0));
@@ -363,11 +362,9 @@ impl GpuTrainingGuard {
 
     /// Zero the accumulator — write zeros to pinned host memory (visible to GPU).
     pub fn reset_accumulators(&mut self) -> Result<(), MLError> {
-        unsafe {
-            std::ptr::write_volatile(self.acc_mapped.host_ptr, 0.0_f32);
-            std::ptr::write_volatile(self.acc_mapped.host_ptr.add(1), 0.0_f32);
-            std::ptr::write_volatile(self.acc_mapped.host_ptr.add(2), 0.0_f32);
-        }
+        self.stream.memset_zeros(&mut self.acc_buf).map_err(|e| {
+            MLError::ModelError(format!("reset acc_buf: {e}"))
+        })?;
         Ok(())
     }
 

crates/ml/src/cuda_pipeline/training_guard_kernel.cu

@@ -85,7 +85,7 @@ extern "C" __global__ void training_guard_check_and_accumulate(
         acc_buf[1] += grad_norm;
     }
     acc_buf[2] += 1.0f;
-    __threadfence_system();  /* ensure CPU visibility of accumulator writes */
+    /* acc_buf is device memory — no threadfence needed (DtoH at epoch boundary). */
 }
 
 /* ------------------------------------------------------------------ */