[기술] GPU와 CUDA (5) - 스레드 레이아웃

스레드 레이아웃

스레드 레이아웃 결정

앞서 CUDA 커널의 레이아웃은 그리드와 블록의 형태로 결정한다고 하였다.
구체적으로 다음의 과정을 따른다.

블록 형태 결정 (즉, 스레드를 어떻게 배치할껀지 결정)
데이터의 크기 및 블록 형태에 따라 그리드 형태 결정

블록 형태는 커널의 알고리즘 특성과 GPU 환경을 고려하여야 한다. 이때 레지스터, 공유 메모리 크기 등도 고려해야 할 요소이다.

큰 벡터의 합을 연산하는 CUDA 커널 (2)

벡터 차원이 1,024보다 크면 블록을 여러 개 지정해야 한다. 하나의 블록이었다면, 각 스레드가 벡터의 첫 번째 원소부터 담당하여 연산한다. 블록이 여러 개라면 각 블록이 벡터의 서로 다른 영역을 연산해야 한다. 각 블록마다 같은 인덱스의 스레드가 존재하므로, 블록 인덱스 고려없이 스레드가 백터 원소를 담당하면, 모든 블록의 동일 인덱스 스레드는 벡터의 같은 원소에 접근한다. 따라서 모든 블록이 벡터의 같은 영역을 처리한다. 원하는 구현은 다음과 같다.

0번 블록이 0 ~ 1023번 데이터를 처리
1번 블록이 1024 ~ 2047번 데이터를 처리
…

// 예를 들어서 스레드 수가 1024이고, 3번쨰 블록의 512번쨰 원소를 연산한다면

tId = blockIdx.x * blockDim.x + threadIdx
element = vector[tId]

이 원리로 앞서 CUDA 커널의 벡터 합 연산을 다시 작성하면

// 단일 블록의 벡터 합 연산
__global__ void vecAdd(int* _a, int* _b, int* _c) {
    int tID = threadIdx.x;
    _c[tID] = _a[tID] + _b[tID];
}


// 여러 블록을 고려하여, 크기가 큰 벡터 연산
__global__ void vecAdd(int* _a, int* _b, int* _c) {
    int tID = blockIdx.x * blockDim.x + threadIdx.x;
    _c[tID] = _a[tID] + _b[tID];
}

위와 같이 수정하면, 스레드가 담당하는 원소의 합을 연산한다. 스레드 레이아웃은 데이터 크기를 한 블록의 크기 (= 스레드 개수)로 나눈 몫으로 블록을 잡아준다. 그런데 예를 들어, 벡터 데이터가 1025개이고, 스레드 수가 512인 경우 블록을 3개로 잡아야 한다. 이렇게 되면 마지막 블록은 스레드 하나만 벡터 연산에 활용하고, 나머지 스레드는 올바른 벡터 인덱스에 접근하지 못한다.

따라서 이런 경우를 대비하여, 스레드가 데이터 인덱스를 벗어나는 예외 케이스를 꼭 고려해야 한다.
이를 반영하면

__global__ void vecAdd(int numThreads, int ...) {
    int tID = blockIdx.x * blockDim.x + threadIdx.x;

    ...

    // 블록 - 스레드 인덱스가 전체 스레드 수를 넘어가면 예외 처리
    if (tID > numThreads) {
        return;
    }

    ...
}

이제까지의 내용을 종합하여, 단일 블록 1024 스레드보다 큰 벡터의 합을 연산하는 커널은

__global__ void vecAdd(int* _a, int* _b, int* _c, int _size) {
    int tID = blockIdx.x * blockDim.x + threadIdx.x;

    ...

    if (tID < _size) {
        _c[tID] = _a[tID] + _b[tID];
    }
}

...

int main() {
    ...

    dim3 dimGrid(ceil(float)NUM_DATA / 256, 1, 1);
    dim3 dimBlock(256, 1, 1);
    vecADd <<<dimGrid, dimBlock>>> (d_a, d_b, d_c, NUM_DATA);

    ...
}

결론

그리드 형태는 블록을, 블록 형태는 스레드를 결정한다.
데이터 크기에 따라 블록 당 최대 스레드 수와 블록 형태를 고려해야 한다.
스레드가 메모리 주소를 잘못 참조하지 않도록, 데이터 형태에 따른 예외케이스를 고려해야 한다.
CUDA 커널 실행 시간은 데이터 전송 시간까지 고려해야 한다.

스레드 인덱싱 1

메모리와 배열

스레드 인덱싱은 CUDA 내장 변수를 사용해, 특정 데이터에 정확히 접근하도록 한다. 1차원 벡터와 논리적 메모리의 매핑은 어렵지 않다. 그런데 2차원 행렬과 논리적 메모리와의 매핑은 다소 다른다. 2차원 행렬에서 1행, 2행, 3행 … 데이터는 논리적 메모리에서 옆으로 계속 나열된다. 따라서 이를 고려하여 메모리 접근을 해야한다. (데이터에 따른 메모리 접근 패턴)

스레드 인덱싱 2

1차원 블록

가장 간단한 경우는 커널 내 스레드가 배열 데이터와 1:1 매핑되는 경우이다.
1차원 블록은 스레드가 x-차원 인덱스만 갖기에, threadIdx.x가 스레드의 글로벌 인덱스와 동일하다.

2차원 블록

블록의 가로는 x-차원이고, 세로는 y-차원이다.
이때 블록 하나는 (threadIdx.x, threadIdx.y) 스레드로 이루어져 있다.
따라서 어떤 행의 하위 블록 스레드에 접근하려면 두 가지를 알아야 한다.

x-차원으로 나열된 블록의 차원 (=blockDim.x)
각 블록의 스레드의 x-차원: threadIdx.x
각 블록의 스레드의 y-차원: threadIdx.y

// 1. 자신이 속한 블록 앞까지의 스레드 수
// = 블록의 x-차원 * 이 x-차원이 세로로 나열된 y-차원 갯수
// 2. 자신이 속한 블록에서의 스레드 인덱스
// = 현재 블록에서 가로로 몇 인덱스에서 자기 자신이 등장하는가?
2D_BLOCK_TID = blockDim.x * threadIdx.y + threadIdx.x

3차원 블록

// 이를 3차원 블록으로 확장하면 다음과 같다.
// 앞에서부터 x-차원, y-차원, z-차원의 스레드 수를 계산
// 현재 2차원 블록에서의 스레드를 계산 (이전과 같음)
TID_IN_BLOCK = (blockDim.x * blockDim.y * threadIdx.z) + 2D_BLOCK_TID

그리드 내 스레드의 전역 번호

그리드 내 블록이 1개라면 TID_IN_BLOCK이 그리드 내의 각 스레드 전역 인덱스이다.
그리드 내 블록이 여러 개라면, 블록 인덱스를 고려하여 쌓아가면 된다.

1차원 그리드

// 블록 하나에 속한 스레드 갯수
NUM_THREAD_IN_BLOCK = blockDim.z * blockDim.y * blockDim.x

// 자신이 속한 블록의 번호
blockIdx.x

// 1차원 그리드에서 스레드 전역 번호
1D_GRID_TID = (NUM_THREAD_IN_BLOCK * blockIdx.x) + TID_IN_BLOCK

2차원 그리드

// 2차원 그리드는 1차원 그리드가 y-차원 방향으로 나열된 경우이다. 
2D_GRID_TID = (blockIdx.y * (gridDim.x * NUM_THREAD_IN_BLOCK)) + 1D_GRID_TID

3차원 그리드

// 2차원 그리드에서 계산한 스레드 인덱스를 그대로 사용할 수 있다.
GLOBAL_TID = (blockIdx.z * (gridDim.y + gridDim.x * NUM_THREAD_IN_BLOCK)) + 2D_GRID_TID

인덱스 상수의 헤더 파일

// 블록 인덱스
#define BID_X blockIdx.x
#define BID_Y blockIdx.y
#define BID_Z blockIdx.z

// 스레드 인덱스
#define TID_X threadIdx.x
#define TID_Y threadIdx.y
#define TID_Z threadIdx.z

// 그리드 차원
#define Gdim_X gridDim.x
#define Gdim_Y gridDim.y
#define Gdim_Z gridDim.z

// 블록 차원
#define Bdim_X blockDim.x
#define Bdim_Y blockDim.y
#define Bdim_Z blockDIm.z

#define TID_IN_BLOCK (TID_Z * (Bdim_Y * Bdim_X) + TID_Y * Bdim_X + TID_X)
#define NUM_THREAD_IN_BLOCK (Bdim_X * Bdim_Y * Bdim_Z)

#define GRID_1D_TID (BID_X * NUM_THREAD_IN_BLOCK) + TID_IN_BLOCK
#define GRID_2D_TID (BID_Y * (Gdim_X * NUM_THREAD_IN_BLOCK) + GRID_1D_TID)
#define GLOBAL_TID (BID_Z * (Gdim_Y * Gdim_X * NUM_THREAD_IN_BLOCK) + GRID_2D_TID)

결론

그리드 형태는 블록을, 블록 형태는 스레드를 결정한다.
데이터 크기에 따라 블록 당 최대 스레드 수와 블록 형태를 고려해야 한다.
스레드가 메모리 주소를 잘못 참조하지 않도록, 데이터 형태에 따른 예외케이스를 고려해야 한다.
CUDA 커널 실행 시간은 데이터 전송 시간까지 고려해야 한다.

스레드 인덱싱 3: 2차원 데이터

몇 가지 특징

행렬을 다룰 떄 사용하는 대표적인 인덱싱 방법은 2차원 스레드 인덱스를 사용하여, 각 스레드가 행렬 원소를 가리키게 하는 것이다. 따라서 2차원 형태의 스레드 레이아웃이 가장 직관적이다.
2차원 스레드 블록을 사용하면, 각 스레드는 (x, y)의 2차원 인덱스를 가진다. 스레드 번호를 매핑할 때에는 x-차원 번호와 y-차원 번호를 각각 행과 열 중 어느 것에 대응할지 결정해야 한다. (둘 다 가능하다.)

1차원 그리드, 2차원 블록 레이아웃

row = threadIdx.x
col = threadIdx.y

// 행의 길이 = blockDim.x
indx = blockDim.x * row + col

2차원 그리드, 2차원 블록 레이아웃

그리드가 2차원이면, 블록은 (x, y) 인덱스를 가진다.
블록이 2차원이면, 스레드는 (x, y) 인덱스를 가진다.

// 블록 인덱싱 복습

// 1차원 블록
// 블록 내에서 몇 row를 지나고 현재 row에서 몇 col을 지나야 있는지 확인
row = threadIdx.y
col = threadIdx.x
index = blockDim.x * row + col

// 2차원 블록
col = blockDim.x * blockIdx.x + threadIdx.x
row = blockDim.y * blockIdx.y + threadIdx.y
index = row * COL_SIZe + col

두 행렬의 합 (그리드 1개, 블록 1개)

__global__ void matAdd_2D_index (float* _dA, float* _dB, float* _dC) {
    unsigned int col = threadIdx.x;
    unsigned int row = threadIdx.y;
    unsigned int index = row * blockDim.x + col;

    _dC[index] = _dA[index] + _dB[index];
}

...
...

// 커널 호출
dim3 BlockDim(COL_SIZE, ROW_SIZE);
matAdd_2D_index <<<1, blockDim>>>(dA, dB, dC);

결론

스레드 레이아웃 및 인덱싱 방법은 절대적인 것이 없다.
데이터 및 알고리즘에 따라 목적에 맞게 설계해야 한다.

대규모 행렬의 합

그리드 2개 이상, 블록 2개 이상

__global__ void MatADd_G2D_B2D
(float* MatA, float* MatB, float* MatC, int ROW_SIZE, int COL_SIZE) {
    unsigned int col = blockDim.x * blockIdx.x + threadIdx.x;
    unsigned int row = blockDim.y * blockIdx.y + threadIdx.y;
    unsigned int index = row * COL_SIZe + col;

    MatC[index] = MatA[index] + MatB[index];
}

...
...

// 커널 호출
// 2차원 블록 크기
dim3 blockDim(32, 32);
// 2차원 그리드 크기
dim3 gridDim(ceil((float)COL_SIZE / blockDim.x), ceil((float)ROW_SIZE / blockDim.y))

MatAdd_G2D_B2D <<<gridDim, blockDim>>> (A, B, C, ROW_SIZE, COL_SIZE);

결론

그리드 크기에서 올림 처리를 하여, 처리되지 않는 데이터가 생기는 것을 방지한다.
올림 처리로 마지막 블록에서 행렬을 벗어난 위치의 메모리에 잘못 접근할 수 있다. 이를 예외처리 한다.

__global__ void MatADd_G2D_B2D
(float* MatA, float* MatB, float* MatC, int ROW_SIZE, int COL_SIZE) {
    unsigned int col = blockDim.x * blockIdx.x + threadIdx.x;
    unsigned int row = blockDim.y * blockIdx.y + threadIdx.y;
    unsigned int index = row * COL_SIZe + col;

    if (col < COL_SIZE && row < ROW_SIZE) {
        MatC[index] = MatA[index] + MatB[index];
    }
    ...
}

...
...

// 커널 호출
// 2차원 블록 크기
dim3 blockDim(32, 32);
// 2차원 그리드 크기
dim3 gridDim(ceil((float)COL_SIZE / blockDim.x), ceil((float)ROW_SIZE / blockDim.y))

MatAdd_G2D_B2D <<<gridDim, blockDim>>> (A, B, C, ROW_SIZE, COL_SIZE);

1차원 그리드, 1차원 블록

1차원 그리드와 1차원 블록을 사용한다는 것은 스레드 번호 중 하나의 차원만 사용한다는 의미이다. 가장 쉬운 방법은 스레드 수가 전체 행렬 크기와 같게 그리드와 블록을 정의한다. 그러나 행렬의 크기가 커지면, 필요한 스레드 수가 급격히 늘어난다.

스레드 x-차원의 전역 번호를 행렬의 열로 매핑하면, 스레드 번호로는 행을 구분할 수 없다. 이는 각 스레드가 담당하는 열의 모든 행을 처리하도록 하는 전략이다. (스레드를 옆으로 행렬의 열만큼 나열하고, 행렬을 위에서 입력하면 스레드 - 열간 1:1 매핑이 되어 행의 모든 성분을 처리한다.)

// 스레드가 열을 담당하여, 그 열의 모든 행 성분을 처리하는 코드
__global__ void MatAdd_G1D_B1D(float* MatA, float* MatB, float* MatC, int ROW_SIZE, int COL_SIZE) {
    // 블록의 열 수: 블록 차원 * 블록 인덱스 + 스레드 인덱스
    unsigned int col = blockDim.x * blockIdx.x + threadIdx.x;
    
    // 열이 COL_SIZE보다 작으면, 해당 열의 행을 모두 순회하여
    if (col < COL_SIZE) {  
        // 해당 열과 해당 행의 메모리 상 인덱스를 계산한 후, 해당 인덱스의 행렬 A, B 값을 C에 누적
        for (int row = 0; row < ROW_SIZE; row++) {
            int index   = row * COL_SIZE + col;
            MatC[index] = MatA[index] + MatB[index];
        }
    }
}

...
...

dim3 blockDim3(32);
dim3 gridDim(ceil((float)COL_SIZe / blockDim.x));
MatAdd_G1D_B1D <<<gridDim, blockDim>>> (A, B, C, ROW_SIZE, COL_SIZE);

결론

그리드 크기에서 올림 처리를 하여, 처리되지 않는 데이터가 생기는 것을 방지한다.
올림 처리로 마지막 블록에서 행렬을 벗어난 위치의 메모리에 잘못 접근할 수 있다. 이를 예외처리 한다.

2차원 그리드, 1차원 블록

이 레이아웃에서는 스레드는 x-차원 인덱스만 가지고, 블록은 x-차원, y-차원 인덱스를 가진다. 따라서 각 차원을 행렬의 행과 열로 매핑할 수 있다. 이때 x-차원 번호는 블록과 스레드가 모두 있다. 따라서 x-차원 전역 인덱스를 만들어서 사용할 수 있다. 행과 매핑되는 y-차원 인덱스는 그대로 사용한다.

// 각 스레드의 행렬 매핑
row = BlockDim.x * blockIdx.x + threadIdx.x
col = blockIdx.y

__global__ void MatAdd_G2D_B1D(float* MatA, float* MatB, float* MatC, int ROW_SIZE, int COL_SIZE) {
    unsigned int col = blockDim.x * blockIdx.x + threadIdx.x;
    unsigned int row = blockIdx.y;
    unsgiend int index = row * COL_SIZE + col;

    if (col < COL_SIZE && row < ROW_SIZE) {
        MatC[index] = MatA[index] + MatB[index];
    }
}

... 
...
// 블록은 1차원이므로 32 설정
dim3 blockDim(32);
// 그리드는 2차원이므로 ceil(전체 행렬 사이즈/블록 차원) = 그리드의 x-차원 수
// y-차원인 행은 행렬의 행 사이즈 그대로 입력
dim3 gridDim(ceil((float)COL_SIZE / blockDim.x), ROW_SIZE);

MatAdd_G2D_B1D <<<gridDim, blockDim>>>(A, B, C, ROW_SIZE, COL_SIZE);

결론

동일한 알고리즘을 다양한 스레드 레이아웃으로 구현 가능하다.
알고리즘에 가장 적합한 스레드 레이아웃을 디자인 해야한다.
스레드 레이아웃따라 알고리즘 속도가 다르다. -> GPU가 CUDA를 처리하는 방법과 관련있다.

스레드 레이아웃#

스레드 레이아웃 결정#

큰 벡터의 합을 연산하는 CUDA 커널 (2)#

결론#

스레드 인덱싱 1#

메모리와 배열#

스레드 인덱싱 2#

1차원 블록#

2차원 블록#

3차원 블록#

그리드 내 스레드의 전역 번호#

1차원 그리드#

2차원 그리드#

3차원 그리드#

인덱스 상수의 헤더 파일#

결론#

스레드 인덱싱 3: 2차원 데이터#

몇 가지 특징#

1차원 그리드, 2차원 블록 레이아웃#

2차원 그리드, 2차원 블록 레이아웃#

두 행렬의 합 (그리드 1개, 블록 1개)#

결론#

대규모 행렬의 합#

그리드 2개 이상, 블록 2개 이상#

결론#

1차원 그리드, 1차원 블록#

결론#

2차원 그리드, 1차원 블록#

결론#