[CUDA] 기초
GPU는 많은 양의 연산을 병렬처할 수 있으며, 이를 지원하기위해 CUDA라는 GPU programming으르 통해 효율적인 연산 병렬화를 지원한다.
행렬 곱셈의 예)
행렬 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이 필요한다.
CUDA 프로그램 구조
CUDA program은 한 개 이상의 영역(phase)로 구성된다. [CPU영역 + GPU영역, nvcc가 이 둘을 분리한다.]
device code
kernel이라 불리는 data 병렬 함수와 관련된 자료구조를 명시하는 keyword를 확장한 상태로 작성한다. kernel은 일반적으로 data 병렬성을 활용하기 위해 많은 수의 thread로 작성된다. 행렬 곱에서는 곱을 하나의 kernel function으로 구현하고(Single Instruction) P의 각 원소 계산을 하나의 thread에 mapping한다(Multiple Thread).
위의 그림은 CPU serial code가 실행되다 중간중간에 kernel function이 시작되는 모습을 보여주며 kernel은 device로 옮겨져 많은 수의 thread(통합해서 grid)를 생성한다.
행렬곱의 예)
Host(CPU)에서의 행렬 곱 연산
1. host의 memory에 matrix M,N,P할당 후 M,N을 읽는다.
2. M x N 계산을 한다. => 다중 루프 연산
3. 계산 결과를 matrix P에 적는다.
row major의 한계점
계산의 해야되는 각 원소들이 떨어져 주소에 mapping되어 locality관점에서 문제가 생긴다.
Device(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
데이터 전송
host ocde로 부터 data를 전송할 수 있는 memory로 device의 global, constant memory가 있다. cuda memory는 programmer가 관리할 수 있도록 API를 제공하고있다.
cudaMalloc()
host code에서 불려지면 대상을 저장하기 위한 global memory를 할당한다.
- 인자 1 : 할당할 대상물을 가리키는 포인터 주소
- 인자 2 : 할당할 대상물의 크기
device global memory 할당 후, host로부터 device로 data를 전송할 것을 요청한다.
cudaMemcpy()
- 인자 1 : 복사 연산의 목적지를 가리키는 포인터
- 인자 2 : 복사할 근원지를 가리키는 포인터
- 인자 3 : 복사할 byte 수
- 인자 4 : 복사에 참여하는 memory의 type (host,device <=> host,device)
CUDA에서의 행렬 곱 예)
kernel함수는 병렬영역에 있는 code를 모든 thread를 이용해 수행하도록 지정한다.
__global__
함수가 kernel임을 지정하며, host fuction에서 호출되며, device에서 thread grid를 생성한다.
__device__
cuda device에서만 호출되며, kerenl 함수, 다른 device함수에서만 호출된다.
Kernel functions and Threading
Threads와 blocks
Thread : devcie에서 programmer에 의해 작성된 kernel program을 수행하는 가장 작은 단위의 주체
Blcok(thread block) : thread의 집합, block 내의 thread들은 하나의 SM을 공유(SM할당 단위)
Grid : block의 집합, 하나의 kerenl에 의해 처리되는 모든 thread의 집합
threads를 controlling하기위한 변수들
gridDim : grid 한 차원의 길이
blockDim : thread block 한 차원의 길이
blockIdx : thread가 소속된 block의 index
threadIdx : thread에 부여된 index
각 thread가 자료구조의 어느 부분을 사용해야 하는것도 지정해야한다. 이 keyword(threadIdx)는 미리 정의된 변수를 지정한다. 이 변수들을 사용해 thread는 수행 중에 thread의 thread의 좌표값을 담고 있는 hardware register로 접근한다.
" thread는 서로다른 threadId를 가진다. "
CPU의 code의 바깥 2개의 for loop는 GPU code의 threadId로 대체된다. loop를 반복할때마다 i,j를 증가시키는 대신 cuda의 threading hardware는 각 thread에 해당되는 threadId를 한꺼번에 생성한다. kernel이 호출되면 kernel은 병렬 thread의 grid에 의해 수행되며, 일반적으로 grid는 수 만 개의 가벼운 thread로 구성된다.
그림 17에서 각 thread block은 4x2x2x의 3차원 thread로 구성되어 있다. 즉, grad 1은 64개의 thread를 갖는다. host code가 kernel을 호출할때, host는 execution configuration변수를 사용해 grid, thread block의 크기를 지정한다.
