junyeokk
Blog
Parallel Computing·2025. 10. 20

커널 실행 설정 - Grid와 Block

왜 스레드를 계층적으로 구성하는가

GPU는 수천 개의 코어를 가지고 있다. 이 코어들에서 수백만 개의 스레드가 동시에 실행된다. 문제는 이 많은 스레드를 어떻게 효율적으로 관리할 것인가다.

CUDA는 스레드를 계층적으로 구성한다. 스레드들을 블록으로 묶고, 블록들을 다시 그리드로 묶는다.

plain
grid (커널 1개당 1개)
├─ block 0
│  ├─ thread 0
│  ├─ thread 1
│  └─ thread 2
├─ block 1
└─ block 2

이렇게 설계한 이유가 있다.

첫째, 하드웨어 매핑이 자연스럽다. GPU는 물리적으로 여러 개의 Streaming Multiprocessor(SM)로 구성된다. 각 블록은 하나의 SM에 할당되어 실행된다. 블록 단위로 스케줄링하면 하드웨어와 소프트웨어 구조가 일치한다.

둘째, 동기화 범위를 제한한다. 수백만 개의 스레드를 전역 동기화하면 성능이 급감한다. CUDA는 블록 내부에서만 동기화를 허용한다(__syncthreads()). 블록 크기를 적절히 제한하면 동기화 비용을 낮출 수 있다.

셋째, 메모리 공유가 가능하다. 같은 블록의 스레드들은 shared memory를 공유한다. 블록 단위로 데이터를 캐싱하면 global memory 접근을 줄여 성능을 높일 수 있다.

커널 실행 문법

커널을 호출할 때 <<<grid, block>>> 문법으로 스레드 구조를 지정한다.

c
__global__ void kernel_func(...) {
    // GPU에서 실행되는 코드
}

// 호출
kernel_func<<<dimGrid, dimBlock>>>(...);
kernel_func<<<8, 16>>>(...);  // 8개 블록, 블록당 16스레드

총 스레드 개수는 grid 크기 × block 크기다. <<<8, 16>>>은 128개 스레드를 생성한다.

스레드가 자신의 위치를 아는 방법

각 스레드는 동일한 코드를 실행한다. 그런데 스레드마다 다른 데이터를 처리해야 한다. 어떻게 구분할까?

CUDA는 스레드마다 고유한 ID를 부여한다.

c
blockIdx.x   // 이 스레드가 속한 블록 번호
threadIdx.x  // 블록 내에서 이 스레드의 번호
blockDim.x   // 블록당 스레드 개수
gridDim.x    // 총 블록 개수

이 변수들을 조합하면 전역 인덱스를 계산할 수 있다.

c
int globalIdx = blockIdx.x * blockDim.x + threadIdx.x;

예를 들어 <<<6, 4>>>로 호출하면 24개 스레드가 생성된다.

plain
block 0:  [0][1][2][3]
block 1:  [4][5][6][7]
block 2:  [8][9][10][11]
block 3:  [12][13][14][15]
block 4:  [16][17][18][19]
block 5:  [20][21][22][23]

blockIdx.x=2, threadIdx.x=3인 스레드의 globalIdx는 2 * 4 + 3 = 11이다.

1D 인덱싱 커널

1차원 배열을 처리하는 가장 기본적인 패턴이다.

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

if (idx < n) 조건이 중요하다. 배열 크기가 블록 크기의 배수가 아닐 수 있다. 예를 들어 1000개 원소를 처리하는데 블록 크기가 256이면 4개 블록(1024스레드)이 필요하다. 마지막 24개 스레드는 범위를 벗어나므로 조건문으로 걸러야 한다.

왜 2D, 3D 구조가 필요한가

데이터가 항상 1차원 배열은 아니다. 이미지는 2D, 볼륨 데이터는 3D다. 데이터 구조와 스레드 구조를 맞추면 인덱싱이 직관적이다.

c
// 1D로 2D 접근 (복잡함)
int row = idx / width;
int col = idx % width;

// 2D 스레드 구조 사용 (직관적)
int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;

2D 이미지 처리에서 (x, y) 좌표로 바로 접근할 수 있다.

2D Layout 예제

c
dim3 gridDim(3, 5);   // 3×5 = 15개 블록
dim3 blockDim(4, 3);  // 블록당 4×3 = 12개 스레드
// 총 180개 스레드

전역 좌표 계산:

c
int gx = blockIdx.x * blockDim.x + threadIdx.x;
int gy = blockIdx.y * blockDim.y + threadIdx.y;

2D 좌표를 1D 배열 인덱스로 변환하려면:

c
int idx = gy * width + gx;  // row-major order

2D Indexing 커널

행렬의 각 원소에 1을 더하는 커널이다.

c
__global__ void addOne2D(int* a, int width, int height) {
    int gx = blockIdx.x * blockDim.x + threadIdx.x;
    int gy = blockIdx.y * blockDim.y + threadIdx.y;
    
    if (gx < width && gy < height) {
        int idx = gy * width + gx;
        a[idx] = a[idx] + 1;
    }
}

x, y 모두 범위 체크가 필요하다. 행렬 크기가 블록 크기의 배수가 아닐 수 있기 때문이다.

3D Layout

3D 구조는 볼륨 데이터(CT 스캔, MRI, 3D 시뮬레이션)에서 사용한다.

c
dim3 gridDim(3, 5, 2);   // 30개 블록
dim3 blockDim(4, 3, 2);  // 블록당 24개 스레드
// 총 720개 스레드

전역 좌표:

c
int gx = blockIdx.x * blockDim.x + threadIdx.x;
int gy = blockIdx.y * blockDim.y + threadIdx.y;
int gz = blockIdx.z * blockDim.z + threadIdx.z;

dim3 타입

CUDA는 3차원 좌표를 위한 전용 타입을 제공한다.

c
dim3 gridDim(6);        // (6, 1, 1) - 1D
dim3 gridDim(3, 5);     // (3, 5, 1) - 2D
dim3 gridDim(3, 5, 2);  // (3, 5, 2) - 3D

생략된 차원은 자동으로 1이 된다. 그래서 단순한 정수도 dim3로 암시적 변환된다.

c
kernel<<<3, 4>>>();           // 동일
kernel<<<dim3(3), dim3(4)>>>(); // 동일

블록 크기 선택 가이드

블록 크기를 어떻게 정해야 할까?

하드웨어 제약이 있다. 블록당 최대 스레드는 1024개다. blockDim.x * blockDim.y * blockDim.z <= 1024를 만족해야 한다.

성능을 위한 권장사항도 있다. 블록 크기는 32의 배수가 좋다. GPU는 32개 스레드를 하나의 warp로 묶어 실행하기 때문이다. 128, 256, 512 같은 값이 일반적이다.

데이터 크기에 맞춰야 한다. 블록 크기가 너무 작으면 스레드 관리 오버헤드가 커지고, 너무 크면 리소스를 낭비한다. 보통 256개 정도로 시작해서 벤치마크로 최적값을 찾는다.

참고 자료