CUDA Programming Model

References

• https://developer.nvidia.com/blog/cuda-refresher-cuda-programming-model/
• CUDA Toolkit Documentation
• https://github.com/nvidia/cuda-samples

Contents

• CUDA kernel and thread hierarchy
• Matrix Addition 예제
• Memory hierarchy

---

이전 CUDA 관련 포스트들에서 CUDA의 간략한 아키텍처와 CUDA 프로그래밍의 기본적인 부분에 대해서 알아봤습니다.

GPU와 CUDA

CUDA 프로그래밍 기초 (예제 : vector addition)

CUDA Thread 구조와 Data Mapping (예제 : 이미지 흑백, Blur 처리)

 

새로운 것들을 더 알아보기 전에, 다시 한 번 CUDA의 스레드 계층과 하드웨어 계층 간의 연관성을 알아보고, 또 메모리 계층은 어떻게 구성되어 있는지 알아봐야겠다고 생각했었는데, NVIDIA 블로그에 CUDA 프로그래밍 모델에 대한 간략한 포스트가 있어서 내용을 가져와 정리해보려고 합니다. 

LINK

CUDA Refresher: The CUDA Programming Model | NVIDIA Developer Blog

 

---

CUDA Programming Model

CUDA 프로그래밍 모델은 어플리케이션과 GPU 하드웨어 사이의 어떠한 관계가 있는지 보여주는 GPU 아키텍처의 추상화를 제공합니다. 간략하게 알아보도록 합시다.

 

CUDA 프로그래밍 모델에서 사용되는 두 가지 키워드, host와 device를 먼저 소개하겠습니다.

host는 시스템에서 사용가능한 CPU를 의미합니다. CPU와 연결된 시스템 메모리를 host memory라고 합니다.

그리고, GPU를 device라고 하며 GPU memory는 device memory라고 지칭합니다.

 

CUDA 프로그램을 실행하기 위한 3가지 스텝은 다음과 같습니다.

• host memory의 input data를 device memory로 복사. host-to-device transfer라고 합니다.
• GPU 프로그램을 load하고 실행 (성능을 위해 data를 on-chip에 caching)
• device memory의 output result를 host memory로 복사. device-to-host transfer라고 합니다.

 

CUDA kernel and thread hierarchy

아래 이미지는 GPU에서 실행되는 함수인 CUDA 커널을 보여주고 있습니다. 프로그램에서 병렬로 실행될 수 있는 부분은 K개의 서로 다른 CUDA 스레드에 의해서 병렬로 K번 실행됩니다.

모든 CUDA 커널은 __global__ 선언 지정자(declaration specifier)로 시작합니다. 그리고 프로그래머는 각 스레드에 유니크한 전역 ID를 내장된 변수를 사용하여 할당할 수 있습니다.

스레드의 그룹은 CUDA 블록(block)이라고 합니다. CUDA 블록은 다시 그리드(grid)로 그룹화됩니다. 한 커널은 스레드의 블록들로 구성된 그리드로 실행이 됩니다.

각 CUDA 블록은 하나의 streaming multiprocessor(SM)에 할당되어 실행되며, GPU의 다른 SMs로 이동할 수 없습니다 (선점, 디버깅, 또는 CUDA 동적 병렬화 중에는 예외입니다.). 하나의 SM은 CUDA 블록에 필요한 리소스에 따라서 여러 개의 블록을 동시에 실행할 수 있습니다. 각 커널은 하나의 device에서 수행되고 CUDA는 여러 커널들이 하나의 device에서 동시에 실행할 수 있도록 지원합니다. 아래 Figure 3은 GPU에서 사용할 수 있는 하드웨어 리소스에 대한 커널 스레드 계층과의 매핑을 보여줍니다.

 

CUDA는 스레드와 블록을 위해서 내장된 3D 변수를 정의합니다. 스레드들은 내장된 3D 변수 threadIdx를 사용하여 인덱싱할 수 있고, 3차원의 인덱싱은 일반적으로 vectors, matrix, volume의 element에 접근할 수 있는 방법을 제공합니다. 스레드와 똑같이, 블록 또한 blockIdx라는 내장된 3D 변수를 통하여 인덱싱할 수 있습니다.

 

여기서 몇 가지 주목할만한 것들이 있습니다.

• CUDA Architecture는 블록당 스레드의 수를 제한합니다. (1024 threads per block limit)
• 스레드 블록의 차원은 커널 안의 내장된 blockDim 변수를 통해서 알아낼 수 있습니다.
• 한 블록 안의 존재하는 모든 스레드는 내재된 함수인 __syncthreads()를 사용하여 동기화할 수 있습니다. __syncthreads() 함수가 호출되면, 호출한 스레드는 블록에 있는 모든 스레드가 해당 지점에 도달할 때까지 기다리게 됩니다. (동기화와 관련된 내용은 따로 다루도록 하겠습니다.)
• 블록당 스레드의 수와 그리드당 블록의 수는 <<<...>>> 문법에 의해서 지정됩니다. (int나 dim3 타입으로)

아래의 Matrix Addition 예제를 살펴보겠습니다. 아래 matrixAdd 커널 함수 코드는 두 행렬을 더하기 위해서 2차원의 blockIdx, threadIdx와 blockDim을 사용하는 방법을 보여주고 있습니다. 아래 예제 코드에서 행렬의 각 요소를 쉽게 인덱싱하기 위해서 2차원의 블록을 사용했습니다. 각 블록에는 256개의 스레드가 있으며, 각각 x, y 방향으로 16개씩 스레드가 존재합니다. 코드에서 필요한 총 블록의 수는 데이터의 크기(여기서는 N x N 행렬)를 각 블록의 크기로 나누어서 계산됩니다.

__global__
void matrixAdd(const float *A, const float *B, float *C, const int M, const int N)
{
    int ROW = blockIdx.x * blockDim.x + threadIdx.x;
    int COL = blockIdx.y * blockDim.y + threadIdx.y;

    if (ROW < M && COL < N) {
        C[(ROW * N) + COL] = A[(ROW * N) + COL] + B[(ROW * N) + COL];
    }
}

int main()
{
	...
    // Matrix addition kernel launch from host code
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks((N + threadsPerBlock.x - 1)/threadsPerBlock.x, (N + threadsPerBlock.y- 1)/threadsPerBlock.y);
    matrixAdd<<<numBlocks, threadsPerBlock>>>(MatA, MatB, MatC);
    ...
}

 

Matrix Addition 전체 코드는 아래 링크를 참조하시기 바랍니다. 코드는 NVIDIA에서 제공되는 샘플을 참고하여 작성하였습니다.

https://github.com/junstar92/parallel_programming_study/blob/master/CUDA/matrixAdd/matrixAdd.cu

GitHub - junstar92/parallel_programming_study

컴파일 후 실행하면 다음과 같은 결과를 얻을 수 있습니다.

 

Memory hierarchy

CUDA를 사용할 수 있는 GPU는 다음과 같은 메모리 계층을 가지고 있습니다.

GPU architecture에서 확인할 수 있는 메모리는 다음과 같습니다.

• Registers - 레지스터는 각 스레드에서 private 합니다. 즉, 한 스레드에서 할당된 레지스터는 다른 스레드에서 볼 수 없고 접근할 수도 없습니다. 레지스터 활용은 컴파일러에 의해서 결정됩니다.
• L1 / Shared memory(SMEM) - 모든 SM은 L1 캐시와 공유메모리로 사용할 수 있는 빠른 on-chip scratchpad memory(SPRAM)을 가지고 있습니다. 하나의 CUDA 블록 안에 있는 모든 스레드들은 서로 shared memory를 공유하고, 하나의 SM 내에서 실행 중인 모든 CUDA 블록들은 SM으로부터 제공되는 physical memory를 공유할 수 있습니다.
• Read-only memory - 각 SM은 하나의 instruction cache, constant memory, texture memory, R0 cache를 가지고 있으며, 커널 코드에서 읽기 전용으로 사용합니다.
• L2 cache - L2 캐시는 모든 SMs 사이에서 공유되고, 모든 CUDA 블록의 모든 스레드들이 이 메모리에 액세스할 수 있습니다. RTX3080의 L2 캐시 사이즈는 5MBytes입니다.
• Global memory - GPU와 DRAM(GPU에 부착된)의 framebuffer 사이즈입니다.

NVIDIA CUDA 컴파일러는 메모리 리소스를 훌륭하게 최적화합니다. 하지만, CUDA 개발자는 필요에 따라 CUDA 프로그램을 최적화하기 위해 메모리 계층을 효율적으로 사용할 수 있습니다.

 

Compute capability

GPU compute capability는 GPU 하드웨어에 의해서 결정되는 일반적인 사양과 사용 가능한 기능을 결정합니다. 이 compute capability는 프로그램에서 런타임에 현재 GPU에서 사용할 수 있는 하드웨어 기능 또는 명령을 결정하는데 사용할 수 있습니다.

모든 GPU에는 X.Y의 형태로 compute capability가 제공되는데, X는 major revision number고 Y는 minor revision number입니다. minor revision number는 동일 아키텍처에서의 새로운 기능의 추가와 같은 개선이 되었다는 것을 의미합니다.

 

CUDA를 지원하는 GPU의 자세한 compute capability는 NVIDIA에서 제공하는 샘플 코드 deviceQuery를 실행하여 확인할 수도 있습니다. 

https://github.com/NVIDIA/cuda-samples/tree/master/Samples/deviceQuery

GitHub - NVIDIA/cuda-samples: Samples for CUDA Developers which demonstrates features in CUDA Toolkit

---

정리

CUDA 프로그래밍 모델에서 host 코드는 CPU에서C/C++ 프로그램을 실행하고, 커널은 물리적으로 분리된 GPU device에서 실행되는 이기종 환경(heterogeneous envrionment)를 제공합니다. 또한 CUDA 프로그래밍 모델은 host와 device가 각각 분리된 메모리 공간(host memory와 device memory)를 사용한다고 가정합니다. 그리고, CUDA PCIe Bus를 통해 host와 device 메모리 간의 데이터 전송을 제공합니다.

 

CUDA에서는 많은 built-in 변수들이 존재하고 쉽게 프로그래밍할 수 있도록 다차원 인덱싱을 가능하게 합니다.

또한, 레지스터, shared memory, L1 cache, L2 cache, globla memory와 같은 다양한 메모리가 있으며, 이러한 메모리를 효율적으로 사용하여 CUDA 프로그램을 최적화할 수 있습니다.