CUDA 스레드 레이아웃 설정 및 커널 호출<<<>>> - 공부하는 도비

CUDA 프로그래밍에서 스레드 레이아웃 설정과 커널 호출은 GPU 병렬 처리를 효과적으로 활용하기 위한 핵심 개념입니다. 

1. 스레드 레이아웃 설정 

GPU 병렬 처리를 위해 CUDA는 스레드(Thread)를 그리드(Grid)와 블록(Block)으로 구성합니다. 

각각의 스레드는 고유한 ID를 가져 작업을 병렬로 수행합니다. 

 

(1) 스레드 구성 요소

- 그리드(Grid) : 여러 블록으로 구성된 구조

- 블록(Block) : 여러 스레드로 구성된 구조

- 스레드(Thread) : 실제로 작업을 수행하는 단위

 

스레드들이 자신이 처리할 데이터가 무엇인지 알기 위해서는 자신이 어떤 블록에 속해 있는지, 또 블록 내 자신의 스레드 번호를 알아야 합니다. 이를 위해서 CUDA는 현재 그리드 및 블록의 형태와 각 스레드가 자신이 속한 블록 번호, 그리고 자기 자신의 스레드 번호를 확인할 수 있는 "내장 변수(built-in variable)"를 제공합니다. 

내장 변수의 값은 커널이 실행될 때 결정되며, 각 스레드는 자신에게 할당된 내장 변수 값을 참조할 수 있습니다. 

 

(2) 스레드 식별을 위한 내장 변수 

- gridDim : 그리드 내의 블록 수(그리드의 크기, 형태)

• gridDim.x, gridDim.y, gridDim.z : 각 차원의 길이는 1이상의 양의 정수로 정의되어야 합니다. (실제로 2차원 그리드지만, 사용하지 않는 차원(ex. z)의 크기는 1로 표현함)
• x 차원의 최대 길이 : 2^31-1(2,147,483,647)
• y 차원과 z 차원의 최대 길이 : 65,535

- blockDim : 블록 내의 스레드 수(블록의 크기, 형태)

• blockDim.x, blockDim.y, blockDim.z : 각 차원의 길이는 1 이상의 양의 정수로 정의되어야 합니다. (실제로 2차원 블록이지만, 사용하지 않는 차원(ex. z)의 크기는 1로 표현함)
• x 차원과 y차원의 최대 길이 : 1,024
• z 차원의 최대 길이 : 64
• 블록 하나가 가질 수 있는 최대 스레드의 수는 1,024개 (블록이 가지는 스레드 수는 x 차원, y 차원, z 차원의 크기를 모두 곱한 것과 같음)

- blockIdx : 현재 그리드 내의 블록 ID (1D, 2D, 3D 가능)

• blockIdx.x, blockIdx.y, blockIdx.z 

- threadIdx : 현재 블록 내의 스레드 ID (1D, 2D, 3D 가능)

• threadIdx.x, threadIdx.y, threadIdx.z 

 

2. 커널 호출 

CUDA에서 GPU에서 실행될 코드를 "커널(kernel)"이라고 부릅니다. 커널은 일반적으로 __global__ 키워드를 사용해 정의합니다. 

 

(1) 커널 함수 정의

__global__ void addVectors(int *a, int *b, int *c, int n) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

 

(2) 커널 호출

커널 호출은 <<<그리드의 형태, 블록의 형태>>> 형태로 정의 합니다.

기본 문법은 다음과 같습니다.

제일 처음 정의한 커널 함수 이름을 적어주고, 그리드와 블록의 형태, 그리고 넘겨줄 매개변수를 작성하면 됩니다. 

addVectors <<<그리드 형태, 블록의 형태>>>(매개변수);

 

제일 처음 GPU를 공부할 때, 별 생각 없이

<<<1, n>>> 과 같은 형태로 커널에서 사용할 스레드의 수를 n개로 지정을 했었는데,

정확한 의미는 (1, 1, 1) 형태의 그리드를 사용하며, 블록의 형태는 (n, 1, 1)로 설정하라는 것입니다. 

<<<1, n>>의 의미>

 

위에서 언급한 내장변수로 각 스레드의 고유한 Id를 정의할 수 있습니다. 

그리드와 블록이 1차원이 아닌 경우 (x, y, z)와 같이 각 차원의 형태(크기)를 <<<>>>에 인자로 전달해 줄 수도 있습니다. 

예를 들어, 그리드가 (2, 3, 4), 블록의 형태가 (5, 6, 7)이라면 실행 구성을 <<<(2, 3, 4), (5, 6, 7)>>> 처럼 설정합니다. 

 

(3) dim3를 사용한 스레드 레이아웃 설정 

스레드 레이아웃은 CUDA가 지원하는 구조체 변수인 dim3을 사용해서도 지정할 수가 있는데,

dim3는 x, y, z 멤버 변수를 가지는 구조체로 각각 x, y, z 차원의 크기를 담는 역할을 합니다. 

보통 <<<>>>안에 스레드 레이아웃을 바로 적기 보다는 dim3 구조체를 사용해서 전달을 합니다. 

dim3 dimGrid(3, 1, 1);
dim3 dimBlock(5, 1, 1);

addVectors <<<dimGrid, dimBlock>>>();

 

위 코드로 생성된 스레드 레이아웃은 다음 그림처럼 표현될 수 있습니다.