Kim Seon Deok
[CUDA] ch1. Introduction to CUDA Programming 본문
* Learn CUDA Programming 책을 바탕으로 하였습니다.
Outline
- High-performance computing의 역사
- Technical requirements
- CUDA 를 사용해 Vector addition 연산
- CUDA 에서 Error reporting
- CUDA 에서 Data type support
High-performance computing의 역사
High-Performance-Computing(고성능 컴퓨팅)은 과거 Mega-Floating Point Operations(MFLOPs)에서부터 PetaFLOP 계산이 가능한 수준으로 향상되었다.
- Floating-Point Operations (FLOPs) per second
프로세서의 이론적 최대 성능(peak compute)을 측정할 때 사용되는 기본 단위, 즉 프로세서의 성능 측정 지표
- Mega FLOPs = 10^6 -> 1초에 백만 개의 부동 소수점 연산 수행
- Peta FLOPs = 10^15 -> 1초에 1조 개의 부동 소수점 연산을 수행
- Instruction-Level Parallelism(ILP)
특정 코드에서 independent한 instruction이 동시에 실행될 수 있는지를 나타냄
CPU는 더 빠른 clock rate를 달성하기 위해 5~15stage로 수행한다.
HPC technology의 변화
HPC technology에서 주요 변화 단계를 epoch로 나눌 수 있다.
- Epoch1 : CRAY-1 -> single vector CPU아키텍처로 최대 160MegaFLOP의 성능을 나타냄
- Epoch2 : CRAY-2 -> 4 core vector CPU아키텍처로 최대 2GigaFLOP의 성능을 나타냄
- Epoch3 : CRAY T3D -> 3D Torus 네트워크 구조를 가지며 1TeraFLOP의 성능을 나타냄, 대역폭은 300MB/s이다.
이후의 기술적 변화는 다음과 같이 나타났다.
- 8비트, 16비트, 32비트로 발전하며, 현재는 64비트 instruction set으로 변화
- ILP 증가
- core 갯수 증가
이러한 기술적 변화는 clock rate의 증가로 인해 support 되었고, 현재 clock rate는 4GHz에 도달했다.
또한 이러한 성능 향상은 semiconductor 산업에 지대한 영향을 미쳤던 두 가지 법칙에 기인한다.
1. Moore's Law : 밀도 높은 intergrated circuit에 포함된 transistor 수가 2년 마다 2배로 증가한다는 법칙
2. Dennard scailing : Moore의 법칙을 지속시키는 scailing 법칙으로 트랜지스터 크기와 전력 밀도 간의 관계에 대한 법칙
Q : 트랜지스터의 갯수 f : operating freqeuncy C : capacitance
v : operating voltage Ileakage : leakage current
무어의 법칙과 덴나드 스케일링은 transistor size를 줄여서 number of transistors per chip을 향상시킴에 있어 cost effectiveness가 증가한다는 것에서 연관이 있다.
트랜지스터 사이즈가 1/S의 비율로 줄어드는 동안 트랜지스터 갯수는 2배로 증가했다. 또한 frequency는 2년 마다 40%씩 증가했다. 하지만 이 규칙은 leakage current로 인해 65nm에 도달한 이후로는 더이상 유지되지 못했다.
leakage current의 영향을 줄이기 위해 switching process에 변화가 생겼는데, processor design에서 전압은 1V로 유지되었고 더 이상 power envelope를 일정하게 유지하는 것이 어려워짐으로 인해 "Powerwall"현상이 발생했다.
덴나드 스케일링이 유지되지 못하게 되면서 프로세서의 크기가 45nm에서 16nm으로 줄어들었음에도 불구하고 enegy/chip size가 3배로 증가되었다. 동시에 아키텍쳐의 pipeline stage는 15stage 이상으로 늘어났다.
instruction pipeline을 유지하기 위해 프로그램의 branch와 memory address를 예측하는 speculation 기술이 사용되었다.
깊은 pipeline stage와 legacy software 작성 방식은 사용되지 않는 일부 트랜지스터를 만들어 내었고, clock cycle가 낭비되게 하였다. 결론적으로 프로그램의 성능 측면에서 개선이 없었다.
이후 GPU가 등장하였고 graphic processing에 사용되었다. 나중에 GPU가 non-graphic으로 활용되면서 General Purpose Computation using GPU(GPGPU)라는 용어를 사용하게 되었다. 대부분의 compute-intensive한 task들은 matrix multiplication과 같은 data-parallel한 특성을 갖고 있다. 당시 GPU를 사용하는 데 유저들이 갖고 있었던 유일한 문제는 GPU를 활용하려면 Graphic pipeline을 이해해야 한다는 점이었다. GPU에서 계산 작업은 shader execution을 중심으로 이루어졌다. High Performance Computing을 수행하려면 이보다 더 일반적인 interface가 필요했고, 이러한 문제점은 2007년 CUDA가 개발되면서 해결되었다.
computer architectuer의 진화는 sequential processing에서 출발해 distributed memory로 나아가면서 programming 모델에 큰 영향을 주었다. GPU 아키텍쳐도 무어의 법칙과 덴나드 스케일링에 영향을 받지만, 트랜지스터를 다르게 할당하여 기존 아키텍쳐보다 높은 성능을 달성한다.
High Performance Computing System의 발전 과정
Heterogeneous Computing
GPU는 코드 중 parallel한 특성을 가진 부분을 가속화해서 application을 가속화하는 데 사용된다.
CPU는 코드 중 latency bound한 부분을 실행한다. 그렇기 때문에 높은 효율을 달성하는 CPU와 높은 throughput을 달성하는 GPU를 결합하면 application의 성능을 향상시킬 수 있다.
- CPU -> latency bound
- GPU -> throughput bound
Amdahl의 법칙은 application의 일부분만 parallelize할 수 있을 때 얻을 수 있는 최대 speedup을 정의할 때 사용된다.
CPU는 latency bound한 코드에 적합하며 GPU는 throughput bound한 코드에 적합하다.
optimization 시 CPU 코드, GPU 코드둘 중 한 부분만 빠르게 실행된다 하면 프로그램에 대한 성능 향상이 이루어지지 않는다. 최적으로 사용되도록 하기 위해선 두 프로세서에서 모두 성능 이득을 얻어야 한다.
이는 곧 최상의 성능을 얻기 위해서는 heterogeneous computing이 필요하다는 것에 근거가 된다.
즉, CPU에서 특정 유형의 작업을 GPU로 offload하는 접근을 heterogeneous computing이라고 한다.
Programming Paradigm
[Computer architecture의 분류]
| single 프로세서 | vector 프로세서 | |
| single instruction | SISD | SIMD |
| multiple instruction | MIMD |
SISD - Single Instruction Single Data(단일 명령 단일 데이터)
ex) single core CPU와 같이 하나의 코어를 갖는 직렬 처리 연산 장치로, 하나의 명령어를 하나의 데이터에 대해 수행하는 가장 간단한 구조
MISD - Multi Instruction Single Data (복수 명령 단일 데이터)
개념적으로만 기술되고, 아직 실현되지 않은 컴퓨터 아키텍쳐이다.
여러 개의 명령어를 하나의 데이터에 동시에 수행하는 구조
SIMD - Single Instruction Multiple Data (단일 명령 복수 데이터)
동일한 제어 명령으로 여러 개의 코어를 제어하는 구조이다.
SISD와의 비교 → 크기가 16인 배열 x와 y의 원소를 더함
SISD → 16개의 원소를 하나씩 처리 → 16번의 반복이 필요
data level parallelism = 동일한 연산을 동시에 여러 데이터에 적용함
SIMD → 4개의 원소를 한번에 처리 → 4번의 반복이 필요
데이터 배열에 연산이 적용된다는 의미에서 vector processor 혹은 array processor라고도 한다.
MIMD - Multiple Instruction Multiple Data (복수 명령 복수 데이터)
명령어와 데이터는 1:1로 연결되며 다수의 SISD가 하나의 칩 안에 들어 있는 구조라 보면 됨.
MIMD는 여러 개의 연산 유닛을 가지므로, 여러 스레드가 병렬로 동시에 작업을 수행하게 된다.
task level parallelism = 각 연산 유닛이 서로 다른 일을 한다. → 각 스레드에 독립된 작업을 분배한다.
ex) multi core CPU, 여러 독립된 프로세서의 집합 = 각 프로세서가 자신만의 control unit과 context를 가짐
여러 개의 명령어를 여러 개의 데이터에 적용하는 컴퓨팅 아키텍쳐
Low latency vs Higher throughput
CPU는 다양한 종류의 cache를 갖추고 있어 cache된 dataset에 대한 low latency를 위해 설계되었다.
많은 수의 트랜지스터가 speculative execution과 out of order execution을 위해 사용되는데, CPU는 높은 clock rate에서 실행되기 때문에 빈번하게 사용되는 data를 cache에 저장하고 다음에 실행할 instruction을 예측하여 latency를 hide를 한다. (if-else문이 없는 application과 같이 instruction pipeline을 쉽게 채울 수 있는 application을 사용해서 latency를 hiding할 수 있다) 결론적으로 CPU는 latency reducing architecture이다.
GPU는 data parallelsim에 최적화되어있으며 latency reducing 혹은 high througphut 아키텍쳐이다.
GPU는 CPU와 달리, 다른 thread의 연산으로 latency를 hide한다. 한 thread가 계산을 하기 위해 data를 기다리는 동안 다른 thread는 execution을 시작할 수 있어, clock cycle 낭비를 줄일 수 있다.
GPU는 ALU에 많은 transistor를 할당하는 반면, CPU는 latency를 줄이기 위해 ALU를 사용한다.
GPU에는 많은 register가 있으며 모든 thread의 context switching 정보는 register에 저장된다.
CPU register는 제한되어 있어, thread와 관련된 정보는 cache와 같은 low memory hierarchy에 저장된다.
따라서 CPU에서 thread 간 context switching 에 걸리는 시간은 GPU에 비해 더 오래 걸린다.
Programming approaches to GPU
CUDA는 NVIDIA에서 개발한 parallel 컴퓨팅 플랫폼 및 프로그래밍 모델 아키텍쳐이며 CPU와 GPU 모두를 고려한 heterogeneous programming model이다.
GPU 아키텍쳐는 기존 라이브러리를 활용해 accelerate하는 방식, openACC directive를 사용하는 방식, C/C++, Python, Fortran 등의 언어를 활용하는 방식을 통해 구현이 가능하다.
CUDA에서는 두 프로세서가 interact한다.
host ≒ CPU로, device function을 호출한다. CPU에서 실행되는 코드는 host code이다.
device ≒ GPU이다.
GPU Architecture
CUDA가 인기를 얻게 된 이유 중 하나는 하드웨어와 소프트웨어가 최상의 성능을 얻기 위해 설계되고 긴밀하게 결합되어 있기 때문이다. 그렇기 때문에 소프트웨어를 통한 CUDA 프로그래밍과 하드웨어 설계 간의 관계를 이해하는 것이 중요하다.
| software | Executes on / as | hardware |
| CUDA thread | CUDA core / SIMD code | |
| CUDA block | streaming multiprocessor | |
| CUDA grid | GPU device |
CUDA thread : CUDA thread는 CPU thread와 다르게 대규모 register 크기와 하드웨어 기반 scheduler (thread scheduler가 register에 존재)를 통해 빠른 context switching을 수행한다. 각 CUDA thread는 동일한 kernel을 실행하고 서로 다른 data에 독립적으로 작동해야 한다.(SIMT)
CUDA block : CUDA thread의 그룹이다. CUDA block은 single streaming multiprocessor(SM)에서 실행된다. 즉 block 내 모든 thread들은 하나의 SM core에서만 실행되고 다른 SM core에서 실행되지 않는다.
CUDA grid : CUDA block의 그룹이다. device에서 실행된다.
Vector addition using CUDA
CUDA를 사용하여 두 vector 간 덧셈연산을 하는 과정은 다음과 같다.
1. host memory 할당 -----> malloc()
2. device memory 할당 -----> cudaMalloc()
3. host memory에서 device memory로 data 전송 --> cudaMemcpy()
4. CUDA function 호출 및 실행 -----> <<< # blocks , # threads >>>
5. synchronize -----> cudaDeviceSynchronous()
6. device memory에서 host memory로 data 전송 --> cudaMemcpy()
7. host memory 해제 -----> free()
8. device memory 해제 -----> cuda
* CUDA에서 모든 kernel 호출은 asynchronous하다. host는 kernel을 호출한 후 free상태가 되어 이후 다음 instruction을 실행한다. host가 device function의 완료를 기다려야 하는 경우, host 코드가 device function의 완료를 기다리도록 하는 API가 필요한데, 그 중 하나가 바로 cudaDeviceSynchronous()이다. 이는 이전에 시작된 모든 device call이 완료될 때까지 기다리게 하는 함수이다.
case 1) creating multiple blocks -> use of several blocks with one thread
__global__ void device_add(int *a, int *b, int *c)
{
c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];
}
...
device_add<<< N, 1 >>> // <<<블록 갯수 , 스레드 갯수>>>위 예제에서는 CUDA block을 사용해 벡터 덧셈을 병렬로 실행한다. 특히 blockIDx.x를 사용해서 각 블록이 다른 element들을 연산하도록 하는데, 전체적으로 각 블록은 병렬로 실행된다.
device_add함수는 N번 호출되어 실행되는 데, 이 때 블록을 통해 병렬로 호출된다.
case 2) creating multiple threads -> use of one block with several threads
__global__ void device_add(int *a, int *b, int *c)
{
c[threadIdx.x] = a[threadIdx.x] + b[threadIdx.x];
}
....
device_add<<< 1, N >>>위 예제에서는 CUDA thread을 사용해 벡터 덧셈을 병렬로 실행한다. 특히 threadIDx.x를 사용해서 각 스레드가 다른 element들을 연산하도록 하는데, 전체적으로 각 스레드는 병렬로 실행된다.
device_add함수는 N번 호출되어 실행되는 데, 이 때 스레드를 통해 병렬로 호출된다.
case3) multiple blocks with multiple threads
key point는 바로 index를 찾는 것 -> threadIdxd와 blockIdx를 할당해서 unique한 ID를 부여해 주어야 함
case 3-1) combining blocks and threads -> 4 blocks with 8 threads
__global__ void device_add(int *a, int *b, int *c)
{
int index = blockIdx.x * blockDim.x + threadIdx.x ;
c[index] = a[index] + b[index];
}
....
device_add<<< 4, 8 >>>
case 3-2) combining blocks and threads -> 8 blocks with 4 threads
__global__ void device_add(int *a, int *b, int *c)
{
int index = blockIdx.x * blockDim.x + threadIdx.x ;
c[index] = a[index] + b[index];
}
....
device_add<<< 8, 4 >>>
thread와 block의 조합으로 thread의 unique한 ID를 계산할 수 있다.
이 때 blockDim이라는 변수가 모든 thread에 주어지는데, 이 변수는 block의 dimension, 즉 block 당 thread 갯수이다.
* thread와 block 이라는 hierarchy가 왜 필요한가?
- CUDA 프로그래밍 모델이 설정해 놓은 제약 때문이다.
- CUDA 유저는 global indexing을 올바르게 수행하기 위해 block과 grid 크기를 찾아야 한다.
block과는 다르게 하나의 block 내 thread들 효율적으로 communicate하고 synchronize하는 매커니즘을 갖고 있다.
다른 block에 속한 thread는 kernel 실행 중에 서로 communicate하고 synchronize하는 것이 불가능하다.
이러한 제약을 통해 scheduler가 각각의 SM에서 block들을 독립적으로 스케줄링할 수 있게 되는 것이다.
thread는 서로 communicate하기 위해 shared memory를 사용한다.
[Full code]
/* Device add using Global memory */
#include "stdio.h"
#include "stdlib.h"
#define N 512
__global__ void device_add(int *a, int *b, int *c){
c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];
}
void fill_array(int *data){
for(int idx = 0 ; idx < N ; idx++)
{
data[idx] = idx;
}
}
void print_output(int *a, int *b, int *c){
for(int idx = 0 ; idx < N ; idx++)
printf("%d + %d = %d\n", a[idx], b[idx], c[idx]);
}
int main(void)
{
int *a, *b, *c;
int *d_a, *d_b, *d_c;
int size = N * sizeof(int);
a = (int *)malloc(size);
fill_array(a);
b = (int *)malloc(size);
fill_array(b);
c = (int * )malloc(size);
cudaMalloc((void *) &d_a, N * sizeof(int));
cudaMalloc((void *) &d_b, N * sizeof(int));
cudaMalloc((void *) &d_c, N * sizeof(int));
cudaMemcpy(d_a, a, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, b, cudaMemcpyHostToDevice);
device_add<<<a, b, c>>>(d_a, d_b, d_c);
cudaMemcpy(c, d_c, cudaMemcpyDeviceToHost);
printf_output(a,b,c);
free(a);
free(b);
free(c);
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
return 0;
}위 코드는 GPU Memory hierarchy 중 Global memory를 이용하는 코드이다.
Launching kernels in multiple dimensions
thread와 block을 1차원, 2차원, 3차원으로 실행할 수 있다.
threadIdx.x -> 1차원에서 kernel 실행
예를 들어 image에서 parallel task를 수행할 때 2차원에서 thread와 block을 실행하게 된다.
Data type support in CUDA
CUDA는 char, float, double, long 등 일반적으로 c/c++에서 지원하는 data type 뿐만 아니라 float2, float4와 같은 vector type도 지원한다.
data type은 naturally align 된 data type을 사용하는 것이 권장되며 크기가 1,2,4,8,16 byte인 data type의 경우 GPU가 single memory instruction을 호출하도록 한다. 만약 data type이 align 되어 있지 않다면 compiler는 여러 instruction을 생성하게 되고, 이는 memory와 memory bus를 비효율적으로 활용하게 된다.
또한 CUDA는 structure와 class 등 복잡한 data structure를 지원한다.
'CUDA Programming' 카테고리의 다른 글
| [CUDA] ch3. CUDA Thread Programming (0) | 2024.01.16 |
|---|---|
| [CUDA] ch2. CUDA Memory Management - Shared memory & Texture memory & Registers (0) | 2023.12.18 |
| [CUDA] ch2. CUDA Memory Management - Global memory (2) | 2023.12.18 |
