From 8e0bcbcad3c16bacfc54a56d3bdd0ebaddaec089 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Mon, 16 Mar 2026 09:38:53 +0100 Subject: [PATCH] fix(cuda): forward-declare warp_reduce_sum_all before BF16 matvec MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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 --- crates/ml/src/cuda_pipeline/common_device_functions.cuh | 3 +++ 1 file changed, 3 insertions(+) diff --git a/crates/ml/src/cuda_pipeline/common_device_functions.cuh b/crates/ml/src/cuda_pipeline/common_device_functions.cuh index eb26c4998..42e99435c 100644 --- a/crates/ml/src/cuda_pipeline/common_device_functions.cuh +++ b/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. *