제2부 CUDA 프로그래밍

2.2 첫 CUDA 커널 II — 인덱싱·경계검사·검증

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

2.1에서 명시적 메모리로 첫 SAXPY를 빌드·실행했다. 이 장은 그 커널을 견고하고 정확하게 만드는 부분 — 경계 검사, grid-stride loop, 통합 메모리, 그리고 정확성 검증·디버깅 — 을 손에 익힌다.

배우는 것:

• 경계 검사 if (i < n)가 왜 선택이 아니라 필수인지.
• cudaMallocManaged(통합 메모리)로 코드를 단순화하되 동기화 의무가 남는 이유.
• grid-stride loop로 임의 크기 배열을 어떤 런치 구성에도 견고하게 처리.
• 명시적 vs 통합 메모리의 트레이드오프와 정확성 검증 규율.

선수 지식: 2.1(빌드·실행·호스트/디바이스 흐름·오류 검사). 실행 구성 설계는 2.3, 메모리 모델 상세는 2.4~2.5.

2.2.2 왜 경계 검사가 필요한가

블록당 스레드 수는 보통 32의 배수(워프 크기 32)로 고른다 — 256이 흔하다. 그런데 배열 길이 N이 256으로 나눠떨어지란 법이 없다. 그러면 마지막 블록의 일부 스레드는 배열 범위 밖을 가리킨다. 이 스레드들이 동기화 없이 쓰기를 하면 메모리 오버런이다. 그래서 커널의 필수 줄:

if (i < N) { /* 실제 작업 */ }

주황색 칸 i = 10, 11은 배열 범위 밖이라 if (i < n) 가드가 제외한다 — 올림 나눗셈으로 항상 grid*block ≥ n이기 때문이다.

(n+block-1)/block 올림 나눗셈은 항상 grid*block >= n이라 i >= n인 스레드가 존재한다. if (i < n) 가드를 빠뜨리면 그 스레드들이 배열 밖을 쓴다(메모리 오염·크래시·불법 접근). 디버그 빌드에선 실패하지 않고 릴리스에서만 깨지기도 한다 — 그래서 compute-sanitizer로 확인하는 습관이 중요하다(2.14).

2.2.3 grid-stride loop — 임의 크기를 견고하게

2.1의 "원소당 스레드 한 개" 버전은 매우 크거나 그리드 한도를 넘는 N에서 깨질 수 있고, 스레드 수를 데이터에 묶어버린다. grid-stride loop는 고정된 수의 스레드가 보폭(stride)만큼 건너뛰며 배열 전체를 순회하게 한다 — 어떤 N에도, 어떤 런치 구성에도 견고하고, 한 스레드가 여러 원소를 처리해 런치 오버헤드도 줄인다.

__global__ void saxpy_gs(int n, float a, const float* x, float* y) {
  int idx    = blockIdx.x * blockDim.x + threadIdx.x;
  int stride = blockDim.x * gridDim.x;            // 그리드 전체 스레드 수
  for (int i = idx; i < n; i += stride) {         // 보폭만큼 건너뛰며 전부 처리
    y[i] = a * x[i] + y[i];
  }
}

이제 그리드를 데이터에 정확히 맞출 필요가 없다 — SM 수에 비례한 고정 그리드로 런치해도 정확하다. 이 패턴을 기본기로 삼는다. 루프 안에서도 i < n 검사(for 조건)가 경계를 지킨다. 그리드 크기 설계·점유율 API는 2.3.

2.2.4 통합 메모리 — 코드 단순화

통합 메모리(cudaMallocManaged)는 호스트·디바이스가 같은 포인터로 접근하는 단일 할당을 만든다. 드라이버가 페이지를 필요한 쪽으로 자동 마이그레이션한다. cudaMemcpy가 사라져 코드가 짧아진다 — 첫 학습·프로토타이핑에 좋다.

#include <cstdio>
#include <cmath>
#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)

__global__ void saxpy(int n, float a, const float* x, float* y) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) y[i] = a * x[i] + y[i];
}

int main(void) {
  const int N = 1 << 20; const float a = 2.0f;
  float *x, *y;                                  // 호스트·디바이스 공용 포인터
  CUDA_CHECK(cudaMallocManaged(&x, N * sizeof(float)));
  CUDA_CHECK(cudaMallocManaged(&y, N * sizeof(float)));
  for (int i = 0; i < N; ++i) { x[i] = 1.0f; y[i] = 2.0f; }  // 호스트가 직접 씀

  int block = 256, grid = (N + block - 1) / block;
  saxpy<<<grid, block>>>(N, a, x, y);
  CUDA_CHECK(cudaGetLastError());
  CUDA_CHECK(cudaDeviceSynchronize());           // 호스트가 읽기 전 반드시 필요

  float max_err = 0.0f;
  for (int i = 0; i < N; ++i) max_err = fmaxf(max_err, fabsf(y[i]-4.0f));
  printf("max error = %f  ->  %s\n", max_err, max_err==0.0f?"PASS":"FAIL");

  CUDA_CHECK(cudaFree(x));  CUDA_CHECK(cudaFree(y));
  return 0;
}

차이의 핵심: 명시적 복사가 없어졌지만 cudaDeviceSynchronize()는 여전히 필수다. 동기화 없이 커널 직후 호스트에서 y[i]를 읽으면 커널이 아직 안 끝났을 수 있다(데이터 레이스). 통합 메모리는 편의이지 성능 보장이 아니다 — 마이그레이션 비용·프리페치 튜닝은 2.5.

2.2.5 트레이드오프 — 명시적 vs 통합 메모리

견해: 처음엔 명시적 버전으로 H→D→커널→D→H 흐름을 확실히 익히고, 그다음 통합 메모리로 단순화하라. 명시적 모델을 건너뛰면 데이터가 "자동으로" 옮겨진다고 착각하기 쉽고, 나중에 성능 디버깅 때 그 착각이 비용이 커진다. 두 모델의 비용·튜닝은 2.5에서 정량 비교.

비엔비디아 정직: 같은 SAXPY를 OpenMP·SYCL·OpenCL·HIP로도 쓸 수 있고 이식성은 그쪽이 유리. CUDA를 택하는 이유는 엔비디아 GPU에서의 성숙한 툴체인·라이브러리 생태계(2.15 이후)다 — 이식성과 생태계·성능 사이의 선택.

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

• [ ] (n+block-1)/block 올림으로 남는 스레드가 생기고 경계 검사가 그것을 막음을 설명한다.
• [ ] grid-stride loop가 그리드 크기를 N에서 분리함을 알고, stride=blockDim.x*gridDim.x를 제시한다.
• [ ] 통합 메모리에서도 cudaDeviceSynchronize가 필수인 이유(데이터 레이스)를 제시한다.
• [ ] "컴파일·실행됨 ≠ 정답"을 알고 호스트 독립 검증을 한다.

결과가 틀리면:

2.2.7 흔한 함정

• 동기화 누락. 커널 런치는 비동기. cudaDeviceSynchronize()(또는 동기 cudaMemcpy) 없이 결과를 읽으면 레이스. "가끔 맞고 가끔 틀린" 버그의 1순위.
• 경계 검사 누락. N이 blockDim 배수가 아니면 마지막 블록 일부가 범위를 넘는다. if (i < n) 필수. compute-sanitizer로 확인.
• 올림 나눗셈 실수. grid = N / block(내림)이면 꼬리 원소가 처리 안 됨. 반드시 (N + block - 1) / block.
• 포인터 혼동. 통합 메모리는 호스트/디바이스 포인터 혼동을 없애지만 동기화 의무는 남는다.
• 검증 생략. 항상 호스트에서 독립 검산.

2.2.8 연습문제

연습 1 (경계 버그 재현). if (i < n)을 지우고 N=1000, block=256으로 실행하라. 무슨 일이 나는가? compute-sanitizer로는?

연습 2 (grid-stride 변환). 2.1의 SAXPY를 grid-stride로 바꾸고, 그리드를 데이터와 무관한 고정값(예: 80)으로 런치해도 정답이 나오는지 확인하라.

연습 3 (동기화 버그 재현). 통합 버전에서 cudaDeviceSynchronize를 지워라. 결과가 어떻게 되나?

2.2.9 요약 · 다음 장

• 커널의 둘째 필수 줄은 경계 검사 if (i < n) — 올림 나눗셈으로 남는 스레드의 메모리 오버런을 막는다.
• grid-stride loop를 임의 크기 처리의 기본 패턴으로 삼는다(그리드 크기를 N에서 분리).
• 통합 메모리는 코드를 단순화하지만 동기화 의무는 남는다. 명시적 모델로 흐름을 익힌 뒤 단순화하라.
• "컴파일·실행 ≠ 정답" — 항상 호스트에서 독립 검증한다.

다음 장: 2.3 실행 구성 — 그리드·블록·grid-stride 설계. <<<grid, block>>>의 두 숫자를 근거를 갖고 고르는 법 — 32 정렬, 블록 크기 선택, 2D 매핑·병합, 다이버전스 회피, 점유율 API를 본다.