[CUDA] 기초

GPU는 많은 양의 연산을 병렬처할 수 있으며, 이를 지원하기위해 CUDA라는 GPU programming으르 통해 효율적인 연산 병렬화를 지원한다. 

 

행렬 곱셈의 예)

[그림 1] 행렬 연산, [그림2] GPU/CPU연산의 차이

행렬 P의 각 원소는 M의 행과 N의 열의 내적연산의 결과이다. 이 과정에서 P의 서로 다른 원소를 얻기위해 내적을 동시에 실행할 수 있다(data 계산의 병렬성). 이 원리를 통해 CUDA device(GPU)는 CPU보다 더 빨리 연산이 가능하다. 

 

pitfall! 

"GPU가 CPU에 비해 연산이 더 빠르다고 CPU가 필요없는 것은 아니다. "

CPU와 GPU는 사용목적과 배경에 따라 서로 다르게 발전했다. GPU는 CPU보다 훨씬 많은 양의 ALU를 보유하고 있어 병렬연산에 용이하지만, 현재까지 CPU(host)없이 GPU는 구동할 수 없다. GPU를 구동하기위해 host는 연산과 데이터를 GPU에 옮겨야하며, GPU의 연산 결과는 추후 host로 리턴해야된다. 이를 위해 CUDA program이 필요한다. 

[그림3] CPU, GPU간 통신 

CUDA 프로그램 구조 

 CUDA program은 한 개 이상의 영역(phase)로 구성된다. [CPU영역 + GPU영역, nvcc가 이 둘을 분리한다.] 

 

device code 

 kernel이라 불리는 data 병렬 함수와 관련된 자료구조를 명시하는 keyword를 확장한 상태로 작성한다. kernel은 일반적으로 data 병렬성을 활용하기 위해 많은 수의 thread로 작성된다. 행렬 곱에서는 곱을 하나의 kernel function으로 구현하고(Single Instruction) P의 각 원소 계산을 하나의 thread에 mapping한다(Multiple Thread).

[그림4] GPU kernel의 thread.(초록 화살표)

위의 그림은 CPU serial code가 실행되다 중간중간에 kernel function이 시작되는 모습을 보여주며 kernel은 device로 옮겨져 많은 수의 thread(통합해서 grid)를 생성한다. 

 

행렬곱의 예)

Host(CPU)에서의 행렬 곱 연산 

[그림5] host에서의 행렬 곱 연산 

1. host의 memory에 matrix M,N,P할당 후 M,N을 읽는다. 

2. M x N 계산을 한다. => 다중 루프 연산 

3. 계산 결과를  matrix P에 적는다.

[그림6] 행렬의 주소 mapping 방식(row major)

row major의 한계점 

계산의 해야되는 각 원소들이 떨어져 주소에 mapping되어 locality관점에서 문제가 생긴다. 

 

Device(GPU)에서의 행렬 곱 

[그림7] GPU에서의 행렬곱

1. M,N,P matrix의 복사본을 할당 후 device로 복사한다. 

2. 행렬 곱셈을 병렬적으로 하는 kernel을 호출한다. 

3. device memory부터 P를 host memory에 복사한다. 

"input data를 device로 전송, device에서 계산 시작, device로부터 결과 취합" 

 

Device memory와 데이터 전송 

CUDA에서 host, device는 별개의 memory space를 갖는다. device에 있는 여러 종류 memory를 할당하고, 이동하고, 사용하도록 하기위한 device memory와 data transfer과정을 살펴본다. 

Device Memory 

Global Memory

 - off chip memory로 host에서의 memory와 동일한 개념 

 - 모든 thread에서 접근 가능 

 - Access latency가 가장 높으므로, 성능 제약 발생 

Register 

 - 가장 빠른 on chip memory 

 - 많은 thread를 관리하기 위해서 대용량 register가 존재 

Shared Memory 

 - 낮은 latency를 제공하는 programmable memory 

 - 동일 thread block내의 모든 thread가 공유 가능 

 - programmer가 사용 여부를 선택 

 - 사용 용도 : 연산과정에서의 중간값 저장, 자주 참조하는 data등 

Constant Memory 

 - size가 작은 read only memory 

 

데이터 전송  

[그림8] host/device data transfer

host ocde로 부터 data를 전송할 수 있는 memory로 device의 global, constant memory가 있다. cuda memory는 programmer가 관리할 수 있도록 API를 제공하고있다. 

cudaMalloc()

host code에서 불려지면 대상을 저장하기 위한 global memory를 할당한다. 

 - 인자 1 : 할당할 대상물을 가리키는 포인터 주소

 - 인자 2 : 할당할 대상물의 크기 

[그림9] cudaMalloc()의 예

device global memory 할당 후, host로부터 device로 data를 전송할 것을 요청한다. 

cudaMemcpy()

 - 인자 1 : 복사 연산의 목적지를 가리키는 포인터 

 - 인자 2 : 복사할 근원지를 가리키는 포인터 

 - 인자 3 : 복사할 byte 수 

 - 인자 4 : 복사에 참여하는 memory의 type (host,device <=> host,device)

[그림10] cudaMemcpy()의 예

CUDA에서의 행렬 곱 예) 

[그림11] CUDA에서의 행렬 곱 예

kernel함수는 병렬영역에 있는 code를 모든 thread를 이용해 수행하도록 지정한다. 

__global__ 

함수가 kernel임을 지정하며, host fuction에서 호출되며, device에서 thread grid를 생성한다. 

__device__

cuda device에서만 호출되며, kerenl 함수, 다른 device함수에서만 호출된다. 

[그림12] device memory allocation

[그림13] device memory allocation

Kernel functions and Threading 

Threads와 blocks

Thread : devcie에서 programmer에 의해 작성된 kernel program을 수행하는 가장 작은 단위의 주체 

Blcok(thread block) : thread의 집합, block 내의 thread들은 하나의 SM을 공유(SM할당 단위) 

Grid : block의 집합, 하나의 kerenl에 의해 처리되는 모든 thread의 집합 

[그림14] GPU구조

threads를 controlling하기위한 변수들 

gridDim : grid 한 차원의 길이 

blockDim : thread block 한 차원의 길이 

blockIdx : thread가 소속된 block의 index 

threadIdx : thread에 부여된 index

[그림15] thread를 제어하기 위한 변수

각 thread가 자료구조의 어느 부분을 사용해야 하는것도 지정해야한다. 이 keyword(threadIdx)는 미리 정의된 변수를 지정한다. 이 변수들을 사용해 thread는 수행 중에 thread의 thread의 좌표값을 담고 있는 hardware register로 접근한다. 

" thread는 서로다른 threadId를 가진다. "

[그림16] CPU,GPU의 행렬 곱 code 비교

CPU의 code의 바깥 2개의 for loop는 GPU code의 threadId로 대체된다. loop를 반복할때마다 i,j를 증가시키는 대신 cuda의 threading hardware는 각 thread에 해당되는 threadId를 한꺼번에 생성한다. kernel이 호출되면 kernel은 병렬 thread의 grid에 의해 수행되며, 일반적으로 grid는 수 만 개의 가벼운 thread로 구성된다. 

[그림17] thread 계층구조

그림 17에서 각 thread block은 4x2x2x의 3차원 thread로 구성되어 있다. 즉, grad 1은 64개의 thread를 갖는다. host code가 kernel을 호출할때, host는 execution configuration변수를 사용해 grid, thread block의 크기를 지정한다. 

[그림18] execution configuration