⑤ CUDA Graphs

한 줄로 — 매번 똑같이 반복되는 GPU 작업 묶음을 "그래프"로 한 번 찍어두고, 그다음부터는 통째로 한 방에 실행하는 기술입니다. ③에서 본 launch 비용과 지터(들쭉날쭉함)를 없애줘요. 조금 고급 주제지만, 비유로 차근차근 가봅시다.

5.1 왜 그래프인가 — "주문서를 매번 새로 쓰는 낭비"

③에서 커널 호출은 "주방에 주문서를 넣는 것"이라고 했죠. 그 주문서를 넣는 행위 자체가 ~µs(마이크로초)의 비용이었습니다. 큰 커널 하나면 무시할 만한데, 작은 커널을 수십 개씩 매번 반복하면 어떻게 될까요?

// 매 반복마다 주문서를 N장 새로 쓰는 셈 → N번의 비용 + 들쭉날쭉(지터)
a<<<...>>>();  b<<<...>>>();  c<<<...>>>();  d<<<...>>>();   // 반복...

매번 똑같은 순서로 a→b→c→d를 돌린다면, "이 순서를 한 번만 적어두고 다음부턴 그걸 통째로 실행"하면 낭비가 사라지겠죠. 그게 CUDA Graphs입니다:

// 한 번 그래프로 찍어두면, 그다음부턴:
cudaGraphLaunch(graphExec, stream);   // 단 1번의 호출로 a→b→c→d 전부 실행

💡 얻는 것 3가지

비용 분할상환 — 주문서 N장 쓰던 걸 1번으로.
CPU 해방 + 지터↓ — CPU는 한 번만 던지고 빠짐, GPU가 알아서 쭉 실행 → 타이밍이 일정해짐.
드라이버가 미리 최적화 — 전체 순서를 미리 아니까, 커널 사이 빈틈을 더 잘 메웁니다.

🎯 언제 쓰고, 언제 쓰지 말아야 하나

판단 기준은 단순합니다 — launch 비용(~수 µs)이 커널 실행 시간 대비 큰가?

딱 맞는 경우 — 수십 µs짜리 짧은 커널 여러 개를 같은 구조로 반복: 딥러닝 학습/추론의 한 스텝(레이어 수십 개), 시뮬레이션의 타임스텝, 신호 처리 파이프라인. 실제로 PyTorch의 torch.cuda.graphs, TensorRT 등이 내부에서 이걸 씁니다.
효과 없는 경우 — 커널 하나가 수 ms 이상 도는 큰 작업(launch 비용이 0.1%도 안 됨), 또는 매번 구조가 바뀌는 워크로드(재빌드 비용이 절약분보다 큼).

즉 그래프는 "GPU를 빠르게"가 아니라 "CPU가 GPU에 일 시키는 비용을 없애는" 최적화입니다. 프로파일러에서 커널 사이사이 GPU가 노는 빈틈(gap)이 보인다면 — 그게 그래프를 쓸 신호예요.

👤 CPU 개발자라면 — CPU로 치면 무엇인가

비유하자면 — 똑같은 함수 호출 시퀀스를 매번 다시 호출하는 대신, 한 번 "녹화"해두고 그 녹화본을 재생하는 것입니다. CPU에선 함수 호출 비용이 거의 0이라 이런 게 필요 없죠. 그런데 GPU는 커널 호출마다 CPU↔GPU 통신 오버헤드(~µs)가 있어서, 호출을 "묶어서 한 번에" 보내면 그 통신 비용이 크게 줄어드는 거예요. CPU 세계의 "배치 처리(batching)"나 "명령 버퍼 기록 후 제출"(그래픽 API의 command buffer와 사실상 같은 개념)에 해당합니다.

5.2 3단계 동작 — 정의 → 인스턴스화 → 실행

그래프는 세 단계로 씁니다. 핵심 아이디어는 "비싼 준비는 딱 한 번, 실행은 싸게 무한 반복"이에요.

단계	무엇을 하나	비유	비용
① 정의 (definition)	작업(노드)과 순서(의존성)를 적어 `cudaGraph_t` 만들기	레시피 적기	1회
② 인스턴스화 (instantiation)	실행 가능한 형태로 컴파일(검증·최적화) → `cudaGraphExec_t`	주방에 레시피 세팅	1회 (비쌈)
③ 실행 (execution)	`cudaGraphLaunch`로 반복 재생	"그거 한 번 더!"	매번 (쌈)

여기서 그래프(graph)는 수학의 그래프 — 작업(노드) + 화살표(의존성)로 이루어진 흐름도(DAG)입니다. 노드 종류엔 커널뿐 아니라 memcpy(복사), memset, host 함수, 다른 그래프, 이벤트 등이 들어갈 수 있어요.

그래프 예 (화살표 = "먼저 끝나야 함"):
   [복사] → [커널 a] → [커널 b] → [복사 결과]
                    ↘ [커널 c] ↗      (a 다음 b·c 병렬 가능)

그림 — 그래프(작업+의존성)를 한 번 정의·인스턴스화해 두고, 실행만 싸게 반복한다

5.3 만드는 법 — 두 가지 방식

① Stream Capture — 가장 쉬움 (기존 코드 재활용)

이미 짜둔 스트림 코드를 Begin/EndCapture로 감싸기만 하면, 그 안의 작업들이 실행되는 대신 "녹화"되어 그래프가 됩니다. 순서(의존성)는 스트림 사용에서 자동으로 알아냅니다.

cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);  // 녹화 시작
a<<<...,stream>>>();          // 이 줄들은 지금 실행되는 게 아니라
b<<<...,stream>>>();          // 그래프에 '기록'됨
cudaStreamEndCapture(stream, &graph);   // 녹화 끝 → graph 완성

② Explicit API — 정밀 제어

노드를 손으로 하나씩 추가하고 화살표(의존성)를 직접 지정합니다. 그래프 구조를 세밀하게 다룰 때 써요.

cudaGraphCreate(&graph, 0);
cudaGraphAddKernelNode(&nA, graph, NULL, 0, &pA);     // 노드 a (의존 없음)
cudaGraphAddKernelNode(&nB, graph, &nA, 1, &pB);     // 노드 b (a에 의존)

5.4 업데이트 — 구조는 그대로, 데이터만 바꾸기

여기가 실전에서 가장 유용합니다. 보통 반복할 때 처리할 데이터(포인터)만 바뀌고 작업 순서는 똑같죠. 그럼 그래프를 매번 새로 만들 필요 없이, 파라미터만 갈아끼우면 됩니다.

준비: 빌드 + 인스턴스화 (딱 1번)
반복: 데이터 포인터만 갱신 → 실행
      └─ cudaGraphExecKernelNodeSetParams(...)  → cudaGraphLaunch(...)

노드 하나만 바꾸기 — cudaGraphExecKernelNodeSetParams 등으로 특정 노드 파라미터 교체
전체를 한꺼번에 — cudaGraphExecUpdate로 (구조가 같은) 새 그래프의 파라미터를 일괄 적용

📌 알아둘 성질

업데이트는 다음 실행(launch)부터 적용되고, 이미 돌고 있는 실행엔 영향을 주지 않습니다. 단 구조(노드·화살표 모양)는 똑같아야 해요 — 노드를 새로 추가하거나 빼려면 그래프를 다시 만들어야 합니다.

5.5 고급 — 동적 흐름과 GPU 자율 실행

Conditional Node (CUDA 12.3+) — 그래프 안의 if / while

그래프는 구조가 고정이라 "조건에 따라 다르게 실행"이 약점이었습니다. 이걸 푸는 게 conditional node — 그래프 안에 IF나 WHILE 상자를 두고, 그 안의 작업을 조건 값에 따라 실행하거나 반복하는 거예요.

커널 안에서 cudaGraphSetConditional로 조건 값을 정함
IF: 조건이 참이면 안의 작업을 1번 / WHILE: 참인 동안 반복 (예: "수렴할 때까지 반복하는 계산"에 딱)

그림 — conditional 노드(보라 점선 상자): 커널이 정한 조건 값에 따라 IF는 body를 1회 실행/건너뛰고, WHILE은 조건이 0이 될 때까지 body를 반복한다

Device Graph Launch (CUDA 12+) — GPU가 직접 그래프 실행

원래 그래프 실행은 CPU가 시켰는데, 이제 GPU 커널이 직접 다른 그래프를 실행할 수 있습니다. GPU가 "다음엔 어떤 그래프를 돌릴지" 스스로 정해서, CPU에 묻는 왕복을 없애는 거죠. 복잡한 단계적 처리를 GPU 안에서 끝낼 수 있습니다.

그림 — 기존엔 그래프 사이마다 CPU↔GPU 경계를 왕복(위)했지만, device launch는 GPU 커널이 직접 다음 그래프(B 또는 C)를 골라 실행한다(아래)

⚠️ 흔한 함정

준비(빌드·인스턴스화)는 딱 한 번만 — 매 반복마다 새로 만들면 오히려 더 느려집니다(가장 흔한 실수).
구조 고정 — 작업 순서를 바꾸려면 재빌드. 동적 분기는 conditional node로 일부만 해결.
녹화(capture) 중엔 일부 동기 호출 금지 — capture 모드 제약을 확인하세요.

✅ ⑤ 장 정리

CUDA Graphs = 정의 → 인스턴스화(딱 1번) → 실행(반복). 매번 반복되는 커널 묶음을 한 번 찍어두고 데이터만 갈아끼우며 통째로 실행 → launch 비용과 지터를 없앱니다. 동적 흐름은 conditional node, CPU 우회는 device launch로 확장돼요. 반복 워크로드가 있다면 꼭 고려할 기법입니다.