CUDA Thread

기존 C에서 "Hello World"를 찍기 위한 코드는 다음과 같다.

int main(void) {
    printf("Hello World!\n");
    return 0;
}

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
 
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
 
#define NUM_DATA 512
 
__global__ void vecAdd(int *_a, int *_b, int *_c)
{
    int tID = threadIdx.x;
    _c[tID] = _a[tID] + _b[tID];
}
 
int main(void)
{
    int *a, *b, *c;
    int *d_a, *d_b, *d_c;
 
    int memSize = sizeof(int)*NUM_DATA;
    printf("%d elements, memSize = %d bytes\n", NUM_DATA, memSize);
 
    a = new int[NUM_DATA]; memset(a, 0, memSize);
    b = new int[NUM_DATA]; memset(b, 0, memSize);
    c = new int[NUM_DATA]; memset(c, 0, memSize);
 
    for (int i = 0; i < NUM_DATA; i++) {
        a[i] = rand() % 10;
        b[i] = rand() % 10;
    }
 
    cudaMalloc(&d_a, memSize);
    cudaMalloc(&d_b, memSize);
    cudaMalloc(&d_c, memSize);
 
    cudaMemcpy(d_a, a, memSize, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, memSize, cudaMemcpyHostToDevice);
    
    vecAdd <<<1, NUM_DATA >>>(d_a, d_b, d_c);
 
    cudaMemcpy(c, d_c, memSize, cudaMemcpyDeviceToHost);
 
    //check results
    bool result = true;
    for (int i = 0; i < NUM_DATA; i++) {
        if ((a[i] + b[i]) != c[i]) {
            printf("[%d] The resutls is not matched! (%d, %d)\n", i, a[i] + b[i], c[i]);
            result = false;
        }
    }
 
    if (result) printf("GPU works well!\n");
 
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
 
    delete[] a;
    delete[] b;
    delete[] c;
 
    return 0;
}

// Kernel call
vecAdd<<<1, NUM_DATA >>>(d_a, d_b, d_c);
cudaDeviceSynchronize(); // synchronization function

CUDA를 활용했을 때와 CPU만을 활용했을 때 차이를 측정해보자.

256크기 측정 결과

실제로 시간 측정을 해보면 CPU 버전보다 훨씬 느린 것을 확인할 수 있다.
데이터의 양이 적어 기본 작업보다 적게 나오는 것일까?

그렇다면 데이터의 양을 늘려보자(256 → 1024)

1024크기 측정 결과

약간 차이가 줄어들긴 했지만 여전히 CPU 버전보다 현저히 느리다.

더 늘린다면 어떻게 될까?

2048 크기 측정 결과

더 늘렸더니 이번엔 정확한 데이터가 나오지 않는다.

왜 이런 현상이 발생할까?

이를 이해하기 위해 CUDA Programming Model에 대한 이해가 필요하다.
우리가 사용하고 있는 GPU는 SIMT Architecture이다.

SIMT Architecture(Single Instruction, Multiple Threads)
Thread들을 그룹 단위로 관리하며, 실행한다.
즉, 모든 thread가 같은 프로그램 코드를 공유하는 것이다.(kernel)

여기서 SIMD와 SIMT의 차이를 생각하자.

SIMD는 하나의 instruction을 여러 개의 data에 적용시키는 것이고,
SIMT는 여러 개의 thread에 적용시키는 것이다.
즉, SIMT에서는 각 작업 항목에 스레드를 개별적으로 분기할 수 있다.

그렇다면, CUDA에서 thread들을 어떻게 관리하는지 계층도를 살펴보자.

CUDA Thread Hierarchy

• Thread
  • Basic processing unit
• Warp
  • 32 Threads
  • Basic exection unit
  • Controlled by the same instructions
• Block
  • Groups of threads
  • Threads in a block have different thread IDs
    • threadIdx
  • Can be 1D, 2D, or 3D
• Grid
  • Groups of blocks
  • Blocks in a grid have different block IDs
    • blockIdx
  • Can be 1D, 2D, or 3D

여기서 주요 개념 용어들을 살펴보자.

grid

block

• gridDim
  • 현재 그리드의 차원
  • Grid 안의 block 수 결정
• blockIdx
  • 현재 스레드의 블록 ID
• blockDim
  • 현재 블록의 차원
  • Block 안의 thread 수 결정
• threadIdx
  • Block 안의 현재 스레드의 ID

이렇게 계층이 정해져 있다는 것은 각 공간마다 허용되는 사이즈가 정해져 있다는 소리이다.
즉, 최대로 사이즈를 얼마까지 잡을 수 있는지 미리 정해져 있다.
이것은 CUDA C Programming Guide에 명시되어 있다.

CUDA C Programming Guide

여기서 우리가 앞서 했던 VectorSum 예제의 문제점을 살펴볼 수 있다.
블록마다 최대 threads의 개수는 1024개로 제한되어 있다.
그런데 우리는 Vector 배열의 크기만큼 threads 수를 설정해 주기 때문에
1024개가 넘어가면 CUDA에서 블록 내에 허용하는 최대 thread 개수를 넘어가기 때문에 오류가 발생하는 것이다.

그렇다면 최대 사이즈를 넘지 않게 Block을 여러 개 사용하거나,
Grid까지 활용하는 방식을 알아보자.

먼저 함수를 통해 thread layout을 지정해줄 수 있다.(그리드와 블록의 차원)

dim3 data type은 x, y, z 필드 값을 통해 접근 가능하다.

dim3 dimGrid(4, 1, 1);
dim3 dimBlock(8, 1, 1);
vecAdd <<<dimGrid, dimBlock>>>(d_a, d_b, d_c);