extern "C" {

__global__ void vecadd_f32(const float* a, const float* b, float* c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

void launch_vecadd_f32(const float* a, const float* b, float* c, int n, void* stream) {
    int block = 256;
    int grid = (n + block - 1) / block;
    vecadd_f32<<<grid, block, 0, (cudaStream_t)stream>>>(a, b, c, n);
}

}