제2부 CUDA 프로그래밍

2.1 첫 CUDA 커널 I — 빌드·실행·호스트/디바이스 흐름

이 장은 단계별 워크스루 실습이다. 빈 디렉터리에서 시작해 한 줄씩 작성하고, nvcc로 빌드하고, 실행해 결과를 검증한다. 이 장은 호스트/디바이스 흐름과 빌드·실행을, 인덱싱·경계검사·grid-stride·검증의 깊은 부분은 2.2로 잇는다.

2.1.1 이 장에서 배우는 것 / 선수 지식

배우는 것:

호스트(CPU)·디바이스(GPU)의 역할 분리, 둘 사이 데이터가 오가는 흐름.
__global__ 커널 작성, blockIdx·blockDim·threadIdx로 전역 인덱스 계산.
cudaMalloc·cudaMemcpy로 명시적 메모리를 다루고, <<<grid, block>>> 런치.
cudaGetLastError·cudaDeviceSynchronize로 오류를 검출하고 동기화한다.

선수 지식: C/C++ 기초(포인터·배열). 실습 환경(드라이버·CUDA Toolkit)은 준비됐다고 가정 — 미설치면 0.3·A.1. GPU 실행 모델(SIMT·워프)의 큰 그림은 1.1~1.2. 통합 메모리·grid-stride·검증의 깊은 부분은 2.2.

이 장의 목표는 최적화가 아니라 동작·검증·디버깅의 루프를 손에 익히는 것이다. 성능 튜닝은 2.3(실행 구성)·2.7(타일링)에서.

2.1.2 왜 호스트/디바이스를 나누는가

CPU와 GPU는 별개 프로세서이고 전통적으로 별개 메모리(시스템 DRAM vs GPU의 HBM·GDDR)를 가진다. CUDA 프로그램은 이 둘을 오케스트레이션한다: CPU(호스트)가 흐름을 지휘하고, 데이터를 GPU(디바이스)로 보내고, "이 함수를 수천 스레드로 실행하라"고 지시(커널 런치)한 뒤 결과를 가져온다.

핵심 비대칭: CPU 코드는 한 번 호출하면 한 번 실행되지만, GPU 커널은 한 번 런치하면 수천 스레드가 같은 코드를 각자의 데이터에 대해 실행한다. 그래서 커널 안에서 가장 먼저 하는 일은 "나는 몇 번 스레드이고, 따라서 배열의 어느 칸을 맡는가"를 계산하는 것이다. 이 한 줄이 CUDA의 핵심이다.

예제는 SAXPY(y[i] = a*x[i] + y[i])다. BLAS 표준 연산이고 각 원소가 독립적이라 "CUDA의 hello world"로 이상적이다.

2.1.3 스레드 인덱싱과 런치 구문

2.1.3.1 스레드 계층 — 그리드·블록·스레드

커널을 런치할 때 두 차원을 지정한다: 블록의 격자(grid) 와 블록당 스레드(block).

각 블록 안 thread는 threadIdx.x = 0..3, block은 blockIdx.x = 0..3이다. 전역 인덱스 i = blockIdx.x * blockDim.x + threadIdx.x이므로 block 2의 thread 1 → i = 2*4 + 1 = 9.

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

threadIdx.x — 블록 안 스레드 번호(0..blockDim.x-1).
blockIdx.x — 그리드 안 블록 번호(0..gridDim.x-1).
blockDim.x — 블록당 스레드 수(런치 때 정함).
gridDim.x — 그리드의 블록 수.

.x .y .z로 2D·3D 격자도 만들지만 1D 배열엔 .x만. 차원 설계는 2.3.

2.1.3.2 런치 구문과 그리드 크기 계산

int block = 256;
int grid  = (N + block - 1) / block;   // N을 덮는 데 필요한 블록 수(올림 나눗셈)
saxpy<<<grid, block>>>(N, a, d_x, d_y);

(N + block - 1) / block는 정수 올림 나눗셈 관용구다. N=1000, block=256이면 4블록 → 1024 스레드, 그중 24개는 경계 검사에 걸려 아무 일도 안 한다(경계 검사는 2.2). <<< >>>는 비동기다 — 런치는 GPU에 작업을 큐잉만 하고 즉시 호스트로 돌아온다. 그래서 결과를 읽기 전에 동기화가 필요하다.

2.1.4 워크스루 — 완전한 SAXPY (명시적 메모리)

빈 디렉터리에서 시작한다. 아래 전체를 saxpy.cu로 저장한다.

#include <cstdio>
#include <cstdlib>
#include <cmath>

// ── 오류 검사 매크로: 모든 CUDA 런타임 호출을 감싼다 ───────────────
#define CUDA_CHECK(call)                                                   \
  do {                                                                     \
    cudaError_t err__ = (call);                                           \
    if (err__ != cudaSuccess) {                                           \
      fprintf(stderr, "CUDA error %s:%d: '%s' returned %s\n",             \
              __FILE__, __LINE__, #call, cudaGetErrorString(err__));      \
      exit(EXIT_FAILURE);                                                  \
    }                                                                      \
  } while (0)

// ── 디바이스 커널: y[i] = a*x[i] + y[i] ───────────────────────────
__global__ void saxpy(int n, float a, const float* x, float* y) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;   // 내 전역 인덱스
  if (i < n) {                                      // 경계 검사 (2.2)
    y[i] = a * x[i] + y[i];
  }
}

int main(void) {
  const int   N = 1 << 20;          // 1,048,576 원소
  const float a = 2.0f;
  const size_t bytes = N * sizeof(float);

  // 1) 호스트 메모리 할당·초기화
  float* h_x = (float*)malloc(bytes);
  float* h_y = (float*)malloc(bytes);
  for (int i = 0; i < N; ++i) { h_x[i] = 1.0f; h_y[i] = 2.0f; }

  // 2) 디바이스 메모리 할당
  float *d_x = nullptr, *d_y = nullptr;
  CUDA_CHECK(cudaMalloc(&d_x, bytes));
  CUDA_CHECK(cudaMalloc(&d_y, bytes));

  // 3) 호스트 → 디바이스 복사
  CUDA_CHECK(cudaMemcpy(d_x, h_x, bytes, cudaMemcpyHostToDevice));
  CUDA_CHECK(cudaMemcpy(d_y, h_y, bytes, cudaMemcpyHostToDevice));

  // 4) 커널 런치
  int block = 256;
  int grid  = (N + block - 1) / block;
  saxpy<<<grid, block>>>(N, a, d_x, d_y);

  // 5) 런치 자체의 오류 확인 (런치 구성·인자 오류)
  CUDA_CHECK(cudaGetLastError());
  // 6) 커널 완료까지 대기 + 실행 중 오류 확인
  CUDA_CHECK(cudaDeviceSynchronize());

  // 7) 디바이스 → 호스트 복사
  CUDA_CHECK(cudaMemcpy(h_y, d_y, bytes, cudaMemcpyDeviceToHost));

  // 8) 검증: 모든 원소가 a*1 + 2 = 4 이어야 한다
  float max_err = 0.0f;
  for (int i = 0; i < N; ++i)
    max_err = fmaxf(max_err, fabsf(h_y[i] - 4.0f));
  printf("N = %d, max error = %f\n", N, max_err);
  printf(max_err == 0.0f ? "RESULT: PASS\n" : "RESULT: FAIL\n");

  // 9) 정리
  CUDA_CHECK(cudaFree(d_x));  CUDA_CHECK(cudaFree(d_y));
  free(h_x);  free(h_y);
  return 0;
}

빌드·실행:

nvcc -O2 -arch=native -o saxpy saxpy.cu
./saxpy
# 예상(개략):  N = 1048576, max error = 0.000000  /  RESULT: PASS

-arch=native는 빌드 머신 GPU에 맞는 아키텍처를 자동 선택(최신 CUDA). 안 되면 세대에 맞춰 명시: Ampere -arch=sm_80, Ada -arch=sm_89, Hopper -arch=sm_90. 타깃·PTX·SASS 정리는 2.12.

2.1.4.1 한 줄씩 — 무슨 일이 일어나는가

호스트 초기화: h_x=1.0, h_y=2.0 → a=2로 정답 2*1+2=4(검증 쉬운 값).
cudaMalloc(&d_x, bytes): GPU 글로벌 메모리에 할당하고 디바이스 포인터를 받는다. d_x는 GPU 주소공간을 가리키므로 호스트 코드에서 역참조하면 안 된다(2.2 함정).
cudaMemcpy(HostToDevice): PCIe(또는 NVLink, 4.6)를 건너 복사. 이 전송이 곧잘 병목(전송 최적화는 2.4~2.5).
런치: 약 N개 스레드를 큐잉. 호출은 즉시 반환(비동기).
cudaGetLastError(): 런치 구성 오류(블록이 너무 큼·잘못된 인자)를 검출한다. 런치는 에러를 던지지 않고 저장만 하므로 명시적으로 읽어야 한다.
cudaDeviceSynchronize(): 커널 끝까지 대기하고 실행 중 비동기 오류(디바이스 측 잘못된 메모리 접근)를 반환. 5·6은 한 쌍 — 하나는 런치 오류, 하나는 실행 오류.
cudaMemcpy(DeviceToHost): 결과 회수.
검증: 전 원소가 4.0인지 최대 오차로. GPU 계산을 "실행됐으니 맞겠지"로 믿지 않고 항상 호스트에서 독립 검증.
정리: cudaFree·free.

2.1.5 중간 점검 (체크포인트)

[ ] 빌드: nvcc -O2 -arch=native -o saxpy saxpy.cu가 오류 없이 실행 파일을 만든다. (실패: nvcc PATH? -arch 값이 GPU와 맞나? 문법?)
[ ] 실행: ./saxpy가 RESULT: PASS.
[ ] 검증 의미: max error = 0이 왜 0인지 설명한다(2*1+2=4, 정수처럼 떨어지는 FP 값이라 오차 없음).
[ ] 인덱스 공식: block=256일 때 block 3의 thread 10이 전역 몇 번인가(=778).

증상별 의심(상세는 2.2):

증상	의심
전부 초기값(2.0) 그대로	커널이 안 실행됐다 → `cudaGetLastError` 확인, 런치 구성·`-arch` 점검
가끔만 PASS, 비결정적	동기화 누락 — 커널 끝나기 전에 결과를 읽음
즉시 크래시·`illegal memory access`	디바이스 포인터를 호스트에서 역참조, 또는 인덱스 오버런
`cudaMalloc` 실패	메모리 부족, GPU/드라이버 미인식(`nvidia-smi` 확인)

디버깅: 빌드에 -lineinfo를 추가하고 compute-sanitizer ./saxpy로 실행하면 잘못된 메모리 접근·레이스를 줄 번호와 함께 검출한다(2.14).

2.1.6 흔한 함정

오류를 안 읽음. 대부분의 CUDA 호출은 예외를 던지지 않고 cudaError_t를 반환한다. 반환값을 무시하면 실패가 조용히 묻힌다. 모든 호출을 CUDA_CHECK로. 런치는 반환값이 없으니 cudaGetLastError()로 따로.
두 오류를 헷갈림. cudaGetLastError는 런치 시점(구성·인자), cudaDeviceSynchronize는 실행 중 오류. 디바이스 측 잘못된 접근은 후자에서만 보인다.
포인터 혼동. 호스트 포인터(h_x)와 디바이스 포인터(d_x)는 다른 주소공간. 디바이스 포인터를 호스트에서 읽으면 illegal memory access/세그폴트.
-arch 불일치. 빌드 타깃이 실제 GPU보다 높으면 no kernel image is available. 세대에 맞춰 -arch(2.12).
검증 생략. "컴파일·실행됐다"는 "정답이다"가 아니다. 항상 호스트에서 독립 검산.

2.1.7 연습문제

연습 1 (요소별 변형). saxpy를 z[i] = x[i]*y[i] + sinf(x[i])로 바꿔라. 출력 버퍼를 적절히 할당하고 호스트에서 검산하라.

정답

출력 z를 따로 cudaMalloc하고 시그니처를 (int n, const float* x, const float* y, float* z)로. 본문 if(i<n) z[i]=x[i]*y[i]+sinf(x[i]);. sinf는 디바이스 호출 가능 단정밀 함수. 검산은 호스트에서 동일 계산 후 차이를 보되 FP 오차 때문에 정확히 0이 아닐 수 있으니 허용오차(예: 1e-5 상대)를 둔다.

연습 2 (오류 검사 체험). cudaMalloc에 말도 안 되는 크기((size_t)1 << 50)를 줘 보라. CUDA_CHECK 없을 때와 있을 때 동작 차이는?

정답

CUDA_CHECK 없으면 d_x가 nullptr인 채 진행돼 커널에서 illegal memory access로 늦게 실패하거나 조용히 틀린다. 있으면 할당 직후 out of memory를 파일·줄 번호와 함께 즉시 보고하고 종료 — 오류를 발생 지점에서 검출하는 가치.

연습 3 (스레드 수의 의미). N=1M, block=256이면 약 100만 스레드. 물리 코어는 수천 개뿐인데 어떻게 실행되는가?

정답

스레드는 워프(32) 단위로 SM 스케줄러가 시간 분할 실행하고, 메모리 대기 동안 다른 워프로 전환해 지연을 은닉한다. "코어당 한 스레드"가 아니라 수많은 스레드를 생성해 점유율을 채우고 지연을 은닉한다(1.2·1.6).

2.1.8 요약 · 다음 장

CUDA 프로그램은 호스트가 데이터를 디바이스로 옮기고, 커널을 수천 스레드로 런치하고, 결과를 회수하는 구조다.
커널의 첫 필수 줄: 전역 인덱스 i = blockIdx.x*blockDim.x + threadIdx.x.
런치는 비동기 — cudaGetLastError(런치 오류)와 cudaDeviceSynchronize(실행 오류)를 쌍으로 쓰고, 결과는 항상 호스트에서 독립 검증.

다음 장: 2.2 첫 CUDA 커널 II — 인덱싱·경계검사·검증. 경계 검사가 왜 필수인지, grid-stride loop로 임의 크기를 견고하게, 통합 메모리로 단순화, 그리고 정확성 검증·디버깅을 손에 익힌다.