Applicaton CUDA의 가속 절차
-
성능 분석
-
병목 탐색
-
알고리즘 분석 및 병렬화 전략 수립
-
데이터 의존성 해결
-
병렬화
-
최적화
NVIDIA 에서 진행된 CUDA 강의를 듣고 정리한 포스트입니다.
NVCC Compiler
NVCC Compiler : GPU Code, CPU Code를 인식, CPU Code는 다른 컴파일러로 넘깁니다.
NVCC Library Partion을 만들고 기존 project에서 확장하는 것이 가능합니다.
Graphic Card Device
CUDA 지원 그래픽카드 : 모든 GPU ( Pascal, Mawell, Kepler, Fermi ... )
DP, 대용량 연산에는 TESLA가 필수적입니다.
Kernal
GPU 코드를 Kernal이라고 합니다.
__global__ void func(...)
{
...
}__global__ 을 이용해서 Kernal임을 명시합니다.
Thread
Thread는 연산 단위 입니다.
CUDA Thread : 순차적으로 연산하는 Thread (index 최대 3차원)
Thread Block : 연산 제어 위한 block (index 최대 3차원)
Grid : Kernal (index 없음)
Warp : CUDA의 최소 제어 Thread묶음 (32개 고정)
계층 구조
-----------------------------------GRID--------------------------------------
-----------ThreadBlock---------------------ThreadBlock------------
------Thread--------Thread-----------Thread-------Thread--------
학습목표
강좌를 마치면 여러분은 다음 사항을 수행할 수 있게 될 것입니다.
- CPU 함수를 호출하고 GPU 커널(kernel)을 구동(launch)하는 C/C++ 프로그램을 작성, 컴파일, 실행하기
- 실행 설정(execution configuration)을 통해 병렬 스레드 계층(thread hierarchy) 제어하기
- 직렬로 실행되는 반복문을 GPU 상에서 병렬로 실행되도록 수정하기
- CPU와 GPU를 위한 메모리를 할당하고 해제하기
- CUDA 코드에서 생성된 오류 처리하기
- CPU-only 애플리케이션 가속화하기
GPU를 위한 애플리케이션 코드 작성하기
CUDA는 다양한 범용 프로그래밍 언어 확장을 지원합니다. 본 강좌에서는 그 중에서 C/C++을 사용합니다. 개발자들은 언어 확장을 통해 소스 코드를 GPU 상에서 손쉽게 실행할 수 있게 됩니다.
아래는 .cu 파일을 보여 줍니다. .cu는 CUDA 가속화 프로그램을 위한 파일 확장자입니다. 여기에서 두 개의 함수가 있는데 첫 번째는 CPU 상에서 실행되고 두 번째는 GPU 상에서 실행될 것입니다. 각각의 함수가 서로 어떻게 다른지, 정의되는 방식과 호출되는 방식에 유념하여 잘 살펴 보세요.
void CPUFunction()
{
printf("This function is defined to run on the CPU.\n");
}
__global__ void GPUFunction()
{
printf("This function is defined to run on the GPU.\n");
}
int main()
{
CPUFunction();
GPUFunction<<<1, 1>>>();
cudaDeviceSynchronize();
}아래에 코드의 중요 부분과 함께 가속화 컴퓨팅에서 사용하는 주요 용어들을 정리했습니다:
__global__ void GPUFunction()__global__키워드는 해당 함수가 GPU에서 실행되며 전역적으로 호출될 수 있음을 의미합니다. 지금 문맥에서 전역적이란 CPU나 GPU, 그 누구에 의해서든 호출될 수 있다는 뜻입니다.- CPU에서 수행되는 코드는 종종 호스트(host) 코드라고 불립니다. GPU에서 수행되는 코드는 디바이스(device) 코드라고 합니다.
- 리턴 타입은
void입니다.__global__키워드로 정의된 함수는 반드시void타입을 리턴해야 합니다.
GPUFunction<<<1, 1>>>();- GPU에서 실행되는 함수를 호출할 때, 통상적으로 이 함수를 커널(kernel)이라고 부르며 호출 대신에 구동(launch)이라는 표현을 씁니다.
- 커널을 구동할 때에는 실행 설정(execution configuration)을 해주어야 합니다. 이것은 커널에 인수를 전달하는 부분 직전에서
<<<...>>>꼴의 문법을 사용하여 이루어집니다. 실행 설정은 본 강좌에서 자세히 살펴 볼 것입니다. 지금은 일단 커널이 (<<<1, 1>>>의 첫 번째 값 1로부터 온) 1 개의 스레드 블록을 사용하고, 스레드 블록당 (<<<1, 1>>>의 두 번째 값 1로부터 온) 1 개의 스레드를 사용한다고만 이해하면 됩니다.
cudaDeviceSynchronize();- 대부분의 C/C++ 코드와는 달리 커널 구동은 비동기적(asynchronous)입니다. 즉, CPU 코드는 커널 구동이 완료되는 것을 기다리지 않고 다음 코드를 실행합니다.
- CUDA 런타임에서 제공되는
cudaDeviceSynchronize의 호출은 호스트(CPU) 코드로 하여금 디바이스(GPU) 코드의 완료를 기다리게 하며, 디바이스 코드가 완료되었을 때 비로소 CPU가 실행을 재개합니다.
hello.cu 연습
#include <stdio.h>
void helloCPU()
{
printf("Hello from the CPU.\n");
}
/*
* Refactor the `helloGPU` definition to be a kernel
* that can be launched on the GPU. Update its message
* to read "Hello from the GPU!"
*/
__global__ void helloGPU()
{
printf("Hello from the GPU!\n");
}
int main()
{
helloCPU();
/*
* Refactor this call to `helloGPU` so that it launches
* as a kernel on the GPU.
*/
helloGPU<<<1, 1>>>();
/*
* Add code below to synchronize on the completion of the
* `helloGPU` kernel completion before continuing the CPU
* thread.
*/
cudaDeviceSynchronize();
}
컴파일
!nvcc -arch=sm_70 -o hello-gpu 01-hello/01-hello-gpu.cu -run
결과
Hello from the CPU.
Hello from the GPU!실행 설정
<<<n, m>>> : n개의 ThreadBlock 아래에 동일한 m개의 Thread가 존재합니다.
->Kernal은 n x m 번 구동됩니다.
01-basic-parallel.cu
#include <stdio.h>
/*
* Refactor firstParallel so that it can run on the GPU.
*/
__global__ void firstParallel()
{
printf("This should be running in parallel.\n");
}
int main()
{
/*
* Refactor this call to firstParallel to execute in parallel
* on the GPU.
*/
firstParallel<<<5, 5>>>();
/*
* Some code is needed below so that the CPU will wait
* for the GPU kernels to complete before proceeding.
*/
cudaDeviceSynchronize();
}
컴파일
!nvcc -arch=sm_70 -o basic-parallel 02-first-parallel/01-basic-parallel.cu -run결과
This should be running in parallel.
This should be running in parallel.
This should be running in parallel.
This should be running in parallel.
This should be running in parallel.
This should be running in parallel.
This should be running in parallel.
This should be running in parallel.
This should be running in parallel.
This should be running in parallel.
This should be running in parallel.
This should be running in parallel.
This should be running in parallel.
This should be running in parallel.
This should be running in parallel.
This should be running in parallel.
This should be running in parallel.
This should be running in parallel.
This should be running in parallel.
This should be running in parallel.
This should be running in parallel.
This should be running in parallel.
This should be running in parallel.
This should be running in parallel.
This should be running in parallel.25회 출력되었다. (5 x 5)
스레드 인덱스와 블록 인덱스 (Thread and Block Indices)
스레드 블록 안에 있는 각 스레드에는 0부터 작하는 인덱스가 할당됩니다. 또한, 각 블록 역시 0부터 시작하는 인덱스를 가지고 있습니다. 스레드가 모여 스레드 블록이 구성되듯이 블록이 모여 그리드(grid)가 됩니다. 그리드는 CUDA 스레드 계층 구조에서 최상위 개체에 해당합니다. 요약하면, CUDA 커널은 하나 이상의 블록으로 이루어진 그리드에서 실행되고, 각각의 블록은 1 개 이상의 스레드로 구성되는데 블록당 스레드의 개수는 동일합니다.
CUDA 커널은 (블록 안에서의) 스레드 인덱스를 나타내는 변수와 (그리드 안에서의) 블록 인덱스를 나타내는 변수를 가지고 있습니다. 이들을 각각 threadIdx.x와 blockIdx.x라고 합니다.
참고) 모든 index는 0부터 시작합니다.
01-thread-and-block-idx.cu
#include <stdio.h>
void helloCPU()
{
printf("Hello from the CPU.\n");
}
/*
* Refactor the `helloGPU` definition to be a kernel
* that can be launched on the GPU. Update its message
* to read "Hello from the GPU!"
*/
void helloGPU()
{
printf("Hello from the GPU!\n");
}
int main()
{
helloCPU();
/*
* Refactor this call to `helloGPU` so that it launches
* as a kernel on the GPU.
*/
helloGPU<<<1, 1>>>();
/*
* Add code below to synchronize on the completion of the
* `helloGPU` kernel completion before continuing the CPU
* thread.
*/
cudaDeviceSynchronize();
}컴파일
!nvcc -arch=sm_70 -o thread-and-block-idx 03-indices/01-thread-and-block-idx.cu -run결과
Success!병렬화 수준을 높이기 위해 블록 차원 사용하기 (Block Dimensions)
하나의 스레드 블록에는 최대 1024 개까지의 스레드만 들어갈 수 있습니다. 가속화 애플리케이션의 병렬화 수준을 높이기 위해서는 다수의 스레드 블록이 서로 협력하도록 해야 합니다.
CUDA 커널은 하나의 블록에 들어있는 스레드의 수를 저장하는 blockDim.x라는 특별한 변수를 가지고 있습니다. 이 변수를 blockIdx.x 및 threadIdx.x와 함께 사용하여 블록과 스레드의 병렬 수행을 조직화함으로써 병렬화 수준을 높일 수 있는데, 이때 threadIdx.x + blockIdx.x * blockDim.x와 같은 수식을 사용합니다. 상세한 예는 다음과 같습니다. (역주: 이는 처리할 데이터를 1 차원으로 보았을 때의 예입니다. 2 차원의 영상 데이터 등을 처리한다고 할 때에는 threadIdx.x + blockIdx.x * blockDim.x 뿐 아니라 threadIdx.y + blockIdx.y * blockDim.y도 함께 사용됩니다.)
실행 설정 <<<10, 10>>>은 블록 당 10 개의 스레드(역주: blockDim.x는 10)를 가진 블록 10 개로 이루어진, 총 100 개의 스레드로 구성된 그리드를 구동합니다. 따라서 0부터 99까지의 인덱스를 가진 값들 중에서 고유한 하나를 스레드에게 주어 계산시킬 수 있을 것입니다. (역주: 즉, 서로 다른 스레드가 같은 인덱스의 값을 계산하면 중복이 되므로 중복이 되지 않도록 각 스레드에게 일감을 배분한다는 뜻입니다.)
-
블록의
blockIdx.x가 0이면blockIdx.x * blockDim.x는 0이 됩니다. 이 0을 가능한threadIdx.x값인 0부터 9의 값에 더하면 100 개의 스레드 그리드 중 0부터 9까지의 인덱스를 얻을 수 있습니다. -
블록의
blockIdx.x가 1이면blockIdx.x * blockDim.x는 10이 됩니다. 이 10을 가능한threadIdx.x값인 0부터 9의 값에 더하면 100 개의 스레드 그리드 중 10부터 19까지의 인덱스를 얻을 수 있습니다. -
블록의
blockIdx.x가 5이면blockIdx.x * blockDim.x는 50이 됩니다. 이 50을 가능한threadIdx.x값인 0부터 9의 값에 더하면 100 개의 스레드 그리드 중 50부터 59까지의 인덱스를 얻을 수 있습니다. -
블록의
blockIdx.x가 9이면blockIdx.x * blockDim.x는 90이 됩니다. 이 90을 가능한threadIdx.x값인 0부터 9의 값에 더하면 100 개의 스레드 그리드 중 90부터 99까지의 인덱스를 얻을 수 있습니다.
blockDim 은 한 블럭에서의 Thread의 개수
blockDim.x : 한 블럭의 x축 방향의 Thread의 개수
kernal이 어느 unit에서 실행되고 있는지 알아내기 위한 코드
int idx = blockDim.x * blockIdx.x + threadIdx.x두 벡터 x, y의 합성 벡터 c를 구하는 소스코드 (iteration 방식)
for ( i=0; i<N; i++)
{
c[i] = x[i] + y[i];
}의 CUDA 병렬 처리 방식
int i = blockDim.x * blockIdx.x + threadIdx.x;
C[i] = X[i] + Y[i];위와 같은 방법으로 loop를 한 번 없애는 것이 가능합니다.
For 반복문 가속하기 (Accelerating For Loops)
CPU-only 애플리케이션의 for 반복문은 가속화하기에 매우 적합한 대상입니다. 반복문으로 각 반복을 순차적으로 실행하는 대신에, 각 반복을 고유의 스레드에서 병렬로 실행되도록 하면 됩니다. 아래 반복문을 보면 반복이 몇 번 실행되는지와 각 반복에서 어떤 일이 일어나는가를 쉽게 알 수 있습니다.
int N = 2<<20;
for (int i = 0; i < N; ++i)
{
printf("%d\n", i);
}이 반복문을 병렬화하기 위해서 두 단계의 과정이 필요합니다.
- 커널은 반복 중 1 회만 수행하도록 작성합니다.
- 커널은 다른 커널의 동작 여부를 알 수 없으므로 정확한 횟수로 커널이 수행되도록 실행 설정을 해야 합니다. 즉, 반복문의 반복 횟수만큼 커널이 수행되어야 합니다.
01-single-block-loop.cu
#include <stdio.h>
/*
* Refactor `loop` to be a CUDA Kernel. The new kernel should
* only do the work of 1 iteration of the original loop.
*/
__global__ void loop()
{
int idx = blockDim.x * blockIdx.x + threadIdx.x;
printf("This is iteration number %d\n", idx);
}
int main()
{
/*
* When refactoring `loop` to launch as a kernel, be sure
* to use the execution configuration to control how many
* "iterations" to perform.
*
* For this exercise, only use 1 block of threads.
*/
int N = 10;
loop<<<1, N>>>();
cudaDeviceSynchronize();
}
컴파일
!nvcc -arch=sm_70 -o multi-block-loop 04-loops/02-multi-block-loop.cu -run결과
This is iteration number 0
This is iteration number 1
This is iteration number 2
This is iteration number 3
This is iteration number 4
This is iteration number 5
This is iteration number 6
This is iteration number 7
This is iteration number 8
This is iteration number 902-multi-block-loop.cu
#include <stdio.h>
/*
* Refactor `loop` to be a CUDA Kernel. The new kernel should
* only do the work of 1 iteration of the original loop.
*/
__global__ void loop()
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
printf("This is iteration number %d\n", i);
}
int main()
{
/*
* When refactoring `loop` to launch as a kernel, be sure
* to use the execution configuration to control how many
* "iterations" to perform.
*
* For this exercise, be sure to use more than 1 block in
* the execution configuration.
*/
int N = 10;
loop<<<2, N/2>>>();
cudaDeviceSynchronize();
}
컴파일
!nvcc -arch=sm_70 -o multi-block-loop 04-loops/02-multi-block-loop.cu -run결과
This is iteration number 0
This is iteration number 1
This is iteration number 2
This is iteration number 3
This is iteration number 4
This is iteration number 5
This is iteration number 6
This is iteration number 7
This is iteration number 8
This is iteration number 9병렬의 순차성??
위의 결과가 0~9까지 순차적이었던 이유는 thread 는 모두 병렬적이지만 호출 규약에 의한 순서 차이에 의해 발생한 현상입니다. 이론상으로는 랜덤하게 출력되었어야만 한다는 것이다. 실제로 block은 이런 차이에 의한 순차적인 현상이 발생할 확률이 더 적습니다.
loop<<<5, 2>>>();
로 바꾸어 실행시켜 보세요.
결과
This is iteration number 8
This is iteration number 9
This is iteration number 6
This is iteration number 7
This is iteration number 0
This is iteration number 1
This is iteration number 2
This is iteration number 3
This is iteration number 4
This is iteration number 5thread 사이에서는 순차적으로 kernal 구동이 끝나는 경향을, block은 랜덤하게 구동이 끝나는 모습을 볼 수 있습니다. 결과에 속으면 안 됨! 항상 cuda는 병렬적입니다.
GPU와 CPU가 사용하는 메모리 할당하기 (Allocating Memory)
버전 6 이상의 최근 버전 CUDA에서는 CPU 호스트와 GPU 디바이스가 사용하는 메모리를 보다 쉽게 할당할 수 있습니다. 가속화 애플리케이션의 최적 성능을 위한 다양한 중급 및 고급 메모리 관리 기법들 이 있기는 하지만, 지금 소개하는 기본 CUDA 메모리 관리 기법만으로도 별다른 어려움 없이 CPU-only 애플리케이션에 비해 훨씬 훌륭한 성능을 얻을 수 있습니다.
메모리를 할당하고 해제하며, 호스트와 디바이스 코드가 참조할 수 있는 포인터를 얻기 위해서는 아래 예제와 같이 malloc과 free를 cudaMallocManaged와 cudaFree로 바꾸어 주기만 하면 됩니다.
// CPU-only
int N = 2<<20;
size_t size = N * sizeof(int);
int *a;
a = (int *)malloc(size);
// Use `a` in CPU-only program.
free(a);
// Accelerated
int N = 2<<20;
size_t size = N * sizeof(int);
int *a;
// Note the address of `a` is passed as first argument.
cudaMallocManaged(&a, size);
// Use `a` on the CPU and/or on any GPU in the accelerated system.
cudaFree(a);
필요한 스레드 수와 블록 설정 간의 불일치 해결하기 (Grid Size Work Amount Mismatch)
때로는 반복문을 병렬화하기 위해 필요한 정확한 스레드 개수를 실행 설정 과정에서 지정할 수 없는 경우가 있을 수 있습니다.
이러한 사례는 최적 블록 크기를 결정할 때 흔히 볼 수 있습니다. 예를 들어, GPU 하드웨어의 특성 때문에 블록에 들어가는 스레드의 수를 32의 배수로 설정해야만 좋은 성능이 나오는 경우가 있습니다. 32의 배수인 256 개의 스레드를 가지는 블록을 여러 개 구동하고 싶고, 1000 개의 병렬 작업을 실행해야 하는 경우를 생각해 봅시다. 이 경우 정확히 1000 개의 스레드를 가지는 블록 수를 정할 수 없습니다. 32의 배수로는 정확히 1000을 만들 수 없기 때문입니다.
이 시나리오는 다음과 같은 방법으로 해결할 수 있습니다.
- 할당된 작업을 수행하는 데 필요한 스레드 수보다 더 많은 수의 스레드를 만들도록 실행 설정을 작성합니다.
- 처리해야 하는 데이터의 총 크기 또는 작업을 위해 필요한 총 스레드 수를
N이라고 하고 이 값을 커널에 인수로 넘겨줍니다. tid+bid*bdim를 이용하여 스레드의 인덱스를 계산한 후, 이 인덱스 값이N을 초과하지 않는지 확인하여 초과하지 않는 경우에만 커널의 관련 작업을 수행합니다.
N 값과 블록 당 스레드 개수를 알고 있고, 그리드의 스레드 개수와 N이 정확히 일치하는 것을 보장할 수 없을 때의 통상적 예제 코드는 아래와 같습니다. 이 코드는 최소 N 개 이상의 스레드가 존재하되, 최대 1 개의 추가적 블록에 해당하는 스레드까지만 만들도록 되어 있습니다.
// Assume `N` is known
int N = 100000;
// Assume we have a desire to set `threads_per_block` exactly to `256`
size_t threads_per_block = 256;
// Ensure there are at least `N` threads in the grid, but only 1 block's worth extra
size_t number_of_blocks = (N + threads_per_block - 1) / threads_per_block;
some_kernel<<<number_of_blocks, threads_per_block>>>(N);위 코드의 실행 설정은 N 보다 큰 개수의 스레드를 가지는 그리드를 만들기 때문에 "추가된" 스레드 중 하나가 실행될 때 인덱스 범위를 벗어나는 데이터 접근을 시도하지 않도록 주의를 기울여야 합니다.
__global__ some_kernel(int N)
{
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < N) // Check to make sure `idx` maps to some value within `N`
{
// Only do work if it does
}
}그리드보다 큰 데이터 집합
최고 성능의 실행 설정을 위한 선택에 의해서이든지 아니면 그저 필요에 의해서든지, 그리드의 스레드 개수가 데이터 집합의 크기보다 작은 경우가 있기 마련입니다. 1000 개의 원소를 가진 배열과 250 개의 스레드를 가진 그리드를 생각해 봅시다. 여기에서 그리드의 각 스레드는 4 번씩 사용되어야 합니다. 이를 달성하기 위한 통상적인 방법은 커널에서 그리드 폭 반복문(grid-stride loop)을 사용하는 것입니다.
그리드 폭 반복문에서 각 스레드는 tid+bid*bdim를 이용하여 고유의 인덱스를 계산합니다. 그리고는 방금 계산한 인덱스 값에 그리드의 스레드 개수를 더하여 다음 계산을 합니다. 이 과정을 배열의 범위를 벗어나기 전까지 반복합니다. 예를 들어 500 개의 원소를 가진 배열과 250 개의 스레드를 가진 그리드가 있을 때, 인덱스 20을 가진 스레드는 다음과 같은 과정을 밟을 것입니다.
- 500 개의 원소를 가진 배열에서 20 번째 원소를 처리합니다.
- 인덱스를 그리드의 폭인 250 만큼 건너뛰어 인덱스 270을 얻습니다.
- 500 개의 원소를 가진 배열에서 270 번째 원소를 처리합니다.
- 인덱스를 그리드의 폭인 250 만큼 건너뛰어 인덱스 520을 얻습니다.
- 520은 인덱스 범위를 초과하므로 스레드는 작업을 멈춥니다.
CUDA는 그리드 안에 있는 블록 수를 나타내는 변수 gridDim.x를 제공합니다. 그리드의 총 스레드 수는 단순히 그리드의 블록 수에 블록 당 스레드 수를 곱한 gridDim.x * blockDim.x로부터 얻을 수 있습니다. 이 내용을 기억하면서 아래 예제를 보세요.
__global void kernel(int *a, int N)
{
int indexWithinTheGrid = threadIdx.x + blockIdx.x * blockDim.x;
int gridStride = gridDim.x * blockDim.x;
for (int i = indexWithinTheGrid; i < N; i += gridStride)
{
// do work on a[i];
}
}03-grid-stride-double.cu
#include <stdio.h>
void init(int *a, int N)
{
int i;
for (i = 0; i < N; ++i)
{
a[i] = i;
}
}
/*
* In the current application, `N` is larger than the grid.
* Refactor this kernel to use a grid-stride loop in order that
* each parallel thread work on more than one element of the array.
*/
__global__
void doubleElements(int *a, int N)
{
for( int i = blockIdx.x * blockDim.x + threadIdx.x; i < N ; i += gridDim.x * blockDim.x)
{
a[i] *= 2;
}
}
bool checkElementsAreDoubled(int *a, int N)
{
int i;
for (i = 0; i < N; ++i)
{
if (a[i] != i*2) return false;
}
return true;
}
int main()
{
/*
* `N` is greater than the size of the grid (see below).
*/
int N = 10000;
int *a;
size_t size = N * sizeof(int);
cudaMallocManaged(&a, size);
init(a, N);
/*
* The size of this grid is 256*32 = 8192.
*/
size_t threads_per_block = 256;
size_t number_of_blocks = 32;
doubleElements<<<number_of_blocks, threads_per_block>>>(a, N);
cudaDeviceSynchronize();
bool areDoubled = checkElementsAreDoubled(a, N);
printf("All elements were doubled? %s\n", areDoubled ? "TRUE" : "FALSE");
cudaFree(a);
}
컴파일
!nvcc -arch=sm_70 -o grid-stride-double 05-allocate/03-grid-stride-double.cu -run결과
All elements were doubled? TRUE오류 처리하기 (Error Handling)
다른 애플리케이션과 마찬가지로 CUDA 역시 오류 처리가 필수적입니다. 대부분이라고까지는 할 수 없지만 많은 CUDA 함수가 cudaError_t 타입의 값을 리턴합니다. (메모리 관리 함수의 예를 확인해 보세요.) 이 값을 이용해서 함수 호출 중에 오류가 발생했는지를 알 수 있습니다. 아래 코드는 cudaMallocManaged를 호출할 때 오류를 처리하는 예제입니다.
cudaError_t err;
err = cudaMallocManaged(&a, N) // Assume the existence of `a` and `N`.
if (err != cudaSuccess) // `cudaSuccess` is provided by CUDA.
{
printf("Error: %s\n", cudaGetErrorString(err)); // `cudaGetErrorString` is provided by CUDA.
}void를 리턴하도록 되어 있는 구동 커널은 cudaError_t 타입의 값을 리턴하지 않습니다. 실행 설정의 오류와 같이 커널 구동 단계에서 발생하는 오류를 점검하기 위해서 CUDA는 cudaGetLastError 함수를 제공하는데 이 함수는 cudaError_t 타입의 값을 리턴합니다. 비동기 에러 처리하기
/*
* This launch should cause an error, but the kernel itself
* cannot return it.
*/
someKernel<<<1, -1>>>(); // -1 is not a valid number of threads.
cudaError_t err;
err = cudaGetLastError(); // `cudaGetLastError` will return the error from above.
if (err != cudaSuccess)
{
printf("Error: %s\n", cudaGetErrorString(err));
}마지막으로, 비동기 커널의 실행 중에 나오는 오류와 같이 비동기적으로 발생하는 오류를 잡아내기 위해서는 뒤이어 실행되는 동기화를 위한 CUDA 런타임 API 호출에서 리턴된 상태값을 검사하는 것이 필수적입니다. 이러한 동기화 함수에는 앞에서 본 cudaDeviceSynchronize가 있는데 이 함수는 이전에 구동된 커널 중 실패한 것들의 오류를 리턴합니다.
실습: 오류 처리 추가하기
01-add-error-handling.cu 프로그램을 컴파일하고 실행하면 배열의 원소를 이배수하는 작업이 성공하지 못했다는 결과를 출력합니다. 하지만 관련된 오류를 출력하지는 않습니다. CUDA 오류를 처리하도록 프로그램을 수정함으로써 어디가 잘못되었는지를 알 수 있게 하고 디버깅을 효과적으로 할 수 있도록 하세요. CUDA 함수 호출 중에 나오는 동기적 오류와 커널 실행 중에 나오는 비동기적 오류 모두를 검사해야 합니다.
01-add-error-handling.cu
#include <stdio.h>
void _CUDACHECK(cudaError_t err)
{
if (err != cudaSuccess)
{
printf("Error: %s\n", cudaGetErrorString(err));
}
}
void init(int *a, int N)
{
int i;
for (i = 0; i < N; ++i)
{
a[i] = i;
}
}
__global__
void doubleElements(int *a, int N)
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int stride = gridDim.x * blockDim.x;
for (int i = idx; i < N + stride; i += stride)
{
a[i] *= 2;
}
}
bool checkElementsAreDoubled(int *a, int N)
{
int i;
for (i = 0; i < N; ++i)
{
if (a[i] != i*2) return false;
}
return true;
}
int main()
{
/*
* Add error handling to this source code to learn what errors
* exist, and then correct them. Googling error messages may be
* of service if actions for resolving them are not clear to you.
*/
int N = 10000;
int *a;
size_t size = N * sizeof(int);
_CUDACHECK(cudaMallocManaged(&a, size));
init(a, N);
//size_t threads_per_block = 2048; thread per block cannot be bigger than 1024 ! !
size_t threads_per_block = 1024;
size_t number_of_blocks = 32;
doubleElements<<<number_of_blocks, threads_per_block>>>(a, N);
_CUDACHECK(cudaGetLastError());
cudaDeviceSynchronize();
_CUDACHECK(cudaGetLastError());
bool areDoubled = checkElementsAreDoubled(a, N);
printf("All elements were doubled? %s\n", areDoubled ? "TRUE" : "FALSE");
_CUDACHECK(cudaFree(a));
}
컴파일
!nvcc -arch=sm_70 -o add-error-handling 06-errors/01-add-error-handling.cu -run결과
All elements were doubled? TRUE학습목표
강좌를 마치면 여러분은 다음 사항을 수행할 수 있게 될 것입니다.
- NVIDIA 커맨드라인 프로파일러(nvprof)를 이용하여 가속화 애플리케이션의 성능 프로파일링하기
- 실행 설정 최적화를 위한 스트리밍 멀티프로세서에 대한 체계적 이해
- 페이지 폴트와 데이터 마이그레이션에 관련된 통합 메모리의 동작 이해
- 페이지 폴트와 데이터 마이그레이션을 감소시켜 성능을 향상시키는 비동기적 메모리 프리패칭 사용
- 반복적 개발 사이클을 적용하여 애플리케이션 가속화와 배치를 신속히 진행하기
NVIDIA 프로파일러를 활용한 최적화
실습: nvprof을 이용하여 애플리케이션 프로파일링 하기
가속화 코드 베이스의 최적화 시도가 실제로 성공했는가를 확인하는 방법 중의 하나로서 성능과 관련된 정량적 정보를 프로파일링하는 방법이 있습니다. nvprof는 이러한 작업을 지원하는 NVIDIA의 커맨드라인 프로파일러입니다. CUDA 툴킷과 함께 배포되는 이것은 가속화 애플리케이션을 프로파일링하는 강력한 도구입니다.
nvprof를 사용하기는 쉽습니다. 가장 기본적인 사용법은 nvcc로 컴파일한 실행 파일의 경로를 전달해 주는 것입니다. nvprof는 해당 애플리케이션을 실행한 후 GPU 활동, CUDA API 호출 기록, 통합 메모리(Unified Memory) 활동 정보 등을 정리하여 출력해 줍니다. 자세한 내용은 강좌를 진행하며 다루도록 하겠습니다.
애플리케이션을 가속화하거나 이미 가속화된 애플리케이션을 최적화할 때에 과학적이고 반복적인 접근법을 취하세요. 애플리케이션 변경 후에는 프로파일링을 수행하고, 기록을 남기고, 수정사항이 성능에 미치는 의미를 적으세요. 보다 이른 단계에, 그리고 자주 이러한 관찰을 수행하세요. 작은 노력들이 모여 성능을 획기적으로 향상시키고 출시를 돕는 경우가 많이 있습니다. 빈번한 프로파일링은 특정한 코드 변경이 실제 성능에 어떠한 영향을 미치는가를 여러분에게 가르쳐 줄 것입니다. 이러한 지식은 코드 베이스를 한참 변경한 후에 수행하는 프로파일링으로는 얻을 수 없는 중요한 자산입니다.
nvprof을 이용하여 애플리케이션 프로파일링 하기
!nvprof ./single-thread-vector-add결과
==231== NVPROF is profiling process 231, command: ./single-thread-vector-add
Success! All values calculated correctly.
==231== Profiling application: ./single-thread-vector-add
==231== Profiling result:
Type Time(%) Time Calls Avg Min Max Name
GPU activities: 100.00% 2.35052s 1 2.35052s 2.35052s 2.35052s addVectorsInto(float*, float*, float*, int)
API calls: 81.83% 2.35047s 1 2.35047s 2.35047s 2.35047s cudaDeviceSynchronize
17.30% 496.87ms 3 165.62ms 20.488us 496.81ms cudaMallocManaged
0.85% 24.326ms 3 8.1087ms 7.4851ms 9.1192ms cudaFree
0.01% 258.15us 94 2.7460us 615ns 70.957us cuDeviceGetAttribute
0.01% 248.00us 1 248.00us 248.00us 248.00us cuDeviceTotalMem
0.00% 123.70us 1 123.70us 123.70us 123.70us cudaLaunch
0.00% 18.788us 1 18.788us 18.788us 18.788us cuDeviceGetName
0.00% 9.1430us 4 2.2850us 701ns 6.5900us cudaSetupArgument
0.00% 4.1340us 1 4.1340us 4.1340us 4.1340us cudaConfigureCall
0.00% 3.6370us 3 1.2120us 626ns 1.6940us cuDeviceGetCount
0.00% 2.0250us 2 1.0120us 688ns 1.3370us cuDeviceGet
0.00% 1.0670us 1 1.0670us 1.0670us 1.0670us cudaGetLastError
==231== Unified Memory profiling result:
Device "Tesla V100-SXM2-16GB (0)"
Count Avg Size Min Size Max Size Total Size Total Time Name
2304 170.67KB 4.0000KB 0.9961MB 384.0000MB 42.24435ms Host To Device
768 170.67KB 4.0000KB 0.9961MB 128.0000MB 11.28947ms Device To Host
768 - - - - 104.4759ms Gpu page fault groups
Total CPU Page faults: 1536스트리밍 멀티프로세서와 워프
CUDA 애플리케이션을 실행하는 GPU는 스트리밍 멀티프로세서(streaming multiprocessor; SM)라는 프로세싱 유닛을 가지고 있습니다. 최대한 많은 수의 병렬처리를 위하여 GPU에 있는 SM 숫자의 배수로 이루어진 블록수를 그리드의 크기로 선택함으로써 성능 이득을 얻을 수 있습니다.
SM은 하나의 블록 안에 있는 32 개의 스레드를 하나의 그룹처럼 다루어 스레드 생성, 관리, 스케쥴링, 실행을 합니다. 이러한 32 개의 스레드로 이루어진 그룹을 워프(warp#Warps))라고 합니다. SM과 warp에 대한 상세 내용은 본 강좌의 범위를 벗어나지만 32의 배수를 스레드 개수로 갖는 블록 크기를 선택함으로써 성능 이득을 얻을 수 있다는 것을 기억하는 것은 중요합니다.
- warp는 스케쥴링의 특성상 순차적으로 kernal이 구동이 끝납니다.
프로그램으로 GPU 디바이스 속성 질의하기
GPU 상의 SM 개수는 GPU 모델에 따라 다르기 때문에, 서로 다른 SM 개수를 가진 GPU 간의 프로그램 이식성을 유지하기 위하여 SM 개수는 코드 베이스에 하드코드되어서는 안됩니다. 이 정보는 프로그램으로 얻어내야 합니다.
아래 예제는 CUDA C/C++에서, SM 값을 포함하여, 현재 활성화된 GPU의 다양한 속성을 알아내는 데 사용되는 C 구조체를 얻는 방법을 보여줍니다.
int deviceId;
cudaGetDevice(&deviceId); // `deviceId` now points to the id of the currently active GPU.
cudaDeviceProp props;
cudaGetDeviceProperties(&props, deviceId); // 'prop' is useful통합 메모리(Unified Memory) 세부 내용
여러분은 cudaMallocManaged 함수를 이용하여 호스트와 디바이스 코드가 사용할 메모리를 할당해 왔습니다. 지금까지 이 함수가 제공하는 자동 메모리 마이그레이션, 쉬운 프로그래밍과 같은 혜택을 이용해 오면서도 cudaMallocManaged가 실제로 할당하는 통합 메모리(unified memeory; UM)에 대한 세부 내용은 신경 쓸 필요가 없었습니다. nvprof는 가속화 애플리케이션의 UM 관리에 대한 상세한 정보를 제공하는데, 이 정보와 아울러 UM의 작동 원리를 보다 잘 이해하게 된다면 가속 애플리케이션 최적화의 추가적인 기회를 얻을 수 있습니다.
통합 메모리(UM) 마이그레이션
UM이 할당될 때, 메모리는 아직 호스트 또는 디바이스에 적재되지 않습니다. 호스트나 디바이스가 그 메모리에 접근하려고 하면 페이지 폴트가 일어나고 이 시점에 호스트와 디바이스는 필요한 데이터를 연속적으로 읽어들여 옵니다. 이것을 메모리 마이그레이션(memory migration)이라고 합니다. 이와 마찬가지로 호스트나 디바이스가 아직 적재되지 않은 메모리에 접근을 시도한다면 페이지 폴트가 일어나고 마이그레이션이 시작됩니다.
페이지 폴트와 요청시 마이그레이션은 가속화 애플리케이션 개발을 쉽게 해주어 큰 도움이 됩니다. 특히 애플리케이션이 실제로 실행되어 데이터를 필요로 할 때까지 어느 데이터가 필요한지 알 수 없는 경우가 있는데, 이와 같이 흩어진 데이터를 다루는 경우, 또는 다수의 GPU가 접근하는 데이터와 같은 경우에 있어서 요청시 마이그레이션은 매우 유용합니다.
하지만 어떤 데이터가 필요한지 미리 알 수 있는 경우와 큰 연속 메모리 영역이 필요한 경우도 많이 있습니다. 이 경우에는 오히려 페이지 폴트와 요청시 마이그레이션의 오버헤드가 큰 부담이 되므로 이러한 오버헤드 비용은 피하는 것이 좋습니다. (오버헤드 : A 작업을 하기 위해 부가적으로 필요한 B 작업을 하는 데 걸리는 비용.)
본 강좌의 나머지 부분은 요청시 마이그레이션을 이해하고 프로파일러 출력에서 이를 확인하는 것에 대해 할애할 것입니다. 이러한 지식을 통하여 여러분은 오버헤드를 피하는 것이 유리한 경우를 이해하고 대처할 수 있게 될 것입니다.
비동기 메모리 프리패칭
페이지 폴트와 호스트에서 디바이스로 또는 디바이스에서 호스트로의 메모리 전달인 요구시 메모리 마이그레이션으로 인한 오버헤드를 줄이는 강력한 기법을 비동기 메모리 프리패칭(asynchronous memory prefetching)이라고 합니다. 이 기법을 이용하면 애플리케이션이 메모리를 사용하려고 하기 전에 프로그래머가 UM을 임의의 CPU나 GPU로 마이그레이션할 수 있는데, 이는 비동기적으로 백그라운드에서 수행됩니다. 이를 통해 줄어든 페이지 폴트 및 요구시 마이그레이션만큼 GPU 커널과 CPU 함수의 성능이 향상됩니다.
프리패칭은 데이터를 큰 덩어리로 마이그레이션하는 경향이 있어 요구시 마이그레이션보다 메모리 전달 빈도가 낮아집니다. 이는 접근할 데이터를 런타임 전에 알 수 있고, 데이터 접근 패턴이 산발적이지 않은 경우에 아주 적합합니다.
CUDA는 cudaMemPrefetchAsync 함수를 이용하여 메모리에서 GPU/CPU로의 비동기적 프리패칭을 손쉽게 처리합니다. 아래 예제는 현재 활성화된 GPU로 데이터를 프리패치한 후, CPU로 프리패치하는 방법을 보여 줍니다.
int deviceId;
// The ID of the currently active GPU device.
cudaGetDevice(&deviceId);
// Prefetch to GPU device.
cudaMemPrefetchAsync(pointerToSomeUMData, size, deviceId);
// Prefetch to host. `cudaCpuDeviceId` is a built-in CUDA variable.
cudaMemPrefetchAsync(pointerToSomeUMData, size, cudaCpuDeviceId); NVIDIA Visual Profiler를 쓰시오.
NVIDIA Visual Profiler를 쓰시오.
마지막 과제
마지막 실습: N-body 시뮬레이터 가속화와 최적화
n-body 시뮬레이터는 일군의 물체가 서로 중력을 미치는 상황에서 각 물체의 운동을 예측합니다.01-nbody.cu는 간단하지만 잘 작동하는 3 차원 n-body 시뮬레이터입니다. 커맨드라인 인수를 이용하여 애플리케이션에 서로 영향을 미치는 물체의 개수를 전달할 수 있습니다.
현재는 CPU-ony 형식이고 4096 개의 물체에 대하여 초당 3000 만 번의 상호작용을 계산할 수 있습니다. 여러분이 할 일은 아래와 같습니다.
- 시뮬레이션의 정확성을 유지하면서 GPU 가속화하세요.
- 반복 작업을 통해 4096(=2<<11) 개의 물체에 대해 300 억 번의 상호작용을 계산하도록 최적화하세요.
- 반복 작업을 통해 ~65,000(=2<<15) 개의 물체에 대해 3250 억 번의 상호작용을 계산하도록 최적화하세요.
완료한 다음에는 이 노트북을 여신 브라우저로 돌아가서 Assess 버튼을 누르세요. 시물레이션 정확도를 유지하면서 제시된 목표를 달성했을 경우 여러분은 CUDA C/C++을 이용한 애플리케이션 가속화 기초에 대한 역량을 인정하는 인증서를 받으실 수 있습니다.
목표
3250 억 번
: 325 billion times
소스 코드
#include <math.h>
#include <stdio.h>
#include <stdlib.h>
#include "timer.h"
#include "check.h"
#define SOFTENING 1e-9f
/*
* Each body contains x, y, and z coordinate positions,
* as well as velocities in the x, y, and z directions.
*/
typedef struct { float x, y, z, vx, vy, vz; } Body;
/*
* Do not modify this function. A constraint of this exercise is
* that it remain a host function.
*/
void randomizeBodies(float *data, int n) {
for (int i = 0; i < n; i++) {
data[i] = 2.0f * (rand() / (float)RAND_MAX) - 1.0f;
}
}
/*
* This function calculates the gravitational impact of all bodies in the system
* on all others, but does not update their positions.
*/
__global__ void bodyForce(Body *p, float dt, int n) {
int stride = blockDim.x * gridDim.x;
int i = blockDim.x * blockIdx.x + threadIdx.x;
for(;i < n; i += stride)
{
float Fx = 0.0f; float Fy = 0.0f; float Fz = 0.0f;
for (int j = 0; j < n; j++) {
float dx = p[j].x - p[i].x;
float dy = p[j].y - p[i].y;
float dz = p[j].z - p[i].z;
float distSqr = dx*dx + dy*dy + dz*dz + SOFTENING;
float invDist = rsqrtf(distSqr);
float invDist3 = invDist * invDist * invDist;
Fx += dx * invDist3; Fy += dy * invDist3; Fz += dz * invDist3;
}
p[i].vx += dt*Fx; p[i].vy += dt*Fy; p[i].vz += dt*Fz;
}
}
__global__ void physics(Body *p, float dt, int n)
{
int stride = blockDim.x * gridDim.x;
for(int i = blockDim.x * blockIdx.x + threadIdx.x;i < n; i += stride)
{
p[i].x += p[i].vx*dt;
p[i].y += p[i].vy*dt;
p[i].z += p[i].vz*dt;
}
}
int main(const int argc, const char** argv) {
/*
* Do not change the value for `nBodies` here. If you would like to modify it,
* pass values into the command line.
*/
int nBodies = 2<<11;
int salt = 0;
if (argc > 1) nBodies = 2<<atoi(argv[1]);
/*
* This salt is for assessment reasons. Tampering with it will result in automatic failure.
*/
if (argc > 2) salt = atoi(argv[2]);
const float dt = 0.01f; // time step
const int nIters = 10; // simulation iterations
int bytes = nBodies * sizeof(Body);
float *buf;
int deviceId;
cudaGetDevice(&deviceId);
int numberOfSMs;
cudaDeviceGetAttribute(&numberOfSMs, cudaDevAttrMultiProcessorCount, deviceId);
cudaMallocManaged(&buf, bytes);
cudaMemPrefetchAsync(buf, bytes, deviceId);
Body *p = (Body*)buf;
/*
* As a constraint of this exercise, `randomizeBodies` must remain a host function.
*/
randomizeBodies(buf, 6 * nBodies); // Init pos / vel data
double totalTime = 0.0;
/*
* This simulation will run for 10 cycles of time, calculating gravitational
* interaction amongst bodies, and adjusting their positions to reflect.
*/
/*******************************************************************/
// Do not modify these 2 lines of code.
for (int iter = 0; iter < nIters; iter++) {
StartTimer();
/*******************************************************************/
/*
* You will likely wish to refactor the work being done in `bodyForce`,
* as well as the work to integrate the positions.
*/
bodyForce<<<numberOfSMs * 32, 256>>>(p, dt, nBodies); // compute interbody forces
/*
* This position integration cannot occur until this round of `bodyForce` has completed.
* Also, the next round of `bodyForce` cannot begin until the integration is complete.
*/
physics<<<numberOfSMs * 32, 256>>>(p, dt, nBodies);
cudaDeviceSynchronize();
/*******************************************************************/
// Do not modify the code in this section.
const double tElapsed = GetTimer() / 1000.0;
totalTime += tElapsed;
}
double avgTime = totalTime / (double)(nIters);
float billionsOfOpsPerSecond = 1e-9 * nBodies * nBodies / avgTime;
#ifdef ASSESS
checkPerformance(buf, billionsOfOpsPerSecond, salt);
#else
checkAccuracy(buf, nBodies);
printf("%d Bodies: average %0.3f Billion Interactions / second\n", nBodies, billionsOfOpsPerSecond);
salt += 1;
#endif
/*******************************************************************/
/*
* Feel free to modify code below.
*/
cudaFree(buf);
}