dist: coalesce grads into buckets for all-reduce (KI-5)

Replace the per-parameter eager all-reduce (~150 tiny serial NCCL calls
for dim512, DDP's dominant cost after T10's batched forward) with a
coalesced bucketed all-reduce: pack grads into a few large contiguous
scratch buffers, all-reduce each bucket once (fused via ncclGroupStart/
End), fold the 1/world average into one per-bucket scale, unpack back.

The packed buffer is the concatenation of the grad tensors, so NCCL's
element-wise sum over a bucket equals the per-tensor sums — bit-identical
to the un-bucketed path; only launch/latency overhead is removed. DDP
cross-rank param identity + loss-match are preserved.

Adds xtrain_cuda::device::copy_d2d (cudaMemcpy D2D) for the pack/unpack.

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

crates/xtrain-cuda/src/device.rs

@@ -14,3 +14,15 @@ pub fn set_device(device: u32) -> Result<()> {
 pub fn synchronize() -> Result<()> {
     error::check(unsafe { ffi::cudaDeviceSynchronize() })
 }
+
+/// Device-to-device copy of `count` bytes (`dst <- src`) on the same GPU. Issued
+/// on the null stream (like every other xtrain kernel), so it orders with the
+/// surrounding work. Used by the DDP bucketed all-reduce to pack/unpack grads
+/// into a flat scratch buffer.
+///
+/// # Safety
+/// `dst`/`src` must point to at least `count` valid bytes of device memory on the
+/// current device, with no overlap.
+pub unsafe fn copy_d2d(dst: *mut u8, src: *const u8, count: usize) -> Result<()> {
+    error::check(unsafe { ffi::cudaMemcpy(dst, src, count, ffi::CUDA_MEMCPY_D2D) })
+}

crates/xtrain-cuda/src/ffi.rs

@@ -5,6 +5,7 @@ pub type CudaStream = *mut c_void;
 
 pub const CUDA_MEMCPY_H2D: i32 = 1;
 pub const CUDA_MEMCPY_D2H: i32 = 2;
+pub const CUDA_MEMCPY_D2D: i32 = 3;
 
 pub const CUDA_SUCCESS: i32 = 0;
 pub const CUDA_ERROR_OUT_OF_MEMORY: i32 = 2;