|cuda| reduction max(parallel reduction, binary reduction)

0. [cuda] - reduction max

누가 볼지 모르겠지만, 다양한 방법의 reduction 수행을 통해 연산을 최대한 가속하는 것에 목표를 두었다.

1. 풀이

1. 가장 기본적으로 nvidia에서 제공하는 atomicMax 활용

모든 threads 에서 병렬적으로 atomicMax를 수행한다.

atomicMax에서는 값이 대입되기 전까지 변수에 접근을 막아 sync를 맞추는 작업이 수행된다.

즉, 결과 자체는 동일하지만 수행시간이 느린 것을 알 수 있었다.

--------------------------

2. shared memory와 atomicMax 동시 활용

block으로 구분되어 병렬적으로 연산이 수행되는데, 각각의 block이 갖는 최대값을 활용해 선별적으로 atomicMax를 수행하였다.

이러한 결과 데이터 병목현상을 막을 수 있었다.

3. 추가적으로 gpu 내부 int 사용시 수행시간의 차이 발생 확인

엄청난 차이가 발생하지는 않았지만, blockIdx, threadIdx를 수행하는 횟수가 1번인 경우에 대해 int에 대입하지 않고 바로 활용하는 것이 미세한 이득이 있는 것을 확인하였다.

4. threads당 workload 증가

이러한 경우 테스트 하려다 pass

---------------------------

5. binary max

기본적인 binary 연산을 통해 block이 수행하는 연산량이 매 수행마다 절반으로 감소한다.

6. unroll

for문을 해체하여 수행시간 감소

7. shared memory에 올리며 바로 비교

thread를 절반으로 줄이는 것이 가능해진다.

정리하자. 실험적으로 binary를 활용하는 경우 수행시간 가속이 이상적으로 이루어지지 않았다.

현재 테스트를 진행한 GPU는 

Turing의 RTX2070을 활용하였는데, atomicMax를 바로 수행하는 것이 연산속도가 짧게 나오는 것을 볼 수 있었다.

하지만, 이와 같은 실험을 linux 환경의 GTX TITAN X에서 수행한 결과 N = 1024 근처인 경우 local < binary < global의 결과를 얻을 수 있었다.

이와 관련한 자세한 것들은 계속 공부해가며 이해하자.

2. 소스코드

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include<iostream>
#include<algorithm>
#include<stdio.h>
#include<Windows.h>
 
#include<iostream>
 
#define THREADS 64
 
using namespace std;
 
__global__ void getMax_global(unsigned int * d_data, unsigned int *d_max, int n);
__global__ void getMax_local(unsigned int * d_data, unsigned int *d_max, int n);
__global__ void getMax_local2(unsigned int * d_data, unsigned int *d_max, int n);
__global__ void getMax_binary(unsigned int * d_data, unsigned int *d_max, int n, const int k);
__global__ void getMax_binary2(unsigned int * d_data, unsigned int *d_max, int n, const int k);
__device__ void warpReduce(volatile unsigned int * s_data, int tid);
 
 
__global__ void getMax_global(unsigned int * d_data, unsigned int *d_max, int n) {
    int gid = (blockIdx.x*blockDim.x + threadIdx.x);
    atomicMax(d_max, d_data[gid]);
}
 
 
//tid, gid를 사용하는 경우
__global__ void getMax_local(unsigned int * d_data, unsigned int *d_max, int n) {
    __shared__ unsigned int s_max;
    int tid = threadIdx.x;
    int gid = blockIdx.x * blockDim.x + threadIdx.x;
    //Shared variable init
    if (tid == 0) {
        s_max = 0;
    }
 
    atomicMax(&s_max, d_data[gid]);
    __syncthreads();
 
    if (tid == 0) {
        atomicMax(d_max, s_max);
    }
}
 
//tid, gid를 사용하지 않는 경우
__global__ void getMax_local2(unsigned int * d_data, unsigned int *d_max, int n) {
    __shared__ unsigned int s_max;
    //Shared variable init
    if (threadIdx.x == 0) {
        s_max = 0;
    }
 
    atomicMax(&s_max, d_data[blockIdx.x * blockDim.x + threadIdx.x]);
    __syncthreads();
 
    if (threadIdx.x == 0) {
        atomicMax(d_max, s_max);
    }
}
 
//tid만 사용하는 경우
__global__ void getMax_local3(unsigned int * d_data, unsigned int *d_max, int n) {
    __shared__ unsigned int s_max;
 
    int tid = threadIdx.x;
    //Shared variable init
    if (tid == 0) {
        s_max = 0;
    }
 
    atomicMax(&s_max, d_data[blockIdx.x * blockDim.x + tid]);
    __syncthreads();
 
    if (tid == 0) {
        atomicMax(d_max, s_max);
    }
}
 
//tid 일부 사용하는 경우
__global__ void getMax_local4(unsigned int * d_data, unsigned int *d_max, int n) {
    __shared__ unsigned int s_max;
 
    int tid = threadIdx.x;
    //Shared variable init
    if (tid == 0) {
        s_max = 0;
    }
 
    atomicMax(&s_max, d_data[blockIdx.x * blockDim.x + threadIdx.x]);
    __syncthreads();
 
    if (tid == 0) {
        atomicMax(d_max, s_max);
    }
}
 
 
 
//more workload
__global__ void getMax_local5(unsigned int * d_data, unsigned int *d_max, int n) {
    __shared__ unsigned int s_max;
 
    int tid = threadIdx.x;
    //Shared variable init
    if (tid == 0) {
        s_max = 0;
    }
 
    atomicMax(&s_max, d_data[blockIdx.x * blockDim.x + threadIdx.x]);
    __syncthreads();
 
    if (tid == 0) {
        atomicMax(d_max, s_max);
    }
}
 
 
 
 
__global__ void getMax_binary(unsigned int *d_data, unsigned int *d_max, int n, const int k) {
    extern __shared__ unsigned int s_data[];
 
    int tid = threadIdx.x;
 
    s_data[tid] = d_data[blockIdx.x * blockDim.x + threadIdx.x];
    __syncthreads();
    for (int s = ((THREADS*k) >> 1); s >= 1; (s >>= 1)) {
        if (tid < s) {
 
            if (s_data[tid] < s_data[tid + s]) {
                s_data[tid] = s_data[tid + s];
            }
        }
        __syncthreads();
    }
 
    if (tid == 0) {
        atomicMax(d_max, s_data[0]);
    }
}
 
 
__device__ void warpReduce(volatile unsigned int * s_data, int tid) {
    if (s_data[tid] < s_data[tid + 32]) s_data[tid] = s_data[tid + 32];
    if (s_data[tid] < s_data[tid + 16]) s_data[tid] = s_data[tid + 16];
    if (s_data[tid] < s_data[tid + 8]) s_data[tid] = s_data[tid + 8];
    if (s_data[tid] < s_data[tid + 4]) s_data[tid] = s_data[tid + 4];
    if (s_data[tid] < s_data[tid + 2]) s_data[tid] = s_data[tid + 2];
    if (s_data[tid] < s_data[tid + 1]) s_data[tid] = s_data[tid + 1];
}
 
//using unroll
__global__ void getMax_binary2(unsigned int * d_data, unsigned int *d_max, int n, const int k) {
    extern __shared__ unsigned int s_data[];
    int tid = threadIdx.x;
 
 
    s_data[tid] = d_data[blockIdx.x * blockDim.x + threadIdx.x];
    __syncthreads();
    for (int s = ((THREADS*k) >> 1); s > 32; (s >>= 1)) {
        if (tid < s) {
            if (s_data[tid] < s_data[tid + s]) s_data[tid] = s_data[tid + s];
        }
        __syncthreads();
    }
 
    if (tid < 32) warpReduce(s_data, tid);
 
    if (tid == 0) {
        atomicMax(d_max, s_data[0]);
    }
}
 
 
 
__device__ void warpReduce2(volatile unsigned int * s_data, int tid) {
    if (s_data[tid] < s_data[tid + 32]) s_data[tid] = s_data[tid + 32];
    if (s_data[tid] < s_data[tid + 16]) s_data[tid] = s_data[tid + 16];
    if (s_data[tid] < s_data[tid + 8]) s_data[tid] = s_data[tid + 8];
    if (s_data[tid] < s_data[tid + 4]) s_data[tid] = s_data[tid + 4];
    if (s_data[tid] < s_data[tid + 2]) s_data[tid] = s_data[tid + 2];
    if (s_data[tid] < s_data[tid + 1]) s_data[tid] = s_data[tid + 1];
}
 
//using more unroll
template <unsigned int blockSize>
__global__ void getMax_binary3(unsigned int * d_data, unsigned int *d_max, int n, const int k) {
    extern __shared__ unsigned int s_data[];
    int tid = threadIdx.x;
 
    s_data[tid] = d_data[blockIdx.x * blockDim.x + tid];
    __syncthreads();
 
    if (blockSize*k >= 1024) {
        if (tid < 512)
            if (s_data[tid] < s_data[tid + 512])
                s_data[tid] = s_data[tid + 512];
        __syncthreads();
    }
    if (blockSize*k >= 512) {
        if (tid < 256)
            if (s_data[tid] < s_data[tid + 256])
                s_data[tid] = s_data[tid + 256];
        __syncthreads();
    }
    if (blockSize*k >= 256) {
        if (tid < 128)
            if (s_data[tid] < s_data[tid + 128])
                s_data[tid] = s_data[tid + 128];
        __syncthreads();
    }
    if (blockSize*k >= 128) {
        if (tid < 64)
            if (s_data[tid] < s_data[tid + 64])
                s_data[tid] = s_data[tid + 64];
        __syncthreads();
    }
 
    if (tid f32) warpReduce2(s_data, tid);
 
    if (tid == 0) {
        atomicMax(d_max, s_data[0]);
    }
}
 
//using First Add During Load
template <unsigned int blockSize>
__global__ void getMax_binary4(unsigned int * d_data, unsigned int *d_max, int n, const int k) {
    extern __shared__ unsigned int s_data[];
    int tid = threadIdx.x;
    int gid = blockIdx.x * blockDim.x * 2 + tid;
    //printf("blockDim : %d\n", blockDim.x);
        
    s_data[tid] = (d_data[gid] > d_data[gid + blockDim.x]) ? d_data[gid] : d_data[gid + blockDim.x];
    __syncthreads();
 
    if (blockSize*k >= 1024) {
        if (tid < 512)
            if (s_data[tid] < s_data[tid + 512])
                s_data[tid] = s_data[tid + 512];
        __syncthreads();
    }
    if (blockSize*k >= 512) {
        if (tid < 256)
            if (s_data[tid] < s_data[tid + 256])
                s_data[tid] = s_data[tid + 256];
        __syncthreads();
    }
    if (blockSize*k >= 256) {
        if (tid < 128)
            if (s_data[tid] < s_data[tid + 128])
                s_data[tid] = s_data[tid + 128];
        __syncthreads();
    }
    if (blockSize*k >= 128) {
        if (tid < 64)
            if (s_data[tid] < s_data[tid + 64])
                s_data[tid] = s_data[tid + 64];
        __syncthreads();
    }
 
    if (tid < 32) warpReduce2(s_data, tid);
 
    if (tid == 0) {
        atomicMax(d_max, s_data[0]);
    }
}
 
 
//고냥 처음 넣을 때, shared memory에 바로 올리지 말고
//비교한 이후 shared memory에 넣어주자.
//이를 위해서 blocksize를 절반을 택하거나
//threads의 크기를 절반을 택하는 두가지 경우를 모두 테스트해보자.
template <unsigned int blockSize>
__global__ void getMax_binary5(unsigned int * d_data, unsigned int *d_max, int n, const int k) {
    extern __shared__ unsigned int s_data[];
    int tid = threadIdx.x;
    int gid = blockIdx.x * blockDim.x + tid;
    
    s_data[tid] = (d_data[gid] > d_data[gid + n]) ? d_data[gid] : d_data[gid + n];
    __syncthreads();
 
    if (blockSize*k >= 1024) {
        if (tid < 512)
            if (s_data[tid] < s_data[tid + 512])
                s_data[tid] = s_data[tid + 512];
        __syncthreads();
    }
    if (blockSize*k >= 512) {
        if (tid < 256)
            if (s_data[tid] < s_data[tid + 256])
                s_data[tid] = s_data[tid + 256];
        __syncthreads();
    }
    if (blockSize*k >= 256) {
        if (tid < 128)
            if (s_data[tid] < s_data[tid + 128])
                s_data[tid] = s_data[tid + 128];
        __syncthreads();
    }
    if (blockSize*k >= 128) {
        if (tid < 64)
            if (s_data[tid] < s_data[tid + 64])
                s_data[tid] = s_data[tid + 64];
        __syncthreads();
    }
 
    if (tid < 32) warpReduce2(s_data, tid);
 
    if (tid == 0) {
        atomicMax(d_max, s_data[0]);
    }
}
 
 
//이건 위랑 동일한데 함수명만 다르게
//threads 수가 절반!
template <unsigned int blockSize>
__global__ void getMax_binary6(unsigned int * d_data, unsigned int *d_max, int n, const int k) {
    extern __shared__ unsigned int s_data[];
    int tid = threadIdx.x;
    int gid = blockIdx.x * blockDim.x + tid;
 
    s_data[tid] = (d_data[gid] > d_data[gid + n]) ? d_data[gid] : d_data[gid + n];
    __syncthreads();
 
    if (blockSize*k >= 1024) {
        if (tid < 512)
            if (s_data[tid] < s_data[tid + 512])
                s_data[tid] = s_data[tid + 512];
        __syncthreads();
    }
    if (blockSize*k >= 512) {
        if (tid < 256)
            if (s_data[tid] < s_data[tid + 256])
                s_data[tid] = s_data[tid + 256];
        __syncthreads();
    }
    if (blockSize*k >= 256) {
        if (tid < 128)
            if (s_data[tid] < s_data[tid + 128])
                s_data[tid] = s_data[tid + 128];
        __syncthreads();
    }
    if (blockSize*k >= 128) {
        if (tid < 64)
            if (s_data[tid] < s_data[tid + 64])
                s_data[tid] = s_data[tid + 64];
        __syncthreads();
    }
 
    if (tid < 32) warpReduce2(s_data, tid);
 
    if (tid == 0) {
        atomicMax(d_max, s_data[0]);
    }
}
 
//use if, else if
template <unsigned int blockSize>
__global__ void getMax_binary7(unsigned int * d_data, unsigned int *d_max, int n, const int k) {
    extern __shared__ unsigned int s_data[];
    int tid = threadIdx.x;
    int gid = blockIdx.x * blockDim.x + tid;
 
    s_data[tid] = (d_data[gid] > d_data[gid + n]) ? d_data[gid] : d_data[gid + n];
    __syncthreads();
 
    if (blockSize*k >= 1024) {
        if (tid < 512)
            if (s_data[tid] < s_data[tid + 512])
                s_data[tid] = s_data[tid + 512];
        __syncthreads();
        if (tid < 256)
            if (s_data[tid] < s_data[tid + 256])
                s_data[tid] = s_data[tid + 256];
        __syncthreads();
        if (tid < 128)
            if (s_data[tid] < s_data[tid + 128])
                s_data[tid] = s_data[tid + 128];
        __syncthreads();
        if (tid < 64)
            if (s_data[tid] < s_data[tid + 64])
                s_data[tid] = s_data[tid + 64];
        __syncthreads();
    }
    else if (blockSize*k >= 512) {
        if (tid < 256)
            if (s_data[tid] < s_data[tid + 256])
                s_data[tid] = s_data[tid + 256];
        __syncthreads();
        if (tid < 128)
            if (s_data[tid] < s_data[tid + 128])
                s_data[tid] = s_data[tid + 128];
        __syncthreads();
        if (tid < 64)
            if (s_data[tid] < s_data[tid + 64])
                s_data[tid] = s_data[tid + 64];
        __syncthreads();
    }
    else if (blockSize*k >= 256) {
        if (tid < 128)
            if (s_data[tid] < s_data[tid + 128])
                s_data[tid] = s_data[tid + 128];
        __syncthreads();
        if (tid < 64)
            if (s_data[tid] < s_data[tid + 64])
                s_data[tid] = s_data[tid + 64];
        __syncthreads();
    }
    else if (blockSize*k >= 128) {
        if (tid < 64)
            if (s_data[tid] < s_data[tid + 64])
                s_data[tid] = s_data[tid + 64];
        __syncthreads();
    }
 
    if (tid < 32) warpReduce2(s_data, tid);
 
    if (tid == 0) {
        atomicMax(d_max, s_data[0]);
    }
}
 
//not use id
template <unsigned int blockSize>
__global__ void getMax_binary8(unsigned int * d_data, unsigned int *d_max, int n, const int k) {
    extern __shared__ unsigned int s_data[];
 
    s_data[threadIdx.x] = (d_data[blockIdx.x * blockDim.x + threadIdx.x] > d_data[blockIdx.x * blockDim.x + threadIdx.x + n]) ? d_data[blockIdx.x * blockDim.x + threadIdx.x] : d_data[blockIdx.x * blockDim.x + threadIdx.x + n];
    __syncthreads();
 
    if (blockSize*k >= 1024) {
        if (threadIdx.x < 512)
            if (s_data[threadIdx.x] < s_data[threadIdx.x + 512])
                s_data[threadIdx.x] = s_data[threadIdx.x + 512];
        __syncthreads();
        if (threadIdx.x < 256)
            if (s_data[threadIdx.x] < s_data[threadIdx.x + 256])
                s_data[threadIdx.x] = s_data[threadIdx.x + 256];
        __syncthreads();
        if (threadIdx.x < 128)
            if (s_data[threadIdx.x] < s_data[threadIdx.x + 128])
                s_data[threadIdx.x] = s_data[threadIdx.x + 128];
        __syncthreads();
        if (threadIdx.x < 64)
            if (s_data[threadIdx.x] < s_data[threadIdx.x + 64])
                s_data[threadIdx.x] = s_data[threadIdx.x + 64];
        __syncthreads();
    }
    else if (blockSize*k >= 512) {
        if (threadIdx.x < 256)
            if (s_data[threadIdx.x] < s_data[threadIdx.x + 256])
                s_data[threadIdx.x] = s_data[threadIdx.x + 256];
        __syncthreads();
        if (threadIdx.x < 128)
            if (s_data[threadIdx.x] < s_data[threadIdx.x + 128])
                s_data[threadIdx.x] = s_data[threadIdx.x + 128];
        __syncthreads();
        if (threadIdx.x < 64)
            if (s_data[threadIdx.x] < s_data[threadIdx.x + 64])
                s_data[threadIdx.x] = s_data[threadIdx.x + 64];
        __syncthreads();
    }
    else if (blockSize*k >= 256) {
        if (threadIdx.x < 128)
            if (s_data[threadIdx.x] < s_data[threadIdx.x + 128])
                s_data[threadIdx.x] = s_data[threadIdx.x + 128];
        __syncthreads();
        if (threadIdx.x < 64)
            if (s_data[threadIdx.x] < s_data[threadIdx.x + 64])
                s_data[threadIdx.x] = s_data[threadIdx.x + 64];
        __syncthreads();
    }
    else if (blockSize*k >= 128) {
        if (threadIdx.x < 64)
            if (s_data[threadIdx.x] < s_data[threadIdx.x + 64])
                s_data[threadIdx.x] = s_data[threadIdx.x + 64];
        __syncthreads();
    }
 
    if (threadIdx.x < 32) warpReduce2(s_data, threadIdx.x);
 
    if (threadIdx.x == 0) {
        atomicMax(d_max, s_data[0]);
    }
}
 
 
 
int main(int argc, char* argv[]) {
    /*if (argc < 2) {
        puts("usage: matmul [N]");
        return 0;
    }*/
 
    //int N = atoi(argv[1]);
    int N = 1024;
    for (int m = 0; m < 3; m++) {
        printf("N: %d\n", N);
        //total size
        size_t sz = sizeof(int) *N;
 
        unsigned int * data = (unsigned int*)malloc(sz);
        for (int i = 0; i < N; i++) {
            data[i] = (unsigned int)(i + 1);
        }
        //Struct for time measure
        /*struct timeval start, end, timer;
        srand(time(NULL));*/
        LARGE_INTEGER freq, start, end;
 
        // H2D memcpy
        unsigned int *d_data;
        cudaMalloc((void **)&d_data, sz);
        unsigned int *d_max;
        cudaMalloc((void **)&d_max, sizeof(unsigned int));
 
        //Declare max variable
        unsigned int max;
 
        /******************************
           CPU
        *******************************/
        QueryPerformanceFrequency(&freq);
        QueryPerformanceCounter(&start);
        //init max
        max = 0;
        for (int i = 0; i < N; i++) {
            if (max < data[i]) {
                max = data[i];
            }
        }
        //Timemeasure end
        QueryPerformanceCounter(&end);
        double diff = (end.QuadPart - start.QuadPart) / (freq.QuadPart / 1000000.0);
        printf("CPU, elapsend time: %lf\n", diff);
        printf("CPU, max value : %d\n", max);
 
        int threads, grid;
        int a = 2;
        for (int k = 0; k < 4; k++) {
            threads = THREADS * a; //variable threads means number of threads per block        
            grid = (N%threads) ? N / threads + 1 : N / threads; // variable grid means number of total blocks
 
            cout << threads << " " << grid << endl;
            /******************************
           Global max
        *******************************/
            for (int i = 0; i < 10; i++) {
 
                max = 0;
                cudaMemcpy(d_data, data, sz, cudaMemcpyHostToDevice);
                cudaMemcpy(d_max, data, sizeof(unsigned int), cudaMemcpyHostToDevice);
                //Time measure start
                QueryPerformanceFrequency(&freq);
                QueryPerformanceCounter(&start);
                getMax_global << <grid, threads >> > (d_data, d_max, N);
                //Wait until thread function(gpu) is over
                cudaThreadSynchronize();
                cudaMemcpy(&max, d_max, sizeof(unsigned int), cudaMemcpyDeviceToHost);
 
                //Time measure end
                QueryPerformanceCounter(&end);
                diff = (end.QuadPart - start.QuadPart) / (freq.QuadPart / 1000000.0);
                printf("Global max, elapsend time: %lf\n", diff);
                printf("Global max, max value : %d\n", max);
 
            }
 
 
            /******************************
               Local max
            *******************************/
 
            //using gid, tid
            for (int i = 0; i < 10; i++) {
                max = 0;
                cudaMemcpy(d_data, data, sz, cudaMemcpyHostToDevice);
                cudaMemcpy(d_max, data, sizeof(unsigned int), cudaMemcpyHostToDevice);
                //Time measure start
                QueryPerformanceFrequency(&freq);
                QueryPerformanceCounter(&start);
                getMax_local << <grid, threads >> > (d_data, d_max, N);
                //Wait until thread function(gpu) is over
                cudaThreadSynchronize();
                cudaMemcpy(&max, d_max, sizeof(unsigned int), cudaMemcpyDeviceToHost);
 
                //Time measure end
                QueryPerformanceCounter(&end);
                diff = (end.QuadPart - start.QuadPart) / (freq.QuadPart / 1000000.0);
                printf("Local max, elapsend time: %lf\n", diff);
                printf("Local max, max value : %d\n", max);
            }
 
            //not use int id
            for (int i = 0; i < 10; i++) {
                max = 0;
                cudaMemcpy(d_data, data, sz, cudaMemcpyHostToDevice);
                cudaMemcpy(d_max, data, sizeof(unsigned int), cudaMemcpyHostToDevice);
                //Time measure start
                QueryPerformanceFrequency(&freq);
                QueryPerformanceCounter(&start);
                getMax_local2 << <grid, threads >> > (d_data, d_max, N);
                //Wait until thread function(gpu) is over
                cudaThreadSynchronize();
                cudaMemcpy(&max, d_max, sizeof(unsigned int), cudaMemcpyDeviceToHost);
 
                //Time measure end
                QueryPerformanceCounter(&end);
                diff = (end.QuadPart - start.QuadPart) / (freq.QuadPart / 1000000.0);
                printf("Local2 max, elapsend time: %lf\n", diff);
                printf("Local2 max, max value : %d\n", max);
            }
 
            //using tid
            for (int i = 0; i < 10; i++) {
                max = 0;
                cudaMemcpy(d_data, data, sz, cudaMemcpyHostToDevice);
                cudaMemcpy(d_max, data, sizeof(unsigned int), cudaMemcpyHostToDevice);
                //Time measure start
                QueryPerformanceFrequency(&freq);
                QueryPerformanceCounter(&start);
                getMax_local3 << <grid, threads >> > (d_data, d_max, N);
                //Wait until thread function(gpu) is over
                cudaThreadSynchronize();
                cudaMemcpy(&max, d_max, sizeof(unsigned int), cudaMemcpyDeviceToHost);
 
                //Time measure end
                QueryPerformanceCounter(&end);
                diff = (end.QuadPart - start.QuadPart) / (freq.QuadPart / 1000000.0);
                printf("Local3 max, elapsend time: %lf\n", diff);
                printf("Local3 max, max value : %d\n", max);
            }
 
            //using part of tid
            for (int i = 0; i < 10; i++) {
                max = 0;
                cudaMemcpy(d_data, data, sz, cudaMemcpyHostToDevice);
                cudaMemcpy(d_max, data, sizeof(unsigned int), cudaMemcpyHostToDevice);
                //Time measure start
                QueryPerformanceFrequency(&freq);
                QueryPerformanceCounter(&start);
                getMax_local4 << <grid, threads >> > (d_data, d_max, N);
                //Wait until thread function(gpu) is over
                cudaThreadSynchronize();
                cudaMemcpy(&max, d_max, sizeof(unsigned int), cudaMemcpyDeviceToHost);
 
                //Time measure end
                QueryPerformanceCounter(&end);
                diff = (end.QuadPart - start.QuadPart) / (freq.QuadPart / 1000000.0);
                printf("Local4 max, elapsend time: %lf\n", diff);
                printf("Local4 max, max value : %d\n", max);
            }
 
 
 
 
 
            /******************************
               Binary max
            *******************************/
            for (int i = 0; i < 10; i++) {
                if (N%THREADS != 0) {
                    puts("N must be times of THREADS when using binary reduction");
                    return 0;
                }
                max = 0;
                cudaMemcpy(d_data, data, sz, cudaMemcpyHostToDevice);
                cudaMemcpy(d_max, data, sizeof(unsigned int), cudaMemcpyHostToDevice);
                //Time measure start
                QueryPerformanceFrequency(&freq);
                QueryPerformanceCounter(&start);
                getMax_binary << <grid, threads, THREADS*a * sizeof(unsigned int) >> > (d_data, d_max, N, a);
                //Wait until thread function(gpu) is over
                cudaThreadSynchronize();
                cudaMemcpy(&max, d_max, sizeof(unsigned int), cudaMemcpyDeviceToHost);
 
                //Time measure end
                QueryPerformanceCounter(&end);
                diff = (end.QuadPart - start.QuadPart) / (freq.QuadPart / 1000000.0);
                printf("Binary max, elapsend time: %lf\n", diff);
                printf("Binary max, max value : %d\n", max);
            }
 
            for (int i = 0; i < 10; i++) {
                if (N%THREADS != 0) {
                    puts("N must be times of THREADS when using binary reduction");
                    return 0;
                }
                max = 0;
                cudaMemcpy(d_data, data, sz, cudaMemcpyHostToDevice);
                cudaMemcpy(d_max, data, sizeof(unsigned int), cudaMemcpyHostToDevice);
                //Time measure start
                QueryPerformanceFrequency(&freq);
                QueryPerformanceCounter(&start);
                getMax_binary2 << <grid, threads, THREADS*a * sizeof(unsigned int) >> > (d_data, d_max, N, a);
                //Wait until thread function(gpu) is over
                cudaThreadSynchronize();
                cudaMemcpy(&max, d_max, sizeof(unsigned int), cudaMemcpyDeviceToHost);
 
                //Time measure end
                QueryPerformanceCounter(&end);
                diff = (end.QuadPart - start.QuadPart) / (freq.QuadPart / 1000000.0);
                printf("Binary max2, elapsend time: %lf\n", diff);
                printf("Binary max2, max value : %d\n", max);
            }
 
            //unroll
            for (int i = 0; i < 10; i++) {
                if (N%THREADS != 0) {
                    puts("N must be times of THREADS when using binary reduction");
                    return 0;
                }
                max = 0;
                cudaMemcpy(d_data, data, sz, cudaMemcpyHostToDevice);
                cudaMemcpy(d_max, data, sizeof(unsigned int), cudaMemcpyHostToDevice);
                //Time measure start
                QueryPerformanceFrequency(&freq);
                QueryPerformanceCounter(&start);
                getMax_binary3<64> << <grid, threads, THREADS*a * sizeof(unsigned int) >> > (d_data, d_max, N, a);
                //Wait until thread function(gpu) is over
                cudaThreadSynchronize();
                cudaMemcpy(&max, d_max, sizeof(unsigned int), cudaMemcpyDeviceToHost);
 
                //Time measure end
                QueryPerformanceCounter(&end);
                diff = (end.QuadPart - start.QuadPart) / (freq.QuadPart / 1000000.0);
                printf("Binary max3, elapsend time: %lf\n", diff);
                printf("Binary max3, max value : %d\n", max);
            }
 
            //more unroll
            for (int i = 0; i < 10; i++) {
                if (N%THREADS != 0) {
                    puts("N must be times of THREADS when using binary reduction");
                    return 0;
                }
                max = 0;
                cudaMemcpy(d_data, data, sz, cudaMemcpyHostToDevice);
                cudaMemcpy(d_max, data, sizeof(unsigned int), cudaMemcpyHostToDevice);
                //Time measure start
                QueryPerformanceFrequency(&freq);
                QueryPerformanceCounter(&start);
                getMax_binary4<64> << <grid, threads, THREADS*a * sizeof(unsigned int) >> > (d_data, d_max, N, a);
                //Wait until thread function(gpu) is over
                cudaThreadSynchronize();
                cudaMemcpy(&max, d_max, sizeof(unsigned int), cudaMemcpyDeviceToHost);
 
                //Time measure end
                QueryPerformanceCounter(&end);
                diff = (end.QuadPart - start.QuadPart) / (freq.QuadPart / 1000000.0);
                printf("Binary max4, elapsend time: %lf\n", diff);
                printf("Binary max4, max value : %d\n", max);
            }
 
 
            //grid 수를 절반으로
            for (int i = 0; i < 10; i++) {
                if (N%THREADS != 0) {
                    puts("N must be times of THREADS when using binary reduction");
                    return 0;
                }
                max = 0;
                cudaMemcpy(d_data, data, sz, cudaMemcpyHostToDevice);
                cudaMemcpy(d_max, data, sizeof(unsigned int), cudaMemcpyHostToDevice);
                //Time measure start
                QueryPerformanceFrequency(&freq);
                QueryPerformanceCounter(&start);
                getMax_binary5<64> << <grid / 2, threads, THREADS*a * sizeof(unsigned int) >> > (d_data, d_max, N / 2, a);
                //Wait until thread function(gpu) is over
                cudaThreadSynchronize();
                cudaMemcpy(&max, d_max, sizeof(unsigned int), cudaMemcpyDeviceToHost);
 
                //Time measure end
                QueryPerformanceCounter(&end);
                diff = (end.QuadPart - start.QuadPart) / (freq.QuadPart / 1000000.0);
                printf("Binary max5(grid/2), elapsend time: %lf\n", diff);
                printf("Binary max5(grid/2), max value : %d\n", max);
            }
 
 
            //threads 수를 절반으로
            for (int i = 0; i < 10; i++) {
                if (N%THREADS != 0) {
                    puts("N must be times of THREADS when using binary reduction");
                    return 0;
                }
                max = 0;
                cudaMemcpy(d_data, data, sz, cudaMemcpyHostToDevice);
                cudaMemcpy(d_max, data, sizeof(unsigned int), cudaMemcpyHostToDevice);
                //Time measure start
                QueryPerformanceFrequency(&freq);
                QueryPerformanceCounter(&start);
                getMax_binary6<64> << <grid, threads / 2, THREADS*a / 2 * sizeof(unsigned int) >> > (d_data, d_max, N / 2, a / 2);
                //Wait until thread function(gpu) is over
                cudaThreadSynchronize();
                cudaMemcpy(&max, d_max, sizeof(unsigned int), cudaMemcpyDeviceToHost);
 
                //Time measure end
                QueryPerformanceCounter(&end);
                diff = (end.QuadPart - start.QuadPart) / (freq.QuadPart / 1000000.0);
                printf("Binary max6(threads/2), elapsend time: %lf\n", diff);
                printf("Binary max6(threads/2), max value : %d\n", max);
            }
 
 
            //use else if
            for (int i = 0; i < 10; i++) {
                if (N%THREADS != 0) {
                    puts("N must be times of THREADS when using binary reduction");
                    return 0;
                }
                max = 0;
                cudaMemcpy(d_data, data, sz, cudaMemcpyHostToDevice);
                cudaMemcpy(d_max, data, sizeof(unsigned int), cudaMemcpyHostToDevice);
                //Time measure start
                QueryPerformanceFrequency(&freq);
                QueryPerformanceCounter(&start);
                getMax_binary7<64> << <grid, threads / 2, THREADS*a / 2 * sizeof(unsigned int) >> > (d_data, d_max, N / 2, a / 2);
                //Wait until thread function(gpu) is over
                cudaThreadSynchronize();
                cudaMemcpy(&max, d_max, sizeof(unsigned int), cudaMemcpyDeviceToHost);
 
                //Time measure end
                QueryPerformanceCounter(&end);
                diff = (end.QuadPart - start.QuadPart) / (freq.QuadPart / 1000000.0);
                printf("Binary max7(else if), elapsend time: %lf\n", diff);
                printf("Binary max7(else if), max value : %d\n", max);
            }
 
 
            //not use id
            for (int i = 0; i < 10; i++) {
                if (N%THREADS != 0) {
                    puts("N must be times of THREADS when using binary reduction");
                    return 0;
                }
                max = 0;
                cudaMemcpy(d_data, data, sz, cudaMemcpyHostToDevice);
                cudaMemcpy(d_max, data, sizeof(unsigned int), cudaMemcpyHostToDevice);
                //Time measure start
                QueryPerformanceFrequency(&freq);
                QueryPerformanceCounter(&start);
                getMax_binary8<64> << <grid, threads / 2, THREADS*a / 2 * sizeof(unsigned int) >> > (d_data, d_max, N / 2, a / 2);
                //Wait until thread function(gpu) is over
                cudaThreadSynchronize();
                cudaMemcpy(&max, d_max, sizeof(unsigned int), cudaMemcpyDeviceToHost);
 
                //Time measure end
                QueryPerformanceCounter(&end);
                diff = (end.QuadPart - start.QuadPart) / (freq.QuadPart / 1000000.0);
                printf("Binary max8, elapsend time: %lf\n", diff);
                printf("Binary max8, max value : %d\n", max);
            }
 
            a *= 2;
        }
    }
    
    
 
    return 0;
}
 
 
 
 
 
 
 
 

3. 참고

https://developer.download.nvidia.com/assets/cuda/files/reduction.pdf?ncid=afm-chs-44270&ranMID=44270&ranEAID=a1LgFw09t88&ranSiteID=a1LgFw09t88-bKHPLM.ndfvRA2bRYR_vrA

질문이나 지적 있으시면 댓글로 남겨주세요~

도움 되셨으면 하트 꾹!