Grid, Block, Thread 좌표계

CUDA kernel launch는 이런 모양이다.

kernel<<<gridDim, blockDim>>>();

1D launch에서는 다음처럼 쓸 수 있다.

kernel<<<3, 4>>>();

의미는:

gridDim.x = 3
blockDim.x = 4

즉 grid 안에 block이 3개 있고, block 하나 안에 thread가 4개 있다.

CUDA가 자동으로 주는 좌표

kernel 안에서는 다음 값들을 사용할 수 있다.

blockIdx.x
threadIdx.x
blockDim.x
gridDim.x

범위는 다음과 같다.

blockIdx.x  = 0 ... gridDim.x - 1
threadIdx.x = 0 ... blockDim.x - 1

kernel<<<3, 4>>>이면:

blockIdx.x  = 0, 1, 2
threadIdx.x = 0, 1, 2, 3

threadIdx는 block마다 다시 시작한다

threadIdx.x는 전체 grid에서의 번호가 아니라, 현재 block 안에서의 번호다.

block 0: thread 0 1 2 3
block 1: thread 0 1 2 3
block 2: thread 0 1 2 3

그래서 blockIdx.x와 threadIdx.x를 함께 써야 전체 위치를 만들 수 있다.

mental model

grid  = blocks의 배열
block = threads의 배열

blockIdx.x는 grid 안에서 block의 위치이고, threadIdx.x는 block 안에서 thread의 위치다.

인터랙티브 설계

• kernel<<<3, 4>>>를 입력하면 3개 block과 각 block의 4개 thread를 그린다.
• gridDim.x와 blockDim.x 슬라이더를 바꾸면 전체 좌표가 다시 계산된다.
• thread 하나를 클릭하면 blockIdx.x, threadIdx.x, blockDim.x, gridDim.x가 표시된다.

확인

• kernel<<<5, 8>>>이면 총 thread 수는 몇 개인가?
• threadIdx.x가 block마다 0부터 다시 시작하는 이유는 무엇인가?
• gridDim.x와 blockDim.x는 누가 정하는가?