Global Index와 Boundary Check

1D CUDA indexing의 핵심 공식은 하나다.

int i = blockIdx.x * blockDim.x + threadIdx.x;

이 i가 전체 배열에서 현재 thread가 처리할 위치다.

왜 이 공식인가

전체 위치는 두 부분으로 나뉜다.

현재 block의 시작 위치
+
block 안에서 thread의 위치

blockDim.x = 4이면 block 하나가 thread 4개를 가진다.

block 0 시작 위치 = 0 * 4 = 0
block 1 시작 위치 = 1 * 4 = 4
block 2 시작 위치 = 2 * 4 = 8

여기에 threadIdx.x를 더한다.

block 2, thread 0 -> 2 * 4 + 0 = 8
block 2, thread 1 -> 2 * 4 + 1 = 9
block 2, thread 2 -> 2 * 4 + 2 = 10
block 2, thread 3 -> 2 * 4 + 3 = 11

그래서:

global index = block start + local thread offset

block 수는 보통 올림 나눗셈으로 계산한다.

int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;

예를 들어 n = 17, threadsPerBlock = 4이면 block은 5개가 필요하다.

총 thread 수 = 5 * 4 = 20
실제 원소 수 = 17
초과 thread = i 17, 18, 19

그래서 kernel 안에서 막아야 한다.

if (i < n) {
    c[i] = a[i] + b[i];
}

이 조건이 없으면 배열 바깥 메모리에 접근할 수 있다.