CUDA

개요

• Host: CPU
• Device: GPU
• device code: code run in GPU (written in c (API), Compile to execute in GPU)
• host code: code run in CPU (written in c)

우선 CPU메모리에서 GPU메모리로 복사가 일어난뒤, (PCI Bus를 통해서) GPU 코드를 로드하고 실행시키게 된다. 코드가 실행되고 결과는 다시 GPU메모리에서 CPU메모리로 Copy되게 된다. 메모리와 디바이스는 서로 메모리가 분리되어 있기 때문에, 메모리의 위치가 중요한 역활을 한다. nvcc는 device code와 host code를 분리하여 서로 다른 컴파일러로 컴파일하고 결과를 하나로 합치는 역활을 한다.

키워드

• __global__: 이 키워드가 붙은 함수는 host에서 호출되어 device에서 실행된다.
• mykernel<<<1,1>>>(): triple brackets은 CPU가 GPU를 호출한다고 명시하는 것을 말한다. __global__을 붙히고 triple brackets를 쓰지 않으면, nvcc는 error을 내보낸다.

메모리 관련 함수

GPU 커널 함수에 넘기는 포인터는 GPU에 할당된 메모리여야 한다. 즉 포이터가 GPU메모리에 위치하는가, CPU메모리에 위치하느냐에 따라서 분리하여 생각해야 된다. 따라서 다음과 같은 함수를 이용해서 메모리를 할당하고 해제하는 것이 필요하다. 여기서 중요한 것은 메모리 할당 함수에 넘기는 포인터는 포인터의 주소값 즉 더블 포인터를 넘겨야 한다. 왜냐하면, 쿠다는 항상 에러 값을 리턴하고자 하기 때문에 기존의 C처럼 할당된 부분이 리턴되는 것이 아니라 포인터의 주소값을 이용해서 포인터 값 그자체를 바꾸기 때문이다.

• cudaMalloc
• cudaFree
• cudaMemcpy

__global__ void add(int *a, int *b, int *c) {
    *c = *a + *b;
}

int main(void) {
    int a, b, c;
    int *da, *db, *dc;
    cudaMalloc((void **)&da, sizeof(int));
    cudaMalloc((void **)&db, sizeof(int));
    cudaMalloc((void **)&dc, sizeof(int));
    // 만약, da에 host에서 접근하면 segmentation fault가 발생한다.
    // 왜냐하면, da는 device에서 사용하는 포인터이기 때문이다. 

    // Host -> Device
    cudaMemcpy(da, &a, sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(db, &b, sizeof(int), cudaMemcpyHostToDevice);

    // Lauch add() kernel
    add<<<1,1>>>(da, db, dc);

    // Device -> Host
    cudaMemcpy(&c, dc, sizeof(int), cudaMemcpyDeviceToHost);

    // Cleanup
    cudaFree(da); cudaFree(db); cudaFree(dc);
}

Block

GPU는 한번에 많은 데이터를 처리할 수 있는데, 이는 Block을 사용하여 이루어진다. 예를 들어서 add가 병렬적으로 처리된다고 할경우 각각의 add는 block이라고 불리운다. 또한 이러한 block들의 묶음을 grid라고 한다.

__global__ void add(int *a, int *b, int *c) {
    c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];
}

int main(void) {
    int *a, *b, *c;
    int *da, *db, *dc;

    a = malloc(N * sizeof(int));
    b = malloc(N * sizeof(int));
    c = malloc(N * sizeof(int));

    cudaMalloc((void **)&da, sizeof(int) * N);
    cudaMalloc((void **)&db, sizeof(int) * N);
    cudaMalloc((void **)&dc, sizeof(int) * N);
    // 만약, da에 host에서 접근하면 segmentation fault가 발생한다.
    // 왜냐하면, da는 device에서 사용하는 포인터이기 때문이다. 

    // Host -> Device
    cudaMemcpy(da, &a, sizeof(int) * N, cudaMemcpyHostToDevice);
    cudaMemcpy(db, &b, sizeof(int) * N, cudaMemcpyHostToDevice);

    // Lauch add() kernel (N block)
    // Lauch N copies of add with add<<<N,1>>>(), each block is distinguished by blockIdx
    add<<<N,1>>>(da, db, dc);

    // Device -> Host
    cudaMemcpy(c, dc, sizeof(int) * N, cudaMemcpyDeviceToHost);

    // Cleanup
    cudaFree(da); cudaFree(db); cudaFree(dc);
}

Threads

각각의 블락은 thread로 분리될 수 있다. 이는 threadIdx로 구분된다. 또한 이는 add<<<1,N>>>()처럼 N개의 스레드를 사용한다고 명시함으로써 사용되게 된다. Warp는 32(64)개의 스레드로 구성됨으로, GPU내의 warp scheduler가 이 스레드를 K개의 warp로 나누어서 각각의 블럭을 처리하게 된다. 기능적으로는 위의 block을 사용한것과 차이가 없지만, 성능의 차이가 나게 된다.

그런데 block으로 나누어진다면, 스레드가 필요없는 것처럼 보인다. 그렇다면 왜 스레드를 사용하여야 하는 것인가? 이는 Communicate와 Synchronize의 측면에서 장점이 있기 때문이다.

그렇다면 과연 블럭과 스레드를 동시에 사용할 수 있는 방법이 있어야 한다.

한 블럭당 M개의 스레드들이 생성된다면, 각각의 스레드에 대한 unique한 인덱스는 다음과 같이 주어진다. (blockIdx와 threadIdx는 0부터 시작한다.)

int index = threadIdx.x + blockIdx.x * M

여기서 M은 built-in variable로 blockDim.x이렇게 주어진다.

__global__ void add(int *a, int *b, int *c, int n) {
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    // 여기서 n을 통해서 index가 N을 넘는 것을 방지한다.
    if (index < n)
        c[index] = a[index] + b[index];
}

#define N (2048*2048)
#define THREADS_PER_BLOCK 512

int main(void) {
    int *a, *b, *c;
    int *da, *db, *dc;

    a = malloc(N * sizeof(int));
    b = malloc(N * sizeof(int));
    c = malloc(N * sizeof(int));

    cudaMalloc((void **)&da, sizeof(int) * N);
    cudaMalloc((void **)&db, sizeof(int) * N);
    cudaMalloc((void **)&dc, sizeof(int) * N);
    // 만약, da에 host에서 접근하면 segmentation fault가 발생한다.
    // 왜냐하면, da는 device에서 사용하는 포인터이기 때문이다. 

    // Host -> Device
    cudaMemcpy(da, &a, sizeof(int) * N, cudaMemcpyHostToDevice);
    cudaMemcpy(db, &b, sizeof(int) * N, cudaMemcpyHostToDevice);

    // Lauch add() kernel (N block)
    // Lauch N copies of add with add<<<N,M>>>(), each block is distinguished by blockIdx
    add<<<N/THREADS_PER_BLOCK,THREADS_PER_BLOCK>>>(da, db, dc, N);

    // Device -> Host
    cudaMemcpy(c, dc, sizeof(int) * N, cudaMemcpyDeviceToHost);

    // Cleanup
    free(a); free(b); free(c);
    cudaFree(da); cudaFree(db); cudaFree(dc);
}

Shared Memory

Shared Memory는 캐쉬와는 다르게 프로그래머가 정확이 어떠한 변수가 메모리 영역에 올라갈지를 결정할 수 있다. Shared Memory영역은 같은 스레드간에 공유될 수 있다. 만약 모든 병렬 처리가 block으로 나누어진다면, 스레드가 필요없는 것처럼 보인다. 그렇다면 왜 스레드를 사용하여야 하는 것인가? 이는 Communicate와 Synchronize의 측면에서 장점이 있기 때문이다. 예를 들어서 1D 배열에서 주위의 변수들의 합을 구하는 프로그램을 짠다고 해보자. 이러할 경우 만약 [-3, 나, +3] 의 영역의 합을 구한다면 같은 메모리 위치에 모두 7번 접근해서 각각의 블럭을 처리해야 한다. 그러나 이러한 일은 매우 비효율 적인 일이다. 따라서 스레드를 이용해서 shared memory를 통해서 이러한 한계를 극복할 수 있다.

// [가장자리 0, 가장자리 2, ..., 가장자리 + RADIUS, 가운데 0, ..., 가운데 BLOCK_SIZE - 2 *  RADIUS, 가장자리 ...]
__global__void stencil_1d(int *in, int *out) {
    // 가장자리를 제외한 부분의 영역을 공유 메모리에 적재
    __shared__ int temp[BLOCK_SIZE + 2 * RADIUS];
    int gindex = threadIdx.x + blockIdx.x * blockDim.x;
    int lindex = threadIdx.x + RADIUS;

    // 가장자리 부분을 공유 메모리에 적재
    temp[lindex] = in[gindx];
    if (threadIdx.x < RADIUS) {
        temp[lindex - RADIUS] = in[gindex - RADIUS];
        temp[lindex + BLOCK_SIZE] = in[gindex + BLOCK_SIZE];
    }
}

Data race

void __syncthreads();

이 함수는 블럭당 모든 스레드를 동기화 시킨다. 즉 RAW, WAR, WAW를 없앨 수 있다.

Host & Device

커널은 비동기적으로 작동하게 된다. 즉 커널과 CPU의 작동은 서로 비동기적이다. 따라서 CPU는 Result를 소비하기 전에 동기화될 필요가 있다.

• cudaMemcpu: 호스트를 잠시 멈추고 GPU가 결과를 가져오기를 기다린후, 결과를 가져온다.
• cudaMemcpyAsync: 비동기적으로 결과를 가져온다.
• cudaDeviceSynchronize: 호스트를 GPU에서 결과 처리가 완료될까지 기다리게 한다.

모든 CUDA API는 Error code를 만들어 낸다. cudaError_t.

• cudaError_t e = ...SomeCUDAOperation...
• cudaGetLastError: 마지막 에러를 가져온다.

CUDA는 여러 GPU중 하나를 선택할 수 있다. 이는 여러 Host가 하나의 GPU를 공유하거나, 여러 GPU를 하나의 Host가 공유하는 상황 모두에서 쓰일 수 있다. 또한 하나의 GPU를 가상머신 처럼 여러 GPU로 쪼개서 사용할 수도 있는데, 이는 클라우드와 같이 하나의 GPU를 공유하는 시스템에서 유용하게 사용될 수 있다.

• cudaGetDeviceCount(int *count)
• cudaSetDevice(int device)
• cudaGetDevice(int *device)
• cudaGetDeviceProperties(cudaDeviceProp *prop, int device)
• cudaMemcpu: GPU와 GPU사이의 통신을 목적으로 사용될 수 있다.

예시

Convolution

__global__ void cuda_conv2d(double* image, double* kernel, double*result, int m, int n, int k) {
    int row_index = threadIdx.y + blockIdx.y * blockDim.y;
    int col_index = threadIdx.x + blockIdx.x * blockDim.x;

    if(row_index < m && col_index < n) {
        for(int i=0;i<k;++i) {
            result[row_index * n + col_index] += image[row_index * k + i] * kernel[n * i + col_index];
        }
    }
}

extern "C" void conv2d(double *image, double *kernel, double *result, int m, int n, int k)
{
    double *d_image, *d_kernel, *d_result;

    cudaMalloc((void**)&d_image, m * k * sizeof(double));
    cudaMalloc((void**)&d_kernel, n * k * sizeof(double));
    cudaMalloc((void**)&d_result, m * n * sizeof(double));

    cudaMemcpy(d_image, image, m * k * sizeof(double), cudaMemcpyHostToDevice);
    cudaMemcpy(d_kernel, kernel, n * k * sizeof(double), cudaMemcpyHostToDevice);

    int row_index = (n - 1) / 32 + 1;
    int col_index = (m - 1) / 32 + 1;

    dim3 dim_block = dim3(32, 32, 1);
    dim3 dim_thread = dim3(row_index, col_index);

    cuda_conv2d<<<dim_thread, dim_block>>>(d_image, d_kernel, d_result, m, n, k);

    cudaMemcpy(result, d_result, m * n * sizeof(double), cudaMemcpyDeviceToHost);

    cudaFree(d_image);
    cudaFree(d_kernel);
    cudaFree(d_result);
}

참고

1. http://haanjack.github.io/cuda/2016/03/27/cuda-prog-model.html
2. CUDA programming NVIDA Manual