다차원 그리드와 데이터 매핑

3.1 다차원 그리드 구성

벡터 덧셈은 데이터가 1차원이라 스레드도 1차원으로 늘어놓으면 끝이었다. 그러나 우리가 다룰 진짜 데이터는 거의 다 2차원 이상이다. 이미지는 (행, 열), 행렬도 (행, 열), 동영상은 (시간, 행, 열), 의료 볼륨은 (z, y, x). 매번 1차원으로 평탄화해 인덱스 계산을 수동으로 하는 것은 번거롭고 실수의 온상이다. CUDA는 그래서 처음부터 그리드와 블록을 1D/2D/3D로 잡을 수 있게 해 두었다.

모양을 지정하는 타입이 dim3다. 1D, 2D, 3D 어느 쪽이든 같은 타입으로 표현한다.

dim3 block(16, 16);              // 16×16 = 256 스레드/블록 (2D)
dim3 grid((W + 15) / 16,
          (H + 15) / 16);        // W×H 픽셀 이미지 덮기

myKernel<<<grid, block>>>(...);

커널 안에서는 threadIdx.x, .y, .z와 blockIdx.x, .y, .z, blockDim.x, .y, .z가 다 살아 있다(쓰지 않는 차원은 1로 셋업된다). 이미지를 다룰 때 (col, row) = (x, y) 매핑을 쓰는 게 관례다. 즉 .x가 가로(열), .y가 세로(행)다. 처음엔 헷갈리지만 한 번 익혀 두면 일관성이 무기가 된다.

왜 굳이 다차원으로 잡을까? 1차원으로도 충분히 일은 된다. 두 가지 이유가 있다. 첫째, 코드 가독성이다. 2D 데이터에 2D 스레드는 머릿속에서 그대로 그려진다. 인덱스 산수 한 번 줄이는 것 이상으로, "내가 픽셀 (col, row)에서 일을 하고 있다"가 자연스럽게 읽힌다. 둘째, 메모리 합치기(coalescing)와의 궁합이다. 같은 워프(연속한 32 스레드)가 메모리에서 연속된 주소를 읽을 때 GPU는 한 번의 트랜잭션으로 묶어 읽는다. 행 우선(row-major) 배열에서 같은 행의 인접 열들이 연속 주소이므로, threadIdx.x가 열에 매핑되면 자연히 합쳐진 접근이 된다. 세부는 6장에서 다루지만, 매핑 관례를 일찍 익히는 게 그 미래를 편하게 한다.

3.2 다차원 데이터에 스레드 매핑

2D 이미지에서 픽셀 (col, row)에 한 스레드씩 매핑하려면 다음 두 줄을 거의 모든 커널 첫머리에 둔다.

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

if (col < W && row < H) {
    int idx = row * W + col;     // 행 우선 1D 평탄화
    out[idx] = doSomething(in[idx]);
}

여기서 row * W + col 패턴이 핵심이다. 메모리상 배열은 결국 1차원으로 쭉 깔려 있다. C/C++의 행 우선 배치에서는 같은 행의 원소들이 연속해 있다. 그래서 위 인덱스 계산이 표준이 된다. 만약 행렬이 column-major(예: BLAS, Fortran 관례)라면 col * H + row가 된다. 어느 쪽을 쓰든 한 프로젝트 안에서 일관되어야 한다 ─ 섞으면 디버깅이 지옥이다.

경계 검사 if (col < W && row < H)도 빠뜨리지 말 것. 그리드를 천장(ceiling)으로 잡았으니 가장자리 블록은 항상 일부 스레드가 범위 밖이다. 이 검사 한 줄이 메모리 침범을 막는다.

3D로 가면 자연스럽게 idx = z * H * W + y * W + x가 된다. 의료 영상의 복셀, 3D 시뮬레이션 격자에서 그대로 쓰인다. 차원이 늘어도 패턴은 같다 ─ 가장 안쪽이 .x, 그 밖이 .y, 가장 바깥이 .z.

그림 3.1 — 너비가 블록 폭의 배수가 아닐 때 마지막 블록 일부는 비활성

3.3 첫 2D 커널: 이미지 블러

3×3 박스 블러를 구현해 보자. 한 픽셀의 새 값은 자기와 8 이웃의 평균(총 9 픽셀의 산술 평균)이다. 입력은 1채널 8비트 그레이스케일 이미지라고 하자.

__global__ void boxBlur3x3(const unsigned char* in,
                           unsigned char* out,
                           int W, int H) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    if (col >= W || row >= H) return;

    int sum = 0, cnt = 0;
    for (int dy = -1; dy <= 1; ++dy) {
        for (int dx = -1; dx <= 1; ++dx) {
            int r = row + dy;
            int c = col + dx;
            if (r >= 0 && r < H && c >= 0 && c < W) {
                sum += in[r * W + c];
                ++cnt;
            }
        }
    }
    out[row * W + col] = (unsigned char)(sum / cnt);
}

// 호스트에서:
//   dim3 block(16, 16);
//   dim3 grid((W + 15) / 16, (H + 15) / 16);
//   boxBlur3x3<<<grid, block>>>(d_in, d_out, W, H);

경계 처리 방식 두 가지를 짚어 두자. 위 코드에서는 유효 이웃의 평균을 쓴다. 즉 가장자리 픽셀은 5~6개의 이웃만으로 평균 낸다. 다른 흔한 선택지는 제로 패딩(범위 밖은 0으로 간주)과 경계 복제(가장 가까운 안쪽 값을 복사)다. 어느 것을 고를지는 응용에 따라 다르다. 여기서는 가장자리 명도가 어두워지지 않게 유효 평균을 썼다.

한 가지 미리 짚을 것. 이 커널의 모든 메모리 접근은 글로벌 메모리(DRAM)로 직행한다. 한 출력 픽셀당 9번 입력을 읽는데, 이웃 픽셀들은 인접 스레드 사이에서 겹친다. 즉 같은 데이터를 여러 번 DRAM에서 끌어온다. 5장에서 공유 메모리 타일링으로 이 중복을 제거할 것이다. 그때까지는 일단 "동작하는 단순 버전"으로 두는 게 정석이다 ─ 최적화 전에 정확성, 정확성 전에 단순함.

3.4 행렬 곱셈: 이 책의 단골 손님

행렬 곱은 GPU 컴퓨팅의 "Hello, World" 너머의 첫 번째 진지한 예제이고, 이 책 전체에 걸쳐 계속 진화한다. 3장에서는 가장 단순한 글로벌 메모리 버전, 5장에서는 공유 메모리 타일링, 6장에서는 메모리 합치기 최적화, 그리고 후반부에서는 텐서 코어 활용까지 ─ 같은 문제를 점점 빠르게 만든다. 그러니 지금 이 단순 버전을 머릿속에 잘 새겨 두자. 이후 모든 비교의 기준선이 된다.

크기 M×K 행렬 A와 K×N 행렬 B의 곱 C = A·B를 구한다. C의 (row, col) 원소는

C[row, col] = Σ_k=0..K−1 A[row, k] · B[k, col]

각 출력 원소가 다른 출력 원소와 독립이라는 점이 보인다. 그러니 한 스레드에 한 C 원소를 맡기는 매핑이 자연스럽다.

__global__ void matMulNaive(const float* A, const float* B, float* C,
                            int M, int N, int K) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;   // C의 행
    int col = blockIdx.x * blockDim.x + threadIdx.x;   // C의 열
    if (row >= M || col >= N) return;

    float acc = 0.0f;
    for (int k = 0; k < K; ++k) {
        acc += A[row * K + k] * B[k * N + col];
    }
    C[row * N + col] = acc;
}

코드는 간결하다. 그러나 성능 분석을 잠깐 해 보자. 출력 원소 한 개를 만들기 위해 A의 한 행(K개)과 B의 한 열(K개)을 글로벌 메모리에서 읽는다. 즉 1번의 출력당 2K번 읽기와 K번 곱셈/덧셈. 산술 강도(FLOP/byte)는 대략 2K FLOP / (8K바이트) = 0.25 FLOP/byte. 너무 낮다. 게다가 인접 출력 원소들 사이에 같은 데이터가 거듭 사용된다 ─ C[row, col]과 C[row, col+1]은 둘 다 같은 A 행을 읽는다. 이 중복이 5장에서 공유 메모리로 제거된다.

그래도 지금 이 버전으로 1024×1024 행렬을 GPU에서 돌리면 CPU naive 버전보다 한참 빠르다. 무식한 병렬성이 이미 큰 일을 하기 때문이다. 단지 그 무식함이 GPU 성능의 5~10%만 쓰고 있을 뿐이다. 우리의 여정은 이 시작점에서 50%, 80%, 그리고 텐서 코어 시대엔 이론치에 근접한 곳까지 가게 된다.

3.5 정리

이 장에서는 1차원 사고에서 다차원 사고로 옮겨갔다. dim3로 그리드와 블록을 잡고, (col, row) = (x, y) 관례로 스레드를 데이터 위에 펼치고, 행 우선 평탄화 row*W + col로 메모리에 접근하는 패턴을 익혔다. 이미지 블러에서는 이웃 픽셀을 읽는 패턴이, 행렬 곱에서는 한 출력 원소가 한 행과 한 열을 모두 훑는 패턴이 나왔다. 두 예제 모두 글로벌 메모리만 쓰는 단순 버전이라 성능은 아직 빛나지 않는다. 4장에서 그 성능이 왜 안 나오는지 ─ GPU 안에서 실제로 어떤 일이 일어나는지 ─ 를 들여다보고, 5장에서 공유 메모리로 드라마틱한 도약을 만든다.