이 글은 CUDA 기반 GPU 병럴처리 프로그래밍 책의 내용을 바탕으로 정리한 글이다.
chapter 4 - 5 에서 핵심은 스레드에 대한 이해다. 기본적이지만 병렬처리를 하려면 반드시 암기하고 있어야 하는 내용이기도 하다. 어느 정도로? 1+1 =? 에 대한 대답정도로.
CUDA 스레드 계층
사실 책에서는 쓰레드, 워프, 블록, 그리드로 나누어 설명한다. 그런데 그렇게 설명을 하면 한 가지 헷갈리는 부분이 생기게 된다. 사실 우리가 코딩에 필요한 부분은 고유 스레드 id를 부여하는 것이다. 그래서 내가 생각하는 계층은 아래와 같이 표현하는 것이 좋고, 실행 단위가 Warp 라는 것을 기억하면 된다.
Tread -> Block -> Grid
참 간단하다. 이 구조를 설명하면, block은 Thread를 가지고 있고 grid는 block을 가지고 있는 형태다. 각각의 요소에 대해 조금 더 설명을 이어가겠다.
Thread
Thread 는 CUDA 연산을 수행하거나, CUDA 코어를 사용하는 기본 단위다. 딱 이 이상도 이하도 아니다. 그냥 thread 하나가 core 하나를 사용한다라고 생각하면 편하다.
Warp ( 실행 단위 )
Block을 설명하기에 앞서 Warp에 대해 설명하고자 한다. Warp는 CUDA의 기본 실행 단위를 말한다. 이게 무슨 말이냐면, 한 warp에 속한 Threads 들은 하나의 제어 장치에 의해 제어가 된다. 이를 좀 더 풀어 말하면, 하나의 명령에 따라 32개의 Thread가 하나의 warp로 묶여, 동시에 움직인다는 말이다. 왜 이걸 설명하느냐, 추후에 병렬 처리 코드 특히, cuda c 프로그램들을 보면 block을 구성하는 단위가 32의 배수다. warp 자체는 indexing에 사용되지는 않지만, 효율적인 프로그래밍을 위해서는 반드시 필요한 개념이다.
Block
block은 warp들의 집합이다. 그런데 이렇게 설명하면, 헷갈리게 된다. 그래서 다시 설명하면, block은 thread들로 구성이 되어 있고, 그 thread들이 32개의 단위로 실행된다고 이해해야 헷갈림이 덜 하다. 블록이 가지는 중요한 사실은 thread별로 고유 ID를 부여한다는 점이다. block 내 thread는 1,2,3 차원의 형태로 배열 될 수 있으며, 한 block가 가질 수 있는 Thread의 수는 1024개이다.( 이를 간과하기 쉽다 ) 또한 CUDA 공식 가이드에 따라 block의 최대 크기는 x, y 차원의 경우 1024, z 차원의 경우는 64 이다.
Grid
grid는 block들의 집합이고, 하나의 grid에 포함된 block들은 자신만의 고유 ID를 가진다 grid 내에서 block은 1,2,3 차원의 형태로 배열되어 질 수 있으며, x-차원의 최대 길이는 2^31-1 이고, y, z 차원의 경우 2^16-1 = 65535 이다. 쉽게 x차원으로는 제한이 없고 y,z 차원으로 제한이 있다고 생각하면 된다.
Thread 계층을 위한 내장 변수
gridDim
gridDim은 그리드의 형태 정보를 담고 있는 structure형 내장 변수이다. 멤버 변수로, x,y,z 를 가지고 있다. 각 차원의 길이는 1이상의 양의 정수로 정의되어야 한다. 잠깐 하나만 짚고 가면, grid는 block들의 집합이었다. 그럼 gridDim에 따라 블록들이 어떻게 배치될지가 결정된다. 그럼 그 형태에 따라 blockIdx를 정할 수 있게 된다. 당연하다고 느껴지겠지만, 인덱싱을 할 때 헷갈릴 우려가 있어서 미리 짚고 넘어간다. 돌아가 커널이 실행될 때 gird및 block의 형태가 결정되며 한 gird내 모든 block은 동일한 형태를 가진다. gridDim 내장 변수는 커널 내 모든 스레드가 공유한다.
blockIdx
blockIdx는 현재 thread가 속한 block의 번호를 담고 있는 struct형 내장 변수이다. grid 내 block들은 서로 다른 block 번호를 가지며, 최대 3차원의 번호를 가질 수 있다. 각 차원은 x,y,z로 표현된다. 여기서 헷갈리지 않아야될 것은 gridDim의 경우 각 차원의 길이가 1이상의 양수로 정의되어야하지만, blockIdx는 차원이 0이 될 수 있다. gridDim과 관련해서 설명하면, gridDim은 block들의 형태를 정의하고, blockIdx로 각 블록을 지칭하게 된다. 시각적으로 표현하면 아래와 같다. 이 예시를 정확히 이해하고 간다면, 아래 나오는 내장 변수에도 헷갈림이 없을 것이다.
blockDim
blockDim은 block의 형태 정보를 담고 있는 structure형 내장 변수이다. 내장 변수로 역시, x, y, z를 가지고 있으며 각 차원의 크기를 담고 있다. gridDim과 마찬가지로 각 차원의 길이는 1 이상의 양의 정수로 정의되어야 한다. 앞서 언급했지만 다시 한번 짚고 넘어가자면, 커널이 실행될 때 grid 및 block의 형태가 결정되며, 한 grid내 모든 block은 동일한 형태를 가진다. 즉, blockDim 역시 모든 thread가 공유하게 된다.
threadIdx
threadIdx는 block 내에 현재 thread가 부여받은 thread 번호를 담고 있는 structure형 내장 변수다. block내 thread들은 서로 다른 번호를 가지며, 최대 3차원의 번호를 가지고 각 차원은 x, y, z로 표현되며 각 차원의 길이는 0 이상이다. 위에서 설명한 gridDim - blockIdx 의 관계와 blockDim - threadIdx의 관계는 같다. 이제 이해를 돕기 위해 하나의 예시를 들고자 한다.
아래의 예시는 물론 차원을 고려하고 있지 않다. 즉 serialize 하게 데이터를 배열했지만, 아래의 예시를 완벽하게 이해한다면, 차원을 늘려도 쉽게 적용가능 할 것이다. 이 예시는 메모리 입장에서 생각했다고 보면 편하다.
예시, gird는 3 x 2의 형태를 가지고 있고, 각 block은 5x2x4 의 형태를 가지고 있다. 이때, blockIdx( 2, 1) 의 threadIdx( 4, 1, 3 ) 에 해당하는 thread에 고유 번호를 부여하고자 한다. blockIdx(0,0)의 threadIdx(0,0,0) 이 0 번의 번호를 부여받고 x부터 y, z 순으로 순차적으로 올라간다고 하면, 해당 thread의 id는 몇 번인가?
- 전략
1. block이 몇 번째 block인지를 파악한다.
2. block 당 thread를 몇 개를 가지고 있는지를 파악한다.
3. 지나온 block에 포함된 thread 수에 현재 block에서 지칭하고 있는 고유 id를 더해준다.
- sol
1. blockIdx는 (2, 1) 이고, 순차적으로 올라감으로 x 로 (0,0), (1,0), (2,0), (0,1), (1,1), (2,1) 이 됨으로 6 번째 block임을 알 수 있다. 위를 수식으로 표현하여 계산하면, ( 2+1 ) x (1 +1 )로 계산할 수 있다.
2. block당 가지고 있는 thread의 개수는 blockDim 을 활용할 수 있으며, 총 5x2x4 = 40 개의 thread를 가지고 있다.
3. 지나온 block은 6개 이고 포함된 thread 개수는 6x40= 240 개이고 , 현재 block에서 지칭하는 thread를 threadIdx를 활용해 계산하면 5x2x4 임으로 40이다. 따라서 240 + 40 = 280 임으로 지칭하는 thread의 고유 id는 280이다.
Thread 레이아웃
thread layout은 thread 배치 형태를 지칭하고 grid와 block의 형태로 정의된다. 즉 gridDim과 blockDim을 활용하게 된다. thread layout은 커널 호출 시 설정하며, kernel을 호출할 때 사용되는 <<< gridDim, blockDim >>> 이라는 실행 구성 문법을 통해 전달된다. 위 예시에서 든 thread layout은 gridDim 3x2x1 이고, blockDim 5x2x4 이다. 이를 코드로 표현하면 아래와 같다.
예시의 thread layout에 대한 코드 표현
Dim3 dimGrid( 3, 2, 1 )
Dim3 dimBlock( 5, 2, 1 )
kernel<<<dimGrid, dimBlock>>>();
쭉 이어가려고 했는데, 그 다음 인덱싱이 grid 내에서 global indexing 하는 부분을 적다보니 너무 길어져서 해당 부분은 다음 글에서 다루겠다.
'CUDA C' 카테고리의 다른 글
| [ CUDA 기반 GPU 병렬처리 프로그래밍 ] chapter 4-5 스레드 계층, 레이아웃과 인덱싱 2 (0) | 2023.11.09 |
|---|---|
| [ CUDA 기반 GPU 병렬처리 프로그래밍 ] chapter 2-3 CUDA 프로그램 구조와 흐름 (0) | 2023.11.06 |
| [ CUDA 기반 GPU 병렬처리 프로그래밍 ] chapter 1 GPGPU 및 병렬처리 개요 (2) | 2023.11.01 |
