cudatraining · 학습 기록

LESSON 01 · 2026.04.18 · T4

vector_add — 첫 커널이 가르쳐 준 것

연산이 더 비싼가, 복사가 더 비싼가. T4 위에서 처음으로 돌려본 커널은 ‘데이터 이동이 병목’이라는 명제를 17배 차이의 숫자로 돌려줬다.

GPU · Tesla T4 · sm_75 N · 67,108,864 결과 · 230.9 GB/s · 72.1%

셋업

로컬이 Apple Silicon 이라 CUDA 를 직접 돌릴 수 없다. GCP 크레딧으로 us-east1-d 에 T4 Spot VM 을 띄우고, Deep Learning VM 이미지로 올린 뒤 원격에서 컴파일했다. 첫 커널은 전형적인 vector add. 단순하지만, 단순해서 배운다.

__global__ void vector_add_kernel(const float* a, const float* b, float* c, size_t n) {
  size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
  size_t stride = gridDim.x * blockDim.x;
  for (size_t i = idx; i < n; i += stride) c[i] = a[i] + b[i];
}

grid-stride loop 을 쓴 이유는 단순하다: 문제 크기 n 이 grid 보다 크면 한 스레드가 여러 원소를 맡는다. launch config 가 n 에 덜 민감해진다.

결과

메인 실행 (block_size = 256):

best kernel time
3.488 ms
effective bandwidth
230.909 GB/s
theoretical (T4)
320.064 GB/s
efficiency
72.145%
H2D copy
44.276 ms
D2H copy
20.439 ms

Block size 를 흔들어봐도 극적이지 않았다.

blockGB/seff
128219.22868.495 %
256230.90972.145 %
512236.59273.920 %

진짜 병목

커널은 3.4 ms. 복사는 합쳐서 60 ms+. 진짜 병목은 커널 안이 아니라 PCIe 위다.

첫 레슨의 교훈은 간단하다. 연산보다 데이터 이동이 더 비싸다. 이 명제가 뒤의 9개 레슨을 모두 규정한다 — pinned memory (02), reduction 트리 (03), tiling (04), fusion (05), Flash Attention (06). 모두 "어떻게 덜 옮길까" 의 변주다.

왜 bytes_moved = n × 4 × 3 인가

한 원소당 float 두 개를 읽고 한 개를 쓴다. 4 바이트 × 3 = 12 바이트/원소. 이걸 시간으로 나눈 게 effective bandwidth. 커널이 "자기가 할 일"을 얼마나 빠르게 하는지만 보는 숫자다. 복사 시간은 여기 포함 안 된다 — 다음 레슨에서 따로 다룬다.

왜 vector add 는 memory-bound 인가

연산 강도(arithmetic intensity)가 1 FLOP 당 12 바이트. T4 의 roofline 꼭짓점은 약 FP32 peak 8.1 TFLOPS / 320 GB/s ≈ 25 FLOP/byte 부근이다. 우리 강도는 그 왼쪽 한참 아래 — 즉 HBM 대역폭이 지붕이다. 커널을 아무리 튜닝해도 bandwidth 가 상한.

다음 레슨

H2D 와 D2H 가 왜 이렇게 비싼가. cudaMallocHost 로 받은 pinned memory 와 그냥 new float[] 로 받은 pageable memory 를 나란히 측정한다.

Index전체 목록으로 Next · 02pageable vs pinned