[컴퓨터구조#15] GPU (1) - 기초, CUDA, GEMM example

참고: Computer Architecture: A Quantitative Approach (5th edition) - 4.1.~4.4.

이번 글에서는 GPU에 대해 알아보자.

목차

1. GPU 아키텍처
2. CUDA
3. GEMM example

---

1. GPU 아키텍처

GPU 미리보기: multi-core, multi-threaded, SIMD-like 아키텍처

프로세싱은 매우 데이터 병렬적인 과정이다. GPU는 highly multithreaded되어 있고, 긴 메모리 지연 시간을 숨기기 위해 쓰레드 스위칭을 사용한다. 또한, 그래픽스 메모리는 wide하고 높은 대역폭을 가지고 있다. 일반 목적 GPU를는 보통 이종 CPU/GPU 시스템에 사용되어 순차 코드는 CPU에서, 병렬 코드는 GPU에서 돌리는 식으로 사용된다. GPU의 프로그래밍 언어/API는 DirectX, OpenGL, Cg, HLSL, CUDA(Compute Unified Device Architecture) 등이 있다.

GPU 아키텍처에서는 높은 메모리 대역폭이 필수적이다. 

예를 들어 엔비디아 Tesla 구조를 통해 살펴보자.

하나의 Streaming Multiprocessor (SM) 안에는 명령어의 fetch와 decode를 담당하는 front end (I-cache, MT Issue, C-Cache), 실행을 담당하는 8개의 Streaming Processor (SP), 그리고 SM 내 PE들이 공유하는 shared memory가 존재한다. 이 Shared memory는 소프트웨어에 의해 관리되는 캐시인 작은 양의 SRAM으로, 외부 DRAM이 존재하는 멀티프로세서와는 다르다.  

• SP: single-precision FP 및 정수를 처리한다. 각 SP는 fine-grained multithreaded되어 있다.
• Warp: 32개의 쓰레드의 집합으로, 병렬적으로 SIMD style(=> warp 내 모든 쓰레드는 같은 명령어를  실행)로 실행된다.
  • 8개의 SP가 존재하므로, 각 SP에서 한 cycle에 한 개의 쓰레드를 실행한다고 하면 한 warp를 처리하는 데 4 clock cycle이 소요된다. (32/8=4)
  • 레지스터, PC 등 하드웨어 컨텍스트 저장에 24 warp(=23*32 thread)가 사용된다.

출처: https://tatourian.blog/2013/09/03/nvidia-gpu-architecture-cuda-programming-environment/

GPU 분류하기

GPU는 SIMD/MIMD 모델에 정확히 들어맞지는 않는다. 말하자면 GPU의 프로그램 모델은 MIMD, 물리적 구현은 vector-like 아키텍처라고 할 수 있다. 따라서, warp 내에서도 모든 쓰레드가 각각의 컨트롤 플로우를 가지며, diverging data control path를 지원한다.

쓰레드에서의 조건부 실행은 MIMD의 illusion을 제공한다. 하지만 성능 저하가 발생할 수 있어 주의해서 general purpose code를 작성해야 한다.

	Static: Discovered at Compile Time	Dynamic: Discovered at Runtime
Insturction-Level Parallelism	VLIW	Superscalar
Data-Level Parallelism	SIMD or Vector	GPU

GPU 아키텍처

그림을 통해 GPU 아키텍처를 다시 정리해 보자.

GPU는 compute-intensive, highly parallel computation 및 그래픽에 특화된 처리 장치이다. GPU 내 트랜지스터는 주로 프로세싱(ALU)에 할당되며, data caching이나 flow control 관련 unit은 매우 제한적이다.

아래 그림에서 보면 CPU에서는 캐시, 컨트롤 unit이 차지하는 비중이 높은 반면 GPU에서는 대부분이 연산 unit(초록색으로 된 core. ALU라고 봐도 됨)이다.

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

 각 코어의 차이점을 다음 표로 정리해 보자.

	CPU-style core	GPU
각 코어를 구성하는 요소들	Fetch/decode Unit ALU *Execution Context	Fetch/decode Unit   (shared) ALU                         (보통 8개) *Execution Context (shared)
	Data Cache OoO control logic branch predictor Memory prefetcher
특징	큰 Fetch/Decode Unit 작은 ALU 큰 Data Cache	하나의 명령어 스트림을 빠르게 작동하도록 만들어주는 요소들을 제거하고, 대신 꼭 필요한 로직을 많이 배치한다.

*Execution Context: PC 등 저장하는 레지스터

GPU에서는 이러한 코어들을 여러 개 배치하여, 각 코어가 동시에 서로 다른 명령어 스트림을 독립적으로 실행할 수 있도록 한다. 만약 16개의 코어가 존재한다면, 한번에 16개의 명령어 스트림 및 (각 코어당 8개의 ALU가 존재하므로 16*8=)128개의 쓰레드를 동시에 실행할 수 있다.

또한, Fetch/Decode Unit 및 Context Data가 공유되도록 하여 많은 ALU 간 명령어 스트림 관리에 필요한 비용/복잡성을 낮춘다. (GPU가 타겟하는 응용은 명령어의 computational density가 높기 때문 → 즉 가능한 많은 쓰레드를 돌리도록 하기 위함.)

하나의 ALU당 하나의 쓰레드가 할당된다. 하나의 쓰레드는 고유한 context 데이터를 갖고 있으며, 이는 Shared Context Data에 쓰레드별로 하나씩 저장된다. 또한, ALU에서 실행되는 모든 쓰레드는 같은 프로그램(명령어) 하에서 돌아가지만, 쓰레드별로 실행 path는 다양하게 diverge할 수 있다는 점을 유의해야 한다.

GPU의 실행 방식에 대해서 생각해보자. 하나의 warp를 실행하는 경우, 처음에 메모리 latency로 인해 stall되는 구간이 존재한다. 따라서 파이프라인처럼 여러 warp들을 순차적으로 실행하여 처음으로 실행한 warp의 stall이 끝날 때까지 계속해서 다음 warp의 실행을 시작한다. 이런 식으로 하드웨어 멀티쓰레딩을 통해 긴 메모리 접근 latency를 숨길 수 있다. 따라서 GPU는 여러 그룹의 쓰레드(warp)를 오버랩할 수 있을 정도로 높은 메모리 bandwidth가 필수적이다. 

이때 계속해서 서로 다른 warp를 실행한다는 것은 context switch에 대한 오버헤드를 고려하게 되는데, GPU에서는 모든 context 정보를 하드웨어에 저장(context 저장에 사용되는 메모리 많음)하므로 별다른 cost 없이 빠르게 쓰레드를 계속 전환할 수 있다. 각 쓰데르가 필요로 하는 context (레지스터) 양에 따라 저장할 수 있는 context의 개수가 달라진다. 더 많은 레지스터를 필요로 할수록 저장할 수 있는 context 개수가 적어지고, 이 경우 latency hiding ability가 더 낮다. 쓰레드가 필요로 하는 레지스터의 양은 CUDA kernel에서 컴파일한 이후에 결정된다.

마지막으로, GPU 아키텍처에서 coherence와 synchronization는 어떻게 되는지 알아보자. 같은 core(=SM) 내 (8개의 ALU)에서는 synchronization이 지원된다. 하지만, 서로 다른 SM에서 돌아가는 쓰레드에 대해서는 synchronization에 대해 지원하지 않는다. 또한, 현재 GPU는 캐시 coherence를 지원하지 않으며(당연하게도 GPU용 캐시 일관성 프로토콜도 역시 존재하지 않는다), 프로그램에서도 하드웨어 차원의 coherence support를 가정하지 않는다.

 

2. CUDA

CUDA는 thread level parallelism을 제공하며, GPU를 host CPU 및 관련 메모리 모델에 대한 highly threaded coprocessor인 GPU에 집중한다. CUDA program은 CPU에서 돌아가는 host code와, GPU 내 CUDA core(=kernel)에서 돌아가는 device code로 나뉜다. CUDA는 thread group의 위계를 사용한다. 이때 thread group이란 프로그래밍 모델의 일부로서 warp보다 큰 단위이다. (warp는 GPU의 스케줄링 단위로, 마이크로아키텍처의 구현 컨셉에 더 가깝다.) CUDA에서 사용되는 shared memory와 barrier synchronization은 같은 thread group 내에서만 허용된다.

CUDA에 쓰이는 용어들을 간단히 정리하면 다음과 같다. Host와 device는 분리된 메모리를 가지며, CUDA program은 host code와 device code를 모두 포함한다.

• Host : 보통 CPU를 뜻함. - ANSI C 로 쓰인 코드.
• Device: 보통 GPU를 뜻함 (data-parallel) - extended ANSI C로 쓰인 코드.
• Kernel: data-parallel function. 커널 호출 = 디바이스 내 lightweight 쓰레드 생성. 쓰레드는 하드웨어를 통해 생성/스케줄됨.

host와 device가 통합된 CUDA는 다음과 같이 돌아간다. host C code에서는 serial하거나 적당히 parallel한 부분을, device SPMD kernel C mode에서는 고도로 parallel한 부분을 다룬다. 

Serial Code (host) → Parallel Kernel (device) KernelA<<<nBlk,nTid>>>(args); → Serial Code (host) → Parallel Kernel (device) KernelA<<<nBlk,nTid>>>(args); → ...

N개의 서로 다른 CUDA 쓰레드에 의해 병렬적으로 N 번 실행되는 벡터 덧셈 코드 예시를 살펴보자.

VecAdd의 세 파라미터 A,B,C는 GPU 메모리 내 cudaMalloc으로 할당된 array에 대한 포인터를 나타내고, main 함수 내 VecADd<<<1, N>>>은 1개의 쓰레드 블록과 N개의 쓰레드를 나타내는 것으로 실행 configuration을 나타낸다.

// Kernel definition
__global__ void VecAdd(float* A, float* B, float* C) {
	int i = threadIdx.x;
    	C[i] = A[i] + B[i];
}

int main() {
	...
    // Kernel invocation with N threads
    VecAdd<<1, N>>>(A,B,C);
}

CUDA 프로그래밍 모델

커널은 쓰레드 블록의 집합인 grid 단위로 실행된다. 보통 하나의 커널은 하나의 grid에 대응한다고 생각하면 된다. 하나의 쓰레드 블록은 최대 512개의 쓰레드로 구성된다. 

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

디바이스와 커널의 차이는 무엇일까?

• 디바이스: CPU 혹은 host의 coprocessor로, 고유의 DRAM 즉 device memory를 갖고 있다. 디바이스에서는 많은 쓰레드를 병렬적으로 돌린다. 보통 GPU를 뜻하지만 다른 종류의 병렬 프로세싱 장치가 될 수 있다.
• 한 응용의 data-parallel한 부분들은 여러 쓰레드에서 돌아가는 디바이스 커널로 표현될 수 있다.
• GPU와 CPU 쓰레드의 차이는 다음과 같다.
  • GPU 쓰레드: 매우 lightweight하며 생성 오버헤드가 거의 없다. GPU의 full efficiency를 위해서 수천 개의 쓰레드가 필요하다.
  • CPU 쓰레드: 멀티코어 CPU에서는 몇 개의 쓰레드만 있어도 된다. 

CUDA 커널은 쓰레드 array로 실행된다. 같은 커널 내 모든 쓰레드는 동일한 코드를 실행하게 되며(SPMD: Single Program Multiple Data), 각 쓰레드는 메모리 주소를 계산하고 컨트롤 결정을 내리는 데 사용되는 쓰레드 id를 가진다. 이 쓰레드 id는 꼭 linear할 필요는 없고 (0,0,0)과 같이 다차원일 수도 있다. 이는 data structure에 따라 실행상 용이한대로 정할 수 있다.

이러한 monolithic한 쓰레드 어레이는 여러 개의 쓰레드 블록으로 나뉠 수 있다. 같은 블록 안에 있는 스레드들은 shared memory, atomic operation, barrier synchronization을 통해 협력한다. 하지만 다른 블록에 있는 쓰레드들은 이를 사용할 수 없다. 하나의 쓰레드 블록은 하나의 SM에서 돌아간다. 한 SM에서 여러 개의 쓰레드 블록이 실행되는 것은 가능하지만, 하나의 쓰레드 블록이 여러 개의 SM으로 span되는 것은 불가능하다.

CUDA 메모리 모델

쓰레드는 고유한 local memory를 가지고, 같은 블록 안에 있는 쓰레드들은 Per-Block shared memory를 통해 소통할 수 있다. 같은 그리드 안에 있는 블록은 global memory를 통해 소통할 수 있다.

Per-CUDA Thread Private Memory → Per-thread local memory, Per-Block Local Memory → Per-block shared memory, GPU Memory → Global memory

Global Memory: 호스트와 디바이스 간의 읽기/쓰기 데이터 통신에 사용된다. 내용물은 모든 쓰레드에게 visible하며, 긴 접근 latency를 가진다. 아키텍처상으로는 GPU의 external DRAM에 위치하게 된다. 

호스트 코드에서 CUDA 디바이스 메모리를 다음과 같이 관리할 수 있다.

• cudaMalloc(): 디바이스의 Global Memory에 오브젝트를 할당한다. 할당된 오브젝트에 대한 포인터 주소와, 할당된 오브젝트 크기를 파라미터로 갖는다.
• cudaFree(): 디바이스 Global Memory에 할당된 메모리를 해제한다. 해제할 포인터를 파라미터로 갖는다.
• 예를 들면 다음과 같이 사용 가능하다. (64*64 single precision float array를 디바이스 스토리지에 할당 및 해제)

TILE_WIDTH = 64;
Float *Md;
int size = TILE_WIDTH*TILE_WIDTH*sizeof(float);

cudaMalloc((void**)&Md, size);
cudaFree(Md);

• cudaMemcpy(): 메모리 데이터 전달에 사용하며, 다음 4가지 파라미터를 갖는다. 또한 비동기적 전송이기 때문에 프로그램의 다른 부분과 동시에 실행이 가능하다.
  • dest 포인터, source 포인터, 복사할 바이트 수, 전달 종류(Host-Host, Host-Device, Device-Host, Device-Device)
• 예를 들면 다음과 같이 사용 가능하다.

cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice); // CPU -> GPU, M -> Md
cudaMemcpy(M, Md, size, cudaMemcpyDeviceToHost); // GPU -> CPU, Md -> M

커널 함수는 다음과 같이 execution configuration을 통해 호출되어야 한다. CUDA 1.0부터 커널 함수는 비동기적이며, blocking을 위해서는 별도의 명시적 synch가 필요하다.

__global__ void KernelFunc(...);
dim3	DimGrid(100, 50);	// 5000 thread blocks
dim3	DimBlock(4, 8, 8);	// 256 threads per block
sizze_t	SharedMemBytes = 64;	// 64 Bytes of shared memory
KErnelFunc<<<DimGrid, DimBlock, SharedMemBytes>>>(...);

 

3. GEMM Exaple

두 행렬 A와 B를 곱해서 결과 C를 만드는 예시(AB=C)를 생각해보자. 하나의 쓰레드는 output C의 한 element를 계산하게 된다. A와 B는 global memory에서 각각 A.width, B.height번 로드된다.

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

Step 1: Simple Host Version in C

i와 j를 통해 C의 element 위치를 지정하고, k를 통해 정해진 row와 column을 곱해 계산한다.

// Matmul on the CPU host in double precision
void MatrixMulOnHost(float* A, float* B, float* C, int A_height, int B_width, int c_width)
{
	for (int i=0; i<A_height; ++i)
    	for (int j=0; j<B_width; ++j) {
        	double sum = 0;
                for (int k=0; k<c_width; ++k) {
                    double a = A[i*A_height+k];
                    double b = B[k*B_width+j];
                    sum += a*b;
                }
            P[i*c_width+j] = sum;
    }
}

Step 2: Input Matrix Data Transfer

host-side code로, 필요한 인풋 데이터를 GPU로 전달하는 부분이다.

void MatrixMulOnDevice(float* A, float* B, float* C, ...)
{
	int size_A = A_width*A_height*sizeof(float);
    int size_B = B_width*B_height*sizeof(float);
    int size_C = A_height*B_width*sizeof(float);
    float* Ad, Bd, Cd;
    
    // Allocate & load A,B to device memory
    cudaMalloc(&Ad, size_A);
    cudaMemcpy(Ad, A, size_A, cudaMemcpyHostToDevice);
    
    cudaMalloc(&Bd, size_B);
    cudaMemcpy(Bd, B, size_B, cudaMemcpyHostToDevice);
    
    // Allocate C on the device
    cudaMalloc(&Cd, size_C);

Step 3: Output Matrix Data Transfer

계산 완료된 output 행렬을 다시 호스트로 전달하는 부분이다.

    // Kernel invocation - step 5
    ...

    // read P from device
    cudaMemcpy(C, Cd, size_C, cudaMemcpyDeviceToHost);

    // Free device matrices
    cudaFree(Ad); cudaFree(Bd); cudaFree(Cd);
}

Step 4: Kernel Function

커널 함수는 각 쓰레드에 적용되어 쓰레드의 behavior를 나타낸다. thread id는 2차원으로 x,y정보를 포함한다.

// Matmul kernel - per thread code
__global__ void MatrixMulKernel(float* Ad, float* Bd, float* Cd, int width_C)
{
	// Cvalue: store the element of the matrix & computed by the thread
    float Cvalue = 0;
    
    for (int k=0; k<width_C; ++k){
    	float Aelement = Ad[threadIdx.y*width_C+k];
        float Belement = Bd[k*width_C+threadIdx.x];
        Cvalue += Aelement*Belement;
    }
    
    Cd[threadIdx.y*width_C+threadIdx.x] = Cvalue;
}

이 코드는 step 1에서의 가장 inner loop에 해당한다.

CPU 코드와는 thread 개수 측면에서 차이가 있다.

• CPU-side parallel program(중간고사 범위): 쓰레드별로 생성/유지 cost가 많이 들어 쓰레드를 그렇게 많이 만들지 않는다. 행렬을 여러 chunk로 나눠 하나의 쓰레드가 각 chunk를 담당하도록 한다. 즉 하나의 쓰레드당 하는 일이 적지 않다.
• GPU-side parallel program: 쓰레드 유지비용이 거의 공짜이므로 쓰레드를 가능한 한 많이 만들 수 있다. output element 하나만 계산하도록 한다. 

Step 5: Kernel Invocation

host side에서, 커널호출하는 부분을 추가한다.

// setup the execution configuration
dim3 dimGrid(1,1);
dim3 dimBlock(width_C, width_C);

// launch the device computation threads
MatrixMulKernel<<<dimGrid, dimBlock>>>(Ad, Bd, Cd, width_C);

 

하지만, 이 예시에서 우리는 더 개선의 여지를 볼 수 있다. 바로 한번에 하나의 스레드 블록만 사용되고 있다는 것이다. 이 실행 예시에서, 각 쓰레드는 Cd의 하나의 element만을 계산한다. 즉, 하나의 쓰레드는 Ad의 row와 Bd의 column을 로드하고, 각 Ad, Bd 쌍에 대해 하나의 곱셈과 하나의 덧셈을 수행한다. 계산 : off-chip 메모리 접근 비율이 1:1로 그다지 높지 않은 것이다. 

이 비율을 높이고 성능을 개선하기 위해, 다음과 같이 Cd를 여러 개의 타일로 나누어 각 쓰레드 블록이 하나의 타일을 계산하도록 만들 수 있다. 쓰레드 블록 내 하나의 쓰레드가 타일의 element 하나를 계산하게 되는 것이다. 이때 타일 크기는 블록 크기와 동일하다.

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

바뀐 방식을 바탕으로, step 4의 커널 함수와 step 5의 커널 호출 코드를 다음과 같이 수정할 수 있다.

Step 4 수정: Kernel Function

__global__ void MatrixMulKernel(float* Ad, float* Bd, float* Cd, int width_C)
{
	// Calculate the row index of the Cd element and A/B
    int Row = blockIdx.y*TILE_WIDTH + threadIdx.y;
    int Col = blockIdx.x*TILE_WIDTH + threadIdx.x;    
    
	// Cvalue: store the element of the matrix & computed by the thread
    float Cvalue = 0;
    
    // each thread computes one element of the block sub-matrix
    for (int k=0; k<width_C; ++k){
        Cvalue += Ad[Row*width_C+k]*Bd[k*width_C+Col];
    }
    
    Cd[Row*width_C+Col] = Cvalue;
}

Step 5 수정: Kernel Invocation

// setup the execution configuration
dim3 dimGrid(B_width/TILE_WIDTH, A_height/TILE_WIDTH);
dim3 dimBlock(TILE_WIDTH, TILE_WIDTH);

// launch the device computation threads
MatrixMulKernel<<<dimGrid, dimBlock>>>(Ad, Bd, Cd, width_C);