【初心者向け】Windows環境でCPUとGPUの配列の総和を算出時間を比較してみた

#include <iostream>
#include <vector>
#include <stdio.h>
#include <numeric>
#include <windows.h>
#include <cuda.h>

#pragma comment(lib, "winmm.lib")

__device__ 
void warpReduce(volatile float* sdata, unsigned int tid)
{
     sdata[tid] += sdata[tid + 32];
     sdata[tid] += sdata[tid + 16];
     sdata[tid] += sdata[tid + 8];
     sdata[tid] += sdata[tid + 4];
     sdata[tid] += sdata[tid + 2];
     sdata[tid] += sdata[tid + 1];
}

__global__
void reduction_sum(float* d_val, float* d_out, int n)
{
    extern __shared__ float sdata[];

    unsigned int tid = threadIdx.x;
    
    sdata[tid] = 0;
    float f = 0.0f;

    for (int i = tid; i < n; i += blockDim.x)
    {
        if ( n <= i )break;
        f += d_val[i];
    }
    sdata[tid] = f;
    
    __syncthreads();
    if (tid < 512) { sdata[tid] += sdata[tid + 512]; } 
    __syncthreads(); 
    if (tid < 256) { sdata[tid] += sdata[tid + 256]; } 
    __syncthreads(); 
    if (tid < 128) { sdata[tid] += sdata[tid + 128]; }
    __syncthreads(); 
    if (tid < 64) { sdata[tid] += sdata[tid + 64]; } 
    __syncthreads(); 
    if (tid < 32) warpReduce(sdata, tid);
    if (tid == 0) d_out[tid] = sdata[0];

}



int main()
{

    std::vector<int> N = { 1024, 524288, 1048576 };

    for (auto n : N)
    {
        std::cout << "---N : " << n << std::endl;
        float* d_val;
        float* d_out;
        float* d_sum;
        std::vector<float> h_val;



        for (int i = 0; i < n; i++) h_val.push_back(float(1));

        cudaMalloc(&d_val, sizeof(float) * n);
        cudaMalloc(&d_out, sizeof(float));
        cudaMemcpy(d_val, h_val.data(), sizeof(float) * n, cudaMemcpyHostToDevice);

        d_sum = (float*)malloc(sizeof(float));




        DWORD start = timeGetTime();

        float h_sum = 0;
        for (auto f : h_val)
        {
            h_sum += f;
        }

        DWORD end = timeGetTime();

        std::cout << "h_val sum = " << h_sum << std::endl;
        std::wcout << "CPU Time[ms] : " << end - start << std::endl;



        LARGE_INTEGER frequency;        // ticks per second
        LARGE_INTEGER t1, t2;           // ticks
        double elapsedTime;


        int block = 1024;
        int gird = 1;
        int sdata_size = 1024;



        QueryPerformanceFrequency(&frequency); // get ticks per second
        QueryPerformanceCounter(&t1); // start timer

        reduction_sum << < gird, block, sizeof(float)* sdata_size >> > (d_val, d_out, n);
        
        QueryPerformanceCounter(&t2); // stop timer

        // compute and print the elapsed time in millisec
        elapsedTime = (t2.QuadPart - t1.QuadPart) * 1000.0 / frequency.QuadPart;

        cudaMemcpy(d_sum, d_out, sizeof(float), cudaMemcpyDeviceToHost);
        printf("\n");
        printf("GPU Time[ms] : %.5f \n", elapsedTime);
        std::cout << "d_val sum = " << d_sum[0] << std::endl;
        printf("\n");

    }


}