Getting Started With CUDA for Python Programmers

CUDA 의 학습 곡선이 가파른 이유는 문법이 아니다. 알고리즘과 GPU 모델 두 개를 동시에 처음 보기 때문이다. Jeremy 의 진단이 정확하다 — “Python 으로 이미 짤 수 있는 알고리즘을, GPU 위에서 다시 짜면서 동시에 indexing/launch/메모리 모델까지 처음 만나니, 어디서 막힌 건지가 안 잡힌다.”

강의가 푸는 두 질문.

1. 알고리즘 자체는 익숙한 Python 위에서 먼저 손에 잡고 — print, breakpoint, 작은 입력을 자유롭게 쓰면서 — 그 위에 “GPU 의 모양” 을 layered 로 얹을 수 있는가.
2. 그렇게 익숙해진 Python kernel 을 거의 mechanical 하게 CUDA C 로 변환 할 수 있는가. (Jeremy 가 ChatGPT 에 그대로 붙여넣어 변환을 추천하는 자리.)

Jeremy 가 pmpp.ipynb 에서 깐 학습 패턴은 RGB→grayscale 과 matmul 두 예제에 똑같이 반복된다. 같은 알고리즘 → 세 단계의 변환: ① 일반 Python loop, ② 한 thread 시점의 “kernel” 함수 + 외부 loop, ③ block-aware 한 “kernel” 함수 + 이중 loop. 이 셋을 끝낸 다음 비로소 CUDA C 로 mechanical 하게 옮긴다.

for i in range(n):
    res[i] = 0.30*x[i] + 0.59*x[i+n] + 0.11*x[i+2n]

— 외부 loop, kernel 개념 없음 —

def k(i, x, out, n):
    out[i] = 0.30*x[i] + ...

run_kernel(k, n, x, res, n)

한 thread 시점만 짠다.
loop 는 runner 가.

def k(bi, ti, bdim, x, out, n):
    i = bi*bdim + ti
    if i < n: out[i] = ...

blk_kernel(k, blocks, 256, x, ...)

CUDA 의 indexing 한 줄을
Python 으로 흉내냄.

— 위 함수를 ChatGPT 에 붙여
변환 →

__global__ void k(uchar* x, ...) {
  int i = blockIdx.x*blockDim.x + threadIdx.x;
  if (i<n) out[i] = 0.30*x[i] + ...
}

이 사다리의 진짜 가치는 L1 과 L2 에서 print, breakpoint 가 그대로 동작한다는 점이다. i 가 어떤 값에서 잘못됐는지, blockidx*blockdim + threadidx 가 어떻게 계산되는지를 Python 안에서 직접 확인할 수 있다. CUDA C 로 옮긴 뒤에는 같은 디버깅이 printf(§09) 또는 cuda-gdb 의 영역.

예제는 (C, H, W) shape 의 RGB tensor 를 (H, W) shape 의 grayscale 로 변환. 가중치는 0.2989·R + 0.5870·G + 0.1140·B 의 luma 공식. notebook 의 첫 셀이 PIL/torchvision 으로 강아지 사진을 가져와 (3, 1066, 1600) tensor 를 만든다.

# pmpp.ipynb cell 9 — naive Python
def rgb2grey_py(x):
    c, h, w = x.shape
    n = h * w
    x = x.flatten()
    res = torch.empty(n, dtype=x.dtype,
                      device=x.device)
    for i in range(n):
        res[i] = (0.2989*x[i]
                + 0.5870*x[i+n]
                + 0.1140*x[i+2*n])
    return res.view(h, w)

이 단계에서 %%time 으로 측정한 시간이 약 3초 — 1066×1600 = 약 1.7M element 의 Python loop. 느린 게 핵심이 아니다, 알고리즘이 정확히 무엇을 계산하는지가 손에 잡히는 게 핵심.

두 번째 단계가 학습 사다리의 진짜 도구다. kernel 함수를 “한 thread 의 시점” 만으로 짜고, loop 는 외부 run_kernel 이 책임진다. 이 함수의 시그니처가 정확히 CUDA kernel 의 메탈 모델이다.

# pmpp.ipynb cell 13–17
def run_kernel(f, times, *args):
    for i in range(times):
        f(i, *args)               # 한 thread 의 호출

def rgb2grey_k(i, x, out, n):
    out[i] = (0.2989*x[i]
            + 0.5870*x[i+n]
            + 0.1140*x[i+2*n])

def rgb2grey_pyk(x):
    c, h, w = x.shape; n = h * w
    x = x.flatten()
    res = torch.empty(n, dtype=x.dtype, device=x.device)
    run_kernel(rgb2grey_k, h*w, x, res, n)   # N번 launch 흉내
    return res.view(h, w)

L0(naive Python) 와 L1(kernel mockup) 의 시간이 비슷하다 — Python loop 가 dominant 하니까. 시간을 줄이는 게 이 단계의 목표가 아니다. thread-시점의 함수를 손에 두는 게 목표.

L1(run_kernel) 은 “한 thread 시점” 만 모델링한다 — block 이 없다. 다음 단계가 block 을 추가한 blk_kernel. CUDA 의 launch shape <<<blocks, threads>>> 와 thread 의 indexing 식 blockIdx*blockDim + threadIdx 를 Python 으로 그대로 흉내낸다.

# pmpp.ipynb cell 21–23
def blk_kernel(f, blocks, threads, *args):
    for i in range(blocks):
        for j in range(threads):
            f(i, j, threads, *args)               # (blockIdx, threadIdx, blockDim)

def rgb2grey_bk(blockidx, threadidx, blockdim, x, out, n):
    i = blockidx * blockdim + threadidx
    if i < n:                                # boundary guard — CUDA 와 동일
        out[i] = 0.2989*x[i] + 0.5870*x[i+n] + 0.1140*x[i+2*n]

def rgb2grey_pybk(x):
    c, h, w = x.shape; n = h*w
    x = x.flatten(); res = torch.empty(n, dtype=x.dtype, device=x.device)
    threads = 256
    blocks = math.ceil(h*w / threads)
    blk_kernel(rgb2grey_bk, blocks, threads, x, res, n)
    return res.view(h, w)

이 시점에서 학습자가 손에 가지고 있는 건 — CUDA 와 정확히 같은 모양의 알고리즘이다. 두 차이만 남았다: ① Python 의 indexing 산술이 GPU 의 32-bit register 산술로, ② Python 의 외부 loop 가 GPU 의 SM 위 병렬 실행으로. 두 변환 모두 mechanical 하다.

CUDA 코드를 끼워넣는 도구는 L001 §05 에서 본 torch.utils.cpp_extension.load_inline 그대로다. Jeremy 는 자주 쓰는 옵션을 묶어 한 줄짜리 wrapper 로 만든다.

# pmpp.ipynb cell 28
def load_cuda(cuda_src, cpp_src, funcs,
              opt=False, verbose=False):
    return load_inline(
        cuda_sources=[cuda_src],
        cpp_sources=[cpp_src],
        functions=funcs,
        extra_cuda_cflags=["-O2"] if opt else [],
        verbose=verbose,
        name="inline_ext",
    )

# 모든 CUDA 셀 머리에 붙는 boilerplate
cuda_begin = r'''
#include <torch/extension.h>
#include <stdio.h>
#include <c10/cuda/CUDAException.h>

#define CHECK_CUDA(x) TORCH_CHECK(x.device().is_cuda(), #x " must be a CUDA tensor")
#define CHECK_CONTIGUOUS(x) TORCH_CHECK(x.is_contiguous(), #x " must be contiguous")
#define CHECK_INPUT(x) CHECK_CUDA(x); CHECK_CONTIGUOUS(x)

inline unsigned int cdiv(unsigned int a, unsigned int b) {
  return (a + b - 1) / b;
}
'''

학습 사다리의 마지막 단. Python rgb2grey_bk 를 그대로 ChatGPT 또는 손으로 CUDA C 로 옮긴다. 변환의 “보일러플레이트” 는 cuda_begin 가 다 써놨고, 함수 본체는 거의 한 글자도 안 바뀐다.

// pmpp.ipynb cell 33 — load_cuda 가 빌드하는 본체
__global__ void rgb_to_grayscale_kernel(
    unsigned char* x, unsigned char* out, int n)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;          // ← 이 한 줄만 추가
    if (i < n)
        out[i] = 0.2989f*x[i] + 0.5870f*x[i+n]
               + 0.1140f*x[i+2*n];
}

torch::Tensor rgb_to_grayscale(torch::Tensor input) {
    CHECK_INPUT(input);                                       // 매크로의 가치
    int h = input.size(1);
    int w = input.size(2);
    auto output = torch::empty({h, w}, input.options());
    int threads = 256;
    rgb_to_grayscale_kernel<<<cdiv(w*h, threads), threads>>>(
        input.data_ptr<unsigned char>(),
        output.data_ptr<unsigned char>(),
        w*h);
    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}

강의 후반부는 같은 학습 패턴을 matmul 에 적용. 입력은 MNIST 의 작은 sample (m1.shape = (5, 784)) 와 random 가중치 (m2.shape = (784, 10)). 출력은 (5, 10). 작은 사이즈를 쓰는 이유 — Python 단계에서 %time 으로 측정 가능한 범위에 두기 위해.

# pmpp.ipynb cell 56 — naive Python matmul
def matmul(a, b):
    (ar, ac), (br, bc) = a.shape, b.shape
    c = torch.zeros(ar, bc)
    for i in range(ar):
        for j in range(bc):
            for k in range(ac):
                c[i,j] += a[i,k] * b[k,j]
    return c

# 2D Python kernel — 한 출력 element 의 시점
def matmul_k(i, j, a, b, c, ac):
    s = 0.0
    for k in range(ac):
        s += a[i,k] * b[k,j]
    c[i,j] = s

그리고 같은 사다리로 CUDA 단계까지 — 2D launch (dim3(16, 16) threads, dim3(cdiv(bc,16), cdiv(ar,16)) blocks) 로 띄우고, kernel 안에서 row = blockIdx.y*blockDim.y + threadIdx.y 와 col = blockIdx.x*blockDim.x + threadIdx.x 두 줄로 좌표를 계산. 나머지는 Python 의 inner loop 가 그대로.

이 단순 CUDA matmul 은 shared memory 를 쓰지 않는다 — 매 element 마다 a 의 한 row, b 의 한 column 을 global memory 에서 다시 읽는다. data reuse 가 보이지만 잡지 않은 상태가 의도적이다. L005 가 이 코드 위에 shared memory tiling 을 한 단 더 얹어 “Going Further” 의 의미가 된다.

notebook 환경에서 CUDA 디버깅을 하면 흔히 만나는 두 문제 — error 가 어디서 났는지 위치가 비동기 launch 때문에 어긋난다, kernel 안의 printf 출력이 stdout 에 안 잡힌다. Jeremy 가 두 가지 환경 설정을 셀 머리에 박는다.

이 강의에서 6개월 뒤 다시 돌아왔을 때 가장 먼저 복원해야 하는 사실들과 — 직접 손에 박아야 하는 코드 자료들.

손에 새기기 — 실습 시퀀스

1. 강아지 사진을 자기 노트북 환경에서 grayscale — pmpp.ipynb 를 그대로 돌려 4단계 (naive → kernel → block kernel → CUDA) 의 결과가 byte 단위로 동일 한지 검증.
2. L0 → L1 변환 직접 해보기 — 다른 알고리즘 (예: per-row max, per-row sum) 을 같은 4단계로 변환. 같은 변수명을 유지하면서.
3. boundary check 일부러 빠뜨려서 본다 — Python 단계에서는 IndexError 가 나고, CUDA 단계에서는 illegal memory access. 같은 버그가 두 환경에서 어떻게 다르게 보이는지.
4. thread 수 sweep — 256 → 64, 128, 512, 1024. %timeit 으로 시간 변화. n_warps 와의 관계 추측.
5. matmul 을 작은 사이즈에서 검증, 큰 사이즈에서 측정 — (5, 784) × (784, 10) 의 결과를 PyTorch torch.matmul 과 torch.allclose. 그 다음 (1024, 1024) × (1024, 1024) 로 키워 시간 비교 — naive 가 cuBLAS 와 얼마나 차이나는가.
6. CUDA_LAUNCH_BLOCKING 의 효과 보기 — 일부러 잘못된 indexing (예: i+5) 을 넣어서 launch 했을 때, BLOCKING=1 과 0 에서 traceback 위치가 어떻게 다른지.
7. printf 디버깅 — kernel 안에 if (i < 5) printf("i=%d, val=%f\n", i, x[i]); 를 박고 wurlitzer 로 볼 수 있는지 확인. 작은 입력에 한정.
8. 한 페이지 plan — 자기가 짜고 싶은 다음 커널을 정하고, “Python 4단계 → CUDA” 의 사다리에 맞춰 계획서 작성.

L003 의 “Python 위에서 알고리즘을 손에 잡는다” 패턴이 시리즈의 여러 강의에서 한 단씩 확장된다.

학습 사다리의 framing 안에서 의도적으로 비워둔 자리들과 — 직접 자기 환경에서 확인해야 손에 박히는 사실들.

• shared memory 가 빠진 matmul 의 시간 — 강의의 단순 CUDA matmul 이 cuBLAS 와 얼마나 차이나는지의 정량은 강의에서 직접 측정하지 않는다. “not a CUDA expert” 라고 Jeremy 도 명시. L005 에서 직접 측정.
• PTX 또는 SASS 직접 보기 — 강의의 사다리에서는 PTX 를 까지 않는다. nvcc -ptx 또는 cuobjdump 로 같은 알고리즘이 어떤 instruction sequence 가 되는지는 별도로 검증.
• LLM 변환의 신뢰성 — Jeremy 가 ChatGPT 변환을 추천하지만, 어디서 실패하는가 는 다루지 않는다. 복잡한 indexing, conditional flow, atomic operation 이 들어간 알고리즘에서 변환의 정확도는 직접 검증 필요.
• CUDA_LAUNCH_BLOCKING=1 의 성능 비용 — “디버깅용” 이라고 했지만, 정량적인 성능 차이는 강의에서 측정하지 않는다. 실제로 얼마나 느려지는지 자기 GPU 에서 측정 필요.
• multi-GPU / distributed 는 빠진다 — 이 강의의 모든 예제는 single GPU. 다음 단계의 학습 사다리는 L017 의 PyTorch distributed 또는 NCCL 강의들.
• Triton 과의 비교 — 강의에서 Triton 은 거의 언급하지 않는다. 같은 학습 사다리가 Triton 에 어떻게 적용되는지는 L014 의 Triton 강의에서.