Shared Code

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

#define N (1 << 20)
#define LIMIT 750000

__global__ void findElements(
    int *arr,
    int n,
    int limit,
    unsigned int *result)
{
    __shared__ unsigned int count;

    if (threadIdx.x == 0)
        count = 0;

    __syncthreads();

    int index = blockIdx.x * blockDim.x + threadIdx.x;
    int jump = blockDim.x * gridDim.x;

    unsigned int myCount = 0;

    for (; index < n; index += jump)
    {
        if (arr[index] > limit)
            myCount++;
    }

    if (myCount > 0)
        atomicAdd(&count, myCount);

    __syncthreads();

    if (threadIdx.x == 0)
        atomicAdd(result, count);
}

int main()
{
    if (cudaSetDevice(0) != cudaSuccess)
    {
        std::cout << "CUDA device initialization failed.\n";
        return 1;
    }

    size_t bytes = N * sizeof(int);

    int *h_arr = (int*)malloc(bytes);
    int *d_arr;
    unsigned int *d_result;
    unsigned int answer = 0;

    cudaMalloc(&d_arr, bytes);
    cudaMalloc(&d_result, sizeof(unsigned int));

    for (int i = 0; i < N; i++)
        h_arr[i] = i;

    cudaMemcpy(d_arr, h_arr, bytes, cudaMemcpyHostToDevice);
    cudaMemset(d_result, 0, sizeof(unsigned int));

    int threads = 256;
    int blocks = 256;

    findElements<<<blocks, threads>>>(
        d_arr, N, LIMIT, d_result
    );

    cudaError_t e = cudaGetLastError();

    if (e != cudaSuccess)
    {
        std::cout << "Error while starting kernel: "
                  << cudaGetErrorString(e) << "\n";
        return 1;
    }

    e = cudaDeviceSynchronize();

    if (e != cudaSuccess)
    {
        std::cout << "CUDA execution error: "
                  << cudaGetErrorString(e) << "\n";
        return 1;
    }

    cudaMemcpy(
        &answer,
        d_result,
        sizeof(unsigned int),
        cudaMemcpyDeviceToHost
    );

    unsigned int expected = N - LIMIT - 1;

    std::cout << "\n========== CUDA RESULT ==========\n";
    std::cout << "Array size       : " << N << "\n";
    std::cout << "Threshold value  : " << LIMIT << "\n";
    std::cout << "Elements matched : " << answer << "\n";
    std::cout << "Correct count    : " << expected << "\n";

    if (answer == expected)
        std::cout << "Status           : SUCCESS\n";
    else
        std::cout << "Status           : INCORRECT\n";

    std::cout << "=================================\n";

    cudaFree(d_arr);
    cudaFree(d_result);
    free(h_arr);

    return 0;
}