지난 글에서는 Host-Device 메모리 전송, cudaMalloc, cudaMemcpy 등을 활용하여 벡터 덧셈 예제를 성공적으로 구동해봤습니다. 이제 GPU 병렬 프로그래밍의 핵심인 스레드, 블록, 그리고 그리드 개념을 파고들 차례입니다. 이 개념들은 GPU 상에서 대규모 병렬 연산을 구성하는 뼈대와 같습니다. 오늘은 스레드와 블록, 그리고 이들을 묶어 전체 문제를 분할하는 방식인 그리드(Grid)에 대해 차근차근 살펴보며, "내가 원하는 만큼 병렬화"하는 방법을 이해해보겠습니다.
스레드(Thread)란?
스레드는 GPU 병렬 연산의 최소 실행 단위입니다. CPU에서도 멀티스레딩을 통해 병렬 처리를 하지만, GPU에서는 수백만 개의 스레드를 가볍게 실행할 수 있습니다. 각각의 스레드는 독립적으로 실행되는 작은 프로그램 흐름이며, 각 스레드는 벡터의 원소 하나를 처리한다거나, 이미지 픽셀 하나를 처리하는 식으로 문제를 잘게 쪼개는 데 사용합니다.
- 예: 10만 개의 데이터를 처리하려면, 10만 개의 스레드를 만들어 각 스레드가 하나씩 처리하도록 할 수 있습니다.
블록(Block)이란?
블록은 여러 스레드를 담고 있는 그룹입니다. 스레드들을 단순히 1차원적으로 늘어놓기보다는 일정한 크기로 묶어 관리하면, GPU 하드웨어가 이를 효율적으로 스케줄링할 수 있습니다.
- 각 블록은 threadIdx.x를 통해 자신의 내부에서 스레드 인덱스를 관리합니다.
- 블록은 1차원, 2차원, 3차원 형태로 스레드를 배치할 수 있습니다. 예를 들어 이미지 처리에서는 2D 그리드를 자주 이용합니다.
- 블록 크기는 하드웨어 제약(예: 한 블록당 최대 1024 스레드) 내에서 유연하게 설정할 수 있습니다.
블록 단위로 스레드를 묶는 이유는 GPU 내부에 블록 단위로 스케줄링되는 "Streaming Multiprocessor(SM)" 단위가 있기 때문입니다. 이를 통해 각 블록이 하드웨어적으로 독립적으로 실행될 수 있어 확장성과 병렬성을 극대화합니다.
그리드(Grid)란?
그리드는 여러 블록을 모아놓은 상위 레벨의 그룹입니다. 즉, 블록들의 모음이 그리드를 이룹니다. 우리가 커널을 호출할 때 <<<블록 수, 블록당 스레드 수>>> 문법을 사용했었는데, 이 때 지정한 "블록 수"가 바로 그리드를 구성하는 블록 개수입니다.
- blockIdx.x를 통해 어떤 블록인지를 식별하고, threadIdx.x를 통해 블록 내 어떤 스레드인지를 식별합니다.
- 전체 스레드 인덱스를 구할 때는 global_index = blockIdx.x * blockDim.x + threadIdx.x 같은 식을 사용합니다.
- 이를 통해 수많은 데이터를 간단한 수식 하나로 각 스레드에 "할당"할 수 있습니다.
정리하면, 그리드 > 블록 > 스레드의 계층 구조를 통해, 문제의 크기에 따라 병렬화를 유연하게 조정할 수 있습니다.
예제: 스레드와 블록 인덱스를 이용해 데이터 인덱싱하기
우리가 앞서 봤던 벡터 덧셈 커널을 다시 살펴봅시다. 커널에서 idx = blockIdx.x * blockDim.x + threadIdx.x; 라는 식을 사용했는데, 이것이 바로 그리드와 블록, 스레드 개념을 접목한 예입니다.
__global__ void vector_add_kernel(const float* a, const float* b, float* c, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
c[idx] = a[idx] + b[idx];
}
}
- blockDim.x는 한 블록 내 스레드의 개수를 의미합니다.
- blockIdx.x는 현재 실행 중인 블록의 인덱스(0부터 시작).
- threadIdx.x는 블록 내 스레드의 인덱스(0부터 시작).
이 식을 통해 전체 스레드 공간을 1차원으로 편평화(Flatten)하여, 각 스레드가 데이터의 특정 부분을 담당하게 합니다.
스레드와 블록 개수 정하기
"그럼 스레드와 블록 개수를 어떻게 정해야 할까?" 하는 궁금증이 생길 수 있습니다. 정답은 문제의 크기와 하드웨어 제약에 따라 유동적이지만, 일반적으로 다음과 같은 과정을 거칩니다.
- 데이터 개수(N) 파악: 처리해야 할 원소 수(N)를 기반으로, 스레드 수를 N 이상 확보해야 모든 원소를 처리할 수 있습니다.
- 블록 크기(threadsPerBlock) 설정: 보통 128, 256, 512와 같은 2의 거듭제곱을 사용합니다. 이는 하드웨어 특성과 align을 고려한 관례적인 선택입니다.
- 블록 수(blocks) 계산: 블록 하나에 threadsPerBlock개의 스레드를 담을 수 있으므로, blocks = (N + threadsPerBlock - 1) / threadsPerBlock로 계산해 N개 이상의 스레드를 확보합니다.
- 2D나 3D 블록/그리드: 이미지나 행렬 같은 2D 데이터에서는 (x,y) 형태의 블록·그리드를 활용할 수 있습니다. 예를 들어 1024x1024 이미지 처리 시 가로세로로 블록을 나누어 처리합니다.
예를 들어, N=100000인 벡터 처리 시, threadsPerBlock = 256이라면, blocks = (100000 + 256 - 1) / 256 = 391 (정수 나눗셈 기준) 정도로 설정할 수 있습니다.
Before vs After: 추상적 개념 → 실용적 이해
이전에는 단순히 “커널을 호출하고, 스레드 인덱스를 구한다” 정도로만 이해했다면 이제는 구조를 명확히 알 수 있습니다.
Before:
- 커널 호출 시 <<<...>>> 문법이 뭔가 이상한 문법이라 생각하고, 스레드 인덱스 계산을 어색해함.
After:
- 스레드(Thread), 블록(Block), 그리드(Grid) 계층 구조를 명확히 이해.
- 데이터 처리 규모에 따라 병렬화 수준을 손쉽게 조절 가능.
- 향후 이미지 처리, 대규모 행렬 연산, 복잡한 3D 연산 등 다양한 문제에 적용 가능.
다음 글 예고
다음 글(#5)에서는 지금까지 배운 것을 종합해 이미지나 2D 데이터 예시로 스레드/블록을 확장하여 생각해보는 과정을 살펴봅니다. 또한 비동기 스트림(Asynchronous Stream)과 간단한 최적화 개념 맛보기를 통해 GPU 프로그램의 잠재력을 더 깊이 파헤칠 준비를 할 겁니다.
유용한 링크 & 리소스
- NVIDIA CUDA Programming Guide: 스레드, 블록, 그리드 구조 공식 문서
- NVIDIA CUDA Samples: 다양한 블록/그리드 설정 예제
- Parallel for-loops with GPUs (NVIDIA 블로그)
- CUDA Occupancy Calculator : 블록 크기 및 하드웨어 리소스 활용도 계산