fix(cuda): forward-declare warp_reduce_sum_all before BF16 matvec

The BF16 matvec helpers (warp_matvec_bf16_shmem, warp_matvec_bf16_broadcast_shmem)
call warp_reduce_sum_all which is defined later in the header. nvcc on H100
rejects this without a forward declaration — broke experience kernel compilation.

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

crates/ml/src/cuda_pipeline/common_device_functions.cuh

@@ -440,6 +440,9 @@ __device__ __forceinline__ void cooperative_load_bias_bf16_to_f32(
         shmem_dst[i] = __bfloat162float(global_src[i]);
 }
 
+/* Forward declaration — defined below, used by BF16 matvec helpers */
+__device__ float warp_reduce_sum_all(float val);
+
 /**
  * Warp-cooperative matrix-vector multiply with BF16 weights, F32 accumulation.
  *