4. CUDA 스레드 블록 아키텍처 (1)

CUDA에서 스레드는 계층 구조로 되어있는데 그리드-블록-스레드 구조로 이루어져있다. 스레드가 모여 블록을 이루고 블록이 모여 그리드를 이루고 이를 그리드 블록 모델이라고 한다.

그리드 블록 모델

스레드를 생성하고 어떻게 병렬처리를 하는지에 대해 알기 위해서는 우선 GPU 아키텍처에 대해 알 필요가 있다. GPU에서는 여러개의 SP가 모여 SM을 이루고 SM들이 모여 TPC를 이룬다. TPC는 GPU 아키텍처마다 쓰는 경우도 있고 안쓰는 경우도 있기 때문에 이는 생략하고 다수의 SP가 SM을 구성하는 것만 일단 기억하도록 하자.

GPU 아키텍처

CUDA 스레드 모델에서 SP는 4개의 스레드를 동시에 실행할 수 있고 하나의 블록은 하나의 SM과 대응하여 동작하게 된다. 동일한 커널 함수를 512개의 스레드를 생성하여 위 그림의 왼쪽 G80 아키텍처에서 실행한다고 하자. G80 아키텍처는 8개의 SP를 가진 SM 16개로 이루어져 있다.

Kernel <<<1, 512 >>> (d_a, d_b, d_c);

이때 커널 함수를 블록 1개에 스레드 512개로 호출한다고 가정하면 1개의 SM에서만 스레드가 생성되어 동작하기 때문에 GPU의 성능을 전부 발휘하지 못한다.

Kernel <<<8, 64>>> (d_a, d_b, d_c);

그렇다면 블록의 개수를 조금 늘려 8개의 블록에서 64개의 스레드를 생성하여 총 512개의 스레드를 생성하여 실행한다면 GPU 절반인 8개의 SM이 동작하게 된다.

Kernel <<<16, 32>>> (d_a, d_b, d_c);

마지막으로 블록의 개수를 16개의 블록에서 32개의 스레드를 생성하여 총 512개의 스레드를 생성하여 실행한다면 16개의 SM 전부 동작하여 효율적으로 GPU의 성능을 끌어낼 수 있다.

 

Turing 아키텍처

하지만 이는 시간이 지나 GPU의 성능과 구조가 바뀌게 되면서 조금 다르게 생각하여야 한다. 필자의 GPU 아키텍처인 Turing 아키텍처만 보더라도 72개의 SM에 SM하나당 64개의 SP를 가진다. 이때 위에서의 조건으로 스레드를 생성하여 실행한다면 56개의 SM이 유휴상태가 된다. NVIDIA에서는 64배수로 스레드를 지정하는 것을 권장하고 있고 추후 그래픽 카드의 발전을 위해 충분히 큰 수를 지정하도록 권장하고 있다.