tensor: add scale elementwise CUDA kernel + FFI

New csrc/ops/elementwise.cu (out[i]=in[i]*alpha), compiled by
xtrain-cuda/build.rs and exposed via launch_scale_f32 FFI, gated behind
not(no_cuda) like the existing vecadd smoke test.

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

crates/xtrain-cuda/build.rs

@@ -28,6 +28,7 @@ fn main() {
         .cudart("shared")
         .flag("-gencode=arch=compute_120,code=sm_120")
         .file("../../csrc/test/vecadd.cu")
+        .file("../../csrc/ops/elementwise.cu")
         .compile("xtrain_cuda_kernels");
 }
 

crates/xtrain-cuda/src/ffi.rs

@@ -24,9 +24,19 @@ unsafe extern "C" {
     pub fn cudaGetErrorString(error: i32) -> *const c_char;
 }
 
-// The vector-add smoke-test kernel, compiled from csrc/test/vecadd.cu by build.rs.
-// Only linked when CUDA is actually compiled (i.e. nvcc was present).
+// GPU kernels compiled from csrc/ by build.rs. Only linked when CUDA is
+// actually compiled (i.e. nvcc was present).
 #[cfg(not(no_cuda))]
 unsafe extern "C" {
+    // Vector-add smoke test (csrc/test/vecadd.cu).
     pub fn launch_vecadd_f32(a: *const f32, b: *const f32, c: *mut f32, n: i32, stream: CudaStream);
+
+    // Elementwise scale: out[i] = in[i] * alpha (csrc/ops/elementwise.cu).
+    pub fn launch_scale_f32(
+        input: *const f32,
+        out: *mut f32,
+        alpha: f32,
+        n: i32,
+        stream: CudaStream,
+    );
 }

csrc/ops/elementwise.cu

@@ -0,0 +1,17 @@
+extern "C" {
+
+// out[i] = in[i] * alpha   (in-place safe: out may alias in)
+__global__ void scale_f32(const float* in, float* out, float alpha, int n) {
+    int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    if (idx < n) {
+        out[idx] = in[idx] * alpha;
+    }
+}
+
+void launch_scale_f32(const float* in, float* out, float alpha, int n, void* stream) {
+    int block = 256;
+    int grid = (n + block - 1) / block;
+    scale_f32<<<grid, block, 0, (cudaStream_t)stream>>>(in, out, alpha, n);
+}
+
+}