Lecture 2 · Ch 1-3 PMPP book

GPU Mode 의 reading-group 흐름은 Programming Massively Parallel Processors (Hwu, Kirk, El Hajj 의 PMPP) 를 같이 읽는 데서 시작했다. L002 는 그 책의 처음 세 챕터를 한 시간으로 깐다 — 처음 책을 펴는 사람에게는 가이드라인을 주고, 이미 읽은 사람에게는 코드 위주의 요약을 준다.

강의가 답하려는 질문은 명확하다.

1. 병렬 사고가 시퀀셜 사고와 어떻게 다른가 — “문제를 어떻게 thread 들에 나눠 줄 것인가” 의 첫 결정.
2. host(CPU) 와 device(GPU) 메모리는 어떻게 분리돼 있고, 어떻게 옮기는가 — cudaMalloc, cudaMemcpy 의 첫 사용.
3. grid · block · thread 의 좌표계가 왜 그런 식으로 설계됐는가 — blockIdx · blockDim · threadIdx 로 indexing 하는 그 익숙한 한 줄의 의미.

이 세 질문을 추상으로 답하지 않고 — Andreas 가 깐 방식은 세 개의 커널을 차례로 보여주는 것이다. vector_addition.cu (1D, 일반 CUDA C 프로그램), RGB→grayscale (2D, PyTorch extension), mean filter (한 thread 가 여러 입력을 읽는 stencil). 각 커널이 직전 커널이 답하지 못한 새 차원을 하나씩 추가한다.

PMPP 가 제일 먼저 못 박는 사실. 한 thread 가 한 출력 element 를 만든다는 분해가 GPU 의 자연스러운 첫 번째 패턴이다. vector add, RGB→grayscale 모두 이 모델 안에 들어간다 — 각 thread 가 자기 좌표에 해당하는 output 하나만 책임진다. 입력끼리 의존성이 없는 (embarrassingly parallel) 문제일 때 이 분해가 그대로 통한다.

그런데 이 framing 안에서 두 가지가 빠르게 따라 나온다.

• thread 끼리 통신하지 않는다 — 적어도 첫 단계에서는. 각자 독립이라는 가정. (이게 깨지는 첫 자리가 뒤의 mean filter 와 §07.)
• 입력 데이터의 모양에 맞춰 launch shape 을 짠다 — vector 면 1D grid, image 면 2D, volume 이면 3D. dim3 가 그 도구.

강의에서 Andreas 가 명시적으로 말한 점 — “모든 알고리즘이 embarrassingly parallel 한 건 아니다. 의존성 있는 알고리즘 (예: prefix sum, reduction) 은 다른 분해가 필요하다.” 이건 L009 Reductions 가 본격 다루는 주제로 넘어간다. L002 의 framing 으로는 “일단 의존성 없는 케이스부터 손에 익힌다” 가 답.

GPU 는 자체 DRAM 을 가지고 있다. CPU 의 가상주소를 GPU 가 그대로 dereference 할 수 없다. 그래서 모든 “naive CUDA” 프로그램의 첫 5줄은 device 쪽 메모리를 잡고 (cudaMalloc), CPU 데이터를 옮기고 (cudaMemcpy H→D), 커널을 launch 하고, 결과를 다시 받고 (D→H), free 의 의식이다. 이 의식의 의미를 한 번 잡고 가야 그 다음의 모든 강의가 의미를 갖는다.

이 흐름이 PyTorch 안에서는 거의 보이지 않는다 — x.cuda() 한 줄이 L1+L2 를, op 호출이 L3 을 자동으로 한다. 하지만 강의 후반의 RGB→grayscale 과 mean filter 가 PyTorch extension 으로 깔리는 이유 중 하나가 이 의식을 손에서 빼주기 때문이다 — data_ptr<float>() 가 device pointer 를 그대로 돌려준다. 직접 짠 vector_add 와의 비교가 학습의 핵심.

CUDA 의 좌표계는 3차원 grid 안에 3차원 block, 그 block 안에 3차원 thread 의 nested 구조다. 처음 보면 복잡한 듯하지만, 하나의 직관적인 설계 결정에서 나온다 — thread 들은 동시에 실행되니까 launch shape 을 만들 때부터 각 thread 가 자기 좌표를 직접 계산할 수 있게 한다. 그게 다.

1D 인덱싱의 표준식은 다음 한 줄.

// PMPP Ch.2 Figure 2.6 의 표준식
int i = blockIdx.x * blockDim.x + threadIdx.x;

해석은 분해된다 — blockIdx.x 는 “나는 몇 번째 block 안에 있는가”, blockDim.x 는 “block 하나의 thread 수”, threadIdx.x 는 “block 안에서 내 자리”. 셋을 곱하고 더하면 전역 thread index. 이 한 줄이 vector add, grayscale, mean filter 의 첫 줄에 모두 똑같이 등장한다.

왜 block 이라는 중간 단계가 있는가 — 이게 PMPP 가 빠르게 답하는 다음 질문. block 안의 thread 들끼리만 fast 한 통신 (shared memory + __syncthreads) 이 보장된다. 다른 block 의 thread 와는 통신하지 않는다. block 단위로 SM(streaming multiprocessor) 위에 스케줄되기 때문이다. block 사이의 실행 순서나 동시성에 대해 어떤 가정도 할 수 없다 — “CUDA 가 block 들을 SM 들에 임의로 배정한다” 가 PMPP 의 명제.

강의의 첫 코드. PyTorch 의식 (load_inline, ninja) 없이, 그냥 nvcc vector_addition.cu 한 줄로 빌드되는 가장 작은 CUDA C 프로그램. 모든 것이 한 페이지에 들어 있어 “CUDA 프로그램이 무엇으로 구성되는가” 를 처음 본다.

// vector_addition.cu — 강의 repo 그대로 (요약)
__global__ void vecAddKernel(float* A, float* B,
                              float* C, int n) {
  int i = threadIdx.x + blockDim.x * blockIdx.x;
  if (i < n) {
    C[i] = A[i] + B[i];
  }
}

void vecAdd(float* A, float* B,
            float* C, int n) {
  float *A_d, *B_d, *C_d;
  size_t size = n * sizeof(float);

  cudaMalloc(&A_d, size);
  cudaMalloc(&B_d, size);
  cudaMalloc(&C_d, size);

  cudaMemcpy(A_d, A, size, cudaMemcpyHostToDevice);
  cudaMemcpy(B_d, B, size, cudaMemcpyHostToDevice);

  const unsigned int numThreads = 256;
  unsigned int numBlocks = cdiv(n, numThreads);

  vecAddKernel<<<numBlocks, numThreads>>>(
      A_d, B_d, C_d, n);
  gpuErrchk(cudaPeekAtLastError());
  gpuErrchk(cudaDeviceSynchronize());

  cudaMemcpy(C, C_d, size, cudaMemcpyDeviceToHost);
  cudaFree(A_d); cudaFree(B_d); cudaFree(C_d);
}

두 번째 커널은 HxW 짜리 RGB 이미지를 grayscale 로 변환. 한 출력 pixel 마다 한 thread 를 띄운다. 입력은 (H, W, 3) bytes, 출력은 (H, W) bytes. 같은 “한 thread = 한 출력” 패턴이지만 — 좌표가 2D 가 된다.

dim3 threads_per_block(16, 16);   // 256 threads, 16x16
dim3 number_of_blocks(
    cdiv(width,  16),
    cdiv(height, 16));

rgb_to_grayscale_kernel<<<
    number_of_blocks, threads_per_block, 0,
    torch::cuda::getCurrentCUDAStream()
>>>(
    result.data_ptr<unsigned char>(),
    image.data_ptr<unsigned char>(),
    width, height
);

int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
if (col < width && row < height) {
  int outOff = row * width + col;
  int inOff  = (row * width + col) * 3;
  unsigned char r = input[inOff + 0];
  unsigned char g = input[inOff + 1];
  unsigned char b = input[inOff + 2];
  output[outOff] = (unsigned char)
      (0.21f*r + 0.71f*g + 0.07f*b);
}

강의에서 이 시점에 흘러가는 부수적인 메시지 — PyTorch 와 함께 쓰는 CUDA 코드의 표준 패턴이 이미 여기서 다 나타난다. data_ptr, current stream, C10_CUDA_KERNEL_LAUNCH_CHECK(). 이 세 줄은 L010 의 production-ready CUDA library 강의까지 그대로 이어진다.

세 번째 커널은 mean filter — 각 출력 pixel 이 자기 주변 (2r+1)×(2r+1) window 의 평균이다. “한 출력당 한 thread” 의 패턴은 그대로지만, 한 thread 가 여러 입력을 읽는다는 새 차원이 추가된다. 이게 stencil 패턴의 가장 단순한 형태이고, 이후 모든 convolution, attention 의 부모 형태다.

// mean_filter_kernel.cu — 강의 repo (요약)
__global__ void mean_filter_kernel(
    unsigned char* output, unsigned char* input,
    int width, int height, int radius)
{
  int col = blockIdx.x * blockDim.x + threadIdx.x;
  int row = blockIdx.y * blockDim.y + threadIdx.y;
  int channel = threadIdx.z;
  int baseOff = channel * height * width;

  if (col < width && row < height) {
    int pixVal = 0, pixels = 0;
    for (int br = -radius; br <= radius; ++br) {
      for (int bc = -radius; bc <= radius; ++bc) {
        int r2 = row + br, c2 = col + bc;
        if (r2 >= 0 && r2 < height &&
            c2 >= 0 && c2 < width) {
          pixVal += input[baseOff + r2 * width + c2];
          pixels += 1;
        }
      }
    }
    output[baseOff + row * width + col] =
      (unsigned char)(pixVal / pixels);
  }
}

CUDA C 의 모든 함수는 셋 중 하나다 — host 에서만 도는 일반 함수 (__host__, 기본값), GPU 위에서만 도는 함수 (__device__), host 에서 launch 되어 GPU 에서 도는 entry-point 함수 (__global__). 이 qualifier 가 nvcc 에게 “이 함수의 binary 를 어디에 둘 것인가” 를 알려준다.

이 qualifier 들이 PMPP 책의 figure 와 직접 매칭된다. __global__ 함수는 GPU entry point — host 코드가 kernel<<<…>>>() 로 부른다. __device__ 함수는 GPU 안에서만 도는 helper. 둘 다 nvcc 가 PTX 로 컴파일해서 GPU executable 안에 넣는다. __host__ 는 일반 g++ 처럼 처리.

강의 후반에 Andreas 가 짧게 던진 한 마디가 시리즈 전체의 운영 원칙이다 — “데이터 모양이 매우 spontaneous 해서, 다른 모양에는 다른 커널이 최적일 수 있다.” 이 한 줄이 GPU Mode 시리즈가 같은 연산 (matmul, attention, reduction) 을 여러 강의에 걸쳐 반복하는 이유.

한 커널이 모든 모양에 답하려면 — autotune 이 답이다. Triton 의 @triton.autotune · CUTLASS 의 instance 기반 generation 이 그 표준 패턴이다. 그리고 PyTorch 자체가 op 마다 여러 구현을 들고 있고 dispatcher 가 모양에 따라 다른 걸 부른다 — L006 의 “dispatcher 통과 후 어떤 CUDA kernel 에 도달하는가” 가 이 사실의 다음 형태.

PMPP Ch.1–3 을 다시 펼치기 전에 머릿속에 있어야 하는 사실들과 — 직접 손에 박아야 하는 코드 자료.

손에 새기기 — 실습 시퀀스

1. vector_addition.cu 빌드 · 실행 — make 또는 nvcc vector_addition.cu -o vec. n=1000 의 결과를 직접 print 로 확인. 5단계 의식의 모든 줄을 손으로 따라 친다.
2. boundary check 일부러 빠뜨려보기 — if (i < n) 를 지우고 빌드해서 무엇이 일어나는지 본다. cuda-memcheck 또는 compute-sanitizer 가 어떻게 알려주는지 확인.
3. RGB → grayscale 직접 빌드 — rgb_to_grayscale.py 의 load_inline 호출을 추적, ./load_inline_cuda/ 디렉터리 안의 자동 생성된 main.cpp 를 직접 읽는다.
4. 2D launch shape 바꿔보기 — 16×16 → 32×32, 8×8, 1×256 으로 바꾸면서 결과 정확성과 시간 변화를 확인. 같은 코드가 launch shape 으로 다르게 동작하는 첫 경험.
5. mean filter radius sweep — radius ∈ {1, 3, 5, 9} 로 돌리면서 시간이 얼마나 늘어나는지 측정. radius 가 늘어나면 한 thread 가 읽는 pixel 수가 (2r+1)². data reuse 가 보이는가.
6. __device__ helper 함수 추가해보기 — luma 계산 부분을 __device__ __forceinline__ unsigned char rgb2y(...) 로 빼고 같은 결과가 나오는지 확인. nvcc 가 inline 시킨다.
7. 결과를 PyTorch 의 torchvision.transforms.Grayscale 과 비교 — 같은 입력에 두 결과가 byte 단위로 정확히 같은지. 다르면 어디가 다른지 (rounding, 채널 순서, 가중치).
8. 한 페이지 plan — “이 강의의 패턴을 LLM 추론의 한 자리에 어떻게 적용할까” 를 한 페이지로 적어본다 — 예: KV cache 의 한 head 에 대해 “한 thread = 한 (batch, seq) 좌표” 로 분해하는 식.

L002 의 framing (grid/block/thread, host↔device 의식, 한 thread = 한 출력) 이 시리즈 전체에 흩어져 있는 모든 강의의 출발 좌표가 된다.

PMPP Ch.1–3 의 한 시간 요약이 의도적으로 넘어간 자리들과 — 자기 GPU 위에서 직접 확인해야 손에 박히는 사실들.

• PMPP 의 “simplified square matmul” 이 강의에서 살짝만 등장한다 — Andreas 가 “책에서는 단순화된 정사각 matmul 도 한 thread = 한 출력 element 패턴으로 보여준다” 고 언급. 직접 책 그림 3.1x 를 펴서 확인 필요.
• Tensor Core 는 의도적으로 빠진다 — PMPP 의 첫 세 챕터는 일반 fp32 unit 만 다룬다. Tensor Core 는 후속 챕터와 L023 등에서 본격적으로. 이 강의의 framing 안에서는 “한 thread 의 단순 곱셈/덧셈” 으로 충분.
• CUDA stream 의 깊은 의미 — getCurrentCUDAStream 이 등장은 했지만 stream 사이의 동시성, 다중 stream 패턴은 빠진다. 후속 강의 (multi-stream reduction in L009) 에서.
• occupancy / register pressure — block 안에 thread 를 얼마나 많이 띄울 수 있나의 한계. 이 강의에서는 단순히 “16×16 = 256 이 흔하다” 정도로 멈춘다. L004 의 본론.
• 실제로 강의에서 빌드 시연을 어디까지 했는지 — 자막상 nvcc 의 직접 실행은 일부만 보였다. 실제 영상에서 어느 커널까지 라이브 빌드를 보였는지 재확인 필요. (확인 필요)
• RGB→grayscale 의 byte 정확도 — 0.21·R + 0.71·G + 0.07·B 가 float 으로 계산되고 byte 로 cast 된다. rounding 이 PyTorch 의 reference 와 정확히 일치하는가? 직접 비교 실습이 필요.