- 오류처리와 전역 메모리 성능 한계

축하합니다! CUDA 시리즈의 1편과 2편을 마치신 여러분은 CUDA 프로그래머가 되어서 이제 여러분은 CUDA 장치가 지원하는 수백 개의 병렬 쓰레드를 이용한 프로그램을 개발할 능력을 갖추셨습니다. 2편에 소개한 incrementArrays.cu 에서 여러분에게 데이터를 장치에 전송, 계산을 수행하기 위한 커널 실행, 결과 추출이라는 기본적인 CUDA 응용프로그램 패턴의 실행 예제를 보여드렸습니다. 기본적으로 incrementArrays.cu 에 여러분의 계산 커널과 데이터(이번 회에서 소개할 내용)를 집어 넣으시면 어떤 응용프로그램으로도 변형시킬 수 있습니다. 그 다음 회에서는 CUDA 비동기형 I/O 와 스트림에 관해 살펴볼 예정입니다.

"당신은 너무 많은 것을 알고 있소!" 라는 우스개는 지금까지 우리가 공부한 바를 잘 나타내고 있습니다. CUDA 의 좋은 점은 여러분의 프로그래밍 계획을 집단 병렬 프로그램으로 변형하는 자연스러운 방법을 제시하고 있습니다. 하지만 탄탄하고 효율적인 프로그램을 만들기 위해서는 더욱 깊은 이해가 필요합니다.

너무 걱정하지 마시고 시도해 보세요! CUDA 는 프로그래밍 툴이고 훌륭한 소프트웨어를 개발하기 위한 구조물이기도 합니다. 그리고 이것을 배우는 방법은 실제로 해 보는 것 뿐이 없습니다. 실제로 이번 장에서 CUDA 의 특징을 짧은 예제와 여러분의 이해를 돕기 위한 인터넷 상의 정보들을 통해서 여러분에게 실제로 수행해볼 수 있는 방법을 제공하겠습니다. CUDA Zone(http://www.nvidia.com/cuda) 은 CUDA 에 관한 정보와 질문답변을 볼 수 있는 포럼 등을 제공하는 좋은 곳입니다. 그리고 여러분이 직접 질문하고 답변을 얻으실 수도 있습니다.

이번 및 다음 몇 호에 걸쳐서 간단한 배열 뒤집기를 통해서 여러분의 이해를 돕고 공유 메모리가 성능에 미치는 영향을 살펴보겠습니다. 오류 찾기와 성능 검사 등을 CUDA 프로파일링 툴을 이용해서 수행하는 법도 다루겠습니다. 다음 호에 첨부된 소스에서 실제로 배열 뒤집기를 공유 메모리에서 어떻게 수행하는지 보실 수 있습니다. reverseArray_multiblock.cu 는 아직 저성능이지만 CUDA 장치의 광역 메모리 내에서 배열을 뒤집는 법을 보여드릴 것입니다. 하지만 이 프로그램을 여러분의 프로그램에 사용하지는 마세요. 광역 메모리는 이런 종류의 프로그램에 사용하기에 비효율적입니다. 그리고 이 프로그램은 비 연동 메모리 억세스를 하고 있어서 광역 메모리 효율이 좋지 않습니다. 가장 최적의 광역 메모리 밴드폭은 메모리 억세스가 동시에 단일 메모리 교환에 연동되었을 때 얻을 수 있습니다. 다음회에서 광역 메모리와 공유 메모리, 그리고 장치의 계산 능력에 따른 다양한 메모리 억세스 연동 방법을 살펴볼 것입니다.

- CUDA 오류 처리

오류를 발견해서 처리하는 내용은 탄탄하고 실용적인 소프트웨어 개발에 필수입니다. 사람들은 프로그램이 죽거나 오동작하는 것에 대해 관대하지 못합니다. 개발자 측면에서 오류 처리 코드를 넣는 것은 귀찮고 성가신 작업입니다. 그리고 말쑥한 코드를 어지럽히고, 예측 가능한 오류 처리 루틴을 짜는 동안 전체 개발 시간을 지연시킬 수도 있습니다. 네, 오류 처리는 별로 반갑지 않은 작업이지만 이것은 여러분을 위해서 하는 작업이 아니지요 (하지만 잘 만든 오류 처리 코드가 나의 소중한 시간을 아껴주기도 합니다), 이 프로그램을 사용할 사람들을 위해서 하는 작업입니다. 만약 프로그램이 죽으면 사람들은 왜 죽었고, 더 중요한 것은, 어떻게 하면 죽는 것을 방지할 수 있느냐를 알고 싶어합니다. 잘 만든 오류 처리 및 복구 코드는 여러분의 프로그램이 대박 나게 할 것임은 주지의 사실입니다. 상업용 소프트웨어 개발자들은 특히 더 기억해 두세요.

CUDA 를 고안할 때에도 오류 처리의 중요성이 충분히 인식되었습니다. 이를 위해서 커널 호출을 제외한 모든 CUDA 호출은 cudaError_t 라는 오류 코드를 리턴합니다. 작업이 성공적으로 이루어졌으면 cudaSuccess 가 리턴되고, 아니면 오류 코드가 리턴됩니다.

오류 코드를 읽기 쉽게 바꿔주는 것은 다음과 같습니다.

char *cudaGetErrorString(cudaError_t code);

C 언어 프로그래머들은 이 방법과, 오류를 나타내는 변수 errno 와 이를 사람이 읽기 편한 메세지로 변환하는 perror 와 sterror 를 사용하는 C 라이브러리와 유사하다는 것을 발견하실 수 있을 것입니다. 이러한 C 라이브러리 기법은 지금껏 수많은 C 코드에서 애용되어 왔으며, 향후 CUDA 소프트웨어에서도 잘 작동될 것은 의심의 여지가 없습니다.

또한 CUDA 에는 cudaGetLastError 메쏘드가 있어서 호스트 쓰레드의 최후 런타임 콜에 대한 오류를 보고합니다. 여기에는 몇 가지 관련 사항이 있습니다.

* 커널 실행의 비동기성으로 인하여 cudaGetLastError 가 명확하지 않을 수 있습니다. 따라서, 커널 호출을 포함해서 지금까지의 모든 호출이 종료될 때까지 대기했다가 어떤 작업이든 오류가 발생했을 때 오류 코드를 리턴하는 cudaThreadSynchronize 를 사용하세요. 다수의 커널 콜을 동시에 발생시키면 모든 커널 콜이 종료되기 전까지 오류 검사를 할 수 없다는 뜻이 되기도 합니다. 아니면 오류를 검사해서 호스트로 리턴하는 특별한 코드를 직접 커널 내에 작성해 주어야 합니다.
* 오류는 올바른 호스트 쓰레드로 보고됩니다. 만약 프로그램이 다수의 CUDA 장치를 사용하는 경우와 같이, 호스트가 여러 쓰레드를 동시에 실행한다면, 오류는 올바른 호스트 쓰레드로 보고됩니다.
* 만약 cudaGetLastError 를 호출할 때까지 여러 개의 오류가 발생했을 경우, 맨 마지막에 발생한 오류만 보고됩니다. 프로그래머는 실시간을 발생하는 오류 정보를 곧바로 처리해 주어야 합니다. 아니면 잘못된 오류 보고를 사용자에게 출력할 수도 있습니다.

- 소스 코드를 봅시다

reverseArray_multiblock.cu 소스코드를 살펴보시면, 프로그램의 구조는 지난 2편에서 보셨던 moveArray.cu 와 매우 유사하다는 것을 볼 수 있습니다. 오류 처리 루틴 checkCUDAError 는 cudaGetLastError 에서 오류가 보고되었을 때 이것을 읽기 편한 메세지로 출력한 후 프로그램을 종료하는 루틴입니다. 보시는 것처럼 프로그램 전체에 걸쳐서 오류 검사를 위해 checkCUDAError 이 적절히 활용되고 있습니다.

reverseArray_multiblock.cu 는 [0..dimA-1] 의 값을 가지는 1차원 정수형 배열 h_a 를 선언합니다. h_a 는 cudaMemcpy 를 이용해서 장치의 전역 메모리에 상주하는 d_a 로 옮겨집니다. 그리고 나서 호스트는 reverseArrayBlock 커널을 호출하여 d_a 로부터 d_b 로 배열 내용을 역순으로 복사합니다. 그런 다음 장치가 올바른 결과 ([dimA-1..0]) 를 만들었는지 호스트에서 검사합니다.

// includes, system
#include
#include

// Simple utility function to check for CUDA runtime errors
void checkCUDAError(const char* msg);


// Part3: implement the kernel
__global__ void reverseArrayBlock(int *d_out, int *d_in)
{
int inOffset = blockDim.x * blockIdx.x;
int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);
int in = inOffset + threadIdx.x;
int out = outOffset + (blockDim.x - 1 - threadIdx.x);
d_out[out] = d_in[in];
}
/////////////////////////////////////////////////////////////////////
// Program main
/////////////////////////////////////////////////////////////////////
int main( int argc, char** argv)
{
// pointer for host memory and size
int *h_a;
int dimA = 256 * 1024; // 256K elements (1MB total)


// pointer for device memory
int *d_b, *d_a;


// define grid and block size
int numThreadsPerBlock = 256;


// Part 1: compute number of blocks needed based on
// array size and desired block size
int numBlocks = dimA / numThreadsPerBlock;


// allocate host and device memory
size_t memSize = numBlocks * numThreadsPerBlock * sizeof(int);
h_a = (int *) malloc(memSize);
cudaMalloc( (void **) &d_a, memSize );
cudaMalloc( (void **) &d_b, memSize );


// Initialize input array on host
for (int i = 0; i < dimA; ++i)
{
h_a[i] = i;
}


// Copy host array to device array
cudaMemcpy( d_a, h_a, memSize, cudaMemcpyHostToDevice );


// launch kernel
dim3 dimGrid(numBlocks);
dim3 dimBlock(numThreadsPerBlock);
reverseArrayBlock<<< dimGrid,
dimBlock >>>( d_b, d_a );


// block until the device has completed
cudaThreadSynchronize();


// check if kernel execution generated an error
// Check for any CUDA errors
checkCUDAError("kernel invocation");


// device to host copy
cudaMemcpy( h_a, d_b, memSize, cudaMemcpyDeviceToHost );


// Check for any CUDA errors
checkCUDAError("memcpy");


// verify the data returned to the host is correct
for (int i = 0; i < dimA; i++)
{
assert(h_a[i] == dimA - 1 - i );
}


// free device memory
cudaFree(d_a);
cudaFree(d_b);


// free host memory
free(h_a);


// If the program makes it this far, then the results are
// correct and there are no run-time errors. Good work!
printf("Correct!
");

return 0;
}
void checkCUDAError(const char *msg)
{
cudaError_t err = cudaGetLastError();
if( cudaSuccess != err)
{
fprintf(stderr, "Cuda error: %s: %s.
", msg,
cudaGetErrorString( err) );
exit(EXIT_FAILURE);
}
}

이 프로그램 구성의 핵심은 배열 d_a 와 d_b 모두 장치의 전역 메모리에 상주한다는 점입니다. CUDA SDK 는 장치의 특성에 대한 약간의 정보를 알 수 있는 bandwidthTest 라는 예제를 제공합니다. 제 시스템에서 전역 메모리 밴드폭은 대략 60 GB/s 를 가리켰습니다. 물론 대단한 값입니다만, CUDA 장치들은 매우 빠른 부동소숫점 계산 기능을 갖춘 128 개의 하드웨어 쓰레드에 데이터를 공급해 주어야 합니다. 32 비트 부동소숫점 값은 4 바이트를 차지하고, 전역 메모리 밴드폭으로 인하여 장치의 계산 능력이 15 GF/s 혹은 사용 가능한 계산 능력의 일부분으로 한정되어 버립니다. (이 결과는 프로그램이 전역 메모리를 읽기만 할 뿐 값을 기록하는 것을 제외했습니다.) 확실히 고성능 프로그램에서는 어떤 방식으로든 데이터를 재사용해야 합니다. 이것이 바로 공유 메모리와 레지스터 메모리의 기능이며, 프로그래머는 이들 메모리를 활용하여 최대 효율을 끌어내야 합니다. 기계에서 부동소숫점 계산 능력과 메모리 밴드폭과의 균형에 대한 이해(혹은 그 외 다른 특성)를 원하시는 분은, 제 글 HPC Balance and Common Sense(http://www.scientificcomputing.com/ShowPR.aspx?PUBCODE=030&ACCT=3000000100&ISSUE=0702&RELTYPE=PR&ORIGRELTYPE=HPCC&PRODCODE=00000000&PRODLETT=E) 를 읽어보세요.

- 공유 메모리 버젼

이번 소스코드는 다음에 소개할 내용이 담긴 코드 arrayReversal_multiblock_fast.cu 입니다. 공유 메모리가 이 문제에 어떻게 적용되는지 보실 수 있도록 이 소스를 지금 제공하겠습니다.

// includes, system
#include
#include

// Simple utility function to check for CUDA runtime errors
void checkCUDAError(const char* msg);


// Part 2 of 2: implement the fast kernel using shared memory
__global__ void reverseArrayBlock(int *d_out, int *d_in)
{
extern __shared__ int s_data[];


int inOffset = blockDim.x * blockIdx.x;
int in = inOffset + threadIdx.x;


// Load one element per thread from device memory and store it
// *in reversed order* into temporary shared memory
s_data[blockDim.x - 1 - threadIdx.x] = d_in[in];


// Block until all threads in the block have
// written their data to shared mem
__syncthreads();


// write the data from shared memory in forward order,
// but to the reversed block offset as before


int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);


int out = outOffset + threadIdx.x;
d_out[out] = s_data[threadIdx.x];
}
/////////////////////////////////////////////////////////////////////
// Program main
/////////////////////////////////////////////////////////////////////
int main( int argc, char** argv)
{
// pointer for host memory and size
int *h_a;
int dimA = 256 * 1024; // 256K elements (1MB total)


// pointer for device memory
int *d_b, *d_a;


// define grid and block size
int numThreadsPerBlock = 256;


// Compute number of blocks needed based on array size
// and desired block size
int numBlocks = dimA / numThreadsPerBlock;


// Part 1 of 2: Compute number of bytes of shared memory needed
// This is used in the kernel invocation below
int sharedMemSize = numThreadsPerBlock * sizeof(int);


// allocate host and device memory
size_t memSize = numBlocks * numThreadsPerBlock * sizeof(int);
h_a = (int *) malloc(memSize);
cudaMalloc( (void **) &d_a, memSize );
cudaMalloc( (void **) &d_b, memSize );


// Initialize input array on host
for (int i = 0; i < dimA; ++i)
{
h_a[i] = i;
}


// Copy host array to device array
cudaMemcpy( d_a, h_a, memSize, cudaMemcpyHostToDevice );


// launch kernel
dim3 dimGrid(numBlocks);
dim3 dimBlock(numThreadsPerBlock);
reverseArrayBlock<<< dimGrid, dimBlock,
sharedMemSize >>>( d_b, d_a );


// block until the device has completed
cudaThreadSynchronize();


// check if kernel execution generated an error
// Check for any CUDA errors
checkCUDAError("kernel invocation");


// device to host copy
cudaMemcpy( h_a, d_b, memSize, cudaMemcpyDeviceToHost );


// Check for any CUDA errors
checkCUDAError("memcpy");


// verify the data returned to the host is correct
for (int i = 0; i < dimA; i++)
{
assert(h_a[i] == dimA - 1 - i );
}


// free device memory
cudaFree(d_a);
cudaFree(d_b);


// free host memory
free(h_a);


// If the program makes it this far, then results are correct and
// there are no run-time errors. Good work!
printf("Correct!
");

return 0;
}
void checkCUDAError(const char *msg)
{
cudaError_t err = cudaGetLastError();
if( cudaSuccess != err)
{
fprintf(stderr, "Cuda error: %s: %s.
", msg,
cudaGetErrorString( err) );
exit(EXIT_FAILURE);
}
}

다음 회에서는 성능 향상을 위한 공유 메모리 사용법을 살펴보겠습니다. 그 때까지 __shared__, __constant__, register memory 와 같은 CUDA 메모리 형태에 대해서 살펴보세요.