제3부 딥러닝 라이브러리

3.3 cuDNN II — 그래프 융합 API·휴리스틱

3.3.1 이 장에서 배우는 것 / 선수 지식

3.2에서 컨볼루션 알고리즘들(implicit GEMM·Winograd·FFT)을 봤다. 이 장은 그것을 고르는 두 방식(heuristics·find)과 워크스페이스 트레이드오프, 그리고 그래프 API(v8+) 로 conv-bias-act 같은 패턴을 한 커널로 융합하는 법을 다룬다. 실무 증상(첫 스텝 지연·형상별 성능 차·버전 회귀·OOM)이 전부 여기서 나온다.

배우는 것:

• 알고리즘 선택을 heuristics(즉시 추정) 와 find(실측 자동튜닝) 로 호출하고, 워크스페이스·워밍업·재현성 트레이드오프를 설명한다.
• cuDNN 그래프 API로 conv-bias-act 융합을 연산 그래프(노드 + 가상 텐서)로 표현한다.
• 입력 형상 의존·워밍업·버전 회귀·워크스페이스 OOM 함정을 진단한다.

선수 지식: 3.2(프리미티브·컨볼루션 알고리즘), 3.1(cuBLASLt 에필로그 융합 — cuDNN 그래프 융합과 대비), 1.9~1.10(HBM 왕복·메모리 바운드), 2.13~2.14(프로파일링). 코드 골격은 개략이며 정확한 API는 설치 버전·cudnn-frontend로 확인.

3.3.2 알고리즘 선택 — heuristics vs find

cuDNN은 "이 형상에 어떤 알고리즘을 쓸까"를 두 방식으로 답한다. 이 둘의 차이가 곧 첫 스텝 지연·워밍업·재현성 문제의 근원이다.

heuristics — 즉시 추정. 라이브러리 내부 휴리스틱 모델(과거 벤치마크·형상 특징 기반)이 실측 없이 추천 알고리즘 목록을 순위와 함께 즉시 반환한다. 빠르다(밀리초 미만), 단 추정이라 항상 최적은 아니다(드문 형상·새 GPU에서 빗나갈 수 있음).

find — 실측 자동튜닝. 후보 알고리즘들을 실제로 실행해 시간을 재고 가장 빠른 것을 고른다. 그 형상·GPU·버전에서 실제 최고를 찾지만 느리다 — 후보마다 워크스페이스를 할당해 한 번씩 실행해 첫 호출이 수백 ms~수 초. 이것이 "첫 스텝이 느린" 정체다. PyTorch torch.backends.cudnn.benchmark = True가 이 find 모드를 켜고 형상마다 한 번 실측해 캐시한다.

어떤 것을 쓰나:

견해: 학습은 거의 항상 benchmark=True가 옳다(형상 고정). 단 워밍업 스텝의 지연을 측정에 포함하지 마라(2.13).

워크스페이스 트레이드오프. 알고리즘마다 임시 메모리(워크스페이스, HBM)를 요구한다. 워크스페이스를 더 주면 더 빠른 알고리즘(FFT·Winograd 변형)이 후보에 들어온다. 적게 주면 implicit GEMM 같은 저메모리 알고리즘만 남는다.

함정: 워크스페이스 한도를 OOM 직전까지 키웠다가 배치 크기를 늘리면 알고리즘 워크스페이스 + 활성화 텐서가 동시에 커져 OOM. 워크스페이스는 "무비용 속도"가 아니라 메모리를 건 베팅이다(특히 학습은 backward 워크스페이스까지).

3.3.3 워크스루 A — 알고리즘 자동선택 (개략)

heuristics로 추천을 받은 뒤 find로 실측해 둘을 비교한다(레거시 스타일 골격, 정확한 API는 설치 버전 헤더로):

// cuDNN convolution forward — 알고리즘 선택 골격 (개략)
cudnnHandle_t h; cudnnCreate(&h);
// 디스크립터: 입력 N=32,C=64,H=W=56 / 필터 K=64,C=64,3×3 / pad=1, stride=1 (NHWC, FP16)
cudnnSetTensor4dDescriptor(xDesc, CUDNN_TENSOR_NHWC, CUDNN_DATA_HALF, 32,64,56,56);
cudnnSetFilter4dDescriptor(wDesc, CUDNN_DATA_HALF, CUDNN_TENSOR_NHWC, 64,64,3,3);
cudnnSetConvolution2dDescriptor(convDesc, 1,1, 1,1, 1,1, CUDNN_CROSS_CORRELATION, CUDNN_DATA_FLOAT); // 누적 FP32
cudnnSetConvolutionMathType(convDesc, CUDNN_TENSOR_OP_MATH); // 텐서코어 허용

// heuristics: 실측 없이 추천 순위
int cnt=0; cudnnConvolutionFwdAlgoPerf_t heur[8];
cudnnGetConvolutionForwardAlgorithm_v7(h, xDesc, wDesc, convDesc, yDesc, 8, &cnt, heur);
// heur[0]=추정 최고. heur[i].algo/.memory(워크스페이스)/.time(추정)

// find: 실제로 실행해서 측정 (워크스페이스 버퍼 필요)
size_t wsBytes = 256ull<<20; void* ws; cudaMalloc(&ws, wsBytes);
int fcnt=0; cudnnConvolutionFwdAlgoPerf_t found[8];
cudnnFindConvolutionForwardAlgorithmEx(h, xDesc,x, wDesc,w, convDesc, yDesc,y, 8, &fcnt, found, ws, wsBytes);
// found[0].algo = 이 형상·GPU·버전에서 실측 최고. found[0].time = 실측 시간

float a=1.f, b=0.f;
cudnnConvolutionForward(h, &a, xDesc,x, wDesc,w, convDesc, found[0].algo, ws, wsBytes, &b, yDesc, y);

관찰: heur[0].algo와 found[0].algo가 일치하지 않을 수 있다(추정 vs 실측). wsBytes를 줄여 다시 실행하면 후보가 줄고 최고 시간이 나빠질 수 있다(워크스페이스↔속도). 입력 H·W를 56→55로 바꾸면 최적 알고리즘이 바뀔 수 있다(형상 의존).

3.3.4 그래프 API (v8+) — 융합을 그래프로 표현

cuDNN v8부터 그래프 API가 도입됐다. "연산 하나에 함수 하나" 대신 연산들을 노드로 잇는 그래프를 만들고, 라이브러리가 그것을 융합 가능한 커널로 컴파일한다. cuBLASLt 에필로그(3.1)를 일반화해 훨씬 넓은 패턴(conv-bias-act, BN-act, SDPA)을 한 커널로 합친다.

왜 융합인가 — HBM 왕복 제거. 비융합은 각 연산이 결과를 HBM에 쓰고 다음 연산이 다시 읽는다:

융합 쪽은 중간 텐서를 HBM에 쓰지 않아 대역폭을 절약한다.

활성·bias·정규화는 메모리 바운드(1.10)다. 이들을 앞단 conv·GEMM에 붙여 중간 텐서 HBM 왕복을 없애면 메모리 대역폭이 병목인 구간에서 큰 이득.

핵심 개념 — 가상 텐서(virtual tensor). 그래프 내부에서만 존재하고 HBM에 구현(materialize)되지 않는 중간 텐서. conv 출력을 bias 노드 입력으로 잇되 "이것은 가상"이라고 표시하면 라이브러리가 메모리에 안 쓰고 융합한다. 융합의 단위가 이 가상 텐서 연결로 표현된다.

라이브러리는 이 그래프를 융합 엔진 후보(heuristics)로 평가해 실행 계획을 세우고, conv+add+relu를 한 커널로 합친다.

이것이 SDPA 융합으로도 확장된다 — MATMUL(QKᵀ) → scale → SOFTMAX → MATMUL(×V)를 한 그래프로 표현하면 FlashAttention 유형 융합 커널로 실행(세부·가용성은 버전·세대 의존, 본격 내용 3.6).

cuBLASLt 에필로그와의 차이:

거칠게: cuBLASLt 에필로그는 "GEMM에 정해진 꼬리를 붙이는" 특수 융합, cuDNN 그래프는 "임의 패턴을 그래프로 표현하는" 일반 융합. 경쟁이 아니라 역할이 다르다(GEMM 본체는 cuBLAS, 컨볼루션 패턴은 cuDNN).

3.3.5 워크스루 B — conv+bias+act 융합 (cudnn-frontend 개략)

C++ 프론트엔드(cudnn_frontend, 헤더-온리)로 그래프를 간결하게(정확한 빌더 이름은 버전 의존):

namespace fe = cudnn_frontend;
auto graph = std::make_shared<fe::graph::Graph>();
graph->set_io_data_type(fe::DataType_t::HALF).set_compute_data_type(fe::DataType_t::FLOAT); // 누적 FP32
auto X = graph->tensor(...{32,64,56,56} NHWC...);
auto W = graph->tensor(...{64,64,3,3}...);
auto Bias = graph->tensor(...{1,64,1,1}...);
auto Yc = graph->conv_fprop(X, W, fe::graph::Conv_fprop_attributes().set_padding({1,1}).set_stride({1,1}));
auto Yb = graph->pointwise(Yc, Bias, fe::graph::Pointwise_attributes().set_mode(fe::PointwiseMode_t::ADD));
auto Y  = graph->pointwise(Yb, fe::graph::Pointwise_attributes().set_mode(fe::PointwiseMode_t::RELU_FWD));
Y->set_output(true);   // 이것만 HBM에 구현, Yc·Yb는 가상 → 융합
graph->validate(); graph->build_operation_graph(handle);
graph->create_execution_plans({fe::HeurMode_t::A});      // heuristics 모드
graph->build_plans(handle, fe::BuildPlanPolicy_t::HEURISTICS_CHOICE);
int64_t wsBytes = graph->get_workspace_size(); void* ws; cudaMalloc(&ws, wsBytes);
graph->execute(handle, {{X,xPtr},{W,wPtr},{Bias,bPtr},{Y,yPtr}}, ws);  // conv+bias+relu가 한 커널로

관찰: Yc·Yb에 set_output(true)를 안 한 것이 핵심 — 가상 텐서로 남아 융합된다. 디버깅 위해 Yc를 출력으로 표시하면 융합이 깨지고 HBM에 쓰인다(중간값 확인용이지만 느려진다). build_plans에 여러 후보를 빌드해 직접 시간을 재면 그래프 API에서 find에 해당하는 autotune이 된다.

수치로 보는 융합 이득(개략). conv 출력 $Y: 32\times64\times56\times56$ FP16 ≈ 12.8 MB. 비융합이면 이 텐서를 bias에서 쓰고·읽고, relu에서 쓰고·읽고 — 중간 텐서로 인한 HBM 트래픽이 대략 $12.8\mathrm{MB} \times4 \approx51 \,\mathrm{MB}$. 융합하면 최종 Y만 쓴다(12.8MB). 대역폭 1.5 TB/s(개략)라면 절약된 ~38MB는 $38\mathrm{e}6/1.5\mathrm{e}12 \approx25 \,\mu\mathrm{s}$. 작아 보여도 활성·정규화처럼 연산량이 적은 메모리 바운드일수록 이 왕복이 전체의 큰 비중이고, 레이어 수백 개의 학습에서 누적이 크다. (실제 트래픽은 캐시·NHWC·구현 의존, 확인 필요.)

3.3.6 중간 점검 (체크포인트)

• [ ] heuristics와 find의 차이를 "첫 스텝 지연"과 연결해 설명하고, 학습 vs 가변 형상 추론에서 무엇을 켤지 고른다.
• [ ] 워크스페이스 한도를 키울 때의 이득(더 빠른 후보)과 위험(메모리·OOM)을 동시에 말한다.
• [ ] conv-bias-act 융합을 그래프(노드 + 가상 텐서)로 그리고, 가상 텐서가 융합의 표현임을 설명한다.

증상→의심: 첫 스텝(또는 형상 바뀐 첫 스텝)만 수십 배 느림 → find/benchmark=True 자동튜닝 비용(정상, 측정 제외). 입력 크기 바꿨더니 갑자기 느림 → 알고리즘 형상 의존. 버전 상향 뒤 특정 레이어가 느림 → 알고리즘 회귀(실측 비교). 배치 키울 때만 OOM → 워크스페이스 + 활성화 동시 증가(한도↓로 저메모리 알고리즘 강제).

3.3.7 트레이드오프

견해: 학습은 프레임워크 + benchmark=True로 충분한 경우가 많다. 그래프 API 직접은 프레임워크가 융합 못 하는 커스텀 패턴이나 라이브러리 저작 시. 추론 배포면 cuDNN 직접보다 TensorRT가 더 넓게 융합·튜닝(5.2).

비NVIDIA 정직: 같은 자리(딥러닝 프리미티브·융합)를 AMD MIOpen·Intel oneDNN이 채운다. 컴파일러 경로(torch.compile·Triton·XLA·TVM)는 cuDNN 호출 대신/병행 자체 생성 커널로 융합한다(손튜닝 라이브러리 vs 컴파일러 자동 코드생성). 어느 쪽이든 "형상 의존·자동튜닝·워밍업·버전 회귀"라는 이 장의 함정은 공통이다.

3.3.8 흔한 함정

• 워밍업을 측정에 포함. find/benchmark가 켜져 있으면 첫 호출(형상마다)이 자동튜닝으로 느리다. 워밍업 N 스텝을 버리고 측정(2.13).
• 버전 회귀. cuDNN 업그레이드는 휴리스틱·기본 알고리즘·융합 엔진을 바꾼다. 보통 빨라지나 특정 형상에서 느려지거나(회귀) 수치가 미세하게 달라질 수 있다. 업그레이드 후 핵심 형상 회귀 테스트.
• 워크스페이스 OOM. 한도를 키우면 빠른 알고리즘이 들어오나 활성화 텐서·다른 버퍼와 합쳐 OOM. 학습은 backward 워크스페이스까지 — forward만 보고 한도 정하면 안 됨.
• 결정론 깨짐. find는 측정 노이즈로 다른 알고리즘을 선택할 수 있고 일부(atomics backward)는 비결정적. 비트 재현이 필요하면 torch.use_deterministic_algorithms(True) — 대신 느려지거나 일부 연산이 제한된다.
• 융합 디버깅 중 가상 텐서를 출력으로 표시. 중간값 확인하려 conv 출력을 output으로 빼면 그 즉시 융합이 풀려 HBM에 쓰인다("디버깅하니 느려졌다"의 정체). 측정은 가상 텐서 그대로 두고.

3.3.9 연습문제

연습 1 (benchmark 옵션). cudnn.benchmark = True를 켰더니 학습 첫 50 스텝이 불규칙하게 느리다가 이후 안정적으로 빨라졌다. 무슨 일이고, 이 환경에서 계속 켜둘지?

연습 2 (워크스페이스). 워크스페이스 한도를 256MB→16MB로 줄였더니 컨볼루션이 느려졌다. 메커니즘은? 무한정 키우면 안 되는 이유는?

연습 3 (가상 텐서). conv→bias→ReLU 그래프에서 중간 텐서를 가상으로 두는 것과 출력으로 구현하는 것의 성능 차이를 대역폭으로. 디버깅 시 왜 함정인가?

연습 4 (두 융합 비교). cuBLASLt 에필로그(3.1)와 cuDNN 그래프 융합은 둘 다 "연산을 붙여 HBM 왕복을 줄인다". 왜 두 라이브러리에 따로 있나? 각자 강한 영역은?

3.3.10 요약 · 다음 장

• 알고리즘 선택은 heuristics(즉시 추정) 와 find(실측 자동튜닝) 두 방식. 학습(형상 고정)은 find/benchmark=True, 가변 형상은 heuristics. 워크스페이스를 더 주면 빠른 후보가 늘지만 메모리를 건다.
• 그래프 API(v8+) 는 연산을 노드·가상 텐서 그래프로 표현해 conv-bias-act·SDPA를 한 커널로 융합, 중간 텐서 HBM 왕복을 없앤다.
• 실무 증상(첫 스텝 지연·형상별 성능 차·버전 회귀·OOM)은 전부 이 메커니즘에서 — 운영 형상으로 실측·워밍업 제외·버전 회귀 테스트·워크스페이스 한도 관리가 처방.

다음 장: 3.4 CUTLASS I — 계층·GEMM 구성. cuDNN·cuBLAS가 닫힌 라이브러리라면, CUTLASS는 GEMM·conv 커널을 직접 조립하는 템플릿이다 — 라이브러리에 없는 융합·형상이 필요할 때.