[Programming / NVIDIA CUDA] 1D 배열 인덱스 변환

CUDA 프로그래밍을 하다보면 2차원 또는 3차원의 Grid, Block 메모리 구조를 1차원의 배열 인덱스로 변환해야 할 경우가 있다.

예를 들어 다음과 같이 block과 thread 메모리 공간을 할당할 수 있다.

dim3 blocks( GridDimX, GridDimY );

dim3 threads( BlockDimX, BlockDimY );

blocks는 Grid 내의 block 메모리 공간의 차원을 정의하며, threads는 Block 내의 thread 메모리 공간의 차원을 정의한다.

두 개 모두 기본적으로는 3차원의 구조를 가지고 있으며, Z에 대한 차원이 정의되지 않은 경우 Z의 차원은 1로 정의된다.

즉, GrdiDimZ = 1, BlockDimZ = 1이다.

다음과 같이 kernel 함수를 정의한다고 하자.

__global__ void my_kernel( float *dst ) {

    // map from threadIdx/BlockIdx to offset
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;

    // put any value
    dst[offset] = offset;

}

my_kernel() 함수는 다음곽 같은 형식으로 호출될 것이다.

my_kernel<<<blocks,threads>>>( &dst );

x, y는 각각 Grid, Block 구조에서 해당 thread의 전역 좌표이다.

예를 들어, blocks과 threads의 차원을 다음과 같이 설정해 보자.

dim3 blocks( 2, 3 );

dim3 threads( 4, 2 );

이 경우, my_kernel() 함수에서 정해지는 각 dimension은 다음과 같이 할당된다 :

gridDim.x = 2

gridDim.y = 3

blockDim.x = 4

blockDim.y = 2

이 구조로 할당된 Grid 및 Block Dimension은 다음 그림과 같은 구조를 갖는다.

위의 그림에서 각 셀(cell)은 thread를 의미한다. 주황색 셀은 Block(blockIdx.x = 1, blockIdx.y = 0) 내에 있는 (threadIdx.x = 0, threadIdx.y = 1) 인덱스의 thread 이며, thread의 전역 인덱스 (x, y) 는 (4, 3)이다.

전역 인덱스 (x, y)는 다음과 같이 계산된 것이다 :

x = threadIdx.x + blockIdx.x * blockDim.x = 0 + 1 * 4 = 4

y = threadIdx.y + blockIdx.y * blockDim.y = 1 + 1 * 2 = 3

그러면 offset은 어떻게 계산되는가?

전체 thread를 아래 그림과 같이 일렬로 배치한다고 생각해 보자.

일렬로 배치되었을 때 주황색 셀에 해당하는 thread는 다음과 같이 계산된다 : 

offset = x + y * blockDim.x * gridDim.x = 4 + 3 * 4 * 2 = 4 + 24 = 28

즉, 해당 thread는 1D 배열로 늘어뜨릴 경우 28번째에 위치하게 된다.

이 공식들을 이용하여 3D 메모리 구조에 대해서도 유사한 방식으로 적용할 수 있다.

x = threadIdx.x + blockidx.x * blockDim.x

y = threadIdx.y + blockidx.y * blockDim.y

z = threadIdx.z + blockidx.z * blockDim.z

offset은 

offset = x + ( y * blockDim.x * gridDim.x ) + ( z * blockDim.x * gridDim.x * blockDim.y * gridDim.y )

과 같이 계산된다.