제2부 CUDA 프로그래밍

2.11 CUDA Graphs

소형 커널을 수백 개 던지는 워크로드(LLM 디코딩 한 스텝, 작은 배치 추론, 반복 루프)에서는 GPU가 일하는 시간보다 CPU가 커널을 던지는 시간이 병목이 되기도 한다. CUDA Graphs는 이 반복되는 launch 시퀀스를 한 번 기록한 DAG로 묶고 그래프 하나만 던져 런치 오버헤드를 한 자릿수 수준으로 줄인다.

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

배우는 것:

• 커널 런치 오버헤드(CPU·드라이버 경로)가 왜 소형 커널 다발에서 병목인지.
• CUDA Graph가 노드·의존성 DAG라는 모델, 생성 두 방식(스트림 캡처·명시 API).
• 인스턴스화 → 반복 launch 라이프사이클, cudaGraphExecUpdate로 재인스턴스화 회피.
• 그래프가 이득인 곳(추론·반복)과 무용한 곳(동적 형상·소수의 큰 커널).

선수 지식: 2.1(첫 커널·launch), 2.3(그리드·블록), 2.10(스트림·이벤트 — 그래프 캡처는 스트림 위에서). 컴파일 타깃은 2.12.

2.11.2 런치 오버헤드 — 무엇이 고비용인가

kernel<<<...>>>(args) 한 줄은 GPU에서 즉시 실행되는 것이 아니다. CPU 쪽에서 인자 마샬링 → 드라이버 진입 → 커맨드 버퍼 기록 → 리소스 확인을 거쳐 GPU 큐에 들어간다. 이 CPU 비용은 커널 크기와 거의 무관하게 고정(수 μs 개략, GPU·드라이버 의존)이다.

L(주황) = 커널 크기와 무관한 고정 CPU 런치 비용, K/k(파랑) = GPU 실행. 큰 커널은 실행이 길어 launch가 묻히고 GPU가 거의 쉼 없이 일하지만, 소형 커널이 많으면 GPU가 다음 launch를 기다리며 버블이 생겨 CPU-bound가 된다.

한 스텝에 소형 커널 200개를 던지면 launch만으로 수백 μs. 추론 토큰마다·학습 스텝마다 반복되면 CPU가 GPU를 못 따라가 CPU-bound가 된다(Nsight Systems 타임라인에서 GPU 행에 규칙적 빈틈 + CPU 행이 꽉 참, 2.13).

대응은 둘: ① 커널을 합친다(fusion — cuDNN 융합 3.3, FlashAttention 3.6). ② launch 시퀀스 자체를 한 번에 던진다(CUDA Graphs). 함께 쓴다.

2.11.3 그래프 모델 — 노드와 의존성 DAG

CUDA Graph는 작업 노드와 의존성 엣지의 DAG다. 스트림이 "선형 큐"라면 그래프는 부분 순서를 직접 표현한다 — 의존 없는 노드는 병렬, 의존 엣지로 연결된 노드만 순서 강제.

노드 종류: 커널(kernel<<<>>>), memcpy(cudaMemcpyAsync), memset, 호스트 콜백(cudaLaunchHostFunc), 자식그래프, 이벤트 노드.

한 추론 스텝의 예시 DAG. layernorm 뒤의 Q gemm과 KV proj는 서로 의존이 없어 병렬 실행이 가능하고, attention이 둘을 모은다.

이 DAG가 한 번 만들어지면 GPU·드라이버가 전체 작업 그래프를 미리 안다. 매 노드마다 CPU가 개입할 필요 없이 드라이버가 의존성을 따라 연쇄 실행한다. N개 launch → 1개 graph launch. 빨라지는 이유:

• CPU 경로 제거: 인자 마샬링·유효성 검사·리소스 바인딩 대부분을 인스턴스화 시점에 한 번 하고 실행 때 재사용.
• 드라이버 미리보기 최적화: 전체 DAG를 알아 노드 간 디스패치 거품 축소.

단, 그래프는 CPU launch 오버헤드와 노드 간 거품을 줄이는 것이지 커널 자체를 빠르게 하지 않는다. GPU가 이미 바쁜(큰 커널 위주) 워크로드에선 이득이 거의 없다(2.11.7).

2.11.4 생성 방식 — 스트림 캡처 vs 명시 API

스트림 캡처 (권장 출발점). 기존 스트림 코드를 그대로 감싸 그래프로 "녹화"한다. cudaStreamBeginCapture와 cudaStreamEndCapture 사이의 비동기 작업이 실행되는 대신 노드로 기록되고 의존성은 스트림·이벤트 관계에서 자동 추론된다.

cudaGraph_t graph;
cudaStreamBeginCapture(stream, cudaStreamCaptureModeThreadLocal);
kernelA<<<g, b, 0, stream>>>(d_in, d_tmp);    // 실행되지 않고 "기록"됨
kernelB<<<g, b, 0, stream>>>(d_tmp, d_out);
cudaMemcpyAsync(h_out, d_out, bytes, cudaMemcpyDeviceToHost, stream);
cudaStreamEndCapture(stream, &graph);          // graph에 DAG가 담김

장점: 기존 코드를 거의 안 바꾼다. 라이브러리 호출(cuBLAS·cuDNN)도 스트림으로 던지면 캡처에 포함된다. 실무 1순위. 다중 스트림 캡처는 보조 스트림이 메인 스트림에 이벤트로 묶여야 의존성이 올바로 추론된다.

명시 API (정밀 제어). cudaGraphCreate 후 cudaGraphAddKernelNode 등으로 노드를 직접 추가하고 의존 엣지를 직접 건다. 구조를 코드 수준에서 완전 제어하거나 프로그램적으로 합성할 때. 단 장황하고 오류 나기 쉽고 라이브러리 호출을 노드로 표현하기 어렵다.

견해: 거의 모든 실무 시작점은 스트림 캡처. 명시 API는 캡처로 표현 안 되는 동적 합성이 필요할 때만.

2.11.5 인스턴스화와 반복 실행 — 라이프사이클

그래프 객체(cudaGraph_t)는 템플릿(설계도)이다. 실행하려면 인스턴스화해 실행본(cudaGraphExec_t)을 만든다 — 유효성 검사·리소스 바인딩·최적화를 하는 고비용 일회성 단계. 이후 launch는 비용이 낮다.

cudaGraphExec_t exec;
cudaGraphInstantiate(&exec, graph, 0);   // 고비용 일회성
for (int it = 0; it < num_iters; ++it)
    cudaGraphLaunch(exec, stream);       // 매번 저비용 단일 launch — DAG 전체 실행
cudaStreamSynchronize(stream);
cudaGraphExecDestroy(exec); cudaGraphDestroy(graph);

핵심: 고비용 작업(인스턴스화)을 루프 밖에서 한 번, 저비용 작업(launch)을 루프 안에서 매번. 그래프가 이득이려면 인스턴스화 비용이 반복 횟수에 걸쳐 분할 상환되어야 한다. 한 번만 실행할 그래프는 만들 가치가 없다.

주의: 인스턴스화는 그래프 구조·노드 파라미터의 스냅샷을 고정한다. 캡처 당시의 포인터·인자가 실행본에 박힌다. 다음 이터레이션에서 버퍼 주소·형상이 바뀌면 새 그래프를 다시 굽거나 업데이트해야 한다.

2.11.6 업데이트 — cudaGraphExecUpdate

반복 루프에서 위상(노드·엣지)은 그대로인데 파라미터만 바뀌는 경우가 흔하다(같은 커널 시퀀스, 입력 버퍼 포인터만 다름). 매번 재인스턴스화하면 비용이 크다. cudaGraphExecUpdate는 기존 실행본에 새 그래프의 파라미터만 적용해 재인스턴스화를 피한다.

// 변경된 인자로 그래프 재캡처(구조 동일, 파라미터만 다름)
cudaStreamBeginCapture(stream, cudaStreamCaptureModeThreadLocal);
kernelA<<<g, b, 0, stream>>>(d_in2, d_tmp);   // 포인터/인자가 바뀜
kernelB<<<g, b, 0, stream>>>(d_tmp, d_out2);
cudaStreamEndCapture(stream, &graph2);
cudaGraphExecUpdateResultInfo info;                  // CUDA 12.x 구조체(.result/.errorNode)
cudaError_t st = cudaGraphExecUpdate(exec, graph2, &info);  // 12.x 3인자형(11.x는 4인자)
if (st != cudaSuccess) {                       // 위상이 바뀌었으면 폴백
    cudaGraphExecDestroy(exec);
    cudaGraphInstantiate(&exec, graph2, 0);
}
cudaGraphLaunch(exec, stream);

성공 조건: 두 그래프의 위상이 동일(노드 수·종류·의존 엣지 같음). 허용 변경은 주로 노드 파라미터(인자 값, grid/block, memcpy 주소·크기). 노드를 더하거나 빼면 실패 → 재인스턴스화로 폴백. 노드 하나만 갱신이면 cudaGraphExecKernelNodeSetParams로 더 정밀하게.

2.11.7 워크스루 — 반복 커널 시퀀스를 그래프화

소형 커널 시퀀스를 반복하는 루프를 ① 매번 launch와 ② 그래프로 던지는 방식으로 비교한다.

// graphs_demo.cu — nvcc -O3 -arch=sm_80 graphs_demo.cu -o graphs_demo
#include <cstdio>
#include <cuda_runtime.h>
#define CK(x) do{cudaError_t e=(x); if(e){printf("err %s @%d: %s\n",#x,__LINE__,cudaGetErrorString(e));return 1;}}while(0)
__global__ void scale(float* a, float s, int n){ int i=blockIdx.x*blockDim.x+threadIdx.x; if(i<n) a[i]*=s; }
__global__ void addc (float* a, float c, int n){ int i=blockIdx.x*blockDim.x+threadIdx.x; if(i<n) a[i]+=c; }
__global__ void relu (float* a, int n){ int i=blockIdx.x*blockDim.x+threadIdx.x; if(i<n) a[i]=a[i]>0.f?a[i]:0.f; }
int main(){
    const int n=4096, STEPS=2000, bs=256, gs=(n+bs-1)/bs;  // 작게 → launch 비용이 두드러짐
    float* d; CK(cudaMalloc(&d,n*sizeof(float))); CK(cudaMemset(d,0,n*sizeof(float)));
    cudaStream_t s; CK(cudaStreamCreate(&s)); cudaEvent_t t0,t1; CK(cudaEventCreate(&t0)); CK(cudaEventCreate(&t1));
    auto one_step=[&](cudaStream_t st){ for(int k=0;k<25;++k){
        scale<<<gs,bs,0,st>>>(d,1.0001f,n); addc<<<gs,bs,0,st>>>(d,0.0f,n); relu<<<gs,bs,0,st>>>(d,n); } };

    // (A) 매번 launch
    CK(cudaEventRecord(t0,s)); for(int it=0;it<STEPS;++it) one_step(s);
    CK(cudaEventRecord(t1,s)); CK(cudaEventSynchronize(t1)); float ms_stream=0; CK(cudaEventElapsedTime(&ms_stream,t0,t1));

    // (B) 스트림 캡처 → 그래프 → 인스턴스화 1회 → 반복 launch
    cudaGraph_t graph; cudaGraphExec_t exec;
    CK(cudaStreamBeginCapture(s, cudaStreamCaptureModeThreadLocal)); one_step(s); CK(cudaStreamEndCapture(s,&graph));
    CK(cudaGraphInstantiate(&exec, graph, 0));
    CK(cudaEventRecord(t0,s)); for(int it=0;it<STEPS;++it) CK(cudaGraphLaunch(exec,s));
    CK(cudaEventRecord(t1,s)); CK(cudaEventSynchronize(t1)); float ms_graph=0; CK(cudaEventElapsedTime(&ms_graph,t0,t1));

    printf("(A) per-launch stream : %.2f ms\n(B) cuda graph : %.2f ms\nspeedup (A/B): %.2fx\n",
           ms_stream, ms_graph, ms_stream/ms_graph);
    return 0;
}

예상(개략 — GPU·드라이버 의존): (A)가 큰 값, (B)가 작은 값, speedup > 1. 커널을 크게(예: n을 수백만으로) 키우면 GPU 실행이 launch를 묻어버려 두 값이 가까워진다 — 그래프 이득이 사라지는 지점을 직접 본다. 2.13에서 Nsight Systems 타임라인으로 (A)의 버블과 (B)의 조밀함을 눈으로 확인.

함정: 위 (B) 시간에 인스턴스화 비용은 빠져 있다(루프 밖). 공정한 평가는 인스턴스화 1회 + N회 launch를 합쳐 N으로 나눈 분할 상환 비용. 반복 N이 작으면 인스턴스화가 안 갚아진다.

2.11.8 어디서 이득이고, 어디서 무용한가

• 이득 큰 곳: 반복 추론 디코딩(LLM 토큰 생성, 같은 형상 소형 커널 반복 — 교과서적 적용), 작은 배치·짧은 커널 추론, 고정 형상 학습 스텝, CPU-bound로 진단된 워크로드.
• 이득 적거나 무용: 소수의 큰 커널(GEMM 위주, GPU가 이미 바쁨), 동적 형상(매번 재캡처 비용), 데이터 의존 제어 흐름, 한 번만 실행.

2.11.9 프레임워크에서의 CUDA Graphs

직접 안 써도 상위 프레임워크가 내부에서 활용한다:

• PyTorch: torch.cuda.CUDAGraph·torch.cuda.graph(...)·make_graphed_callables. 입력을 정적 버퍼(static input)에 복사해 넣는 패턴 — 입력 주소가 고정돼야 한다는 제약이 2.11.5의 "박힌 포인터"와 같은 이유.
• 추론 엔진: TensorRT-LLM·Triton(5.7·5.11)이 디코딩 루프 반복 launch를 그래프로 묶어 토큰 지연을 줄인다.
• 컴파일러: torch.compile/CUDA Graph 모드가 추적 그래프를 graph launch로 전달한다.

공통 전제: 형상이 안정적이고 입출력 버퍼가 고정되어야 문제없이 적용. 동적인 부분은 패딩·버킷팅으로 형상을 고정하거나 그래프 밖에.

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

• [ ] 커널 launch의 CPU·드라이버 고정 비용이 왜 소형 커널 다발에서 병목인지 설명한다.
• [ ] cudaGraph_t(설계도)와 cudaGraphExec_t(실행본)를 구분하고 인스턴스화가 고비용 일회성임을 안다.
• [ ] 스트림 캡처로 기존 코드를 거의 안 바꾸고 그래프를 만들고, 형상이 같으면 cudaGraphExecUpdate로 재인스턴스화를 피한다.

증상: 그래프를 썼는데 안 빨라짐 → 커널이 이미 커서 launch가 안 묻혔거나 반복 N이 작아 인스턴스화 못 갚음(2.13으로 GPU 빈틈 확인). 캡처가 에러 → 캡처 중 금지된 동기 호출(cudaMemcpy 동기형·cudaStreamSynchronize). 둘째 이터레이션부터 틀린 결과 → 박힌 옛 포인터(업데이트 누락).

2.11.11 흔한 함정

• 캡처 제약 위반. 캡처 구간에서 동기 API(동기 cudaMemcpy·cudaStreamSynchronize·일부 할당)를 부르면 캡처 무효화. 모든 작업을 비동기·스트림 버전으로. 메모리는 캡처 전 할당 또는 cudaMallocAsync.
• 인스턴스화에 박힌 포인터. 그래프는 캡처 시점 인자(포인터)를 고정한다. 버퍼 주소가 바뀌면 옛 주소로 실행 → 같은 버퍼 재사용 또는 업데이트. 업데이트 누락은 가장 흔한 정답 깨짐.
• 위상 변화에 업데이트 적용. 노드를 더하거나 빼고 cudaGraphExecUpdate를 부르면 실패. 실패 코드를 무시 말고 재인스턴스화로 폴백.
• 동적 형상을 그래프 하나로. 시퀀스 길이·배치가 매번 바뀌면 형상별 그래프 풀(자주 쓰는 형상마다 캡처해 캐시)이나 패딩으로 형상 고정.
• 그래프로 커널 내부까지 빨라지길 기대. 그래프는 launch·디스패치 오버헤드만 줄인다. 메모리·연산 바운드면 그쪽 최적화(1.10)가 먼저.

2.11.12 연습문제

연습 1 (이득 소멸). 데모에서 n을 4096 → 4,000,000으로 키우고 재측정하라. speedup이 어떻게 변하며 왜?

연습 2 (손익분기). 인스턴스화 비용을 측정에 포함시켜 손익분기 반복 N*(그래프가 per-launch보다 빨라지는 최소 N)를 추정하라.

연습 3 (포인터 교체). 둘째 이터레이션에서 입력 버퍼를 다른 주소(d2)로 바꿔야 한다. 재인스턴스화 없이 처리하는 두 방법은?

연습 4 (LLM 디코딩). CUDA Graph가 LLM 디코딩에 잘 맞는 이유와 가변 시퀀스 길이를 다루는 두 전략은?

2.11.13 요약 · 다음 장

• 커널 launch에는 커널 크기와 무관한 CPU·드라이버 고정 비용이 있어, 소형 커널이 많으면 GPU가 기아 상태가 되고 CPU-bound가 된다.
• CUDA Graph는 반복 launch 시퀀스를 노드·의존성 DAG로 한 번 기록해 두고 단일 graph launch로 던져 그 오버헤드를 걷어낸다.
• 생성은 스트림 캡처(실무 1순위)와 명시 API. 실행은 인스턴스화(고비용 1회) → 반복 launch(비용이 낮다), 비용은 반복으로 분할 상환. 형상 같으면 cudaGraphExecUpdate.
• 이득은 반복·고정 형상·소형 커널 다발에서 크고, 큰 커널·동적 형상·1회 실행에선 작거나 무용.

다음 장: 2.12 컴파일 흐름. nvcc가 소스를 PTX·SASS로 내리는 경로와 -arch의 의미 — 이 장에서 빌드할 때 쓴 아키텍처 타깃을 제대로 이해한다.