cuDNN 완전 해부: 딥러닝 연산이 GPU에서 빛처럼 빠른 이유

Naive CUDA 컨볼루션:       ~5 TFLOPS 달성 (H100 이론값의 0.5%)
cuDNN 최적화 컨볼루션:     ~900 TFLOPS 달성 (H100 이론값의 90%)
차이: ~180배

import torch

# 이 코드는 내부적으로 cuDNN을 호출
conv = torch.nn.Conv2d(64, 128, kernel_size=3, padding=1).cuda()
x = torch.randn(32, 64, 56, 56, device='cuda')
y = conv(x)  # → cudnnConvolutionForward() 호출

# cuDNN 사용 여부 확인
print(torch.backends.cudnn.enabled)   # True
print(torch.backends.cudnn.version())  # 예: 8906

출력 픽셀 (i, j)에 대해:
  for c_out in range(C_out):
    for c_in in range(C_in):
      for kh in range(K):
        for kw in range(K):
          output[c_out, i, j] += input[c_in, i+kh, j+kw] * filter[c_out, c_in, kh, kw]

총 연산 수: N × C_out × C_in × K^2 × H_out × W_out

입력 특징맵 [3×3 이미지, 3×3 커널]:

원본 입력:            im2col 변환 결과:
┌─────────┐           ┌────────────────────────────┐
│ 1  2  3 │  im2col   │ 1  2  4  5  (패치 0, 위치 0,0) │
│ 4  5  6 │ ────────→ │ 2  3  5  6  (패치 1, 위치 0,1) │
│ 7  8  9 │           │ 4  5  7  8  (패치 2, 위치 1,0) │
└─────────┘           │ 5  6  8  9  (패치 3, 위치 1,1) │
                      └────────────────────────────┘
각 행 = 수용 영역(receptive field) 하나의 모든 픽셀값

그 다음: output = filter_matrix × im2col_matrix
          (표준 GEMM → Tensor Core 완전 활용 가능!)

# im2col이 어떻게 동작하는지 Python으로 보는 예시
import numpy as np

def im2col(input, kernel_h, kernel_w, stride=1, pad=0):
    N, C, H, W = input.shape
    out_h = (H + 2*pad - kernel_h) // stride + 1
    out_w = (W + 2*pad - kernel_w) // stride + 1

    # 패딩 적용
    img = np.pad(input, [(0,0),(0,0),(pad,pad),(pad,pad)], mode='constant')

    # im2col 행렬 생성: (C*kernel_h*kernel_w) × (N*out_h*out_w)
    col = np.zeros((N, C, kernel_h, kernel_w, out_h, out_w))
    for j in range(kernel_h):
        jj = j + stride * np.arange(out_h)
        for i in range(kernel_w):
            ii = i + stride * np.arange(out_w)
            col[:, :, j, i, :, :] = img[:, :, jj[:, None], ii[None, :]]

    col = col.transpose(0, 4, 5, 1, 2, 3).reshape(N*out_h*out_w, -1)
    return col
# cuDNN의 실제 im2col은 이보다 훨씬 최적화되어 있음

일반 3×3 컨볼루션 (2×2 출력):
- 입력 패치: 4×4 = 16개 원소
- 커널: 3×3 = 9개 원소
- 출력: 2×2 = 4개 원소
- 필요한 곱셈 수 (naive): 4 × 9 = 36번

Winograd F(2×2, 3×3) 변환 후:
- 변환된 입력: 4×4 = 16개 원소 (선형 변환, 곱셈 없음)
- 원소별 곱셈: 16번 (4×4)
- 역변환: 4×4 → 2×2 (선형 변환, 곱셈 없음)
- 필요한 곱셈 수: 16번
- 절감: 36 → 16 = 2.25배 감소

Y = A^T [(G × g × G^T) ⊙ (B^T × d × B)] A

여기서:
d = 입력 타일 (4×4)
g = 3×3 커널
B, G, A = 고정된 변환 행렬 (사전 계산됨)
⊙ = 원소별 곱셈 (하다마르 곱)

핵심: G × g × G^T는 커널에 대해 한 번만 계산 (추론 시 미리 계산 가능)
      B^T × d × B는 각 입력 타일에 대해 계산
      둘을 원소별로 곱하면 36번 대신 16번의 곱셈으로 동일한 결과

시간 영역 컨볼루션:   O(N × K^2)  (N = 출력 크기, K = 커널 크기)
주파수 영역 컨볼루션: O(N × log N) (FFT 후 원소별 곱셈)

K가 클수록 FFT 방식이 유리:
K=3:  9 vs log(N) ≈ 7  → 큰 차이 없음
K=11: 121 vs log(N) ≈ 7 → FFT가 17배 적은 곱셈

# 이 한 줄이 실제로 하는 일은?
torch.backends.cudnn.benchmark = True

# benchmark=True의 효과를 직접 확인
import torch
import time

torch.backends.cudnn.benchmark = False
model = resnet50().cuda()
x = torch.randn(32, 3, 224, 224).cuda()

# 첫 번째 실행 (알고리즘 휴리스틱 선택)
t0 = time.time()
for _ in range(10): y = model(x)
torch.cuda.synchronize()
print(f"benchmark=False: {(time.time()-t0)*1000:.0f}ms")

torch.backends.cudnn.benchmark = True
# 첫 번째 실행에서 벤치마킹 (느림!)
t0 = time.time()
y = model(x)  # 이 한 번이 오래 걸림 (~수 초)
torch.cuda.synchronize()
print(f"첫 번째 실행 (벤치마킹 포함): {(time.time()-t0)*1000:.0f}ms")

# 이후 실행 (캐시된 최적 알고리즘 사용)
t0 = time.time()
for _ in range(10): y = model(x)
torch.cuda.synchronize()
print(f"benchmark=True (워밍업 후): {(time.time()-t0)*1000:.0f}ms")
# 보통 20-40% 빠름

y = gamma * (x - mean(x)) / sqrt(var(x) + eps) + beta

Naive 구현의 메모리 왕복:
1. x 로드 (HBM에서 읽기)
2. mean 계산 → HBM에 저장
3. x 로드 + mean 로드 (2회 HBM 읽기)
4. variance 계산 → HBM에 저장
5. x 로드 + mean 로드 + var 로드 (3회 HBM 읽기)
6. 정규화 계산 → HBM에 저장
7. ReLU: HBM 읽기 → 계산 → HBM 쓰기

총: HBM 접근 약 10회

Fused BN+ReLU 커널:
1. x를 레지스터/공유 메모리에 로드 (HBM 읽기 1회)
2. 레지스터에서 mean, variance 계산 (HBM 접근 없음)
3. 레지스터에서 정규화 (HBM 접근 없음)
4. 레지스터에서 ReLU (HBM 접근 없음)
5. 결과를 HBM에 저장 (HBM 쓰기 1회)

총: HBM 접근 2회 (5배 감소!)

# PyTorch에서 fused BN+ReLU 사용
import torch.nn as nn

# 기본 방식: Conv → BN → ReLU (3개 별도 커널)
class NaiveBlock(nn.Module):
    def __init__(self, c):
        super().__init__()
        self.conv = nn.Conv2d(c, c, 3, padding=1)
        self.bn = nn.BatchNorm2d(c)
        self.relu = nn.ReLU()

    def forward(self, x):
        x = self.conv(x)
        x = self.bn(x)    # 별도 cuDNN 커널
        x = self.relu(x)  # 별도 elementwise 커널
        return x

# cuDNN fused: torch.backends.cudnn.benchmark=True 환경에서
# cuDNN이 자동으로 conv+bn+relu를 하나의 커널로 최적화

# cuDNN Graph API (PyTorch 2.0+ torch.compile과 연계)
# Conv + BN + ReLU + Residual Add를 단일 커널로 처리
# torch.compile()이 자동으로 이런 퓨전을 수행
model = torch.compile(model)  # 내부적으로 cuDNN Graph API 활용

# Standard Attention — 메모리 비효율적
def standard_attention(Q, K, V, scale):
    # S = Q × K^T : (N, N) 크기의 어텐션 점수 행렬 생성
    S = torch.matmul(Q, K.transpose(-2, -1)) * scale  # N^2 메모리!
    # 소프트맥스
    P = torch.softmax(S, dim=-1)
    # V와 곱하기
    O = torch.matmul(P, V)
    return O

# 문제: S는 O(N^2) 메모리를 차지함
# N=8192 (문서 길이), FP16 기준:
# S 크기 = 8192 * 8192 * 2 bytes = 128MB (헤드 1개당!)

표준 어텐션 메모리 흐름:
Q, K, V → HBM 읽기
S = QK^T → HBM 쓰기 (N^2 크기!)
softmax(S) → HBM 쓰기
P × V → HBM 쓰기
총 HBM 접근: O(N^2)

FlashAttention 메모리 흐름:
Q, K, V를 타일 단위로 로드 → 공유 메모리
공유 메모리에서 어텐션 부분합 계산 (N^2 행렬 생성 없음)
최종 출력만 HBM에 저장
총 HBM 접근: O(N)  ← N^2에서 N으로!

FlashAttention 타일링:

Q를 블록 Q_1, Q_2, ... Q_Tc로 분할
K와 V를 블록 K_1, V_1, ... K_Tr, V_Tr로 분할

for i in range(Tc):
  Q_i를 SRAM에 로드
  O_i = 0, l_i = 0, m_i = -inf  (소프트맥스 통계 초기화)

  for j in range(Tr):
    K_j, V_j를 SRAM에 로드

    # SRAM에서만 계산 (HBM 접근 없음)
    S_ij = Q_i × K_j^T  (타일 크기의 어텐션 점수)
    m_ij = max(m_i, rowmax(S_ij))  (수치 안정성)
    P_ij = exp(S_ij - m_ij)

    # 온라인 소프트맥스 업데이트 (Numerically stable!)
    O_i = diag(exp(m_i - m_ij)) × O_i + P_ij × V_j
    l_i = exp(m_i - m_ij) × l_i + rowsum(P_ij)
    m_i = m_ij

  # 최종 정규화 후 HBM에 저장
  O_i = diag(l_i)^(-1) × O_i → HBM 쓰기 1회

NCHW 레이아웃 (batch × channel × height × width):
배치=1, 채널=3(RGB), 4×4 이미지:

메모리 배치:
[R00 R01 R02 R03 | R10 R11 ... | R30 R31 R32 R33 |
 G00 G01 G02 G03 | G10 G11 ... | G30 G31 G32 G33 |
 B00 B01 B02 B03 | B10 B11 ... | B30 B31 B32 B33]

→ 같은 채널의 픽셀들이 연속적으로 배치

NHWC 레이아웃 (batch × height × width × channel):
메모리 배치:
[R00 G00 B00 | R01 G01 B01 | R02 G02 B02 | R03 G03 B03 |
 R10 G10 B10 | R11 G11 B11 | ...
 R30 G30 B30 | R31 G31 B31 | R32 G32 B32 | R33 G33 B33]

→ 같은 공간 위치의 모든 채널이 연속적으로 배치

# PyTorch에서 NHWC 사용
x_nchw = torch.randn(32, 64, 56, 56, device='cuda')

# NHWC로 변환 (channels_last 포맷)
x_nhwc = x_nchw.to(memory_format=torch.channels_last)

# channels_last 형식으로 모델 변환
model = model.to(memory_format=torch.channels_last)

# 이 상태에서 forward pass는 cuDNN이 NHWC 커널을 자동 선택
output = model(x_nhwc)

# 벤치마크 비교
import time
model_nchw = resnet50().cuda()
model_nhwc = resnet50().cuda().to(memory_format=torch.channels_last)
x = torch.randn(64, 3, 224, 224).cuda()
x_nhwc = x.to(memory_format=torch.channels_last)

# 실측 결과 (H100 기준):
# NCHW: ~12ms/batch
# NHWC: ~9ms/batch  (약 25% 빠름)

원본 모델 (ONNX/PyTorch)
      ↓
  그래프 분석
      ↓
  레이어 퓨전
  ┌─────────────────────────────────────┐
  │ Conv → BN → ReLU → Conv → BN → ReLU│
  │       ↓ 퓨전 후                     │
  │     단일 최적화 커널                 │
  └─────────────────────────────────────┘
      ↓
  정밀도 선택 (FP32 → FP16 → INT8)
      ↓
  커널 자동선택 (입력 크기별 벤치마킹)
      ↓
  최적화된 실행 엔진

import tensorrt as trt

# INT8 캘리브레이터 설정
class MyCalibrator(trt.IInt8EntropyCalibrator2):
    def __init__(self, data_loader, cache_file):
        super().__init__()
        self.data_loader = iter(data_loader)
        self.cache_file = cache_file
        self.batch_allocation = None

    def get_batch_size(self):
        return 32

    def get_batch(self, names):
        try:
            batch = next(self.data_loader)[0].numpy()
            # 대표 데이터셋으로 activation 범위 캘리브레이션
            # TensorRT가 각 레이어의 FP32 범위를 INT8로 매핑
            if self.batch_allocation is None:
                self.batch_allocation = cuda.mem_alloc(batch.nbytes)
            cuda.memcpy_htod(self.batch_allocation, batch)
            return [int(self.batch_allocation)]
        except StopIteration:
            return None

# 빌더 설정
builder = trt.Builder(logger)
config = builder.create_builder_config()
config.set_flag(trt.BuilderFlag.INT8)  # INT8 활성화
config.int8_calibrator = MyCalibrator(calibration_loader, 'cache.bin')

# PyTorch Profiler로 LLM 연산 추적
import torch
from torch.profiler import profile, ProfilerActivity

model = GPT2Model.from_pretrained('gpt2').cuda()
input_ids = torch.randint(0, 50257, (1, 512)).cuda()

with profile(
    activities=[ProfilerActivity.CPU, ProfilerActivity.CUDA],
    with_stack=True
) as prof:
    output = model(input_ids)

# 상위 CUDA 커널 출력
print(prof.key_averages().table(sort_by="cuda_time_total", row_limit=10))

GPT-2 Forward Pass 내부 CUDA 커널 (512 토큰, batch=1):

연산              커널                          시간   비율
─────────────────────────────────────────────────────────
Linear (QKV)     cublasSgemm (Tensor Core)      2.1ms  35%
Attention        flash_attn_fwd_kernel           1.4ms  23%
Linear (output)  cublasSgemm (Tensor Core)      0.9ms  15%
LayerNorm        layer_norm_kernel               0.3ms   5%
GELU             vectorized_elementwise_kernel   0.2ms   3%
Residual Add     vectorized_elementwise_kernel   0.1ms   2%
Embedding        Embedding_cuda                  0.1ms   2%
기타                                            0.9ms  15%

총 시간: ~6ms (H100 기준)