#include "Website_VectorAdd.h"

__global__ __launch_bounds__(1024)
/**
  hoisted when extracting vector_add
*/
static void
__hoisted_vector_add_0(uint32_t n, float *a, float *b, float *out)
{
    if (1024U * blockIdx.x + threadIdx.x < n)
        out[1024U * blockIdx.x + threadIdx.x] =
            a[1024U * blockIdx.x + threadIdx.x] +
            b[1024U * blockIdx.x + threadIdx.x];
}

void Website_VectorAdd_vector_add(uint32_t n, float *a, float *b, float *out)
{
    cudaStream_t s = KPR_FRESH_STREAM();
    KPR_KCALL(__hoisted_vector_add_0, n / 1024U + (uint32_t) (n % 1024U != 0U),
        1024U, 0U, s, n, a, b, out);
    MUST(cudaStreamSynchronize(s));
    MUST(cudaStreamDestroy(s));
}