CUDA Shared Memory Dynamic Allocation (동적할당)

- Shared Memory 동적할당 (Dynamic Allocation)

Shared Memory 는 선언하면서 초기화도 할 수가 없으며, 배열 선언시 배열의 개수는 Constant 여야 한다.

이럴 경우 extern 으로 Shared Memory를 선언하면서 배열의 크기를 지정하지 않을 수 있다.

[-] Collapse

extern __shared__ float sfArray [];

단 이럴 경우에는 반드시 kernel 함수를 호출할 때 Shared Memory의 크기를 지정해 줘야 한다.
(당연히 Shared Memory의 크기는 Byte 단위로 지정한다.)
왠만하면 GPU 스펙의 Shared Memory 최고크기를 지정해 놓으면 된다. (GTX 770의 경우 65535)

[-] Collapse

Cuda_TestKernel<<<NUM_BLOCK, NUM_THREAD, SIZE_SHARED_MEMORY>>>(dev_fArray);

1개의 Kernel 함수에 다수의 extern shared array가 있을 경우 그 시작번지수는 모두 동일 하다.

예를 들어서 1000개의 float, 300개의 int, 500개의 float이 필요할 경우 아래와 같이 구현이 가능하다.
[-] Collapse

__shared__ float sfDa;
extern __shared__ float sfArray [];
__shared__ float *sfCs;
__shared__ int *siDs;
__shared__ float *sfEs;

if (threadIdx.x == 0)
{
    sfDa = (float)blockIdx.x;
    sfCs = (float *)sfArray;
    siDs = (int *)&sfCs[1000];
    sfEs = (float *)&siDs[300];
}
__syncthreads();