② CUDA 프로그래밍

①장의 하드웨어 그림을 이제 실제 코드로 옮깁니다. 겁먹지 마세요 — CUDA는 C/C++에 몇 가지 문법만 더한 것이고, 핵심은 딱 하나, "어떤 데이터를 어떤 일꾼(스레드)이 처리할지 정하는 일"입니다.

2.1 프로그래밍 모델 — Grid / Block / Thread

CUDA에서 병렬 작업은 3단 계층으로 조직됩니다. ①에서 본 계층과 똑같아요 — 그게 그대로 하드웨어(SM, warp)에 매핑되거든요.

• Thread(스레드) — 일꾼 한 명. 보통 데이터 원소 하나를 처리합니다.
• Block(블록) — 스레드 묶음(한 조보다 큰 단위). 같은 블록끼리는 공용 작업대(shared memory)를 함께 쓰고 서로 기다려줄(동기화) 수 있어요. 한 블록은 통째로 한 SM에서 돌아갑니다.
• Grid(그리드) — 블록 전부의 묶음. 커널 하나를 실행하면 grid 전체가 GPU에 펼쳐집니다.

그럼 "수천 명의 일꾼 중에 나는 몇 번 일꾼이지?"를 어떻게 알까요? CUDA가 각 스레드에게 좌표를 알려줍니다 — blockIdx(내가 몇 번째 블록), threadIdx(블록 안에서 내 위치), blockDim(블록 크기). 이걸 조합해 전역 번호를 계산합니다:

// "나는 전체에서 몇 번 일꾼인가"를 계산하는 마법의 공식
int idx = blockIdx.x * blockDim.x + threadIdx.x;
//        ↑ 몇 번째 블록    × 블록 크기  + 블록 안 내 자리

// 예) 블록 크기 256일 때, 3번 블록의 5번 스레드라면:
//     idx = 3 * 256 + 5 = 773  → "나는 773번 일꾼, 773번 데이터를 처리"

2.2 첫 커널 — 벡터 덧셈으로 감 잡기

가장 기본 예제, 두 배열을 더하는 C = A + B를 CPU와 GPU로 비교해볼게요. 차이를 보면 "아하" 하실 거예요.

// ===== CPU 버전: 루프로 하나씩 순서대로 =====
void vecAdd_cpu(float* A, float* B, float* C, int n) {
    for (int i = 0; i < n; i++)   // i를 0,1,2,... 차례로
        C[i] = A[i] + B[i];
}

// ===== GPU 버전: 일꾼 한 명이 원소 하나 담당 (루프가 사라짐!) =====
__global__ void vecAdd(float* A, float* B, float* C, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;  // 내 번호 = 내 데이터 칸
    if (i < n)                  // 일꾼이 데이터보다 많을 수 있으니 범위 체크
        C[i] = A[i] + B[i];
}

두 가지만 보세요. ① __global__ — 이 표시가 "이 함수는 GPU에서 돌고, CPU가 불러낸다"는 뜻입니다(=커널). ② 루프가 사라졌어요. CPU는 i를 0부터 n까지 돌리지만, GPU는 n개의 일꾼이 각자 자기 i 하나씩 동시에 처리하니 루프가 필요 없죠.

호스트(CPU) 쪽 — 커널을 실제로 부르는 코드

커널을 짰으면, CPU 쪽에서 ① GPU 메모리 확보 → ② 데이터 복사(CPU→GPU) → ③ 커널 실행 → ④ 결과 회수(GPU→CPU) → ⑤ 정리, 이 순서로 부릅니다:

int n = 1<<20;                  // 약 100만 개
size_t bytes = n * sizeof(float);

// ① GPU(device) 메모리 확보 — 'd'는 device용이라는 관습적 표시
float *dA, *dB, *dC;
cudaMalloc(&dA, bytes);  cudaMalloc(&dB, bytes);  cudaMalloc(&dC, bytes);

// ② CPU(host) 데이터를 GPU로 복사
cudaMemcpy(dA, hA, bytes, cudaMemcpyHostToDevice);
cudaMemcpy(dB, hB, bytes, cudaMemcpyHostToDevice);

// ③ 커널 실행! 일꾼을 몇 명, 몇 조로 만들지 정함
int threads = 256;                            // 블록당 일꾼 수 (보통 128~512)
int blocks  = (n + threads - 1) / threads;   // 필요한 블록 수 (올림 나눗셈)
vecAdd<<<blocks, threads>>>(dA, dB, dC, n);    // ← 이 <<< >>> 가 CUDA 고유 문법

// ④ 결과를 GPU에서 CPU로 회수  ⑤ 메모리 해제
cudaMemcpy(hC, dC, bytes, cudaMemcpyDeviceToHost);
cudaFree(dA);  cudaFree(dB);  cudaFree(dC);

<<<blocks, threads>>>가 CUDA에서만 보이는 실행 구성(launch configuration) 문법입니다. "블록 몇 개, 블록당 스레드 몇 개로 일꾼을 펼쳐라"는 뜻이에요. 이 한 줄이 실제로 GPU 안에서 어떻게 처리되는지는 ③ 실행 모델에서 낱낱이 파헤칩니다.

2.3 실전 패턴 — 2D 그리드와 grid-stride loop

벡터 덧셈으로 기본 형태를 봤으니, 실전 코드에서 정말 자주 만나는 두 가지 패턴을 더 익혀둡시다.

패턴 1: 2D 그리드 — 이미지·행렬 다루기

이미지나 행렬처럼 2차원 데이터를 다룰 땐, 인덱스를 1차원으로 욱여넣는 대신 블록과 그리드를 2차원으로 만들 수 있습니다. dim3 타입으로 x·y(·z) 크기를 지정하면 돼요:

__global__ void brighten(unsigned char* img, int W, int H) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;   // 내 픽셀의 가로 좌표
    int y = blockIdx.y * blockDim.y + threadIdx.y;   // 내 픽셀의 세로 좌표
    if (x < W && y < H)                              // 양쪽 다 범위 체크!
        img[y * W + x] = min(img[y * W + x] + 30, 255);
}

// 호스트 쪽: 16×16 = 256 스레드짜리 타일로 이미지를 덮음
dim3 threads(16, 16);
dim3 blocks((W + 15) / 16, (H + 15) / 16);
brighten<<<blocks, threads>>>(dImg, W, H);

"스레드 한 명 = 픽셀 한 개"라는 사고법은 그대로고, 좌표가 (x, y) 두 개가 됐을 뿐입니다. 메모리는 어차피 1차원이라 y*W + x로 접는데, x가 연속이 되게(이웃 스레드 = 이웃 열) 잡는 게 중요해요 — 이게 ⑥장에서 볼 메모리 병합의 조건입니다.

패턴 2: Grid-Stride Loop — CUDA의 표준 관용구

지금까지는 "데이터 개수 = 스레드 개수"로 맞췄지만, 스레드 수보다 데이터가 많아도 되는 더 유연한 형태가 있습니다. 각 스레드가 자기 칸 하나만 처리하고 끝나는 대신, 전체 스레드 수만큼 건너뛰며 여러 칸을 처리하는 거예요:

__global__ void vecAdd(float* A, float* B, float* C, int n) {
    int stride = blockDim.x * gridDim.x;             // 전체 스레드 수
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n; i += stride)
        C[i] = A[i] + B[i];                          // 내 칸 → 한 바퀴 건너 다음 칸 → ...
}

// 일꾼 수를 데이터가 아니라 'GPU 크기'에 맞춰 고정할 수 있음
vecAdd<<<32 * numSMs, 256>>>(dA, dB, dC, n);   // n이 얼마든 동작!

NVIDIA가 공식적으로 권하는 패턴이고, 장점이 셋입니다:

• 크기 독립 — n이 백만이든 10억이든 같은 launch 구성으로 동작. "블록 수 한도 초과" 걱정도 없음.
• launch 구성을 GPU에 맞춤 — 블록 수를 "SM 수 × 몇 배"로 잡아 occupancy를 딱 맞게 채울 수 있음.
• 스레드당 일이 많아짐 — 스레드 생성·인덱스 계산 같은 고정 비용이 분산되고, 루프 안에서 명령 수준 병렬성(ILP)도 생김.

루프 안의 접근(i, i+stride, …)도 매 반복에서 warp 단위로 보면 연속 주소라, 병합도 그대로 유지됩니다. "커널은 grid-stride loop로 시작하라"를 기본기로 삼아도 좋아요.

2.4 메모리 관리 — 세 가지 방법

방법 1: 명시적 복사 (위에서 본 방식)

cudaMalloc으로 GPU 메모리를 잡고, cudaMemcpy로 직접 복사. 가장 기본이고, 무슨 일이 일어나는지 명확합니다.

방법 2: Unified Memory — 편한 자동 방식

float *A;
cudaMallocManaged(&A, bytes);   // CPU·GPU가 '같은 포인터'로 접근
// 명시적 cudaMemcpy 없이 CPU에서 쓰고 GPU에서 읽기 가능
// → 시스템이 필요할 때 알아서 데이터를 옮겨줌 (편하지만 성능 튜닝은 별도)

입문할 때 편합니다. 단, 뒤에서 자동으로 복사가 일어나므로 성능이 항상 최적은 아니에요.

방법 3: Pinned Memory — 빠른 전송용

에러 체크는 꼭 습관으로

CUDA 함수는 실패해도 조용히 넘어갈 때가 많아, 버그 잡기가 어렵습니다. 그래서 보통 매크로로 감싸 모든 호출을 체크합니다:

#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 호출을 이렇게 감싸기

한 가지 함정이 더 있습니다 — 커널 launch는 에러 코드를 반환하지 않아요(③장에서 볼 "비동기"라서). 그래서 커널은 이렇게 따로 체크합니다:

kernel<<<blocks, threads>>>(...);
CUDA_CHECK(cudaGetLastError());        // launch 자체의 에러 (잘못된 구성 등)
CUDA_CHECK(cudaDeviceSynchronize());   // 실행 중 에러 (잘못된 메모리 접근 등) — 디버깅 시

launch 구성이 틀린 경우(블록당 스레드 한도 초과 등)는 cudaGetLastError에서, 커널 안에서 터진 에러(배열 범위 밖 접근 등)는 나중에 동기화하는 지점에서 잡힙니다. "에러가 난 줄"과 "에러가 보고되는 줄"이 다를 수 있다는 걸 기억하세요.

2.5 동기화 — 같은 블록 일꾼들의 협력

같은 블록(같은 SM에 있는 일꾼들)은 서로 도울 수 있습니다. 단 다른 블록끼리는 커널 도중에 직접 협력할 수 없어요(서로 다른 공장에 있으니까요). 협력에 쓰는 두 도구:

① __syncthreads() — "다 모일 때까지 기다려"

블록 안 모든 스레드가 이 줄에 도착할 때까지 서로 기다리는 약속 장소(배리어)입니다. 공용 작업대(shared memory)에 데이터를 같이 채울 때 꼭 필요해요 — "다 채운 다음에 쓰자"를 보장하니까요.

__shared__ float tile[256];      // 블록 공용 작업대
tile[threadIdx.x] = A[idx];        // ① 각자 자기 칸을 채움
__syncthreads();                 // ② 모두 채울 때까지 대기 (이게 없으면 사고!)
// ③ 이제 tile 전체를 안전하게 함께 사용 가능
float left = tile[threadIdx.x - 1];

② 아토믹(atomic) 연산 — "겹치지 않게 안전하게 더하기"

여러 일꾼이 같은 메모리 칸을 동시에 갱신하면 값이 꼬입니다(예: 합 구하기). atomicAdd 같은 아토믹 연산은 "한 번에 한 명씩 안전하게" 갱신해줍니다. 단 경쟁이 심하면 줄서기(직렬화)가 생겨 느려지니 남용은 금물.

atomicAdd(&sum, value);   // 여러 스레드가 sum을 안전하게 누적

2.6 컴파일과 실행 — 직접 돌려보기

CUDA 코드는 확장자를 .cu로 저장하고, NVIDIA의 컴파일러 nvcc로 컴파일합니다. nvcc가 알아서 코드를 둘로 갈라 — CPU 부분은 일반 C++ 컴파일러로, 커널 부분은 GPU 코드로 — 컴파일한 뒤 하나의 실행 파일로 묶어줘요:

nvcc vecadd.cu -o vecadd     # 컴파일 (CUDA Toolkit 설치 시 포함)
./vecadd                     # 실행 — 보통의 프로그램과 똑같이

nvcc -arch=native vecadd.cu -o vecadd   # 지금 꽂힌 GPU 세대에 최적화
nvcc -arch=sm_86  vecadd.cu -o vecadd   # 특정 세대(CC 8.6 = RTX 30) 타깃
nvcc --ptxas-options=-v vecadd.cu       # 커널의 레지스터·shared 사용량 출력 (튜닝 때 유용)

-arch가 바로 ①장의 compute capability입니다. 지정 안 하면 기본 타깃으로 컴파일되는데, 일단 입문 단계에선 옵션 없이 시작해도 됩니다. 내 GPU 확인은 nvidia-smi.

%%writefile vecadd.cu
// (여기에 위의 CUDA 코드 붙여넣기)

!nvcc vecadd.cu -o vecadd && ./vecadd