[CUDA] GPGPU 그리고 CUDA 프로그래밍 기초

GPGPU  - General purpose comupting on Grapthics processing units(그래픽 프로세서에서 general purpose 

computing 가능하다)

[ CUDA란? ]

CUDA는 Nvidia가 만든 병렬 컴퓨팅 플랫폼 및 API 모델이다. 소프트웨어 개발자들은 CUDA를 통해 General한 Purpose를 위한 GPU사용이 가능해졌는데, 이러한 것을 GPGPU(General-Purpose Computing on Graphics Processing Units)라고 한다. CUDA 플랫폼은 GPU의 가상 명령어셋과 병령 처리 요소들을 사용할 수 있도록 만들어주는 소프트웨어 레이어이며, NVIDIA가 만든 CUDA 코어가 장착된 GPU에서 작동한다.

2006년 11월, G80 아키텍처와 함께 처음 발표되었으며, 초기에는 C,C++만 지원했지만 지금은 Fortran이나 C# 등 다양한 언어에서 사용아 가능하며 현재는 10.0 버전까지 나왔다. 또한 OpenACC나 OpenCL과 같은 프레임워크도 지원을 하고 있다. 이전의 Direct3D와 OpenGL을 활용한 Graphics Programming은 상당히 진보된 기술을 요구하였지만, CUDA는 보다 쉽게 GPU의 자원을 사용할 수 있다.

 

[ 용어 정리 ]

• Host: CPU and its memory(Host Memory)

• Device: GPU and its memory(Device Memory)

[ 데이터 흐름 ]

CUDA를 이용해 GPU 프로그래밍을 하면 처리되는 과정은 다음과 같다.

1. Data가 CPU에서 GPU로 복사된다
2. GPU에서 커널 함수를 실행하여 처리한다.
3. 결과를 GPU에서 CPU로 복사한다.

기본적으로 GPGPU HW구조는 다음과 같다.

----------------------------

CPU   |||||| GPU

DRAM <-> GDRAM (여기서 memory copy가 일어남..)

이제 이러한 개념을 바탕으로 예제를 작성해보도록 하자.

---

2. Hello World with CUDA 

---

[ HelloWorld ]

일반적으로 우리가 C언어를 시작하면 아래와 같이 HelloWorld를 출력하는 예제로 시작을 한다.

int main( void ) {
	printf( "Hello, World!\n" );
	return 0;
}

 

[ HelloWorld with CUDA ]

위와 같은 기초적인 C언어 코드에 GPU에서 실행되는 디바이스 코드를 추가하면 다음과 같다.

// __global__ 키워드를 붙이면 Device에서 작동된다.
__global__ void kernel( void ) {

}

int main( void ) {
	kernel<<<1,1>>>();
	printf( "Hello, World!\n" );
}

 

디바이스 코드는 nvcc(NVIDIA's Compiler)가 컴파일을 하는데, 디바이스 코드를 위해서는 몇가지 키워드가 필요하다. __global__ 키워드는 Host를 통해 호출되어 Device에서 작동되는 함수를 가리키며, nvcc는 source file을 host components와 device components로 나눈다. nvcc는 kernel()과 같은 device component에 해당하는 함수를 처리하고, host의 compiler는 main()과 같은 host component에 해당하는 함수를 처리한다.

 

예를 들어 위와 같은 test.cu 파일을 컴파일하여 실행한다고 하면 아래와 같은 명령어로 컴파일하고 실행해주어야 한다.

// cuda 파일은 확장자가 cu이며 nvcc로 컴파일해주어야 한다.
nvcc -o test test.cu

// 실행 방법은 동일하다
./test

 

<<< BlockNum, ThreadNum >>> 는 host를 통해 device code를 호출함을 의미하며 이에 해당하는 파라미터에 대해서는 아래에서 자세히 알아보도록 하겠다.

 

 

[ Block and Thread ]

Block과 Thread는 GPU(디바이스) 코드를 병렬로 처리하기 위한 단위로, 1개의 블록은 N개의 쓰레드로 구성될 수 있다. 블록의 개수를 조절하는 것은 1번째 파리미터이며, 쓰레드의 개수를 조절하는 것은 2번째 파라미터이다.

// add 함수를 1번 실행
add<<< 1, 1 >>>( dev_a, dev_b, dev_c );

// N개의 블록으로 1개의 스레드를 통해 실행(Parallel Block)
add<<< N, 1 >>>( dev_a, dev_b, dev_c );

// 1개의 블록으로 N개의 스레드를 통해 실행(Parallel Thread)
add<<< 1, N >>>( dev_a, dev_b, dev_c );

 

 

 

예를 들어 4개의 블록에 8개의 쓰레드를 사용하면 아래와 같은 구조가 된다.

blockIdx.x를 사용하면 블록의 인덱스로 접근하여 서로 다른 블록을 처리할 수 있고, threadIdx.x를 통해서 스레드의 인덱스로 접근할 수 있다.

 

만약 1개의 블록당 M개의 쓰레드가 있는 경우라면 각 블록 안의 쓰레드의 unique index는 아래와 같이 계산할 수 있다. 여기서 blockDim.x는 각 블럭이 몇 개의 쓰레드를 갖는지를 나타내는 변수이다.

// 고유한 인덱스 = 블록 안에서의 쓰레드 번호 + 블록 번호 * 블록당 쓰레드의 개수
int index = threadIdx.x + blockIdx.x * blockDim.x(= M);

 

예를 들어 3번째 블록의 45번째 쓰레드의 index는 5 + 2 * 8 즉, 21이 된다.

 

이제 위와 같은 개념을 활용해 다음 포스팅에서 조금 더 복잡한 예제를 보도록 하자.

 

 

---

일단 한국 산업구조상 , High performace computing을 필요로 하는 분야는 별로없다.

이미 Powerful 한 performace 가 필요하다면 cloud에서 여러개를 Launching 하면 , 비싼 Engineer 없이 해결 가능한 부분도 없으므로 많이 필요로 하는 분야는 Edge Device만 남게 된다.

 

핵심 !! HOST MEMORY 와 DEVICE MEMORY 의 cudaMemcpy를 최소화해야함..

->이를 위한 syntax들을 제공하는데, cudaStream 을 알고있으면, 커널들을 task로 생각하며 프로그래밍 가능-> GPU 프로그래밍은 기본적으로 SIMD(Single Instruction Multiple Data) 컨셉을 가진다 

WARP <- 1. warp 16 cuda threads

하나의 WARP, KERNEL

SIMD 

if()

{

////

}

else

{

////

}

 

branch사용가능하게됨

 

이 경우는 크게 Medical Device, 자율주행 , 방위산업분야를 예로 들수 있다.