C++ Cuda parallel reduction strategy

 

코드 >>

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "cuda.h"
#include <stdlib.h>
#include <cstring>
#include <time.h>
#include <stdio.h>
#include <cmath>
#include <iostream>
#include "common.h"
 
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
 
inline void gpuAssert(cudaError_t code, const char* file, int line, bool abort = true)
{
    if (code != cudaSuccess)
    {
        fprintf(stderr, "GPU 경고: %s %s %d\n", cudaGetErrorString(code), file, line);
        if (abort) exit(code);
    }
}
int reduction_cpu(int* input, int size)
{
    int sum = 0;
    for (int i = 0; i < size; i++)
    {
        sum += input[i];
    }
    return sum;
}
 
void compare_results(int gpu_result, int cpu_result)
{
    printf("GPU 결과: %d, CPU 결과: %d \n", gpu_result, cpu_result);
    if (gpu_result == cpu_result)
    {
        printf("GPU와 CPU의 결과 같음 \n");
        return;
    }
 
    printf("GPU와 CPU의 결과 다름 \n");
}
 
__global__ void reduction_neighbored_pairs(int* input,
    int* temp, int size)
{
    int tid = threadIdx.x;
    int gid = blockDim.x * blockIdx.x + threadIdx.x;
 
    if (gid > size)
        return;
 
    for (int offset = 1; offset <= blockDim.x / 2; offset *= 2)
    {
        if (tid % (2 * offset) == 0)
        {
            input[gid] += input[gid + offset];
 
        }
    }
 
    if (tid == 0)
    {
        temp[blockIdx.x] = input[gid];
    }
}
 
int main()
{
    printf("Start neighbored pairs reduction kernel. \n");
    int size = 1 << 27; // 128 MB of data
    int byte_size = size * sizeof(int);
    int block_size = 128;
 
    int* h_input, * h_ref;
    h_input = (int*)malloc(byte_size);
    memset(h_input, 0, byte_size);
    int cpu_result = reduction_cpu(h_input, size);
 
    dim3 block(block_size);
    dim3 grid(size / block.x);
 
    printf("grid.x : %d, block.x : %d\n",
        grid.x, block.x);
 
    int temp_array_byte_size = sizeof(int) * grid.x;
    h_ref = (int*)malloc(temp_array_byte_size);
 
    int* d_input, * d_temp;
 
    gpuErrchk(cudaMalloc((void**)&d_input, byte_size));
    gpuErrchk(cudaMalloc((void**)&d_temp, temp_array_byte_size));
    
    cudaMemset(d_temp, 0, temp_array_byte_size);
    cudaMemcpy(d_input, h_input, byte_size, cudaMemcpyHostToDevice);
 
 
    reduction_neighbored_pairs << <grid, block >> > (d_input, d_temp, size);
 
    gpuErrchk(cudaDeviceSynchronize());
 
    cudaMemcpy(h_ref, d_temp, temp_array_byte_size, cudaMemcpyDeviceToHost);
 
    
 
    int gpu_result = 0;
    
    for (int i = 0; i < grid.x; i++)
    {
        gpu_result += h_ref[i];
    }
 
    compare_results(gpu_result, cpu_result);
 
    cudaFree(d_input);
    cudaFree(d_temp);
    free(h_ref);
    free(h_input);
 
 
    gpuErrchk(cudaDeviceReset());
    return 0;
}