Flattening과 Stride

행렬은 사람이 보기에는 2D다.

A =
[ 0  1  2
  3  4  5
  6  7  8 ]

하지만 메모리 안에서는 보통 한 줄로 저장된다.

A memory = [0, 1, 2, 3, 4, 5, 6, 7, 8]

그래서 (row, col) 좌표를 1D index로 바꿔야 한다.

int idx = row * width + col;

width 6 height 4

2D logical matrix

1D physical memory

왜 row * width인가

width = 3이라고 하자.

(row=0, col=0) -> idx = 0 * 3 + 0 = 0
(row=0, col=1) -> idx = 0 * 3 + 1 = 1
(row=0, col=2) -> idx = 0 * 3 + 2 = 2

(row=1, col=0) -> idx = 1 * 3 + 0 = 3
(row=1, col=1) -> idx = 1 * 3 + 1 = 4

row * width는 이전 row들이 차지한 원소 수다. col은 현재 row 안에서의 위치다.

stride

한 row 아래로 내려가려면 메모리에서 width만큼 이동한다.

stride of row = width

그래서 width는 단순한 matrix 크기가 아니라, row 방향 이동 간격이기도 하다.

1D grid로 2D matrix를 처리할 수도 있다

matrix add를 꼭 2D grid로 해야 하는 것은 아니다. 전체 원소 수를 N = height * width로 보고 1D로 처리할 수도 있다.

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

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

2D grid는 row/col 구조를 보존해서 이해하기 좋고, 1D grid는 elementwise 작업에 단순하다.

확인

• width = 6, row = 2, col = 3이면 idx는 몇인가?
• row * width는 무엇을 세는가?
• 2D matrix add를 1D grid로도 처리할 수 있는 이유는 무엇인가?