0. 시작하며
안녕하세요, 가톨릭 대학교 박사 과정 이제영입니다.
오랫동안 blog를 방치하고 있었는데, 가끔은 제가 공부했던 것이나 생각나는 것을 정리해두면 좋을 것 같아 새로운 형태로 다시 시작하게 되었습니다.
오늘은 Cupy를 이용하여 CUDA Custom Kernel을 작성하는 연습을 진행하다가 CUDA에서 헷갈렸던 부분을 다시 정리해봅니다.
개인적으로 정리하는 것이니 제가 이렇게 이해했다 정도로 참고해주시면 감사하겠습니다.
1. Thread, Block, Grid와 GPU
CUDA를 사용한 GPU 프로그래밍시 thread와 block, grid에 대한 이야기는 항상 등장합니다.
CUDA를 배울땐, thread가 모여서 block, block이 모여서 grid 라고하며 보통 CUDA 프로그래밍에서 Grid는 작업을 구성하는 가장 상위 단위로, 하나의 GPU는 여러 Grid를 동시에 실행할 수 있습니다.
그리고 일반적으로 thread와 block를 내가 할당할수 있다는 것과 어떻게 할당 해야 하는지,
각 단위의 의미와 한번에 병렬 처리 하여 연산한다는 것 정도를 집고 주로 Kernel Function을 작성하는 예제로 넘어가게 됩니다.
제가 갑자기 헷갈렸던 것은 GPU의 스펙과 thread/block과의 관계였습니다.
제가 Ambassador로 활동하고 있는 NVIDIA DLI과정 중 Accelerating Data Engineering Pipelines Optimization 과정에선 Hardware System으로서의 Data Engineering Pipelines 파트가 있습니다.
이 파트에서 GPU의 CUDA Core와 Streaming Multiprocessor(SM)의 언급과 쉽게 설명하기 위해 정확하진 않지만, 쉽게 말하자면, 한번에 연산가능한 CUDA Core가 thread의 개수를 의미하는 것이고 SM이 Block의 개수를 의미하는 것이다.라고 강의자료에 되어있습니다.
그런데, 실제로 Block개수와 Thread의 개수를 GPU에 작성되어있는 CORE개수와 SM개수만큼 할당하고 연산할 수 있는 것일까요??
제가 간단하게 검색해보기론 이러한 부분에 대한 이야기가 자세히 나와있지 않은것 같아서 정리해보고자 합니다.
2. 간단한 예시 RTX 3090
이에대해 파악하기 위해 우선 간단하고 유명한 GPU인 RTX 3090중 CUDA 프로그래밍 특히 이번 포스트에서 알아보고자 하는 부분을 이해하기 위한 스펙을 적어보도록하겠습니다.
| 항목 |
세부내용 |
항목 |
세부내용 |
| CUDA core 개수 |
8704개 |
블록당 최대 쓰레드 수의 곱 |
1024개 |
| SM 개수 |
68개 |
워프 당 실행 쓰레드 수 |
32개 |
| SM당 CUDA 코어 수 |
128개 |
SM당 최대 워프 수 |
64개 |
| SM당 최대 쓰레드 수 |
2048개 |
SM당 최대 블록 수 |
32개 |
table 1. RTX 3090 스펙
3. GPU와 CUDA core
일반적으로 GPU를 잘쓴다는 것은 무엇일까요? 여러가지 의미가 있겠지만, 보통 GPU를 잘 다룬다는 것은 Utilization은 최대한 100%로 유지하면서, task를 처리하는 것일 겁니다. 즉 GPU의 사용량은 최대로(효과적으로) 유지하면서 나의 작업을 빠르게 처리하는 것이 보통 GPU를 잘 다룬다는 것일 겁니다.
그리고 실은(어찌보면 당연한 것이겠지만) 우리가 할당할 수 있는 thread의 개수와 block의 개수는 GPU별로 다르며, 또한 한번에 연산가능한 숫자와 한 Kernel function이 할당 받을 수 있는 thread와 block의 개수도 GPU의 제원을 통해 어느정도 확인 가능 합니다.
따라서 우리는 우리가 사용하고자하는 GPU의 사양을 명확하게 알고, CUDA 프로그래밍을 시도해야겠습니다.
일단 가장 먼저 봐야하는 것은 당연하게도 CUDA core의 개수인데요, CUDA Core의 개수는 말 그대로 CUDA core의 총 개수를 의미하는 것이기 때문에, SM당 CUDA코어 수와 SM개수를 곱하는 것으로 구할 수 있습니다. 또한 기본적으론 CUDA core의 개수가 우리가 사용하고자 하는 GPU의 한번에 연산 가능한 thread의 총 개수라고 보면 될 것 같습니다 :)
여기서 한가지 의문이 드는데요, 한번에 연산 가능한 thread의 총 개수가 RTX 3090의 개수라고하면, 할당 가능한 총 쓰레드의 개수는 다를까요??
정답은 다르다입니다. 이를 위해서 확인 해봐야 하는 것은 SM당 최대 쓰레드 수인데요. GPU에서 쓰레드를 할당 한다는 것은, GPU를 사용해 한번에 연산한다는 것과는 다른 의미였습니다…
4. thread 할당과 warp 연산
그러면 한번에 연산 가능한 최대 개수가 8704개라는 것과 할당 가능한 thread의 개수가 차이가 나는 이유가 무엇일까요?? 이는 CUDA 코어와 쓰레드의 차이에 대해 이해 해야합니다.
그냥 쉽게 말하자면, CUDA 코어는 실제 연산을 수행하는 물리적 계산 장치이고, 쓰레드는 작업을 나누기 위한 논리적 작업 단위이기 때문입니다.
어찌보면 모두가 알고 있는 것인데, 제대로 생각해본적이 없는 것 같습니다 :)
그렇다면 우리가 thread를 할당할 때 최대 몇개까지 할당할 수 있을까요??
정답은 $2048 * 68 = 139264 \; 개입니다. 그렇다면 왜 2048개의 thread까지 할당할 수 있을까요?? 이는 GPU에서의 warp연산과 스케줄링을 이해 해야합니다.
GPU에서 연산을 처리할 때, SM은 thread를 Warp(32개의 쓰레드 단위)로 묶어 실행합니다. SM은 한 번에 4개의 Warp를 병렬로 실행하며, 나머지 Warp는 대기 상태로 스케줄링됩니다.
이는 무슨 이야기냐면, SM은 쿠다 코어를 32개를 1개로 묶어 동시에 연산을 시키게 되는데요, RTX 3090은 SM 안에 쿠다 코어가 128개 씩 존재하니, 총 4개의 warp가 동시에 실행된다고 볼 수 있습니다.
마지막으로 SM당 할당 가능한 최대 warp의 숫자는 64개이므로 $64 * 32 = 2048 \; 개까지 할당 가능합니다.
5. thread 할당과 block 할당
마지막으로 block에 대한 이야기인데요, 일반적으로 CUDA를 배운다고 하면, thread의 묶음을 block로 배우곤합니다.
물론 맞는 이야기입니다만, 오늘은 GPU에서 할당되는 연산과 엮어서 같이 볼 예정이므로, block당 thread의 최대개수와 SM당 block의 최대개수를 동시에 보면서 이야기 하겠습니다.
우선 기본적으로 block은 thread의 묶음이 맞습니다. 그래서 우리가 CUDA에서 맨처음 하는 것은 thread 개수와 block 개수를 정하게 되죠, 기본적으로 CUDA는 받은 thread개수와 block개수를 SM이 최대한 공평하게 나누어 가지려고 시도합니다. 지금의 경우에는 68개의 SM이 최대한 공평하게 나누어 가지게되겠네요.
이 때, SM은 최대 32개의 block까지 할당 가능합니다. 또한 한 block당 thread의 최대 개수는 1024개까지입니다. 이는 1차원 thread가 기준으로 만약 2차원 이상이라면 각 차원의 곱이 해당 개수까지만 할당 가능합니다. 이 점을 주목하며 한번 생각 해보겠습니다.
만약 CUDA를 아래와 같이 코딩한다면 GPU에는 어떻게 할당이 될까요?
#include <cuda_runtime.h>
#include <stdio.h>
__global__ void printThreadInfo() {
// Global ID 계산
int globalIdx = blockIdx.x * blockDim.x + threadIdx.x;
// 블록 ID와 쓰레드 ID 출력
printf("BlockIdx.x: %d, ThreadIdx.x: %d, Global ID: %d\n", blockIdx.x, threadIdx.x, globalIdx);
}
int main(){
int threadsPerBlock = 128;
int numBlocks = 4;
printThreadInfo<<<numBlocks, threadsPerBlock>>>();
}
이경우 CUDA는 4개의 SM에 Block을 하나씩 할당하여 128개의 thread씩, 총 512개의 thread를 할당할 것입니다.
또한 RTX 3090기준으론 하나의 SM에 128개의 쿠다 코어가 존재하니 GPU는 한번에 512개의 값을 연산할 수 있을 것입니다.
그렇다면 남은 SM은 어떻게 될까요??
6. block할당과 GPU성능
이 때 CUDA 프로그래밍의 묘미가 발생하게 됩니다.
실은 위의 방식은 예제이기 때문에, 간단하게 이야기했지만, 현재의 상태에선 남은 64개의 SM은 활용하고 있지 않습니다.
즉 우리가 효과적으로 GPU를 활용하기 위해선, 어떻게하면 SM을 효율적으로 사용할 수 있을지를 고민해야합니다.
지금의 경우, 각 SM은 한 번에 128개의 쓰레드만 처리할 수 있으므로 4개의 블록이 128개의 쓰레드를 포함할 때, GPU는 한 번에 512개의 쓰레드를 실행할 수 있습니다. 나머지 SM은 비활성 상태로 남게 됩니다.
또 SM의 개수는 68개이지만, block는 32개씩할당이 가능합니다. 연산만 놓고 봤을땐, 한번에 $128 * 68 = 8704\;번 계산되겠죠? 즉 GPU가 한번 딸깍하면 8704번 덧셈합니다.
우리는 이런 환경에서 어떻게 데이터를 잘게 쪼개어서 효과적으로 연산할지를 고민해야합니다.
여기서 추가적으로 고려 해봐야하는 것도 있습니다. 워프의 기본단위는 32입니다. 즉 하나의 thread는 32개씩 연산이됩니다. 32의 배수로 할당하는 것이 좋을까요?
혹은 128개의 thread를 68개의 block으로 할당할수도, 32개의 thread를 272개의 block에 할당할 수도 있을 것입니다. 두개의 차이는 무엇일까요? 둘다 RTX3090은 한번에 연산할 것입니다.
block의 개수는 68개의 SM에맞추어 배수로 설정하는 것이 좋을까요?? thread의 처리와 block의 처리중 어느것을 우선하여 처리할까요? 예를들면, 272개의 block에 128개의 thread를 할당하면 어떻게 될까요??
많은 질문이 새롭게 떠오르면 즐거운 밤입니다.
감사합니다.
Jeyoung Lee
Ph D. Candidate at The Catholic of University of Korea, Computer Vision researcher
Comments