[CUDA] 프로세서

CUDA Thread Hierachy와 CUDA Processor간의 관계

CUDA Processor와 CUDA의 논리적인 구조는 다음과 같은 상관 관계가 있습니다.

 

CUDA Architecture	CUDA Logical Architecture
CUDA Processor	CUDA Kernel/Grid
CUDA Multi-Processor	복수의 CUDA Block
CUDA Core	CUDA Thread

CUDA Thread Index 구성

CUDA Kernel을 개발하실때 주지해야 하는 사항은, 여러분이 짜는 CUDA 코드는 병렬적으로 처리되는 코드이지만, CUDA Thread 입장에서는 오직 하나의 Thread만이 동작하는 코드입니다. 따라서 CUDA Thread는 고유의 Index를 갖고 있으며, 이를 알아내기 위해서는 CUDA Thread Indexing Keyword에 대해서 알고 계셔야 합니다.

CUDA Programming 모델을 설명하면서 CUDA Thread Index를 위한 키워드 들을 잠깐 설명했었습니다.

 

키워드	설명	차원
gridDim	Kernel의 block의 수	x, y
blockIdx	CUDA block의 인덱스	x, y
blockDim	CUDA block의 크기	x, y, z
threadIdx	CUDA 쓰레드의 인덱스	x, y, z
warpSize	CUDA Instruction이 동시에 처리가능한 CUDA 쓰레드 수

10 x 18 크기의 공간을 [4 x 4] 크기의 타일로 덮는다고 했을때 몇개의 타일이 필요할까요? 이때 Tile은 쪼갤 수 없으며, 공간을 덮기만 하면 됩니다.

이것을 우리는 block의 크기로 먼저 나눌 것입니다.

단순히 전체 갯수 나누기 block의 크기로 계산하면, 10×18/4/4=11.2510×18/4/4=11.25개가 필요하므로 필요한 타일의 수는 11개 또는 12개라고 하실지도 모르겠습니다. 하지만 타일은 쪼갤 수 없다고 하였으므로 가로 5개, 세로 3개로 해서, [10 x 18]의 공간을 [4 x 4] 크기의 타일로 덮으려면 15개의 타일이 필요합니다.

 

CUDA는 개념적으로 위 그림과 같이, A, B 등으로 라벨링한대로 지정된 CUDA Block의 크기에 맞춰서 데이터를 동시에 처리해줍니다. 그러다보면 C와 같이 데이터 경계를 넘어갈 수도 있습니다. CUDA에서는 이것을 CUDA Index 구성을 통해 우리가 원하는 형태로 CUDA Thread를 데이터에 매핑할 수 있습니다.

만약 타일의 크기가 [8 x 2]라고 한다면 어떻게 될까요. 2x9=182x9=18개가 될 것입니다.

 

CUDA의 Grid의 크기를 계산하는 방법도 이와 비슷합니다. 처리해야하는 데이터를 적당한 크기로 자르는 것이죠. 만약 [1920 x 1080] 크기의 데이터를 [16 x 16] 크기의 CUDA Block으로 나눠준다면 총 CUDA block의 수는 120 x 68개가 필요하게 됩니다.

 

이 경우에 대하여,

 width=1920,

height=1080,

blockWidth=16,

blockHeight=16

 

이라고 할 수 있습니다. CUDA Keyword로는 다음과 같이 정리됩니다.

GridDim.x	68
GridDim.y	120
BlockDim.x	16
BlockDim.y	16

 

여기서 우리는 BlockDim의 값을 [16, 16]이라고 정해줬듯이, Grid Dimension도 정해주어서 CUDA가 적절한 수의 CUDA Block을 생성할 수 있도록 해주어야하며, 데이터마다 CUDA block을 Mapping시켜서 처리하는 경우 다음과 같이 코드를 구성하시면 됩니다. 다만 아래 코드는 식을 보여드리기 위한 것이며, block의 크기는 앞서 CUDA Programming Model에서 설명했던 것 처럼 Kernel를 호출할때 제어 파라미터로 넘겨주셔야 합니다.

blockIdx 키워드는 우리가 앞서 쪼갰던 CUDA block의 위치를 의미하며, CUDA Thread는 자신이 속한 blockIdx의 값을 알고 있으므로 그냥 참조하시면 됩니다.

threadIdx는 CUDA block 내에서 존재하는 CUDA thread의 Index입니다. blockIdx와 마찬가지로 각각의 CUDA Thread는 자신만의 CUDA thread Index를 알고 있는데, 이는 상대적인 위치라고 할 수 있습니다. 다만, 이 때문에, CUDA Thread는 어떤 위치에 있는 데이터를 처리해야하는지 알기 위해선 절대적인 인덱스 주소값을 알아야 합니다.

 

그래서 이전에 CUDA Programming 모델링에 대해서 설명드릴때 아래 표를 보여드리면서 CUDA Thread의 인덱싱은 이대로 한다고 말씀드렸던 것입니다.

인덱싱	종류코드
1차원 데이터(x only)	blockDim.x * blockIdx.x + threadIdx.x
2차원 데이터(x, y)	blockDim.x * blockIdx.x + threadIdx.x, blockDim.y * blockIdx.y + threadIdx.y
2차원 데이터(x only)	N.x * blockDim.y * blockIdx.y + N.x * threadIdx.y + blockDim.x * blockIdx.x + blockIdx.x

 

 

EXPERIENTIA DOCET| CUDA 프로세서에 대한 이해