CUDA 프로그래밍 - CUDA 스레드 계층

 

CUDA 프로그래밍 - CUDA 스레드 계층

 

목차

• CUDA 스레드 계층 구조
• 스레드 계층 내장 변수
• 그리드, 블록의 최대 크기
• 스레드 구조, 커널 호출

 

 

안녕하세요. 지난 포스팅에서 CUDA 프로그램을 활용해 벡터의 합을 구하는 예제를 다뤘습니다. 호스트와 디바이스 메모리 간 데이터 복사, 각 원소의 합을 구하는 커널 함수를 작성했습니다. 또한, 연산 시간을 측정하기 위한 간단한 코드도 작성해 봤습니다. 해당 내용이 궁금하시면 지난 글을 참고해 보시길 추천드립니다.

 

 

CUDA 프로그래밍 - 벡터의 합 구하기

 

 

오늘은 이어서 CUDA의 스레드 계층에 대해 알아보겠습니다. CUDA에서 병렬 연산을 하기 위해 어떤 구조로 스레드를 구성하는지, 프로그램 내에서 정보를 어떻게 확인하는지 알아보겠습니다. 스레드 계층은 CUDA의 대규모 병렬 연산을 수행하기 위해 반드시 알아야 하는 부분입니다. 오늘 포스팅에서 소개해 드리는 스레드 계층에 대해 이해한다면 앞으로 CUDA 프로그램을 작성하는 데 도움이 될 것입니다.

 

 

CUDA 스레드 계층 구조

GPU는 SIMT(Single Instruction Multiple Thread) 구조로서 하나의 명령어가 하나의 스레드 그룹을 제어합니다. GPU는 수백에서 수천 개의 연산 코어를 가지고 있고, 이러한 구조적 특성을 가진 GPU의 대규모 병렬 처리 능력을 효율적으로 활용하기 위해 수천에서 수십만 개의 스레드를 사용해야 합니다. 수천 개 이상의 스레드를 효율적으로 사용하기 위해 CUDA는 스레드를 계층적으로 관리합니다.

 

CUDA의 스레드 계층(CUDA thread hierarchy)은 스레드, 워프, 블록, 그리드 네 개의 계층으로 구성되어 있습니다. 스레드 계층 구조의 각 요소에 대해 하나씩 살펴보겠습니다.

 

CUDA 스레드 계층 구조

 

 

스레드(Thread)

CUDA 스레드 계층 구조에서 가장 작은 단위는 스레드입니다. 스레드는 CUDA 스레드 계층 중 가장 낮은 계층으로 CUDA에서 연산을 수행하거나 CUDA 코어를 사용하는 기본 단위입니다. 커널 코드는 모든 스레드에 공유되며 각 스레드가 독립적으로 커널 코드를 수행합니다.

 

워프(Warp)

워프는 32개의 스레드를 하나로 묶은 것을 말하고, CUDA의 기본 수행 단위입니다. 기본 수행 단위라는 것은 한 워프에 속한 스레드들은 하나의 제어 장치에 의해 제어된다는 것을 의미합니다. GPU의 SIMT 구조에서 멀티 스레드의 단위가 되는 것이 워프입니다. 하나의 명령(instruction)에 따라 32개의 스레드가 동시에 동작하고, CUDA 프로그램에서 매우 중요한 개념입니다.

 

블록(Block)

블록은 워프보다 상위 스레드 그룹으로 워프들의 집합입니다. 블록에 포함된 각 스레드는 자신만의 고유한 스레드 번호(thread ID)를 가지며, 동일한 블록 안에는 동일한 번호를 갖는 스레드는 없습니다. 반면, 서로 다른 블록에 포함된 스레드들은 같은 스레드 번호를 가질 수 있습니다. 따라서, 원하는 스레드에 정확하게 연산 명령을 전달하기 위해 블록 번호, 스레드 번호를 모두 사용해야 합니다.

 

2차원 블록 내 스레드

 

 

블록 내 스레드는 1차원, 2차원 또는 3차원 형태로 배치될 수 있습니다. 스레드 번호 또한 배치에 따라 3차원 번호를 가질 수 있습니다.

 

그리드(Grid)

CUDA 스레드 계층 구조에서 가장 상위 단계는 그리드입니다. 그리드는 여러 개 블록들의 그룹입니다. 하나의 그리드에 포함된 블록들은 서로 다른 자신만의 고유한 블록 번호(block ID)를 가지고, 블록과 마찬가지로 그리드 내 블록 또한 1차원, 2차원, 또는 3차원 형태로 배치될 수 있습니다.

 

 

 

커널이 호출되면 그리드가 생성됩니다. 하나의 그리드는 하나의 커널과 1:1로 대응되고, 해당 커널을 수행할 스레드를 생성합니다.

 

 

스레드 계층 내장 변수

CUDA 스레드 계층 구조에서 그리드 내 블록이 1~3차원 형태로 배치되고, 블록 안의 스레드 또한 1~3차원 형태로 배치됩니다. 각 스레드가 자신이 처리할 데이터가 어떤 것인지 알기 위해서 자신이 어떤 블록에 속해 있는지, 블록 내 자신의 스레드 번호는 무엇인지 알아야 합니다.

 

CUDA는 그리드 및 블록의 형태, 각 스레드가 자신이 속한 블록의 번호, 자신의 스레드 번호를 확인할 수 있는 내장 변수(built-in variable)를 제공합니다. 내장 변수의 값은 커널이 실행될 때 결정되고, 각 스레드는 자신에게 할당된 내장 변수 값을 참조할 수 있습니다.

 

gridDim, blockIdx 내장 변수

 

gridDim

gridDim은 그리드의 형태 정보를 담고 있는 구조체형 내장 변수입니다. x, y, z 멤버가 각 차원의 크기를 담고 있습니다. 각 차원의 크기는 1 이상의 정수로 정의되어 있습니다. 위 그림은 x차원의 크기 3, y 차원의 크기 3, z 차원의 크기 1인 그리드입니다. 따라서, gridDim.x, gridDim.y, gridDim.z는 각각 3, 3, 1의 값을 가집니다.

 

blockIdx

blockIdx는 현재 스레드가 속한 블록의 번호를 담고 있는 구조체형 내장 변수입니다. 그리드 내 블록들은 서로 다른 번호를 가지며, 그리드의 형태에 따라 최대 3차원의 번호를 가질 수 있습니다. 블록 번호는 0부터 시작하는 인덱스 번호를 가지고 있습니다. 위 그림에서 왼쪽으로부터 3번째, 위로부터 3번째 위치한 블록 각 차원의 인덱스 blockIdx.x, blockIdx.y, blockIdx.z는 2, 2, 0입니다.

 

blockDim

blockDim은 블록의 형태 정보를 담고 있는 구조체형 내장 변수입니다. gridDim과 마찬가지로 x, y, z 멤버가 각 차원의 크기를 담고 있습니다. 각 차원의 크기는 1 이상의 정수로 정의되어 있습니다. 커널이 실행될 때, 그리드, 블록의 형태가 결정되고, 한 그리드 내의 모든 블록은 동일한 형태를 가집니다. 아래 그림은 x차원의 크기 4, y차원의 크기 3, z차원의 크기 1인 블록입니다. 따라서, blockDim.x, blockDim.y, blockDim.z는 각각 4, 3, 1의 값을 가집니다.

 

blockDim, threadIdx 내장 변수

 

threadIdx

threadIdx는 블록 내에서 현재 스레드가 부여받은 스레드 번호를 담고 있는 구조체형 내장 변수입니다. 한 블록 내 스레드들은 서로 다른 스레드 번호를 가지며 블록의 형태에 따라 최대 3차원의 번호를 가질 수 있습니다. 스레드 번호는 0부터 시작하는 인덱스 번호를 가지고 있습니다. 위 그림에서 왼쪽에서 4번째, 위로부터 3번째 위치한 스레드 각 차원의 인덱스 threadIdx.x, threadIdx.y, threadIdx.z는 3, 2, 0입니다.

 

스레드 번호와 워프의 구성

워프는 연속된 32개 스레드로 구성됩니다. 스레드의 연속성은 threadIdx의 x-차원, y-차원, z-차원 순으로 결정됩니다. 즉, (0, 0, 0) ~ (31, 0, 0) 번 스레드가 하나의 워프를 구성합니다. 만약 x-차원의 길이가 워프의 크기보다 작다면 y-차원 번호가 낮은 순으로 연속성을 가집니다. 예를 들어 x-차원의 크기가 1이고, y-차원의 길이가 32보다 크다면 (0, 0, 0) ~ (0, 31, 0) 번 스레드가 하나의 워프를 이루게 됩니다.

 

메모리 접근 패턴은 커널의 성능에 큰 영향을 미치는 요소 중 하나입니다. 메모리 접근 패턴을 제대로 활용하려면 워프를 구성하는 스레드를 정확하게 인지하는 것이 중요합니다. 또한, 워프 수준에서 스레드들 사이의 작업을 분배할 때도 중요한 역할을 합니다.

 

 

그리드, 블록의 최대 크기

아래의 그림은 엔비디아 CUDA 공식 가이드에 있는 compute capability 표의 일부입니다. 해당 표는 그리드, 블록의 최대 크기를 보여줍니다.

 

Technical Specifications per Compute Capability(CUDA C++ Programming Guide)

 

그리드의 크기

그리드는 최대 3차원 형태로 구성할 수 있고, x-차원의 최대 길이는 $2^{31}-1$입니다. 그리드 내 블록의 수가 $2^{31}-1$만큼 큰 경우는 거의 없기 때문에 x-차원의 크기에는 제한이 없다고 봐도 무방합니다. 반면 y-차원, z-차원의 경우, 최대 크기가 65,535로 제한됩니다.

 

블록의 크기

블록은 최대 3차원 형태로 구성할 수 있고, x-차원과 y-차원의 최대 크기는 1,024이고, z-차원의 최대 크기는 64입니다. 블록에 비해 최대 크기가 제한적이기 때문에 최대 크기를 기억해 두는 것이 좋습니다. 그리고 한 가지 중요한 점은 블록 하나가 가질 수 있는 최대 스레드의 수는 1,024개라는 것입니다. 블록의 최대 스레드 수를 간과하고 설계하게 되면 커널 실행이 되지 않기 때문에 명심해야 합니다.

 

 

스레드 구조, 커널 호출

커널을 수행할 그리드와 블록의 형태를 지정하는 방법에 대해 알아보겠습니다. 커널 함수를 호출할 때, <<<>>> 문법을 사용했습니다. 이 문법은 스레드 레이아웃(thread layout)을 설정하는 문법으로 그리드와 블록의 형태를 정의합니다. 아래와 같이 커널에서 사용하고자 하는 그리드, 블록의 형태를 전달합니다.

 

Kernel<<<그리드 형태, 블록 형태>>>();

 

 

앞선 예제들에서 <<<1, n>>>의 형태로 커널 함수를 호출했습니다. 해당 스레드 레이아웃의 의미는 (1, 1, 1) 크기의 그리드와 (n, 1, 1) 크기의 블록을 사용한다는 것입니다. 즉, x-차원의 길이가 n인 1차원 블록을 스레드 레이아웃으로 설정하는 것입니다.

 

그리드와 블록이 1차원이 아닌 경우, (x, y, z)와 같이 각 차원의 크기를 <<<>>>의 인자로 전달할 수 있습니다. 예를 들어, 그리드의 크기가 (3, 2, 1), 블록의 크기가 (6, 5, 4)인 경우, <<< (3, 2, 1), (6, 5, 4) >>>와 같이 스레드 레이아웃을 설정할 수 있습니다.

 

스레드 레이아웃은 CUDA가 지원하는 구조체 변수형인 dim3를 사용해서 지정할 수 있습니다. dim3는 x, y, z 멤버 변수를 가지는 구조체로 x-차원, y-차원, z-차원 크기를 설정합니다. CUDA 프로그램 작성 시에는 일반적으로 (3, 2, 1)과 같이 직접 숫자를 전달하기보다 dim3를 활용해 스레드 레이아웃을 전달합니다. 아래의 예제는 dim3를 활용해 커널 함수를 호출하는 예제입니다.

 

 

#include <stdio.h>

#include "cuda_runtime.h"
#include "device_launch_parameters.h"


__global__ void CheckIndex(void) {
  printf("threadIdx:(%d,%d,%d) blockIdx:(%d,%d,%d) blockDim:(%d,%d,%d) gridDim:(%d,%d,%d)\n",
    threadIdx.x, threadIdx.y, threadIdx.z,
    blockIdx.x, blockIdx.y, blockIdx.z,
    blockDim.x, blockDim.y, blockDim.z,
    gridDim.x, gridDim.y, gridDim.z);
}

int main() {
  dim3 dim_block(3, 1, 1);
  dim3 dim_grid(2, 2, 1);

  printf("dim_block:(%d,%d,%d) dim_grid:(%d,%d,%d)\n",
    dim_block.x, dim_block.y, dim_block.z,
    dim_grid.x, dim_grid.y, dim_grid.z);

  CheckIndex<<<dim_grid, dim_block>>>();

  cudaError_t cudaStatus = cudaDeviceSynchronize();
  if (cudaStatus != cudaSuccess) {
    printf("cudaDeviceSynchronize() failed: %s\n", cudaGetErrorString(cudaStatus));
  } else {
    printf("cudaDeviceSynchronize() success\n");
  }

  return 0;
}

 

 

위의 코드처럼 스레드 레이아웃은 dim3 구조체를 활용해 설정합니다. dim_grid, dim_block을 설정하고, 설정된 그리드와 블록의 사이즈를 출력합니다. 이후, CheckIndex() 커널 함수에 전달해 그리드, 블 각 차원의 크기와 현재 블록, 스레드의 인덱스 정보를 출력합니다. 위 프로그램의 출력 결과는 다음과 같습니다.

 

 

CUDA 예제 코드 출력 결과

 

 

예제에서 그리드의 크기는 (3, 1, 1), 블록의 크기는 (2, 2, 1)로 설정되었습니다. 따라서, 생성되는 스레드의 개수는 총 12개입니다. 출력된 결과를 보시면, 모든 스레드는 동일한 그리드, 블록의 크기를 가지고 있습니다. 스레드 레이아웃, 즉, 그리드와 블록의 형상은 커널 실행 시 결정되고, 모든 스레드가 같은 값을 공유하는 것입니다.

 

위 결과처럼 스레드가 같은 인덱스 정보를 가질 수 있습니다. 단, 스레드가 같은 인덱스 정보를 갖더라도 블록의 인덱스는 다릅니다. 이처럼 스레드 인덱스뿐만 아니라 블록의 인덱스를 함께 고려해야 합니다.

 

 

마치며

금일 포스팅에서 CUDA의 스레드 계층에 대해 알아보았습니다. CUDA의 스레드 계층은 어떻게 구성이 되어있고, 각 스레드 계층의 정보를 프로그램에서 확인하기 위한 내장 변수에 대해 알아보고, 실제 커널 함수 호출을 통해 확인해 봤습니다. 이어지는 포스팅에서는 오늘 설명드린 스레드 레이아웃 정보를 활용해서 연산하는 프로그램을 작성해 보겠습니다.

 

오늘 포스팅도 도움이 되셨기를 바라면서 글 마치겠습니다. 고맙습니다.

 

 

Reference

1. CUDA 기반 GPU 병렬 처리 프로그래밍 - 기초부터 성능 최적화 전략까지

2. CUDA C++ Programming Guide