딥러닝과 CUDA

16.1 배경: 머신러닝, 그리고 GPU의 등장

머신러닝(machine learning)은 데이터에서 패턴을 추출하는 알고리즘 묶음이다. 크게 세 가지 학습 방식으로 나뉜다. 지도 학습(supervised)은 입력-정답 쌍으로 학습한다. 비지도 학습(unsupervised)은 정답 없이 구조를 찾는다. 강화 학습(reinforcement)은 보상을 최대화하는 정책을 찾는다. 이 책에서는 GPU의 가속 효과를 가장 극적으로 보여 주는 지도 학습, 그중에서도 신경망 기반 모델에 초점을 맞춘다.

지도 학습은 다시 회귀(regression, 연속값 예측)와 분류(classification, 이산 라벨 예측)로 나뉜다. 어느 쪽이든 본질은 함수 근사다. 입력 x를 출력 y로 보내는 함수 f_θ(x)를 데이터로부터 학습한다. θ는 모델의 파라미터.

가장 단순한 단위가 퍼셉트론(perceptron)이다. 입력 벡터 x ∈ ℝⁿ에 가중치 벡터 w ∈ ℝⁿ을 곱한 뒤 바이어스 b를 더하고, 비선형 활성화 함수 σ를 통과시킨다.

y = σ(w · x + b)

여기서 w · x는 내적(dot product). 활성화 σ는 시그모이드, ReLU, tanh 등. ReLU(x) = max(0, x)가 단순함과 안정적인 그래디언트 덕분에 현대 신경망의 표준이 되었다.

퍼셉트론을 한 줄로 쌓아 두면 레이어(layer)가 되고, 여러 레이어를 직렬로 연결하면 다층 퍼셉트론(MLP)이다. 한 레이어의 출력 h = σ(W h^(prev) + b)는 결국 행렬-벡터 곱과 활성화의 조합이다. 깊이가 L이면 L번의 행렬 곱이 차례로 일어난다.

학습이란 손실 함수 L(θ)를 최소화하도록 θ를 조정하는 과정이다. 가장 표준적인 방법이 경사 하강(gradient descent)이다. θ를 그래디언트 ∇_θL의 반대 방향으로 조금씩 이동시킨다.

θ ← θ − η ∇_θL

여기서 η는 학습률(learning rate). ∇_θL를 효율적으로 계산하기 위해 신경망은 역전파(backpropagation)를 쓴다. 체인 룰에 따라 출력층에서 입력층 방향으로 그래디언트를 전파한다. 한 레이어의 그래디언트는 다음 레이어로부터 받은 ∂L/∂h와 자기 가중치 W로부터 계산된다.

16.2 합성곱 신경망(CNN)의 구성 요소

이미지를 다루려면 MLP만으로는 모자라다. 28×28 이미지에서도 입력이 784차원, 천만 픽셀 이미지면 입력만 천만 차원이라 가중치가 폭발한다. 또한 픽셀의 위치 정보(공간적 인접성)를 살려야 한다. 그래서 합성곱 신경망(Convolutional Neural Network, CNN)이 등장했다. 핵심 아이디어는 작은 필터(kernel)를 이미지 전체에 슬라이딩하며 같은 가중치를 공유하는 것이다.

CNN의 표준 구성 요소는 다음과 같다.

• 합성곱 레이어(conv layer): 입력 특징맵에 K개의 필터를 컨볼루션한다. 출력은 K-채널 특징맵.
• 활성화(ReLU): 원소별 비선형. 음수를 잘라낸다.
• 풀링(pooling): 작은 윈도(예: 2×2) 안에서 max나 average를 취해 공간 해상도를 절반으로 줄인다. 파라미터는 없고 계산량을 줄이며 약한 변형 불변성을 준다.
• 완전 연결층(FC, fully connected): 마지막에 행렬 곱으로 클래스 점수를 낸다. 즉 MLP 한 두 층.
• 소프트맥스(softmax): 분류용. 클래스 점수를 확률 분포로 변환.

한 forward 패스는 conv → ReLU → conv → ReLU → pool → … → FC → softmax 식으로 진행된다. backprop은 정확히 이 흐름을 거꾸로 따라간다. softmax의 그래디언트로 시작해, FC를 통과하며 W^T 곱셈으로 그래디언트를 전파, conv 레이어에서는 입력에 대한 그래디언트(컨볼루션의 transpose)와 가중치에 대한 그래디언트(또 다른 컨볼루션) 두 가지를 동시에 계산한다.

그림 16.1 — CNN의 forward와 backward 흐름.

16.3 conv 레이어의 단순 CUDA 커널

이제 conv 레이어 하나를 GPU에 올려 보자. 입력 특징맵 X는 [C, H, W] 텐서다. C는 입력 채널 수(컬러 이미지면 3), H×W는 공간 해상도. 필터는 [K, C, R, S] 텐서다. K는 출력 채널 수, R×S는 필터 크기(보통 3×3). 출력 Y는 [K, H', W']이다. H' = H − R + 1 (패딩 0 가정).

한 출력 원소 Y[k, y, x]는 다음과 같이 계산된다.

Y[k, y, x] = sum_{c=0..C-1}
             sum_{r=0..R-1}
             sum_{s=0..S-1}
               W[k, c, r, s] * X[c, y+r, x+s]

가장 자연스러운 매핑은 출력 한 원소 = 한 스레드다. 7장에서 본 일반 컨볼루션의 직접적 일반화에 해당한다. 7장과 다른 점은 채널 차원 C에 대한 합산이 추가됐다는 점뿐이다.

__global__ void convForward(
    const float* X,   // [C, H, W]
    const float* W,   // [K, C, R, S]
    float* Y,         // [K, H', W']
    int C, int H, int W_,
    int K, int R, int S,
    int Hout, int Wout) {
  int x = blockIdx.x * blockDim.x + threadIdx.x;
  int y = blockIdx.y * blockDim.y + threadIdx.y;
  int k = blockIdx.z;
  if (x >= Wout || y >= Hout) return;

  float acc = 0.0f;
  for (int c = 0; c < C; ++c) {
    for (int r = 0; r < R; ++r) {
      for (int s = 0; s < S; ++s) {
        float xv = X[(c * H + (y + r)) * W_ + (x + s)];
        float wv = W[((k * C + c) * R + r) * S + s];
        acc += xv * wv;
      }
    }
  }
  Y[(k * Hout + y) * Wout + x] = acc;
}

이 커널은 깔끔해 보이지만 효율이 낮다. 두 가지 큰 문제가 있다. 첫째, 각 스레드가 채널 C번, 필터 R×S번을 순차적으로 읽으므로 메모리 트래픽이 출력 원소당 C·R·S번이다. 둘째, 인접 스레드들이 같은 입력 원소를 여러 번 다시 읽는다. 7장에서 본 셰어드 메모리 타일링이 여기서도 도움이 된다. 한 블록이 출력 타일을 담당하면서 입력 타일과 가중치를 셰어드 메모리에 한번 적재한다.

그러나 셰어드 타일링조차 conv의 진짜 잠재력을 끌어내기에는 부족하다. 더 강력한 접근이 GEMM 변환이다.

16.4 conv를 GEMM으로 변환하기 (im2col)

딥러닝 라이브러리가 가장 자주 쓰는 트릭이 im2col(image-to-column)이다. 핵심 아이디어는 이렇다. 컨볼루션의 본질은 결국 곱셈-누산이고, 이걸 큰 행렬 곱(GEMM)으로 다시 쓸 수 있다면 cuBLAS가 수십 년간 갈고닦은 GEMM 커널을 그대로 재사용할 수 있다.

변환 절차는 다음과 같다.

1. 입력 X [C, H, W]를 펼쳐서 행렬 X_col [C·R·S, H'·W']을 만든다. 출력 픽셀마다 그 위치에서 필터가 보는 R×S 윈도를 한 열로 쌓는다.
2. 가중치 W [K, C, R, S]를 그냥 [K, C·R·S]로 reshape한다. W_mat.
3. Y_mat = W_mat × X_col → 결과는 [K, H'·W']. reshape하면 Y [K, H', W'].

그림 16.2 — im2col로 conv를 GEMM으로 환원.

// im2col 커널 스케치
__global__ void im2col(
    const float* X,    // [C, H, W]
    float* X_col,        // [C*R*S, H'*W']
    int C, int H, int W_,
    int R, int S, int Hout, int Wout) {
  int col = blockIdx.x * blockDim.x + threadIdx.x;  // 출력 픽셀 인덱스
  int row = blockIdx.y * blockDim.y + threadIdx.y;  // (c, r, s) 통합 인덱스
  int nCols = Hout * Wout, nRows = C * R * S;
  if (col >= nCols || row >= nRows) return;

  int x = col % Wout;
  int y = col / Wout;
  int s = row % S;
  int r = (row / S) % R;
  int c = row / (R * S);

  X_col[row * nCols + col] = X[(c * H + y + r) * W_ + x + s];
}

im2col이 끝나면 cuBLAS의 sgemm 한 줄로 conv가 끝난다.

// im2col + GEMM 호출
im2col<<<..., ...>>>(X, X_col, C, H, W_, R, S, Hout, Wout);
cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N,
            Hout * Wout,  // M
            K,            // N
            C * R * S,    // K
            &alpha,
            X_col, Hout * Wout,
            W_mat, C * R * S,
            &beta,
            Y_mat, Hout * Wout);

이 접근의 단점은 메모리다. 입력 X의 크기가 C·H·W라면 X_col은 C·R·S·H'·W'다. 즉 R·S 배 정도 부풀려진다. 3×3 필터면 9배. 큰 모델, 큰 배치에서는 이 메모리 폭발이 부담이다. 그래도 GEMM의 압도적인 속도 덕분에 (cuBLAS sgemm은 V100에서 100 TFLOPS 근처까지 도달한다) 대부분의 경우 net 이득이 크다.

im2col의 변형으로 implicit GEMM이 있다. X_col을 실제로 메모리에 만들지 않고, GEMM 커널 안에서 인덱스 계산만으로 X에서 직접 가져오는 방식이다. 메모리 폭발을 피하면서 GEMM의 캐시 친화적 패턴을 거의 살릴 수 있다. cuDNN의 IMPLICIT_PRECOMP_GEMM 알고리즘이 이쪽이다.

16.5 cuDNN: 알고리즘 자동 선택과 텐서 코어

실전에서 conv 커널을 직접 쓰는 일은 드물다. NVIDIA가 제공하는 cuDNN(Deep Neural Network library)이 거의 모든 케이스에서 사람의 손으로 짠 것보다 빠르기 때문이다. cuDNN은 한 conv 연산에 대해 여러 알고리즘을 내장한다.

• IMPLICIT_GEMM: 메모리 폭발 없는 implicit im2col + GEMM.
• IMPLICIT_PRECOMP_GEMM: 인덱스 계산을 미리 해 두는 변형.
• GEMM: 일반 im2col 명시적.
• FFT: 큰 필터에 유리한 푸리에 도메인 곱.
• WINOGRAD: 작은 필터(3×3)에 특화된 곱셈 횟수 감소 알고리즘.
• WINOGRAD_NONFUSED: 위와 비슷하나 변환 단계가 분리.

입력 크기, 필터 크기, 배치 크기, 데이터 타입에 따라 어느 알고리즘이 가장 빠른지가 달라진다. cuDNN은 cudnnFindConvolutionForwardAlgorithm 같은 자동 튜너를 제공해서, 프로그램 시작 시 후보들을 짧게 벤치마크해 최적을 고른다. PyTorch에서 torch.backends.cudnn.benchmark = True가 바로 이걸 켜는 스위치다.

그림 16.3 — 프레임워크 → cuDNN → 텐서 코어로 이어지는 호출 스택.

또 다른 결정적 요소가 텐서 코어(Tensor Core)다. Volta 세대(V100)에서 처음 등장한 이 전용 행렬 곱 유닛은 한 사이클에 4×4×4 행렬 곱-누산을 처리한다. FP32로는 일반 SM의 4배, FP16/BF16/TF32 같은 저정밀도에서는 8~16배의 처리량을 낸다. cuDNN은 GEMM 단계에서 텐서 코어를 자동으로 호출한다. 그래서 데이터 타입을 FP16이나 BF16으로만 바꿔도 conv가 4~8배 빨라진다. 이를 혼합 정밀도(mixed precision) 학습이라고 부른다.

// PyTorch 측 사용 예 (참고)
// torch.set_float32_matmul_precision('high')   # TF32 켬
// model = model.half()                          # FP16
// with torch.autocast(device_type='cuda'):
//     y = model(x)

프레임워크 사용자 입장에서 보이는 건 nn.Conv2d 한 줄이지만, 그 한 줄 뒤에는 cuDNN의 알고리즘 선택, GEMM 변환, 텐서 코어 호출이 일어난다. 이 책의 챕터 7~14에서 다룬 모든 패턴(타일링, 셰어드 메모리, 코얼레싱, 워프 셔플, 스캔, 리덕션 등)이 그 안에서 사용되고 있다. 직접 짜진 않더라도 그 원리를 이해하는 것은 모델 성능 튜닝과 디버깅에 결정적이다.

16.6 정리

딥러닝과 GPU의 결혼은 우연이 아니었다. 신경망이 요구하는 핵심 연산 — 행렬 곱과 컨볼루션 — 이 GPU의 SIMD 아키텍처와 거의 맞춤형으로 어울렸기 때문이다. 한 conv 레이어를 살펴보면 7장의 컨볼루션, 12장의 머지(softmax 같은 reduction), 14장의 희소 곱셈(sparse attention), 그리고 BLAS의 GEMM이 모두 등장한다. 즉 이 책 전체가 딥러닝 한 모델의 forward 한 번에 응축되어 있다.

사용자 관점에서는 cuDNN과 cuBLAS, 그리고 PyTorch/TF/JAX 같은 프레임워크가 거의 모든 디테일을 가려 준다. 그러나 한 단계만 깊게 내려가도 — 모델이 학습 중 OOM이 나거나, 예상보다 느리거나, 새 GPU 세대에서 기대만큼 가속이 안 될 때 — 결국 우리가 알아야 하는 건 셰어드 메모리, 코얼레싱, 텐서 코어 alignment, 메모리 트래픽 같은 본 책의 주제들이다. 알고리즘과 하드웨어의 만남이 곧 성능이고, 성능이 곧 모델 규모이며, 모델 규모가 곧 능력이다.