CUDA C / C++은 대규모 병렬 어플리케이션을 작성하는 가장 강력한 방법중 하나입니다.
강력한 C++ 프로그래밍 언어를 사용하여 GPU에서 실행되는 수천 개의 병렬 스레드로 가속화되는 고성능
알고리즘의 어플리케이션을 개발할 수 있습니다.
딥 러닝으로 알려진 인공 지능의 지속적인 혁신을 뒷받침하는 라이브러리와 프레임워크를 포함하여
많은 개발자가 이러한 방식으로 계산 및 대역폭을 많이 사용하는 응용 프로그램을 가속화했습니다.
CUDA가 지원되는 GPU를 가지고 있는 컴퓨터 또는 클라우드의 GPU 인스턴스가 필요합니다.
CUDA Toolkit이 설치되어 있는 상태여야 합니다. (GPU 클라우드는 Jupyter notebook 인터페이스를 쓰겠네요...)
자 이제 시작해 봅시다!!
Starting Simple
아래의 샘플 코드는 각각 백만개 원소를 가진 두개의 배열을 더하는 간단한 코드입니다.
#include <iostream>
#include <math.h>
// function to add the elements of two arrays
void add(int n, float *x, float *y)
{
for (int i = 0; i < n; i++)
y[i] = x[i] + y[i];
}
int main(void)
{
int N = 1<<20; // 1M elements
float *x = new float[N];
float *y = new float[N];
// initialize x and y arrays on the host
for (int i = 0; i < N; i++) {
x[i] = 1.0f;
y[i] = 2.0f;
}
// Run kernel on 1M elements on the CPU
add(N, x, y);
// Check for errors (all values should be 3.0f)
float maxError = 0.0f;
for (int i = 0; i < N; i++)
maxError = fmax(maxError, fabs(y[i]-3.0f));
std::cout << "Max error: " << maxError << std::endl;
// Free memory
delete [] x;
delete [] y;
return 0;
}
일단, 이 C++ 코드를 컴파일 하고, 실행해 봅시다.
코드를 add.cpp 로 저장하고 C++ 컴파일러로 컴파일 합니다.
$g++ add.cpp -o add
Linux 에서는 g++ 또는 cc
Windows 에서는 msvc.exe
Mac 에서는 clang++ 를 사용합니다.
그리고, 아래와 같이 실행 합니다.
$./add
Max error:0.000000
예상했던 것과는 같게...
합하는 동안의 오류가 0.00000 으로 출력되고 종료합니다.
그럼, 이제는 백만번의 더하기를 병렬로 해봅시다. (매우 쉬움)
먼저 더하기 하는 함수를 GPU가 실행할 수 있는 CUDA의 커널이라는 기능으로 바꿔야 합니다.
이렇게 하려면 함수에 지정자 __global__을 추가하기만 하면 됩니다.
이렇게 함으로써 CUDA C 컴파일러에게 이것이 GPU에서 실행되고 CPU 코드에서 호출될 수 있는
함수임을 알려줍니다.
// CUDA Kernel function to add the elements of two arrays on the GPU
__global__
void add(int n, float *x, float *y)
{
for (int i = 0; i < n; i++)
y[i] = x[i] + y[i];
}
이렇게 global 지정자가 붙어 있는 함수를 커널이라고 한다.
GPU에서 실행되는 코드는 디바이스 코드라 하고
CPU에서 실행되는 코드는 호스트 코드라고 하여 구분합니다.
Memory Allocation in CUDA
GPU에서 계산하려면 GPU에서 액세스할 수 있는 메모리를 할당해야 합니다.
CUDA의 통합 메모리는 시스템의 모든 GPU와 CPU에서 액세스할 수 있는
단일 메모리 공간을 제공하여 이를 쉽게 만듭니다.
통합 메모리에 데이터를 할당하려면 호스트(CPU) 코드 또는 디바이스(GPU) 코드에서
액세스할 수 있는 포인터를 반환하는 cudaMallocManaged()를 호출합니다.
데이터를 해제하려면 cudaFree()를 호출하면 됩니다.
위의 코드에서는
new -> cudaMallocManaged()
delete [] -> cudaFree()
로 교체하면 됩니다.
// Allocate Unified Memory -- accessible from CPU or GPU
float *x, *y;
cudaMallocManaged(&x, N*sizeof(float));
cudaMallocManaged(&y, N*sizeof(float));
...
// Free memory
cudaFree(x);
cudaFree(y);
마지막으로 일반 함수에서 커널로 바꾼 add() 커널을 시작하도록 바꿔보자!
CUDA 커널의 실행에는 삼중꺽쇠 사이에 <<< >>> 실행 설정 파라메터를 사용하여 지정됩니다.
이 삼중꺽쇠 구문을 원래 함수의 인자들 목록 앞에 추가해 주기만 하면 됩니다.
add<<<1,1>>>(N, x, y);
참 쉽죠잉~!
삼중 꺽쇠사이의 파라메터에 대해 설명하기 전에
이 호출은 한개의 GPU 스레드로 실행된다는 것만 알고 넘어 갑시다!
※ One More Thing
당연한 이야기지만 결과에 액세스하기 전에 커널의 실행이 완료될 때까지 CPU는 기다려야 합니다.
(CUDA 커널 실행은 호출하는 CPU 스레드가 멈춰서 기다리지 않기 때문에...)
이렇게 기다리려면 CPU에서 최종 오류 검사를 수행하기 전에 cudaDeviceSynchronize()를 호출하면 됩니다.
최종 수정된 코드는 아래와 같습니다.
#include <iostream>
#include <math.h>
// Kernel function to add the elements of two arrays
__global__
void add(int n, float *x, float *y)
{
for (int i = 0; i < n; i++)
y[i] = x[i] + y[i];
}
int main(void)
{
int N = 1<<20;
float *x, *y;
// Allocate Unified Memory – accessible from CPU or GPU
cudaMallocManaged(&x, N*sizeof(float));
cudaMallocManaged(&y, N*sizeof(float));
// initialize x and y arrays on the host
for (int i = 0; i < N; i++) {
x[i] = 1.0f;
y[i] = 2.0f;
}
// Run kernel on 1M elements on the GPU
add<<<1, 1>>>(N, x, y);
// Wait for GPU to finish before accessing on host
cudaDeviceSynchronize();
// Check for errors (all values should be 3.0f)
float maxError = 0.0f;
for (int i = 0; i < N; i++)
maxError = fmax(maxError, fabs(y[i]-3.0f));
std::cout << "Max error: " << maxError << std::endl;
// Free memory
cudaFree(x);
cudaFree(y);
return 0;
}
이 코드는 CUDA C 파일이라는 표시로 add.cu로 저장합니다.
컴파일은 CUDA C++ 컴파일러인 nvcc 로 아래와 같이 컴파일 합니다.
$ nvcc add.cu -o add_cuda
$./add_cuda
Max error:0.000000
이것은 단순화된 단일 스레드에 대한 첫번째 예제일 뿐이며, 단일 스레드에서만 정확 합니다.
이코드는 실행되는 모든 스레드에서 전체 배열에 대한 더하기를 진행하기 때문이죠.
또한, 이 코드를 여러 스레드로 실행했을경우, 동일한 위치를 읽고 쓸수 있기 때문에
race condition(스레드간 경쟁)상태가 될수 있습니다.
Note: Windows의 경우 빌드시 프로젝트 속성의 구성을 플랫폼을 x64로 지정해야 합니다.
Profile it!
커널실행에 걸리는 시간을 측정하는 간단한 방법은
CUDA Toolkit에 포함된 Commandline GPU Profiler인 nvprof로 커널을 실행하는 것입니다.
아래와 같이 nvprof ./add_cuda 로 실행합니다.
$ nvprof ./add_cuda
==3355== NVPROF is profiling process 3355, command:./add_cuda
Max error:0
==3355==Profiling application:./add_cuda
==3355==Profiling result:
Time(%) Time Calls Avg Min Max Name
100.00% 463.25ms 1 463.25ms 463.25ms 463.25ms add(int,float*,float*)
...
위의 결과는 단일 호출로 nvprof의 출력예 입니다.
NVIDIA Tesla K80 GPU 에서는 약 0.5초가 소요되고
3년 된 Macbook Pro의 NVIDIA GeForce GT 740M에서는 거의 같은 시간이 걸립니다.
이제, 병렬처리로 더욱 빠르게 만들어 보겠습니다.
Picking up the Threads
하나의 스레드로 동작하는 커널을 병렬로 만드는 방법입니다.
핵심은 CUDA의 <<<1, 1>>> 구문에 있습니다.
이것은 실행설정 파라메터로 GPU에서 실행하는 데 사용할 병렬 스레드 수를 CUDA 런타임에 알려줍니다.
여기에는 두 개의 매개변수가 있지만 두 번째 매개변수인 스레드 블록의 스레드 수를 변경하여 시작하겠습니다.
CUDA GPU는 크기가 32의 배수인 스레드 블록을 사용하여 커널을 실행하므로
256개의 스레드를 선택하는 것이 적절한 크기로 보입니다.
add<<<1,256>>>(N, x, y);
이 커널 실행코드를 바꾸는 것만으로 실행을 하게 되면, 이전에 설명 했듯이 병렬 스레드 간에 계산을 분산하지
않아서 모든 스레드가 전체 배열의 원소에 대해 한번씩 계산을 수행합니다.
CUDA C++는 커널이 실행 중인 스레드의 인덱스를 얻을 수 있도록 하는 키워드를 제공합니다.
특히, threadIdx.x는 블록 내 현재 스레드의 인덱스를 포함하고 blockDim.x 블록의 스레드 수를 포함합니다.
이 두개의 미리 정의된 매개변수에 의해 배열내의 요소들이 중복으로 연산 수행에 이용되지 않도록
커널의 루프부분을 수정합시다.
__global__
void add(int n, float *x, float *y)
{
int index = threadIdx.x;
int stride = blockDim.x;
for (int i = index; i < n; i += stride)
y[i] = x[i] + y[i];
}
add 함수의 기능은 바뀌지 않았습니다. (전체 배열내의 요소를 더하는...)
indedx = 0 으로 stride = 1 로 설정하면 첫번째 소스와 동일한 동작을 합니다.
이 코드를 add_block.cu 로 저장하여 nvcc로 컴파일 하고 nvprof에서 다시 실행합니다.
결과는 아래와 같습니다.
Time(%) Time Calls Avg Min Max Name
100.00% 2.7107ms 1 2.7107ms 2.7107ms 2.7107ms add(int,float*,float*)
이 결과는 엄청난 속도의 향상(463ms -> 2.7ms) 이지만, 1스레드에서 256스레드로 바꾸었기 때문에
그리 놀라운 일은 아닙니다.
NVIDIA Tesla K80 GPU 는 Macbook pro GPU(3.2ms) 보다 빠릅니다.
계속해서 더 빠른 퍼포먼스를 얻기 위한 방법을 알아 봅시다.
Out of the Blocks
CUDA GPU에는 Streaming Multiprocessors(SM)로 그룹화된 많은 병렬 프로세서가 있습니다.
각 SM은 여러 개의 동시 스레드 블록을 실행할 수 있습니다.
예를 들어 Pascal GPU 아키텍처 기반 Tesla P100 GPU에는
각각 최대 2048개의 활성 스레드를 지원할 수 있는 56개의 SM이 있습니다.
이 모든 스레드를 최대한 활용하려면 여러 스레드 블록으로 커널을 시작해야 합니다.
이제 실행 설정 파라메터의 첫번째 매개변수가 스레드의 블록의 수를 지정한다는 것을 짐작 하겠죠?
또한, 병렬 스레드 블록이 그리드로 알려진 것을 구성합니다.
처리할 N개의 요소와 블록당 256개의 스레드가 있으므로 최소한 N개의 스레드를 얻으려면 블록 수를 계산하면 됩니다.
편하게 하기 위해 단순히 N을 블록 크기로 나눕니다(N이 blockSize의 배수가 아닌 경우 반올림에 주의)
int blockSize =256;
int numBlocks =(N + blockSize -1)/ blockSize;
add<<<numBlocks, blockSize>>>(N, x, y);
또한 스레드 블록의 전체 그리드를 고려하도록 커널 코드를 업데이트해야 합니다.
CUDA는 그리드의 블록 수를 포함하는 gridDim.x와
그리드에 있는 현재 스레드 블록의 인덱스를 포함하는 blockIdx.x를 제공합니다.
그림1 은 blockDim.x, gridDim.x 및 threadIdx.x를 사용하여 CUDA에서 배열(1차원)로
인덱싱하는 접근 방식을 보여줍니다.
아이디어는 각 스레드가 블록의 시작 부분에 대한 오프셋을 계산하고
(블록 인덱스 곱하기 블록 크기: blockIdx.x * blockDim.x) 블록 내에
스레드의 인덱스(threadIdx.x)를 더하여 인덱스를 얻는다는 것입니다.
코드 ockIdblx.x * blockDim.x + threadIdx.x는 관용적으로 쓰이는 CUDA 코드 입니다.
__global__
void add(int n, float *x, float *y)
{
int index = blockIdx.x * blockDim.x + threadIdx.x;
int stride = blockDim.x * gridDim.x;
for (int i = index; i < n; i += stride)
y[i] = x[i] + y[i];
}
이렇게 업데이트된 커널은 stride를 그리드의 총 스레드 수로 설정합니다. (blockDim.x * gridDim.x)
CUDA에서 이러한 유형의 루프는 grid-stride loop 라고 합니다.
파일을 add_grid.cu로 저장하고 nvcc로 컴파일하고 nvprof로 실행합니다.
Time(%) Time Calls Avg Min Max Name
100.00% 94.015us 1 94.015us 94.015us 94.015us add(int,float*,float*)
8배 빨라진 결과를 보여줍니다~!
Summing Up
다음은 Tesla K80 및 GeForce GT 750M에서 세 가지 버전의 add() 커널 성능에 대한 요약입니다.
위의 표에서 보는것과 같이 GPU에서 매우 높은 대역폭을 얻을 수 있습니다.
이 예제에서는 아주 기본적이면서 제한적인 대역폭만을 사용했지만,
GPU는 고밀도행렬, 선형대수, 딥러닝, 이미지 프로세싱, 신호처리, 물리 시뮬레이션 등과 같이
컴퓨팅 파워를 많이 사용하는 계산에서도 탁월합니다.
NVIDIA 개발자 싸이트나 싸이트내의 NVIDIA 개발자 블로그에는
CUDA C++ 및 기타 GPU 컴퓨팅 주제에 대한 다양한 콘텐츠가 있으니 더 깊이 있는 학습을 위해
이 싸이트들을 이용하기를 권장 드립니다.