2D Matrix Add

1D에서는 전체 배열 위치 i 하나만 만들었다. 2D matrix에서는 row와 col을 만든다.

int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;

보통 CUDA 2D indexing에서는:

x축 = column
y축 = row

로 둔다.

width 6 height 4

2D logical matrix

1D physical memory

2D block

dim3 threadsPerBlock(3, 2);

이건 block 하나가 다음 모양이라는 뜻이다.

3 columns x 2 rows = 6 threads

(0,0) (1,0) (2,0)
(0,1) (1,1) (2,1)

2D grid

전체 matrix가 width x height이면 필요한 block 수는 축별로 올림 나눗셈한다.

dim3 blocksPerGrid(
    (width + threadsPerBlock.x - 1) / threadsPerBlock.x,
    (height + threadsPerBlock.y - 1) / threadsPerBlock.y
);

width = 6, height = 4, threadsPerBlock = (3, 2)이면:

blocksPerGrid = (2, 2)

kernel 안의 작업

if (row < height && col < width) {
    int idx = row * width + col;
    C[idx] = A[idx] + B[idx];
}

2D 좌표로 thread 위치를 정하지만, 실제 접근은 idx로 한다.

확인

• dim3 threadsPerBlock(3, 2)는 block 하나에 thread가 몇 개인가?
• row는 왜 y 축으로 계산하는가?
• col은 왜 x 축으로 계산하는가?
• row < height && col < width가 필요한 이유는 무엇인가?