Shared Code

#include <iostream>
#include <cstdlib>
#include <cuda_runtime.h>

#define ARRAY_SIZE (1 << 20)
#define THRESHOLD 750000

__global__ void count_above_threshold_kernel(
    const int* d_array,
    int size,
    int threshold,
    unsigned int* d_global_count)
{
    __shared__ unsigned int block_count;

    int local_id = threadIdx.x;
-
    if (local_id == 0)
    {
        block_count = 0;
    }

    __syncthreads();

    int g_id = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;

    int local_match_count = 0;

    while (g_id < size)
    {
        if (d_array[g_id] > threshold)
        {
            local_match_count++;
        }

        g_id += stride;
    }

    if (local_match_count > 0)
    {
        atomicAdd(&block_count, local_match_count);
    }

    __syncthreads();

    if (local_id == 0 && block_count > 0)
    {
        atomicAdd(d_global_count, block_count);
    }
}

int main()
{
    int target_device = 0;

    cudaError_t err = cudaSetDevice(target_device);

    if (err != cudaSuccess)
    {
        std::cerr << "Failed to set CUDA device ["
                  << cudaGetErrorString(err) << "]\n";
        return -1;
    }

    size_t array_bytes = ARRAY_SIZE * sizeof(int);

    int* h_array = (int*)malloc(array_bytes);
    unsigned int h_global_count = 0;

    int* d_array;
    unsigned int* d_global_count;

    cudaMalloc(&d_array, array_bytes);
    cudaMalloc(&d_global_count, sizeof(unsigned int));

    for (int i = 0; i < ARRAY_SIZE; i++)
    {
        h_array[i] = i;
    }

    cudaMemcpy(
        d_array,
        h_array,
        array_bytes,
        cudaMemcpyHostToDevice
    );

    cudaMemset(
        d_global_count,
        0,
        sizeof(unsigned int)
    );

    int threads_per_block = 256;
    int blocks_per_grid = 256;

    count_above_threshold_kernel<<<
        blocks_per_grid,
        threads_per_block
    >>>(
        d_array,
        ARRAY_SIZE,
        THRESHOLD,
        d_global_count
    );

    err = cudaGetLastError();

    if (err != cudaSuccess)
    {
        std::cerr << "Kernel launch error: "
                  << cudaGetErrorString(err) << std::endl;
        return -1;
    }

    err = cudaDeviceSynchronize();

    if (err != cudaSuccess)
    {
        std::cerr << "Kernel execution error: "
                  << cudaGetErrorString(err) << std::endl;
        return -1;
    }

    cudaMemcpy(
        &h_global_count,
        d_global_count,
        sizeof(unsigned int),
        cudaMemcpyDeviceToHost
    );

    unsigned int expected_count =
        (ARRAY_SIZE - 1) - THRESHOLD;

    std::cout << "Total elements found = "
              << h_global_count << std::endl;

    std::cout << "Expected elements = "
              << expected_count << std::endl;

    if (h_global_count == expected_count)
    {
        std::cout << "Verification = PASSED"
                  << std::endl;
    }
    else
    {
        std::cout << "Verification = FAILED"
                  << std::endl;
    }

    cudaFree(d_array);
    cudaFree(d_global_count);

    free(h_array);

    return 0;
}