지난 1편에서는 첫 번째 CUDA 예제 moveArray.cu 를 소개하면서 CUDA 툴을 이용하여 프로그램을 만들고 실행하는 것을 연습하였습니다. C 프로그래머에게 그 프로그램은 단순히 CUDA API 를 호출하여 메모리를 할당하여 CUDA 장비와 데이터를 교환하는 것이었습니다. 툴을 이용하여 컴파일하고 실행하는 데 특별히 까다로운 부분은 없었습니다.

이번에는 몇 줄을 더 추가하여 CUDA 장치에서 간단한 계산을 수행하는 예제를 만들어 보겠습니다. 부동소수점 배열의 숫자를 하나씩 증가시키는 예제입니다. 하지만 이 예지는 CUDA 사용에 필요한 기본 구조 대부분의 내용 (데이터를 CUDA 장치로 전송, 계산을 수행, 결과를 얻어냄) 을 보여줄 것입니다.

좀 더 자세한 내용에 들어가기에 앞서, 먼저 다음을 이해하여야 합니다.

* 커널은 무엇인가? 커널은 CUDA 장치를 - 다수 쓰레드에서 병렬로 동시에 - 실행하기 위하여 호스트에서 호출하는 함수이다.
* 커널은 어떻게 호출하는가? 여기에는 커널의 이름과 실행 환경 설정 방법이 필요하다. 이번 글에서 실행 환경은 그룹 내 병렬 쓰레드 숫자와 그룹의 갯수이다. 환경 설정은 매우 중요한 내용이며 향후 글에서 좀 더 깊이 다룰 것이다.
* 호스트와 커널을 싱크로나이즈하는 방법

Listing One 의 예제에는 호스트 루틴 incrementArrayOnHost 와 커널 incrementArraysOnDevice 가 있습니다.

호스트 함수 incrementArrayOnHost 는 각각의 배열 번호에서 그 내용을 하나씩 증가시키는 간단한 내용입니다. 이 함수는 프로그램 마지막에서 커널이 CUDA 장치에서 올바르게 작업을 수행했는지를 검사하기 위해 작성되었습니다.

그 다음은 CUDA 커널인 incrementArrayOnDevice 입니다. CUDA 는 C 언어에 몇 가지 확장 기능을 제공합니다. 함수 형 지정자인 __global__ 은 CUDA 장치에서 실행되는 커널을 선언합니다. 커널은 호스트에서만 호출할 수 있습니다. 커널의 리턴은 반드시 void 로 선언합니다.

커널 incrementArrayOnDevice 는 incrementArrayOnHost 와 동일한 계산을 수행합니다. incrementArrayOnDevice 를 보시면, 루프가 없다는 것을 보실 것입니다. 왜냐하면 이 함수는 CUDA 장치 내의 쓰레드에서 동시에 실행되는 함수이기 때문입니다. 각각의 쓰레드는 고유 ID 가 있어서 배열 내 다른 인덱스를 계산하거나, 혹은 변수 인덱스가 벗어났을 때 대기할 것인지를 판단하는 데 사용됩니다. 이렇게 레지스터 변수 idx 에 담긴 고유 ID 를 이용하면 함수를 간단하게 만들 수 있습니다. idx 는 1 씩 증가시킬 배열 요소를 가리키는 인덱스로 활용됩니다. 쓰레드 개수가 배열보다 클 수 있으므로, 먼저 idx 를 커널로 전달되는, 배열 크기를 지정하는 변수인 N 과 비교해서 작업을 수행할 것인지를 판단합니다.

커널 호출과 환경 설정은 어떻게 할까요? 프로그램 흐름은 main 으로부터 순차적으로 시작해서 Part 2 of 2 라는 문장이 담긴 주석문 바로 다음 줄까지 계속됩니다.

// incrementArray.cu
#include
#include
#include
void incrementArrayOnHost(float *a, int N)
{
int i;
for (i=0; i < N; i++) a[i] = a[i]+1.f;
}
__global__ void incrementArrayOnDevice(float *a, int N)
{
int idx = blockIdx.x*blockDim.x + threadIdx.x;
if (idx}
int main(void)
{
float *a_h, *b_h; // pointers to host memory
float *a_d; // pointer to device memory
int i, N = 10;
size_t size = N*sizeof(float);
// allocate arrays on host
a_h = (float *)malloc(size);
b_h = (float *)malloc(size);
// allocate array on device
cudaMalloc((void **) &a_d, size);
// initialization of host data
for (i=0; i // copy data from host to device
cudaMemcpy(a_d, a_h, sizeof(float)*N, cudaMemcpyHostToDevice);
// do calculation on host
incrementArrayOnHost(a_h, N);
// do calculation on device:
// Part 1 of 2. Compute execution configuration
int blockSize = 4;
int nBlocks = N/blockSize + (N%blockSize == 0?0:1);
// Part 2 of 2. Call incrementArrayOnDevice kernel
incrementArrayOnDevice <<< nBlocks, blockSize >>> (a_d, N);
// Retrieve result from device and store in b_h
cudaMemcpy(b_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost);
// check results
for (i=0; i // cleanup
free(a_h); free(b_h); cudaFree(a_d);

< Listing One: incrementArray.cu >

이제 incrementArrayOnDevice 를 CUDA 장치에 등록하게 됩니다. 여기서 CUDA 커널의 비동기 실행을 하는 C 언어 확장 기능이 소개됩니다. 먼저 커널의 이름과 실행 환경 설정을 <<< 과 >>> 사이에 넣습니다. nBlocks 와 blockSize 두 개의 변수가 실행 환경 설정에 지정되었습니다. 커널로 전달되는 인수들은 기본 C 함수 인수 전달 방식과 같이 ( ) 사이에 넣습니다. 이 예제에서 배열을 가리키는 전역 메모리 포인터 a_d 와 배열 길이를 나타내는 N 이 커널에 전달됩니다.

CUDA 장치가 아무 일도 하지 않고 있으므로 커널은 지정된 실행 환경 설정과 함수 전달 인수에 따라 즉시 실행하게 됩니다. 한편 호스트는 커널이 실행되자마자 다음 줄로 진행됩니다. 이 시점에서 CUDA 장치와 호스트는 각각의 루틴을 동시에 수행합니다. incrementArray.cu 의 경우 호스트는 곧바로 cudaMemcpy 를 호출하게 되는데, 이 함수는 장치 내 모든 쓰레드가 종료(incrementArrayOnDevice 에서 리턴)되어 바뀐 배열값이 호스트로 돌아올 때까지 기다립니다. 프로그램은 호스트의 incrementArrayOnHost 에서 순차적으로 만들어진 값과 CUDA 가 병렬로 incrementArrayOnDevice 에서 만든 값을 순차적으로 비교한 다음 종료합니다.

커널 실행시 실행 환경 설정 ( <<< >>> 사이에 지정한 nBlocks 와 blockSize ) 변수들 몇 가지 지정되었습니다. 이 변수들은 모든 커널에 유효하니다. nBlocks 와 blockSize 설정은 매우 세련된 방식인데, 왜냐하면 이렇게 함으로써 프로그램 재 컴파일 없이 개발자들이 하드웨어의 필요 한도를 지정할 수 있기 때문입니다. 이것은 CUDA 를 이용한 상업용 프로그램 개발에 핵심적인 요소입니다.

향후 살펴볼 것입니다만, 블럭 내 쓰레드들은 서로 통신 및 동기화할 수 있습니다. 이는 놀라운 소프트웨어 특성입니다만 그만큼 비싼 하드웨어가 필요합니다. 향후 더 비싼 장치들은 값싼 구형 장치들보다 더 많은 블럭 당 쓰레드 개수를 지원할 것입니다. 그리드 추상화 기법은 하드웨어 세대와 가격대에 무관하도록 프로그램 재 컴파일 없이 이와 같은 내용을 반영할 수 있게끔 합니다. 실제적으로 그리드는 동일한 차원과 크기의 블럭을 다루는 동일한 커널을 일괄 호출하며, 단일 커널 호출 당 실행되는 쓰레드 개수인 nBlocks 만큼 효과적으로 증대시킬 수 있습니다. 값싼 장치들은 한 개 혹은 수 개의 쓰레드만을 동시 처리할 수 있을 것이고, 향후 개발될 비싼 장치들은 더 많은 쓰레드를 호출할 수 있을 것입니다. 그리드 추상화 기법으로 소프트웨어를 구성할 경우, 동시 실행 가능한 쓰레드 갯수와 블럭 내 서로 협력해야 하는 최대 쓰레드 갯수, 두 가지 변수의 트레이드 오프에 대한 균형을 고려해야 합니다. 이 두 종류 변수의 영향을 인식하여야 합니다. 물론 알고리듬에 따라 필요한 요소가 달라지겠습니다만, 가능하다면 많은 쓰레드 블럭 갯수를 지정하도록 하십시오.

CUDA 장치의 커널 내에는 커널 호출시 환경 설정으로 전달되는 예약 변수들이 몇 종류 있습니다.

* blockIdx: 그리드 내 블럭 인덱스
* threadIdx: 블럭 내 쓰레드 인덱스
* blockDim: 블럭 내 쓰레드 갯수

이 변수들은 정수형 변수를 포함한 구조체입니다. 예를 들어 블럭은 3차원이므로 정수 x, y, z 를 갖습니다. 한편 그리드는 2차원이므로 정수 x, y 를 갖습니다. 이번 예제에서는 CUDA 장치로 전달되는 배열이 1차원이므로 x 만 사용됩니다. 2차원, 3차원 기능이 어떻게 사용될 것인지는 향후 다루겠습니다.

우리 예제에서 커널은 예약 변수들을 이용하여 쓰레드 인덱스 idx 를 다음과 같이 계산합니다.

int idx = blockIdx.x * blockDim.x + threadIdx.x;

nBlocks 는 그리드 내 블럭 개수, blockSize 는 블럭 내 쓰레드 개수입니다. 이번 예제에서는 호스트 코드에서 커널 호출 직전에 초기화됩니다.

int blockSize = 4;
int nBlocks = N/blockSize + (N%blockSize == 0?0:1);

N 이 blockSize 로 딱 떨어지게 나눌 수 없을 경우 nBlocks 계산시 블럭 하나를 더 추가합니다. 그러므로 경우에 따라서 블럭 내 몇 쓰레드는 아무 일도 하지 않게 됩니다.

확실히 이번 예제는 문제를 쉽게 하기 위하여 쓰레드 블럭 네 개 안에 맞도록 배열 크기를 쓰레드 갯수보다 적게 인위적으로 만든 것입니다. 확실히 심하게 간략화한 것이지만, 쉬운 코드를 이용하여 커널이 incrementArrayOnDevice 를 호출하는 방법을 살펴보겠습니다.

중요한 점은, 각각의 쓰레드는 장치 내에서 변수 a_d 를 모두 억세스할 수 있다는 점입니다. 커널 실행시 고유한 데이터 분리는 없습니다. 커널 프로그램을 짤 때 커널을 인식하고 계산의 병렬적 요소를 지정하는 것은 프로그래머에게 달려 있습니다.

그림 1 은 배열 내에서 어떻게 idx 를 계산하고 a_d 를 참조하는지 나타냅니다. (이 내용이 이해가 안 되신다면, printf 구문을 incrementArrayOnDevice 에 첨가하여 idx 변수와 변수 계산에 필요한 관련 값들을 출력해 보세요. 출력을 보시려면 컴파일 시 "make emu=1" 로 컴파일하시고 실행하세요. printf 출력을 보시려면 에뮬레이터 실행파일의 path 를 확실히 지정해 주셔야 합니다.)

다시 말씀드립니다만 커널 호출은 비동기식입니다. 커널 실행과 동시에 프로그램은 호스트 CPU 로 리턴됩니다. 이전 CUDA 호출이 종료되어 있다면 커널은 CUDA 장치에서 실행될 것입니다. 비동기식 커널 호출은 호스트와 장치간 동시 계산을 가능하게 하는 훌륭한 방식입니다. 이 예제에서 incrementArrayOnHost 를 incrementArrayOnDevice 호출 이후로 바꾸어 호스트와 장치가 계산을 중복 수행하게 함으로써 더 나은 성능을 이끌어낼 수 있습니다. 커널 내 작업에 필요한 시간에 따라, 호스트와 디바이스가 동시에 계산을 수행하게 하는 것이 가능합니다.

다음 글을 보실 때까지

* N 과 nBlocks 값을 변경해 보세요. 만약 이 값이 장치의 한계를 넘었을 때 어떤 일이 일어나는지 확인해 보세요.
* 크기가 다른 배열을 처리할 수 있는 루프를 넣는 방법을 생각해 보세요.
* CUDA 장치 메모리 종류를 확인해 보세요. (전역 메모리, 레지스터, 공유 메모리, 상수 메모리) CUDA occupancy calculator 를 살펴보시고 와 nvcc 의 -cubin 이나 --ptxas-option=-v 옵션으로 커널 내 레지스터 사용 갯수를 지정해 보세요.