[GPGPU Series 4] GPGPU Application Thread Hierarchy

이번 글은 CUDA로 작성된 GPGPU Application의 Thread Hierarchy에 대한 설명을 정리하였다. GPGPU Application에서 Thread Hierarchy와 GPU Warp (Wavefront)의 개념을 이해하면 GPGPU의 개념의 절반 이상을 이해했다고 할 수 있다. 다시 말해서 Thread Hierarchy와 Warp 개념은 GPU에서 가장 중요한 부분이다.

보통 CUDA로 작성된 코드는 GPU에서 실행되는 함수와 CPU에서 실행되는 코드로 구분된다. CUDA로 작성된 코드 중 GPU에서 실행되는 함수는 보통(항상?) 아래와 같은 형태로 작성한다. GPGPU에 대한 글을 모두 작성하고 나면 CUDA로 Matrix Multiplication 코드 예제를 작성할 계획이다. CUDA 코드를 전혀 모르더라도 너무 신경 쓰지 않고 읽어도 이해하는데 큰 문제는 없을 것 같다.

__global__ void functionA(..){
	...
}

__global__ void functionA(..){
	...
}

CPU에서 해당 함수를 호출할 때는 아래와 같이 코드를 작성한다.

void main(){
	...
	dim3 dimBlock(16, 16);
	dim3 dimGrid(32, 32);
	functionA <<<dimGrid, dimBlock>>> (...);
	functionB <<<dimGrid, dimBlock>>> (...);
	...
}

dim3는 X, Y, Z 값을 나타내는 Dimension을 줄임말이라고 생각하면 된다. dimBlock(16, 16)은 x = 16, y = 16, z = 1로 정의한 것이다. dimGrid(32, 32)도 비슷하게 x = 32, y = 32, z = 1로 정의한 것이다. functionA, functionB 에 <<<…>>> 안에 표시된 부분이 Thread Hierarchy를 정하는 핵심적인 부분이다. 간단히 설명하면 위의 예제 코드는 총  16 x 16 (256)개의 Thread로 구성된 32 x 32 (1,024) Thread Block (TB)를 생성하여 functionA와 functionB 코드를 실행한다는 의미이다.

그림 1: 16 x 16 Thread Block 예제 그림

dimBlock(16, 16)을 설정하여 functionA와 functionB를 호출하면 그림 1과 같이 각 Thread Block (TB or Cooperative Thread Array (CTA))은 총 256개의 Thread로 구성된다. TB의 각 Thread는 고유의 Index (ID) 번호를 가진다. 제일 왼쪽 위에 위치한 Thread는 threadIdx.x = 0, threadIdx.y = 0, threadIdx.z = 0과 같이 고유의 번호를 가진다. 그림에서는 편의를 위해서 Thread ID (0, 0, 0)으로 표시하였다. TB의 Thread는 각 고유의 ID를 가짐으로 TB의 어느 위치에 존재하는지 알 수 있다. CUDA로 작성된 GPGPU Application이 GPU에서 실행되는 원리를 설명할 때는 Thread ID는 크게 중요하지 않다. 하지만 CUDA 코딩을 할 때는 상당히 중요한 부분이다.

그림 2: 32 x 32 Grid 예제 그림

dimGrid(32, 32)을 설정하면 functionA와 functionB을 호출하면 그림 2와 같이 Grid는 총 32 x 32 (1024) 개의 TB을 생성하게 된다. Grid의 TB 역시 고유한 Index 번호를 가진다. 예를 들어 제일 왼쪽 위에 위치한 TB는 blockIdx.x = 0, blockIdx.y = 0, blockIdx.z = 0가 된다. 그림에서는 편의를 위해서 TB (0, 0, 0)으로 표시하였다. 결과적으로 TB는 Grid의 어디에 위치하는지 알 수 있다.

현재까지 주어진 정보를 조합하면 functionA는 총 1,024개의 TB를 생성하고 각 TB는 256 Thread를 생성한다. 결과적으로 총 262,144개의 Thread를 생성하여 functionA 함수 연산을 수행하는 것이다. 이렇게 실행되는 한개의 Function을 보통 Kernel이라고 한다. 위 예제 코드의 경우 functionA(), functionB() 두개의 함수가 존재하기 때문에 한 Application은 총 2개의 Kernel로 구성되어 있다. 만약 functionA, functionB, …, functionZ까지 있다면 총 26개의 Kernel로 구성된 Application이 된다.

그림 3: Application Hierarchy 예제

그림 3은 Application Hierarchy를 보여준다. 각 Application은 1개 또는 그 이상의 Kernel로 구성되어 있으며, 각 Kernel은 한 개 또는 그 이상의 TB로 구성되어 있다. 각 TB는 1개 또는 그 이상의 Thread로 구성되어 있다. Application의 Kernel을 구성하는 Thread Block, Thread 크기는 항상 같을 필요가 없다. 다시 말해서 functionA, functionB, …, functionZ 모두 다른 Grid, TB 크기를 가질수 있다. Application이 어떻게 구성되어 있는지 작성하였다. 다음 장에서는 Kernel의 TB가 실행되는 순서에 대해서 설명할 예정이다.

2 thoughts on “[GPGPU Series 4] GPGPU Application Thread Hierarchy”

  1. 잘 읽고 있습니다. 그런데 코드 예제에서 `functionA <<>> (…);`부분에서 두 변수 모두 dimGrid로 작성하신게 오타인가 하여 댓글 남깁니다.

    Reply
    • 안녕하세요. 글 읽어주셔서 감사합니다.
      오타가 맞습니다. 코드 예제 부분 수정하였습니다.

      Reply

Leave a Reply to jinu Cancel reply