perf: streams / drop per-op sync

Default-stream kernels run in order and every host read goes through a
stream-ordered cudaMemcpy (to_device), so the per-op cudaDeviceSynchronize
after each kernel was pure overhead — remove all 21 in tensor.rs. Host
data is still correctly ordered by the D2H memcpy that reads it.

Also zero op-output buffers with cudaMemset (device-side, async) instead of
a blocking H2D memcpy of a host zero buffer on every allocation — that
copy was itself a hidden per-op sync point.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

crates/xtrain-cuda/src/ffi.rs

@@ -19,6 +19,7 @@ unsafe extern "C" {
     pub fn cudaMalloc(devptr: *mut *mut u8, size: usize) -> i32;
     pub fn cudaFree(devptr: *mut u8) -> i32;
     pub fn cudaMemcpy(dst: *mut u8, src: *const u8, count: usize, kind: i32) -> i32;
+    pub fn cudaMemset(devptr: *mut u8, value: i32, count: usize) -> i32;
 
     // --- Error ---
     pub fn cudaGetErrorString(error: i32) -> *const c_char;

crates/xtrain-cuda/src/memory.rs

@@ -46,6 +46,12 @@ impl GpuBuffer {
             ffi::cudaMemcpy(dst.as_mut_ptr(), self.ptr, dst.len(), ffi::CUDA_MEMCPY_D2H)
         })
     }
+
+    /// Set every byte of the buffer to `value` on the device (no host copy).
+    /// Used to zero op-output buffers without a blocking H2D memcpy of zeros.
+    pub fn memset(&mut self, value: u8) -> Result<()> {
+        error::check(unsafe { ffi::cudaMemset(self.ptr, value as i32, self.len) })
+    }
 }
 
 impl Drop for GpuBuffer {