- 공유 메모리의 이해 (2)

지난 4회에서 실행 모델과 커널 실행 환경 설정이 레지스터의 갯수와 공유 메모리같은 지역 멀티프로세서 자원에 영향을 미치는지 살펴보았습니다. 이번에는 메모리 성능과 reverseArray_multiblock_fast.cu 내에서 공유 메모리 사용에 대해 살펴보겠습니다.

- CUDA 메모리 성능

지역 메모리와 광역 메모리에는 캐시가 설정되어 있지 않아서, 지역 혹은 광역 메모리 억세스 시 반드시 메모리 참조가 발생합니다. 그렇다면 각각의 메모리 종류에 따른 억세스에 드는 시간은 어느 정도일까요?

“워프” 에서 메모리 명령 하나가 수행되려면 멀티프로세서 네 클럭 사이클이 소요됩니다. 지역 혹은 광역 메모리 참조에는 400 에서 600 클럭 사이클 분량의 메모리 레이턴시가 초래됩니다. 예를 들어 아래에 있는 짜투리 코드를 실행하기 위해서는 광역 메모리에서 읽기 명령 네 클럭, 공유 메모리로 쓰기 명령 네 클럭, 그리고 광역 메모리로부터 부동소숫점 값을 읽어들이는 데 400 에서 600 클럭 사이클이 소요됩니다. 참고로 __device__ 변수 지정자는 해당 변수가 광역 메모리에 위치하게끔 지정합니다. (그 외의 변수 특성은 CUDA 프로그래밍 가이드 4.2.2.1 장을 참조하세요.) __device__ 형으로 지정된 변수는 호스트 코드에서 참조할 수 없습니다.

__shared__ float shared[32];
__device__ float device[32];
shared[threadIdx.x] = device[threadIdx.x];

억세스에 걸리는 시간이 100 에서 150 배 차이가 나는 광역 메모리 억세스를 최소화하고 지역 멀티프로세서 메모리 내 데이터를 재사용하도록 프로그램을 짜야만 합니다. CUDA 에는 쓰레드 스케쥴러가 잘 되어 있어서 광역 메모리 억세스 지연은 실행 환경 설정에서 큰 블럭을 할당함으로써 숨기고 커널 내에서는 레지스터, __shared__, __constant__ 형의 변수를 되도록 사용합니다.

공유 메모리는 칩 안에 내장되어 있어서 광역 메모리보다 훨씬 빠르고 최적화 기능으로 뱅크 충돌을 방지합니다. 어떤 문서에서는 공유 메모리가 레지스터 억세스만큼 빠르다고 언급할 만큼 공유 메모리 속도는 빠릅니다. 하지만 최근 CUBLAS 나 CUFFT 의 속도 향상은 공유 메모리를 덜 사용하게끔 해서 얻어진 것입니다. 그러니 되도록이면 레지스터를 이용해 주세요. CUDA 공유 메모리는 동일한 크기의 메모리 모듈인 메모리 뱅크로 나뉘어집니다. 메모리 뱅크는 32비트 크기의 데이터(정수나 단밀도 부동소숫점)를 연달아 가지고 있어서 연속적인 쓰레드의 연속적인 배열 억세스에 빠른 속도를 보입니다. 뱅크 충돌은 메모리 억세스가 동일한 뱅크에(동일한 어드레스 혹은 동일한 뱅크에 매핑된 다른 어드레스) 다발적으로 발생할 경우 일어납니다. 이럴 경우 하드웨어는 메모리 명령을 순차적으로 처리하게끔 하여, 모든 메모리 참조가 끝날 때까지 쓰레드를 강제로 대기시킵니다. 모든 쓰레드가 동일한 공유 메모리 어드레스를 읽어들여야 한다면 브로드캐스트 구조가 자동으로 실행되어 순차 실행을 방지합니다. 공유 메모리 브로드캐스트는 다수의 쓰레드가 동시에 데이터를 읽어들이는 효율적인 방법입니다. 공유 메모리를 사용해야 할 때 이 기능을 활용하는 것이 좋습니다.

향후 뱅크 충돌 문제에 대해 깊이 다루어 보겠습니다. 지금은 reverseArray_multiblock_fast.cu 에서는 연속적인 쓰레드가 연속적인 값을 참조하므로 뱅크 충돌이 없다고 간주합니다.

A quick summary of local multiprocessor memory types with read/write capability follows:

* Registers:
o The fastest form of memory on the multi-processor.
o Is only accessible by the thread.
o Has the lifetime of the thread.
* Shared Memory:
o Can be as fast as a register when there are no bank conflicts or when reading from the same address.
o Accessible by any thread of the block from which it was created.
o Has the lifetime of the block.
* Global memory:
o Potentially 150x slower than register or shared memory -- watch out for uncoalesced reads and writes which will be discussed in the next column.
o Accessible from either the host or device.
o Has the lifetime of the application.
* Local memory:
o A potential performance gotcha, it resides in global memory and can be 150x slower than register or shared memory.
o Is only accessible by the thread.
o Has the lifetime of the thread.

지역 멀티프로세서 메모리의 종류와 읽기/쓰기 성능에 대한 간추린 내용입니다.
* 레지스터:
- 멀티프로세서에서 가장 빠른 메모리
- 해당 쓰레드에서만 억세스 가능
- 해당 쓰레드 내에서 데이터 유지
* 공유 메모리:
- 뱅크 충돌이 없을 시 레지스터만큼 빠른 속도
- 설정된 블럭 내 쓰레드는 모두 억세스 가능
- 블럭 내 데이터 유지
* 광역 메모리:
- 레지스터나 공유 메모리보다 대략 150 배 느림 – 비 연동 억세스시 주의, 다음 장에서 다룰 예정
- 호스트와 장치 모두 억세스 가능
- 응용프로그램 내 데이터 유지
* 지역 메모리:
- 속도 저하의 주범, 광역 메모리 내에 존재하므로 레지스터나 공유 메모리보다 대략 150배 느림
- 해당 쓰레드에서만 억세스 가능
- 쓰레드 내 데이터 유지

- 공유 메모리 사용시 주의할 점

* 성능 저하를 일으키는 공유 메모리 뱅크 충돌 주의
* 커널 내에서 동적으로 할당되는 변수 시작위치는 동일합니다. 따라서 여러 개의 동적 공유 메모리 배열을 사용하려면 수동으로 오프셋을 지정해야 합니다. 예를 들어, 배열 a 와 b 를 공유 메모리에 동적 할당하려면 다음과 같이 해 주어야 합니다.

__global__ void kernel(int aSize)
{
extern __shared__ float sData[];
float *a, *b;

a = sData;
b = &a[aSize];

- 레지스터/지역 메모리 사용시 주의할 점

* 경우에 따라 레지스터 메모리가 지역 메모리에 설정될 수도 있으며, 이 경우 속도 저하의 원인이 됩니다. nvcc 에 –ptxas-options=-v 옵션으로 컴파일한 후 ptx 어셈블리 코드 혹은 출력에서 lmem 을 살펴보세요.
* 배열 내용이 컴파일시 상수 인덱스로 참조되면 배열은 레지스터에 상주할 수 있으나 변수 인덱스로 참조되면 레지스터에 상주할 수 없습니다. 개발시 배열 내용을 느린 광역 메모리 대신 레지스터 메모리에 상주토록 하기 위하여 루프 언롤링 기법을 사용해야 하는 어려움이 있습니다. 하지만 언롤링을 사용하면 레지스터 사용 빈도를 늘리게 되므로, 변수들을 지역 메모리에 저장해야 하고, 결국 루프 언롤링의 이득을 상쇄시켜 버립니다. nvcc 에 –maxrregcount=value 옵션을 넣어줌으로써 컴파일러가 더 많은 레지스터를 사용하게 할 수 있습니다. (최대 레지스터 카운트는 128 입니다.) 레지스터 사용과 쓰레드 생성에는 트레이드 오프 관계가 있어서, 메모리 속도 지연 감소를 방해합니다. 프로그램 구조에 따라서 이 옵션은 자원 부족을 일으켜 커널을 실행하지 못하게 되기도 합니다.

- 공유 메모리 커널

reverseArray_multiblock.cu 와 reverseArray_multiblock_fast.au 모두 동일한 작업을 수행합니다. 정수값 [0..dimA-1] 을 가지는 1차원 정수 배열 h_a 를 만듭니다. 배열은 장치 내 cudaMemory 로 옮겨지고 호스트는 reverseArrayBlock 커널을 실행하여 배열 내 값의 위치를 뒤바꿉니다. cudaMemory 는 호스트와 장치간 데이터 교환에 사용되어 장치가 정확한 값 ([dimA-1..0]) 을 만들었는지 검사를 수행합니다.
차이점은 reverseArray_multiblock_fast.cu 는 커널 속도 향상을 위하여 공유 메모리를 사용하지만, reverseArray_multiblock.cu 는 모두 전역 메모리만 사용합니다. 두 프로그램의 실행속도를 측정해서 성능 차이를 확인해 보세요. 그리고, reverseArray_multiblock.cu 는 전역 메모리를 비효율적으로 사용합니다. 이후에 CUDA 프로파일러를 이용해서 성능을 진단하고 고치는 것을 다룰 것이며, 새로운 10 종의 아키텍쳐가 많은 경우 이런 최적화 필요성을 줄여줌으로써 성능을 향상시키는지 알아볼 것입니다.

// 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 the 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 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);
}
}

실행시 소요되는 공유 메모리를 결정하려면 호스트와 장치 코드에 약간의 설정이 필요합니다. 이번 예제에서, 커널 내 블럭당 필요한 공유 메모리 (바이트) 는 호스트에서의 실행 환경 설정의 세 번째 옵션 인수에 넣습니다. (호스트 쪽 설정은 만약 공유 메모리가 커널 실행시 지정될 때에만 필요합니다. 컴파일시 값이 고정될 때에는 호스트 쪽 설정이 필요 없습니다.) 기본적으로 실행 환경 설정에서는 공유 메모리가 필요없는 것으로 간주합니다. 예를 들어, arrayReversal_multiblock_fast.cu 의 호스트 코드에서 다음의 짜투리 코드는 블럭 내 쓰레드 갯수와 동일한 값을 가지는 정수형 배열을 공유 메모리에 할당합니다.

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

reverseArrayBlock 커널을 보면, 공유 메모리는 다음과 같이 정의됩니다.

extern __shared__ int s_data[];

Note that the size is not indicated in the kernel -- rather it is obtained from the host through the execution configuration.
Until the next column on profiling, I recommend looking at the reverseArray_multiblock.cu. Do you think there is a performance problem in accessing global memory? If you think there is a problem, try to fix it.

커널 내에서 배열 크기가 지정되지 않고, 실행 환경 설정으로 호스트에서 결정됩니다.

다음 편 profiling 을 들어갈 때까지, reverseArray_multiblock.cu 를 살펴보세요. 만약 전역 메모리 억세스에 속도 문제가 있다고 생각되시면 한 번 고쳐보도록 하세요.