[SlideShare] Parellel Programming & CUDA (병렬프로그래밍과 CUDA)

병렬프로그래밍과 Cuda from Seok-joon Yun

CUDA+발표자료+대본.pdf

CUDA Compiler, Linker 의 Option Setting

CUDA의 경우 Property Sheet 로 Compliler, Linker의 설정값 저장이 되지 않는다.
따라서 CUDA를 적용한 Project 마다 따로 설정을 해 줘야 한다.

먼저 Solution Exploer Window (솔루션 탐색기) 에서 해당 Project에서 우클릭 한뒤 Properties (속성) 을 누른다.

그럼 아래 그림과 같이 가장 밑에 CUDA Linker, CUDA Runtime API 가 보인다.

만약 보이지 않으면 CUDA Project 설정 페이지를 참조하기 바란다.
 ( http://icysword.blog.me/140201834456 의 3번 항목 )

참고로 Geforce GTX 770 의 설정값이다. 다른 GPU의 경우 해당 GPU의 Spec에 맞게 수정해서 설정하기 바란다.

1. Win32 - Release

CUDA Linker -> General -> Target Machine Platform : x86
                       GPU -> GPU_Architecture (1) : sm_30
                       Host -> Runtime Library : Multi-Threaded DLL (/MD) (MFC Project 설정과 동일하게 맞춤)
CUDA Runtime API -> GPU -> GPU Architecture (1), (2) : sm_30
                                Host -> Target MAchine Platform : x86
                                            Runtime Library : Multi-Threaded DLL (/MD) (MFC Project 설정과 동일하게 맞춤)


2. Win32 - Debug

CUDA Linker -> General -> Target Machine Platform : x86
                       GPU -> Generate GPU Debug Information : 예 (-G)
                                   GPU_Architecture (1), (2) : sm_30
                       Host -> Runtime Library : Multi-Threaded Debug DLL (/MDd) (MFC Project 설정과 동일하게 맞춤)
CUDA Runtime API -> GPU -> Generate GPU Dubug Information : 예 (-G)
                                           Generate Line Number Informarion : 예 (-lineinfo)
                                           GPU Architecture (1) : sm_30
                                Host -> Target MAchine Platform : x86
                                            Generate Host Debug Information :  예 (-D_NEXUS_DEBUD -g)
                                            Optimization : Disabled (/Od)
                                            Runtime Library : Multi-Threaded Debug DLL (/MDd) (MFC Project 설정과 동일하게 맞춤)

3. x64

위 설정에서 Target Machine Platform 만 x64 로 변경

CUDA Random

- CUDA 에서 Random 이용하기 

1. Host 함수에서 Device용 curandState Pointer를 생성하여 할당하기

2. Kernel 함수로 curandState 를 초기화 : 병렬화 할 크기만큼 Seed 생성 (1번만 하면 됨)

3. Kernel 함수로 curandState를 이용하여 random 생성

[-] Collapse

#include <cuda.h>
#include <curand_kernel.h>

#include <stdio.h>
#include <iostream>

using namespace std;

__global__ void setup_kernel(curandState *state, int seed)
{
    int id = threadIdx.x + blockIdx.x * blockDim.x;
    curand_init(seed,id,0, &state[id]);
}

__global__ void generate_kernel(curandState *state, float* result)
{
    int id = threadIdx.x + blockIdx.x * blockDim.x;
    result[id] = curand_normal(&state[id]);
}

int main()
{
    float* dev_result;
    float* host_result;

    curandState *devStates;

    cudaMalloc((void**)&devStates, 1024 * 1000 * sizeof(curandState));
    cudaMalloc((void**)&dev_result, 1024 * 1000 * sizeof(float));

    int a = 0;
    setup_kernel<<<1024,1000>>>(devStates, a);

    while(true)
    {
        generate_kernel<<<1024,1000>>>(devStates, dev_result);

        host_result = new float[1024 * 1000];
        cudaMemcpy(host_result, dev_result, 1024 * 1000 * sizeof(float), cudaMemcpyDeviceToHost);

        float sum = 0.0f;
        float sumSq = 0.0f;
        int count = 1024 * 1000;
        for (int i = 0; i < count; i++)
        {
            sum += host_result[i];
            sumSq += host_result[i] * host_result[i];
        }

        float avg = sum / count;
        float stDev = ( sumSq - (sum * sum / count) ) / ( count - 1);

        cout << "Avg : " << avg << " stDev : " << stDev << "   100번째 숫자 : " << host_result[100] << endl;
        getchar();
    }

    delete [] host_result;
    cudaFree(devStates);
    cudaFree(dev_result);

    return 0;
}

- 대표적인 CUDA용 Random 함수 몇가지

curand_uniform(&state) : 0.0f ~ 1.0f 사이의 실수
curand_normal(&state) : 0.0f ~ 1.0f 사이의 정규분포된 실수
curand_log_normal(&state, mean, stddev) : Log-normally 분산 실수
curand_poisson(&state, lambda) : 포아송분포의 unsigned 정수값

CUDA : Parallel Nsight를 이용한 Debugging 및 Performance Analysis

* 필자의 환경 
 - Win 7 Pro. 64bit
 - Visual Studio 2008 SP1
 - CUDA 5.5, Nsight VS 3.1
 - Geforce GTX 770

- Nsight 환경설정

1. Nsight Monitor 실행
  (필자 : 모든 프로그램\NVIDIA Corporation\Nsight Visual Studio Edition 3.1\Nsight Monitor)

2. 트레이에서 Nsight Monitor 을 우클릭 한후 Options.. 선택

3. General -> Microsoft Display Driver ->WDDM TDR enabled : False 로 설정

4. WPF용 D3D 가속을 중지하기
 - Nsight 가 설치된 폴더로 이동
   (필자 : C:\Program Files (x86)\NVIDIA Corporation\Nsight Visual Studio Edition 3.1\Host\Common)

 - DisableWpfHardwareAcceleration.reg 더블클릭하여 레지스트리에 등록

5. 컴퓨터 재부팅

 - Debugging 하기

1. 솔루션 탐색기에서  CUDA를 Debugging 할 Project에서 우클릭 한후 Nsight User Properties  선택

2. Launch -> Launch Oprion -> Connection name : localhost 로 설정

3. Nsight Monitor 실행

4. Visual Studio 상단 메뉴에서 Nsight 선택 후 Enalbe CUDA Memory Checker 활성

5. 이제 원하는 곳에 Break Point를 적용한 후 Nsight -> Start CUDA Debugging 을 눌러서 Debugging이 가능합니다.

 - CUDA Performance Analysis

1. Visual Studio 상단 메뉴에서 Nsight 선택 후 Start Performance Analysisr 선택

2. Activity1.nvact 화면이 생성되는데 가장 아래로 스크롤하여 CUDA 선택을 하면 Application Control의 Lanuch 가 활성화 됨.

3. Lanuch버튼을 클릭하여 어플리케이션을 실행

4. 어플리케이션 종료 후 결과 보고 화면이 화면에 출력됨

CUDA Rolling을 이용한 Thread Optimization

- Rolling을 이용한 Thread 최적화 작업

CUDA로 연산을 병렬화로 구현하는 경우 GRID x THREAD 로 병렬화를 구분한다.

GRID 는 SM (Stream Multiprocessor)에 할당이 되고,
THREAD는 SM 안에 있는 SP (Stream Processor) 즉, CUDA Core에 할당된다.
SM안에 있는 SP 들은 공유 메모리 (Shared Memory)를 이용하여 빠른 속도로 서로 Data를 주고 받을 수 있다.

만약 2000 x 1000 의 작업을 병렬화로 구현해야 하는 경우 GRID : 2000, THREAD : 1000 개로 병렬화를 하면
가장 소스코드도 직관적으로 읽기 쉽고, 구현하기도 편하다.

SM 은 원래 서로 계산값의 영향을 미치지 않는 경우가 많아서 2000개로 나누어도 상관이 없다.
어차피 GPU는 SM 을 자신이 동시에 실행할 수 있는 갯수만큼 실행을 한 뒤 다음 SM들을 실행할 것이다.

하지만 THREAD의 경우는 Thread 간의 동기화 (Synchronize) 를 하는 Code도 들어가야 하며, 물리적으로 SM 안의 CUDA Core의 수가 넘어가는 Thread를 할당 할경우 Context Switching 도 발생하게 된다.
더군다나 GPU 가 한번에 실행시키는 Thread의 수는 생각보다 크지가 않다.

(GTX 770의 경우 동시 실행하는 SM은 8개, SM당 CUDA Core는 192개, WARP 는 32 이다.
 즉, SM 당 물리적으로 할당하는 Thread 수는 192 x 4 = 768개이며, 동시에 실행 가능한 Thread는 한번에 32개이다.)

그러므로 성능을 생각해서는 Thread의 수는 각 GPU의 WARP 만큼 하는것이 가장 이상적이다.
하지만, 굳이 성능을 위해서 머리아프게 Rolling을 해가면서 Code를 복잡하게 짜야할까 ?
굳이 성능이 아니더라도, 호환성을 위해서도 필요하다.
예를 들어서 위의 예제의 경우 Thread의 수를 1000개로 할당했는데, 다른 GPU의 경우 Thread 할당수가 200개 밖에 되지 않는 다면, 해당 Code로 실행을 하면 원하는 결과가 나오지 않는다.

아래의 Code를 참조하여 별로 어렵지 않게 Thread 수를 GPU의 WARP 만큼으로 설정을 하고 원래의 Code를 Rolling Code로 변환이 가능하다.
[-] Collapse

__global__ void Cuda_Test(float* i_fArray)
{
    int m = threadIdx.x - const_iNum[0] / 2;
    float fAng = (float)m * const_fPi[1] / const_iNum[1];
    i_fArray[threadIdx.x] = cosf(fAng) * cosf(fAng) + sinf(fAng) * sinf(fAng) - cosf(fAng);
}

[-] Collapse

__global__ void Cuda_Test(float* i_fArray)
{
    int idx = threadIdx.x;
    while (idx < const_iNumSensor[0])
    {
        int m = idx - const_iNum[0] / 2;
        float fAng = (float)m * const_fPi[1] / const_iNum[1];
        i_fArray[idx] = cosf(fAng) * cosf(fAng) + sinf(fAng) * sinf(fAng) - cosf(fAng);

        idx += blockDim.x;
    }
    __syncthreads();
}

과정을 말로 설명하자면 다음과 같다.

1. index를 하나 선언하여 기존의 threadIdx.x 의 자리에 넣는다.
2. while 문의 조건으로 index 가 기존 Thread의 크기보다 작을 경우 실행하도록 설정한다.
3. while 문 마지막에 index에 새로 설정한 Thread의 크기만큼 더한다.

4. while 문 뒤에 다른 연산이 더 있다면 __syncthreads(); 를 사용하여 Thread를 동기화 한다.

CUDA Shared Memory Dynamic Allocation (동적할당)

- Shared Memory 동적할당 (Dynamic Allocation)

Shared Memory 는 선언하면서 초기화도 할 수가 없으며, 배열 선언시 배열의 개수는 Constant 여야 한다.

이럴 경우 extern 으로 Shared Memory를 선언하면서 배열의 크기를 지정하지 않을 수 있다.

[-] Collapse

extern __shared__ float sfArray [];

단 이럴 경우에는 반드시 kernel 함수를 호출할 때 Shared Memory의 크기를 지정해 줘야 한다.
(당연히 Shared Memory의 크기는 Byte 단위로 지정한다.)
왠만하면 GPU 스펙의 Shared Memory 최고크기를 지정해 놓으면 된다. (GTX 770의 경우 65535)

[-] Collapse

Cuda_TestKernel<<<NUM_BLOCK, NUM_THREAD, SIZE_SHARED_MEMORY>>>(dev_fArray);

1개의 Kernel 함수에 다수의 extern shared array가 있을 경우 그 시작번지수는 모두 동일 하다.

예를 들어서 1000개의 float, 300개의 int, 500개의 float이 필요할 경우 아래와 같이 구현이 가능하다.
[-] Collapse

__shared__ float sfDa;
extern __shared__ float sfArray [];
__shared__ float *sfCs;
__shared__ int *siDs;
__shared__ float *sfEs;

if (threadIdx.x == 0)
{
    sfDa = (float)blockIdx.x;
    sfCs = (float *)sfArray;
    siDs = (int *)&sfCs[1000];
    sfEs = (float *)&siDs[300];
}
__syncthreads();

CUDA : Performance Optimization Example

[-] Collapse

// Step 1 Interleaved Addressing with Divergent Branching

__global__ void reduce0(int *g_idata, int *g_odata)
{
    extern __shared__ int sdata[];

    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

    sdata[tid] = g_idata[i];
    __syncthreads();

    for (unsigned int s = 1; s < blockDim.x; s *= 2)
    {
        if ((tid % (2*s)) == 0)
            sdata[tid] += sdata[tid + s];
        __syncthreads();
    }

    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}

// Step 2 Interleaved Addressing with Bank Collision
// if 문의 % 연산자 제거

__global__ void reduce1(int *g_idata, int *g_odata)
{
    extern __shared__ int sdata[];

    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

    sdata[tid] = g_idata[i];
    __syncthreads();

    for (unsigned int s = 1; s < blockDim.x; s *= 2)
    {
        int index = 2 * s * tid;
        if (index < blockDim.x)
            sdata[index] += sdata[index + s];
        __syncthreads();
    }

    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}

// Step 3 Solve Bank Collision
// Data 전송 Bank를 인접하게 위치하도록 조정

__global__ void reduce3(int *g_idata, int *g_odata)
{
    extern __shared__ int sdata[];

    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

    sdata[tid] = g_idata[i];
    __syncthreads();

    for (unsigned int s = blockDim.x/2; s > 0; s >>= 1)
    {
        if (tid < s)
            sdata[tid] += sdata[tid + s];
        __syncthreads();
    }

    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}

// Step 4 첫번째 요소 로딩과 동시에 합 계산
// Shared Memory에 최소 값을 초기화 할때 2개의 값을 넣으면서, Kernel 실행 수를 절반으로 줄임
// N개의 Data 생성시 N / 2 개의 Thread 생성

__global__ void reduce4(int *g_idata, int *g_odata)
{
    extern __shared__ int sdata[];

    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

    sdata[tid] = g_idata[i] + g_idata[i+blockDim.x];
    __syncthreads();

    for (unsigned int s = blockDim.x/2; s > 0; s >>= 1)
    {
        if (tid < s)
            sdata[tid] += sdata[tid + s];
        __syncthreads();
    }

    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}

// Step 5 Last Warp Unrolling
// 마지막 Warp 부분의 rolling 제거

__device__ void warpReduce(volatile int* sdata, 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 reduce5(int *g_idata, int *g_odata)
{
    extern __shared__ int sdata[];

    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

    sdata[tid] = g_idata[i] + g_idata[i+blockDim.x];
    __syncthreads();

    for (unsigned int s = blockDim.x/2; s > 32; s >>= 1)
    {
        if (tid < s)
            sdata[tid] += sdata[tid + s];
        __syncthreads();
    }

    if (tid < 32) warpReduce(sdata, tid);

    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}

// Step 6 All Code Unrolling
// 모든 rolling 제거 단 block size는 512 이하로 한다는 전제

template <unsigned int blockSize>
__device__ void warpReduce2(volatile int* sdata, int tid)
{
    if (blockSize >= 64) sdata[tid] += sdata[tid + 32];
    if (blockSize >= 32) sdata[tid] += sdata[tid + 16];
    if (blockSize >= 16) sdata[tid] += sdata[tid + 8];
    if (blockSize >= 8) sdata[tid] += sdata[tid + 4];
    if (blockSize >= 4) sdata[tid] += sdata[tid + 2];
    if (blockSize >= 2) sdata[tid] += sdata[tid + 1];
}

template <unsigned int blockSize>
__global__ void reduce6(int *g_idata, int *g_odata)
{
    extern __shared__ int sdata[];

    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

    sdata[tid] = g_idata[i] + g_idata[i+blockDim.x];
    __syncthreads();

    if (blockSize >= 512) {
        if (tid < 256) sdata[tid] += sdata[tid + 256];
        __syncthreads();
    }

    if (blockSize >= 256) {
        if (tid < 128) sdata[tid] += sdata[tid + 128];
        __syncthreads();
    }

    if (blockSize >= 128) {
        if (tid < 64) sdata[tid] += sdata[tid + 64];
        __syncthreads();
    }

    if (tid < 32) warpReduce2<blockSize>(sdata, tid);

    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}

// Step 7 Algorithm Cascading
// 호출하는 Thread 수를 N / log N 으로 생성하도록 하여 Thread 할당하는 시간을 줄임

template <unsigned int blockSize>
__global__ void reduce7(int *g_idata, int *g_odata, unsigned int n)
{
    extern __shared__ int sdata[];

    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

    sdata[tid] = 0;
    while (i < n)
    {
        sdata[tid] += g_idata[i] + g_idata[i + blockSize];
        i += blockDim.x;
    }
    __syncthreads();

    if (blockSize >= 512) {
        if (tid < 256) sdata[tid] += sdata[tid + 256];
        __syncthreads();
    }

    if (blockSize >= 256) {
        if (tid < 128) sdata[tid] += sdata[tid + 128];
        __syncthreads();
    }

    if (blockSize >= 128) {
        if (tid < 64) sdata[tid] += sdata[tid + 64];
        __syncthreads();
    }

    if (tid < 32) warpReduce2<blockSize>(sdata, tid);

    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}