이 글은 CUDA 기반 GPU 병럴처리 프로그래밍 책의 내용을 바탕으로 정리한 글이다.
저번 글에 이어서 오늘은 인덱싱에 대해 다루고자 한다. 여담이지만, 상용화 단계에 들어가기 위해 성능을 올리는 모델이 있다. 매번 느끼는 거지만, 데모 수준과 상용화 수준은 너무나도 큰 차이가 존재한다. 매우 지루한 과정이긴 하지만, 이번 역시 해내야겠지. 이번 상용화 솔루션이 나오고 나면 다음번에는 아키텍쳐를 조금 바꿔서 기존 상용화 상품의 모델을 바꾸는 일을 할 거다. 이제 슬슬 기초부터 탄탄히 잡고 가야할 시긴데, exponential하게 실력이 늘던 때와는 양상이 조금 달라 지루하기도 하지만 여튼 해야한다. 각설하고 인덱싱에 대해 한번 정리하겠다.
메모리 속 배열의 모습
우선 메모리에 대한 이해가 조금 필요하다. 각 스레드가 원하는 데이터에 접근하도록 하려면 메모리 속 배열 저장의 형태를 기억하는 것이 좋다. 1차원 부터 시작하여, 3차원까지 확장해 나가겠다.
우리가 1차원 배열을 생각할 때, 논리적 구조의 형태는 위와 같다. 물론 1차원까지는 메모리 속 형태 역시도 같다. 문제는 2 차원부터인데, 메모리는 serial 한 모습을 유지하는 반면, 논리적 구조의 모습이 달라진다.
2 차원에서 메모리 구조를 생각하면 자연스레 우리는 행렬의 형태로 데이터를 그리게 된다. 이는 그림에서 위쪽 그림의 모습이다. 그러나, 메모리는 앞서도 언급했지만 차원의 개념이 없다. 저장 공간이 일렬로 나열된다. 저장 규칙은 낮은 차원에서 높은 차원의 순서로 저장한다. 풀어서 더 설명하자면 2차원은 x, y 차원으로 나타내며, 위 그림의 예시는 ( 4 x 3 ) 형태이고 (x0, y0) , (x1, y0), (x2, y0), ... , (x3, y2) 의 순서로 저장됨을 의미한다. 즉 y차원이 돌기 위해서는 x차원의 데이터를 모두 순회해야 한다. 이를 코드에서 배열을 빌려와 다시 한번 설명하자면, arr[0][0], arr[0][1], ... , arr[2][3] 까지 간다는 의미이다. 우리가 표기하는 방식과 배열을 표현하는 방식이 뒤집어져 있어 조금은 헷갈릴 여지가 있다. 우리는 좌표를 ( x, y, z, ... ) 순서로 적고 컴퓨터는 arr[...][z][y][x] 로 표현한다고 기억하면 된다. 3 차원은 어떻게 나타나는지만 보겠다
메모리 속 형태를 한번 더 풀어서 설명하자면 2차원에서 데이터를 채우는 방식과 마찬가지로, y0,z0 에서 x방향으로 데이터를 모두 채우고, y1 방향으로 증가후 y1, z0에서 x 방향으로 데이터를 채운다. 이렇게 y 방향으로의 증가가 끝나면, z 방향으로 증가하시키고 앞의 과정을 반복하는 형태로 데이터를 채우게 된다.
스레드의 전역 번호
스레드 계층의 구조를 다시 한번 상기하고 가면, thread, block, grid 형태로 이루어져 있다. 이 중에서 thread의 layout을 잡아주는 내부 변수는 gridDim과 blockDim이었다. 그리고 각각은 3차원까지 표현이 된다. 즉 가장 큰 부분에서 들어가면 gird가 생성되고 그 안에서 block의 형태는 3차원까지 생성 될 수 있다. 그후 각 block은 thread 의 형태를 결정하며, 각 block안에서 thread는 3차원까지 생성 될 수 있다. 이는 kernel이 호출 될 때 생성되고, 한 명령어당 32개의 thread가 동시에 수행되며, 각 thread는 하나의 cuda core를 사용한다.
이 과정을 아래서부터 다시 한번 상기하면, 한 block이 thread를 가지는 데 그 형태는 3차원까지이고, 각 block 안에서 thread id는 고유하다. 그리고 이 block들은 gird 안에서 3차원 까지의 형태를 가질 수 있고, 각 block의 고유 id를 부여받는다. 그럼 우리가 thread 1개 당 고유 id를 부여하기 위해서는 block 내에서 thread의 id를 부여하고 block id를 활용하여 실제 전역 id를 부여해야한다.
1, 2차를 생략하고 3차원으로 바로 가겠다.
blockDim 이라는 변수는 block내의 thread 의 형태를 정의했고 그 형태를 활용하여 block내의 thread의 개수를 알 수 있었다. 차원은 x, y, z 였고, 저장 규칙에 따라 x 탐색 후 y 증가 z 고정 -> y 증가가 끝나면, z 증가 후 다시 앞의 과정 반복이었다.
이제 순서에 맞춰 인덱싱하는 작업을 해보겠다.
1. z 고정, x 탐색 후 y 증가 -> 이 과정은 z가 고정되어 있음으로 2 차원에서 thread id를 부여하는 것과 같은 말이다. 따라서 이 과정을 2D_BLOCK_TID 로 정의 하겠다.
-> 우리가 알고 싶은 것은 z차원이 고정되어 있는 상황에서, 특정 (x,y) 의 고유id 이다.
1-1. y 차원까지 고정되었을 때, x의 thread id는 threadIdx.x 로 표현된다.
1-2. x 차원이 가지고 있는 thread의 수는 blockDim.x 이다.
1-3. 특정 x 차원에서의 thread id는 threadIdx.x 로 표현되고 특정 y 차원에서 threadid는 threadIdx.y로 표현된다.
1-4. blockDim.x * threadIdx.y 는 특정 (x,y) 까지 오기 전 thread의 총 수를 나타낸다
1-5. 따라서 특정 (x,y)의 threadID, 즉 2D_BLOCK_TID = blockDim.x * threadIdx.y + threadIdx.x 이다
2. 이제 차원의 제한을 풀겠다.
2-1 z 차원 하나가 가지고 있는 총 thread의 수는 blockDim.x*blockDim.y 이다.
2-2 특정 z 차원은 threadIdx.z로 표현된다.
2-3 특정 z 차원이 가지고 threadID는 ( blockDim.x * blockDim.y ) * threadIdx.z + 2D_BLOCK_TID 가 된다.
-> 즉 우리가 구하는 TID_IN_BLOCK = ( blockDim.x * blockDim.y ) * threadIdx.z + 2D_BLOCK_TID 이다
3. 스레드의 전역 번호
-> 우리가 하고자 하는 것은 gird ( x, y, z ) 의 block ( x, y, z ) 일 때 전역 thread ID를 부여하는 것이다.
3-1 y차원과 z차원이 고정되어 있다고 하면 1차원의 grid 내 block을 생각할 수 있다.
3-2 block 하나에 속한 스레드의 개수는 blockDim.z * blockDim.y * blockDim.z 이다
3-3 gird의 y,z 차원이 고정되어 있을 때, 특정 x차원에서 자신이 속한 block의 번호는 blockIdx.x 이다
3-4 따라서 1D_GRID_TID = ( blockDim.z*blockDim.y*blockDim.x)*blockIdx.x + TID_IN_BLOCK이 된다.
3-5 gird의 y,z 차원이 고정되었다고 하면, x차원만을 고려하게 되고 이는 1차원 grid라고 볼 수 있다.
3-6 1차원 grid( x 차원 ) 가 가지고 있는 thread의 수는 gridDim.x * NUM_THEAD_IN_BLOCK이다.
3-6-1 NUM_THEAD_IN_BLOCK = blockDim.z*blockDim.y*blockDim.x 이었다.
3-6-2 따라서 1차원 grid가 가지고 있는 thread의 수는 gridDim.x*(blockDim.z*blockDim.y*blockDim.z) 이다
3-7 gird 에서 z차원을 고정 시킨 상태에서, 특정 gird y차원에서 자신이 속한 block의 번호는 blockIdx.y로 표현된다
3-8. 따라서 특정 grid의 y차원이 있는 thread의 수( 2D_GRID_TID )는 ( grid x차원이 가지고 있는 수 ) * ( blockIdx.y ) + 1D_GRID_TID 이다
3-8-1 grid x 차원이 가지고 있는 thread 수는 (blockDim.z*blockDim.y*blockDim.x)*gridDim.x 였다.
3-8-2 따라서 특정 y차원 전까지 가지고 thread 수는 (blockDim.z*blockDim.y*blockDim.x)*girdDim.x*blockIdx.y 이다
3-9 이를 3차원으로 확장하면, GLOBAL_TID = (blockDim.z*blockDim.y*blockDim.x)*gridDim.x*gridDim.y*blockIdx.z + 2D_GRID_TID 이다.
다소 복잡할 수 있다. 그러나 생각해보면 같은 과정을 반복한 것에 불과하다. 이는 layout과 그를 활용한 indexing 임으로 항상 숙지해야 한다.
'CUDA C' 카테고리의 다른 글
| [ CUDA 기반 GPU 병렬처리 프로그래밍 ] chapter 4-5 스레드 계층, 레이아웃과 인덱싱 1 (0) | 2023.11.08 |
|---|---|
| [ CUDA 기반 GPU 병렬처리 프로그래밍 ] chapter 2-3 CUDA 프로그램 구조와 흐름 (0) | 2023.11.06 |
| [ CUDA 기반 GPU 병렬처리 프로그래밍 ] chapter 1 GPGPU 및 병렬처리 개요 (2) | 2023.11.01 |
