이질 데이터 병렬 컴퓨팅과 CUDA 첫걸음

2.1 데이터 병렬성: "같은 일을 각자 자기 자리에서"

요리에 비유해 보자. 한 사람이 양파 천 개를 까야 한다고 하면 끔찍하다. 하지만 천 명이 모여 한 사람이 양파 한 개씩 맡으면, 시간은 양파 한 개 까는 시간으로 줄어든다. 이게 데이터 병렬성(data parallelism)의 본질이다. 같은 종류의 작업을 서로 독립적인 데이터 조각에 동시에 적용하는 것. 양파끼리 서로 의존하지 않고, 까는 절차는 모두 같다. 그래서 사람이 1000명이든 100명이든 같은 코드를 돌릴 수 있다.

컴퓨터 그래픽스가 GPU의 고향인 이유가 여기에 있다. 풀HD 화면은 약 207만 픽셀이다. 한 픽셀에 어떤 셰이더 함수를 적용한다고 할 때, 픽셀들은 서로를 신경 쓸 필요가 없다. 207만 번 함수를 호출하는 게 아니라 207만 개의 가상 일꾼을 동시에 풀어 한 번에 끝내는 것 ─ GPU가 잘하는 일이다.

실용 예 두 개를 들어 보자.

컬러 → 그레이스케일. 입력 RGB 이미지의 각 픽셀에 대해 gray = 0.299·R + 0.587·G + 0.114·B를 계산해 1채널 이미지를 만든다. 픽셀 i의 결과는 픽셀 j에 영향을 주지 않는다. 임베러싱하게(embarrassingly) 병렬이다.

이미지 블러. 한 픽셀의 새 값이 자기 주변 3×3 또는 5×5 이웃의 평균이라고 하자. 출력 픽셀끼리는 여전히 독립이다(입력에서만 읽고 출력에는 쓰니까). "한 픽셀당 한 스레드" 매핑이면 끝이다. 다만 이웃을 9번 읽어야 하니 메모리 접근 패턴은 단순 그레이스케일보다 풍부해진다 ─ 이게 5장에서 다룰 공유 메모리 타일링의 출발점이다.

2.2 CUDA C 프로그램의 구조: host와 device

CUDA C 프로그램은 두 종류의 코드가 한 파일 안에 공존한다. 호스트 코드(host code)는 CPU에서 도는 평범한 C/C++다. 메인 함수, 파일 입출력, 메모리 할당, 커널 호출 같은 일이 여기에 들어간다. 디바이스 코드(device code)는 GPU에서 도는 코드로, "커널(kernel)"이라 불리는 함수와 디바이스 보조 함수들이다. nvcc 컴파일러는 이 두 가지를 분리해 각각 다른 백엔드에 보낸다 ─ 호스트 코드는 g++/clang 같은 일반 컴파일러로, 디바이스 코드는 NVIDIA의 PTX/SASS 백엔드로.

함수가 어디에서 도는지 표시하는 키워드 세 가지를 외워 두자.

• __global__: 호스트가 호출하고 디바이스에서 도는 함수. 즉 커널이다. 반환형은 반드시 void.
• __device__: 디바이스에서 호출하고 디바이스에서 도는 함수. 커널 안에서 부르는 보조 함수.
• __host__: 호스트에서 호출하고 호스트에서 도는 함수 ─ 그냥 보통 C 함수와 같다. 생략 시 기본값. __host__ __device__를 둘 다 붙이면 두 군데서 모두 컴파일된다.

실행 모델은 SPMD(Single Program, Multiple Data)다. 같은 커널 함수를 수많은 스레드가 동시에 자기 데이터 인덱스를 가지고 실행한다. 양파 까기를 천 명이 하는 그림이 그대로다. 다만 모두가 동일한 코드를 보고 있다는 점이 핵심이다. 분기는 가능하지만 비싸다(4장에서 다룬다).

2.3 첫 커널: 벡터 덧셈

가장 단순한 데이터 병렬 작업으로 시작하자. 길이 N인 두 배열 A, B를 더해 C에 넣는다.

// vec_add.cu  ─ 호스트 + 디바이스가 한 파일에
#include <stdio.h>
#include <cuda_runtime.h>

#define N (1 << 20)   // 1,048,576 원소

__global__ void vecAdd(const float* A, const float* B, float* C, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        C[i] = A[i] + B[i];
    }
}

int main(void) {
    size_t bytes = N * sizeof(float);

    // 호스트 메모리 준비
    float *hA = (float*)malloc(bytes);
    float *hB = (float*)malloc(bytes);
    float *hC = (float*)malloc(bytes);
    for (int i = 0; i < N; ++i) { hA[i] = 1.0f; hB[i] = 2.0f; }

    // 디바이스 메모리 할당
    float *dA, *dB, *dC;
    cudaMalloc(&dA, bytes);
    cudaMalloc(&dB, bytes);
    cudaMalloc(&dC, bytes);

    // 호스트 → 디바이스
    cudaMemcpy(dA, hA, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(dB, hB, bytes, cudaMemcpyHostToDevice);

    // 커널 발사
    int threads = 256;
    int blocks  = (N + threads - 1) / threads;
    vecAdd<<<blocks, threads>>>(dA, dB, dC, N);

    // 디바이스 → 호스트
    cudaMemcpy(hC, dC, bytes, cudaMemcpyDeviceToHost);

    printf("hC[0]=%f, hC[N-1]=%f\n", hC[0], hC[N-1]);  // 3.0, 3.0 기대

    cudaFree(dA); cudaFree(dB); cudaFree(dC);
    free(hA); free(hB); free(hC);
    return 0;
}

이 30여 줄 안에 CUDA 프로그램의 기본 골격이 다 들어 있다. (1) 호스트·디바이스 메모리를 따로 잡는다. (2) 호스트에서 디바이스로 데이터를 복사한다. (3) 커널을 <<<blocks, threads>>> 구문으로 발사한다. (4) 결과를 다시 호스트로 가져온다. (5) 양쪽 메모리를 정리한다. "올리고 — 돌리고 — 받고 — 치우기" 4박자다.

2.4 디바이스 메모리와 데이터 전송

호스트 RAM과 GPU의 DRAM(보통 HBM 또는 GDDR)은 물리적으로 다른 칩이고, 그 사이는 PCIe 버스로 연결되어 있다(또는 NVLink, Grace-Hopper 같은 직접 연결도 있지만 이건 나중에). 그래서 호스트에서 만든 포인터는 디바이스에서 그대로 쓸 수 없고, 그 반대도 마찬가지다. 두 세상은 명시적으로 데이터를 주고받아야 한다.

핵심 API 세 개:

• cudaMalloc(void** devPtr, size_t size): 디바이스 DRAM에 size 바이트를 잡고 그 포인터를 *devPtr에 넣는다. 주의할 점은 첫 인자가 포인터의 포인터라는 것 ─ 그래서 cudaMalloc(&dA, bytes)로 부른다.
• cudaMemcpy(void* dst, const void* src, size_t size, cudaMemcpyKind kind): kind는 cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToDevice, cudaMemcpyHostToHost 중 하나.
• cudaFree(void* devPtr): 디바이스 메모리 해제. 호스트 메모리는 평범하게 free.

실제 코드에서는 모든 CUDA API 호출의 반환값을 검사해야 한다. 안 그러면 메모리가 부족하거나 잘못된 인자를 줬을 때 오류가 조용히 묻힌다. 다음과 같은 매크로 한 벌을 갖고 다니면 편하다.

#define CUDA_CHECK(call)                                              \
    do {                                                              \
        cudaError_t _e = (call);                                      \
        if (_e != cudaSuccess) {                                      \
            fprintf(stderr, "CUDA error %s:%d: %s\n",                 \
                    __FILE__, __LINE__, cudaGetErrorString(_e));      \
            exit(1);                                                  \
        }                                                             \
    } while (0)

// 사용
CUDA_CHECK(cudaMalloc(&dA, bytes));
CUDA_CHECK(cudaMemcpy(dA, hA, bytes, cudaMemcpyHostToDevice));

2.5 커널 함수와 스레딩 모델

커널이 발사되면 GPU는 어떻게 일을 분배할까? CUDA의 답은 2계층 위계다. 가장 위에 그리드(grid)가 있고, 그리드는 블록(block)으로 쪼개져 있고, 각 블록은 스레드(thread)로 쪼개져 있다. 블록과 스레드는 1D, 2D, 3D 어떤 모양으로도 잡을 수 있다(다음 장에서 본격적으로 다룬다). 지금은 1D만 보자.

커널 안에서 자신이 누구인지 아는 데 쓰이는 세 가지 내장 변수가 있다.

• threadIdx.x: 블록 안에서 자기 스레드의 번호 (0부터 blockDim.x−1까지).
• blockIdx.x: 그리드 안에서 자기 블록의 번호 (0부터 gridDim.x−1까지).
• blockDim.x: 한 블록에 들어 있는 스레드 수.

여기에 gridDim.x(그리드 안 블록 수)를 더하면 끝이다. 자기 전역 인덱스는 거의 항상 다음 패턴으로 나온다.

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

예를 들어 blockDim.x = 256이고 N = 1,048,576이면 4,096개의 블록이 필요하다. 블록 0의 스레드 0은 i=0, 블록 0의 스레드 255는 i=255, 블록 1의 스레드 0은 i=256 … 이렇게 모든 i가 정확히 한 번씩 한 스레드에 매핑된다.

그런데 blockDim.x가 N을 정확히 나누지 않을 수 있다. 가령 N=1000, blockDim.x=256이면 4블록을 쏜다(=1024 스레드). 마지막 블록의 마지막 24개 스레드는 i가 1000~1023이 되어 배열 범위를 벗어난다. 그래서 커널 안에서 항상

if (i < n) { /* 안전한 작업 */ }

같은 경계 검사를 둔다. 이 한 줄을 빼먹으면 메모리 침범으로 프로그램이 무작위로 망가진다.

2.6 커널 호출: <<<grid, block>>> 구문

CUDA C가 평범한 C와 다르게 보이는 가장 눈에 띄는 부분이 바로 이 삼중 꺽쇠 구문이다.

kernel<<<gridDim, blockDim, sharedMemBytes, stream>>>(args...);

네 인자 중 뒤의 두 개는 선택사항이다. sharedMemBytes는 동적 공유 메모리 크기(5장에서 다룬다), stream은 비동기 실행 큐를 명시하는 핸들이다(20장에서 다룬다). 처음엔 앞의 두 개만 쓰면 된다.

중요한 두 가지를 다시 강조한다.

첫째, 호출 자체가 비동기다. 호스트는 kernel<<<...>>>() 줄을 만나는 즉시 반환되어 다음 줄로 간다. GPU는 큐에 들어온 커널을 자기 페이스로 처리한다. 이 비동기성은 의도적인 설계다. 호스트가 다음 데이터 청크를 준비하는 동안 GPU가 이미 받은 청크를 처리할 수 있다 ─ 즉 파이프라이닝이 가능해진다.

둘째, 그리드와 블록 크기는 발사할 때마다 바꿀 수 있다. 같은 커널을 작은 데이터엔 작게, 큰 데이터엔 크게 쏘면 된다. 다만 한 블록의 최대 스레드 수, 한 그리드의 최대 블록 수 같은 하드웨어 한계는 있다(블록당 1024 스레드, 그리드 차원당 2^31−1 정도). 이 한계는 디바이스 속성으로 조회할 수 있다(4장).

2.7 컴파일: nvcc의 흐름

CUDA 소스(.cu)는 nvcc로 컴파일한다. 명령은 단순하다.

nvcc -O3 -arch=sm_80 vec_add.cu -o vec_add

-arch=sm_80은 컴퓨트 캐퍼빌리티 8.0 (Ampere, 예: A100)에 맞춘다는 뜻이다. 자기 GPU에 맞춰 sm_75(Turing), sm_86(RTX 30), sm_89(RTX 40), sm_90(Hopper) 등으로 바꾸면 된다. 모르면 -arch=sm_60 정도로 시작해도 동작한다.

nvcc 내부에서 일어나는 일은 대략 이렇다. (1) 소스를 호스트 코드와 디바이스 코드로 분리한다. (2) 호스트 코드는 시스템 C++ 컴파일러(g++, clang 등)로 보낸다. (3) 디바이스 코드는 NVIDIA의 cicc/ptxas를 거쳐 PTX(가상 ISA)와 SASS(실제 GPU 기계어)로 변환된다. (4) 결과를 한 실행 파일에 묶는다. 실행 시 GPU 드라이버가 자기 칩에 맞는 SASS를 골라 로드한다. PTX가 들어 있으면 호환되지 않는 신형 GPU에서 JIT 컴파일도 가능하다.

2.8 정리

이 장에서 다룬 4박자 ─ 메모리 잡기, 올리기, 커널 발사, 받기/치우기 ─ 가 모든 CUDA 프로그램의 기본 형태다. 이후 장에서 다루는 모든 최적화는 이 골격 위에 장식이 붙는 식이다. 벡터 덧셈은 산술 강도가 매우 낮아서(원소 한 개 더하려고 12바이트를 읽는다) GPU에서 사실 빛나지 않는 예제다. 그러나 모델 자체를 익히는 데는 이만한 게 없다. 다음 장에서는 한 차원을 추가해서 이미지와 행렬 같은 진짜 다차원 데이터를 다룬다.