벡터가 뭐고 GPU가 왜 사랑하는가

GPU 모델 전체를 노출시키는 가장 작은 워크로드

벡터는 array야. 같은 길이 벡터 둘을 더하면: 모든 i에 대해 w[i] = u[i] + v[i]. 원소당 산술 op 한 번, 원소 사이 의존성 없음 — 교과서적 embarrassingly parallel 정의.

여기서 'embarrassing'은 칭찬이야. Embarrassingly parallel 워크로드는:

원소 사이 통신 없음 — 각 출력은 자기 입력에만 의존.
Trivially load-balanced — 모든 원소 비용 동일.
Linear 메모리 access — fully coalesced, full bandwidth.
동기화 0 — __syncthreads 필요 없음.

그래서 여기서 시작하는 거. 이후 모든 워크로드 (matrix add, GEMM, attention, softmax)가 이 골격 상속받고 복잡함 추가: 2-D 인덱싱, shared-memory 재사용, reduction, mask. 골격부터 마스터.

왜 GPU가 특히 사랑하냐 — vector add는 L2 cache 한 번도 안 쓰고 DRAM bandwidth 포화. '내 메모리 서브시스템 살아 있나?' 의 가장 단순한 테스트 — 손코딩 커널이 여기서 이론치 80% 치면 toolchain이랑 launch geometry 건강.

Code

골격 — 수학 한 줄, 무한 scaling·cuda

// CUDA
__global__ void vector_add(
    const float* __restrict__ A,
    const float* __restrict__ B,
    float*       __restrict__ C,
    size_t N)
{
    size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) C[i] = A[i] + B[i];
}

// __restrict__는 nvcc에 세 포인터가 alias 안 한다고 알려줘.
// 그러면 원래 불법인 reorder가 unlock 됨.

Bandwidth 포화 — back-of-envelope 기대치·python

# RTX 4090: peak DRAM ~ 1.0 TB/s
# Vector add of N FP32 원소:
#   read:   2 * N * 4 byte (A + B)
#   write:  1 * N * 4 byte (C)
#   total:  3 * N * 4 byte
# Peak BW에서 시간: bytes / BW

N = 1 << 26   # 64M 원소 ≈ 768 MB 총 traffic
bytes_ = 3 * N * 4
time_s = bytes_ / 1.0e12
print(f'Vector add {N:,}: ~{time_s*1000:.2f} ms at peak BW')
# 건강한 커널은 이거 80–90% 도달. 50% 미만 → bandwidth-bound인데
# launch geometry 잘못: block 너무 적음, coalescing 나쁨, 등.

벡터가 뭐고 GPU가 왜 사랑하는가

GPU 모델 전체를 노출시키는 가장 작은 워크로드

Code

External links

Exercise

Progress

댓글 0