제2부 CUDA 프로그래밍

2.10 스트림·이벤트·비동기 오버랩

GPU가 페타플롭을 제공해도 PCIe 링크는 수십 GB/s다. 큰 입력을 H2D 복사 → 커널 → D2H 복사 순으로 직렬 처리하면, GPU는 복사하는 동안 유휴 상태이고 복사 엔진은 계산하는 동안 유휴 상태다. 이 장은 그 직렬 의존을 깨는 도구 — 스트림으로 겹치고(overlap), 이벤트로 측정·동기화하고, 큰 입력을 청크로 잘라 복사와 계산을 파이프라인으로 흘려보내는 법을 익힌다.

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

배우는 것:

• 기본 스트림과 명시 스트림의 순서·동시성 규칙, 기본 스트림이 왜 암묵적으로 직렬화하는지.
• cudaMemcpyAsync와 핀드 메모리가 왜 한 묶음으로 동작하는지.
• 복사·연산 오버랩으로 처리량이 얼마나 오르는지를 타임라인으로 추론하고 상한을 계산한다.
• cudaEvent로 GPU 구간을 재고, cudaStreamWaitEvent로 스트림 간 의존성을 표현한다.
• 거짓 의존성·핀드 누락 함정을 피한다.

선수 지식: 2.1~2.3(커널·실행 구성), 2.4~2.5(메모리 모델·복사·pinned), 2.8~2.9(동기화 — 단 여기선 호스트↔디바이스 비동기). 커널 런치가 비동기, cudaMemcpy가 기본 블로킹임을 안다고 가정.

2.10.2 왜 스트림인가 — 직렬의 낭비

GPU 작업은 세 종류의 물리적으로 별개인 엔진이 처리한다: Copy Engine(H2D), Copy Engine(D2H), Compute(SM). 데이터센터 GPU는 양방향 복사 엔진이 별도라 복사와 계산, 양방향 복사까지 동시 진행 가능하다.

순진한 직렬 코드는 이 세 자원을 한 번에 하나씩만 쓴다:

직렬은 H2D·SM·D2H 세 자원이 한 시점에 하나씩만 일해 총 시간 = T_h2d + T_kernel + T_d2h(세 자원이 직렬로 줄섬)이다. 오버랩은 청크별 H2D→커널→D2H를 여러 스트림에 흘려 세 자원이 동시에 실행되어 총 시간 ≈ max(ΣH2D, ΣK, ΣD2H) + 채우기/비우기로 줄어든다.

직렬이 합이라면 오버랩은 최대에 가깝다. 세 단계가 균형이면 이론상 약 3배(2.10.5).

스트림이란: GPU에 제출된 작업(복사·커널)을 담는 순서 있는 큐다. 같은 스트림 안 작업은 제출 순서대로 실행되고(보장), 다른 스트림 작업은 순서 보장 없이 자원이 허락하면 동시에 실행된다. 오버랩은 "여러 스트림"으로 만든다.

2.10.3 기본 스트림 vs 명시 스트림

cudaStream_t s1, s2;
cudaStreamCreate(&s1); cudaStreamCreate(&s2);
kernel<<<g, b, 0, s1>>>(...);                       // 커널: 4번째 인자가 스트림
cudaMemcpyAsync(dst, src, n, cudaMemcpyHostToDevice, s2);  // 복사: ...Async(..., stream)
cudaStreamDestroy(s1); cudaStreamDestroy(s2);

기본 스트림의 함정 — 암묵적 동기화. 기본(레거시) 동작에서 기본 스트림(stream 0)에 제출된 작업은 다른 모든 블로킹 스트림에 배리어처럼 작동한다. 오버랩하려고 스트림을 만들어 놓고 중간에 무심코 cudaMemcpy(Async 아님·스트림 미지정 = 기본 스트림 + 블로킹) 한 줄을 삽입하면 파이프라인 전체가 직렬화된다.

기본 스트림(0)에 개입한 작업이 양옆 스트림을 모두 직렬화해, 실제 실행은 A → 기본 → B 순서가 된다.

해법: (1) 기본 스트림에 아무것도 제출하지 않고 모든 작업을 명시 스트림에만. (2) per-thread default stream을 켠다(nvcc --default-stream per-thread) — 각 호스트 스레드의 기본 스트림이 일반 스트림처럼 동작해 다른 스트림을 막지 않는다.

규칙 한 문장: 같은 스트림 안은 제출 순서 = 실행 순서, 다른 스트림 간은 순서·동시성 보장 없음(의존성이 필요하면 이벤트로, 2.10.6). H2D→커널→D2H를 같은 스트림에 넣으면 그 청크 내부 순서는 추가 비용 없이 지켜지고, 청크끼리는 다른 스트림이라 겹친다.

2.10.4 비동기 복사와 핀드 메모리

오버랩의 전제는 복사가 비동기여야 한다는 것이다. cudaMemcpyAsync는 작업을 스트림에 넣고 즉시 반환 — 단 조건이 있다.

OS의 일반 호스트 메모리는 페이지블 — OS가 물리 페이지를 스왑·이동할 수 있다. GPU DMA 엔진은 물리 주소로 직접 읽으므로, 페이지블에서 복사하면 드라이버가 내부 핀드 staging으로 한 번 복사한 뒤 DMA를 건다. 이 staging이 동기적이라 cudaMemcpyAsync라도 사실상 블로킹이 된다(2.5).

핀드(page-locked) 메모리는 OS가 스왑·이동 못 하게 고정한 호스트 메모리다. DMA 엔진이 물리 주소를 신뢰할 수 있어 복사 엔진이 독립적으로 일하고 SM·다른 복사와 겹친다.

float *h_pinned; cudaMallocHost(&h_pinned, bytes);   // 핀드
// ... 사용 ...
cudaFreeHost(h_pinned);

비동기 호출은 즉시 반환하지만 작업이 끝난 것은 아니다. 호스트에서 결과를 읽기 전 완료를 기다려야 한다:

2.10.5 워크스루 — 단일 vs 다중 스트림 오버랩

// overlap.cu — nvcc -O2 overlap.cu -o overlap
#include <cstdio>
#include <cuda_runtime.h>
#define CHECK(x) do{cudaError_t e=(x); if(e){printf("err %s:%d: %s\n",__FILE__,__LINE__,cudaGetErrorString(e));return 1;}}while(0)
__global__ void work(const float* in, float* out, int n) {
    int i = blockIdx.x*blockDim.x + threadIdx.x;
    if (i < n) { float v = in[i];
        #pragma unroll
        for (int k=0;k<64;++k) v = v*1.0000001f + 1e-6f;
        out[i] = v; }
}
int main() {
    const int N=1<<24, bytes=N*sizeof(float), NSTREAM=4, chunk=N/NSTREAM, cbytes=chunk*sizeof(float);
    float *h_in,*h_out; CHECK(cudaMallocHost(&h_in,bytes)); CHECK(cudaMallocHost(&h_out,bytes));
    for(int i=0;i<N;++i) h_in[i]=1.0f;
    float *d_in,*d_out; CHECK(cudaMalloc(&d_in,bytes)); CHECK(cudaMalloc(&d_out,bytes));
    cudaStream_t st[NSTREAM]; for(int s=0;s<NSTREAM;++s) CHECK(cudaStreamCreate(&st[s]));
    dim3 blk(256); cudaEvent_t t0,t1; float ms; CHECK(cudaEventCreate(&t0)); CHECK(cudaEventCreate(&t1));

    // 타이밍: 이벤트를 널 스트림(0)에 기록 → 기본(널) 스트림이 모든 블로킹 스트림과 동기화하므로
    //         t0~t1이 전 스트림 작업을 감싸는 전역 경계가 된다. 이 관용구는 레거시 기본 스트림에 의존하므로
    //         이 데모는 --default-stream per-thread로 빌드하지 않는다(그러면 stream 0이 st[]와 동기화 안 됨).
    // (A) 단일 스트림, 일괄 처리 (직렬)
    CHECK(cudaEventRecord(t0,0));
    CHECK(cudaMemcpyAsync(d_in,h_in,bytes,cudaMemcpyHostToDevice,st[0]));
    work<<<(N+blk.x-1)/blk.x,blk,0,st[0]>>>(d_in,d_out,N);
    CHECK(cudaMemcpyAsync(h_out,d_out,bytes,cudaMemcpyDeviceToHost,st[0]));
    CHECK(cudaEventRecord(t1,0)); CHECK(cudaEventSynchronize(t1)); CHECK(cudaEventElapsedTime(&ms,t0,t1));
    printf("[A] single-stream  : %.2f ms\n", ms);

    // (B) 다중 스트림 파이프라인 (오버랩): 청크별 H2D→커널→D2H를 같은 스트림에, 청크끼리는 다른 스트림에
    CHECK(cudaEventRecord(t0,0));
    for (int s=0;s<NSTREAM;++s) {
        int off=s*chunk;
        CHECK(cudaMemcpyAsync(d_in+off,h_in+off,cbytes,cudaMemcpyHostToDevice,st[s]));
        work<<<(chunk+blk.x-1)/blk.x,blk,0,st[s]>>>(d_in+off,d_out+off,chunk);
        CHECK(cudaMemcpyAsync(h_out+off,d_out+off,cbytes,cudaMemcpyDeviceToHost,st[s]));
    }
    CHECK(cudaEventRecord(t1,0)); CHECK(cudaEventSynchronize(t1)); CHECK(cudaEventElapsedTime(&ms,t0,t1));
    printf("[B] %d-stream overlap: %.2f ms\n", NSTREAM, ms);
    return 0;
}

예상(개략 — 환경 의존): [B]가 [A]보다 빠르다. 가속비는 커널·복사 시간의 균형에 달렸다. 구체 수치는 지어내지 않고 nsys로 직접 본다(2.13):

nsys profile --stats=true ./overlap
# Nsight Systems GUI에서 H2D/Compute/D2H 레인이 겹친 모습 확인

제출 패턴: 위는 스트림별로 H2D→K→D2H를 일괄 제출하는 depth-first. 모든 H2D→모든 K→모든 D2H 순회하는 breadth-first도 있다. 구형은 breadth-first가 거짓 의존성을 줄여 더 잘 겹쳤지만, 최신 하드웨어·드라이버에선 차이가 대체로 사라졌다(시점·하드웨어 의존). 오버랩이 기대만큼 안 나오면 제출 순서를 바꿔 보는 것은 여전히 유효한 디버깅.

2.10.6 오버랩의 산수와 이벤트

2.10.6.1 가속의 상한

전체 H2D H, 커널 K, D2H D:

• 직렬: $T_{\mathrm{serial}} = H + K + D$
• $N$청크 오버랩(이상적): $T_{\mathrm{overlap}} \approx (H+K+D)/N + (N-1)/N \cdot \max(H,K,D)$ → $N$ 크면 $\max(H,K,D)$

완전 균형 H=K=D=T: T_serial=3T, N 크면 T_overlap→T → 이론상 최대 3배(단계가 셋). 불균형이면 병목이 정한다: 커널이 압도적(K≫H,D)이면 $T_{\mathrm{overlap}}\approx K$ — 복사를 아무리 겹쳐도 커널 시간 밑으로 못 내려가고, 오버랩의 가치는 "복사 시간을 추가 비용 없이 숨기는 것".

청크를 무한히 잘게 분할해도 안 되는 이유: 작아지면 (1) 런치·복사 호출 오버헤드 누적, (2) 그리드가 작아 SM 미충전, (3) 작은 복사는 PCIe 대역폭을 못 살림. 현실적 N은 보통 2~8, 최적값은 프로파일러로.

2.10.6.2 이벤트 — 타이밍과 의존성

cudaEvent는 스트림의 특정 지점에 기록하는 타임스탬프 마커다. 호스트 타이머(std::chrono)로 GPU를 재면 안 된다 — 비동기라 런치까지의 시간만 잰다.

cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop);
cudaEventRecord(start, stream);
kernel<<<g,b,0,stream>>>(...);
cudaEventRecord(stop, stream);
cudaEventSynchronize(stop);                 // stop 기록까지 호스트 대기
float ms; cudaEventElapsedTime(&ms, start, stop);   // 두 이벤트 사이 GPU 경과(ms)

스트림 간 의존성 — cudaStreamWaitEvent는 GPU측에서 표현하고 호스트를 막지 않는다(이것이 cudaStreamSynchronize와 결정적 차이):

cudaEvent_t doneA; cudaEventCreateWithFlags(&doneA, cudaEventDisableTiming);
prodKernel<<<g,b,0,sA>>>(d_buf);   cudaEventRecord(doneA, sA);
cudaStreamWaitEvent(sB, doneA, 0); // B는 doneA가 지나야 진행 (GPU가 대기, 호스트는 안 막힘)
consKernel<<<g,b,0,sB>>>(d_buf);

이것으로 DAG 형태 작업 의존성을 스트림+이벤트로 그린다. (반복적이고 복잡하면 매번 제출 오버헤드가 커지는데, 한 번 캡처해 재생하는 것이 CUDA Graphs — 2.11.)

2.10.7 거짓 의존성과 오버랩이 안 되는 이유

스트림을 잘 나눴는데 안 겹친다면 대개 거짓 의존성이다:

1. 기본 스트림 오염: 중간에 기본 스트림 작업·블로킹 cudaMemcpy가 개입해 직렬화.
2. 핀드 누락: 페이지블이라 ...Async가 사실상 동기.
3. 자원 충돌: 단방향 복사 엔진뿐인 GPU에서 H2D·D2H가 같은 엔진을 다툼(엔진 수는 세대 의존, asyncEngineCount 확인).
4. 호스트 동기화 개입: 루프 안 매 청크 cudaStreamSynchronize/cudaDeviceSynchronize로 파이프라인을 매번 비움. 동기화는 전부 제출한 뒤 마지막에 한 번.
5. 암묵 동기화 API: cudaMalloc·cudaFree·cudaMemset(비Async) 등이 디바이스 전역 동기화 유발. 핫 루프 안에서 할당·해제 금지 — 버퍼는 루프 밖에서 미리.
6. 커널이 SM 포화: 한 청크 커널이 전체 SM을 채우면 다음 스트림 커널이 동시 실행할 빈 SM이 없다(복사-커널 오버랩은 되지만 커널-커널은 안 됨).

디버깅: nsys profile -o trace ./app → Nsight Systems에서 H2D/Compute/D2H 레인이 겹치는지, 안 겹치면 사이에 동기화 마커가 있는지 본다.

2.10.8 스트림 우선순위·MPS 개요

• 스트림 우선순위: 지연 민감 작업(추론 요청)이 배치와 GPU를 공유할 때, cudaStreamCreateWithPriority로 우선권을 준다. 단 선점이 아니라 스케줄링 힌트 — 이미 실행 중인 저우선 커널을 즉시 쫓아내지 않는다. 강한 보장은 MPS·MIG(4.10).
• MPS(Multi-Process Service): 스트림이 한 프로세스 안 동시성 도구라면, MPS는 여러 프로세스의 커널이 동시에 SM을 공유하게 한다. 작은 작업을 내보내는 프로세스가 여럿일 때 이용률↑.

견해: 한 앱의 처리량이 목표면 스트림 오버랩이 1차. 여러 프로세스가 GPU를 못 채우면 MPS, 강한 격리·QoS면 MIG(4.10).

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

• [ ] "오버랩이 왜 빠른가"를 직렬(합) vs 오버랩(최대) 타임라인으로 그린다.
• [ ] 핀드 메모리가 오버랩의 전제임을, 페이지블이면 ...Async가 사실상 동기임을 안다.
• [ ] 균형 모델에서 가속 상한이 ×3(단계 셋)이고 병목 단계가 율속함을 설명한다.
• [ ] cudaStreamWaitEvent가 호스트를 안 막는다는 점을 cudaStreamSynchronize와 구분한다.

2.10.10 흔한 함정

• 기본 스트림 암묵 동기화. 명시 스트림만 쓰거나 --default-stream per-thread. 중간 블로킹 cudaMemcpy 한 줄 조심.
• 핀드 누락. 페이지블 + ...Async = 사실상 동기. cudaMallocHost로.
• 이벤트 타이밍 실수. cudaEventElapsedTime 전 stop 동기화 필수. 호스트 타이머로 비동기 GPU 재지 말 것.
• 완료 전 결과 읽기. ...Async·커널은 즉시 반환. 읽기 전 cudaStreamSynchronize.
• N 과대. 청크가 너무 작으면 오버헤드·SM 미충전으로 오히려 느려진다.
• 핫 루프 안 cudaMalloc/cudaFree. 암묵 동기화. 버퍼는 루프 밖에서 미리.

2.10.11 연습문제

연습 1 (핀드 제거). overlap.cu에서 cudaMallocHost를 malloc으로 바꾸고 빌드·실행하라. [B]가 [A]에 가까워지는가? nsys로 무엇이 달라지나?

연습 2 (가속 산수). 균형 H=K=D=T에서 N=2, N=4의 가속비를 $T_{\mathrm{overlap}} \approx(H+K+D)/N + (N-1)/N\cdot\max$로 구하라(직렬 3T).

연습 3 (의존성). 스트림 A에서 생산, B에서 소비. 호스트를 막지 않고 의존성을 거는 코드는? cudaStreamSynchronize(sA)와 무엇이 다른가?

연습 4 (오버랩 진단). 다중 스트림인데 nsys에서 Compute는 겹치는데 H2D/D2H가 전혀 안 겹친다. 원인 두 가지와 확인법은?

2.10.12 요약 · 다음 장

• 스트림은 순서 있는 작업 큐 — 같은 스트림은 순서 보장, 다른 스트림은 동시성 허용. 오버랩은 "여러 스트림"으로.
• 전제는 비동기 복사 + 핀드 메모리. 페이지블이면 ...Async가 사실상 동기.
• 직렬은 자원 시간의 합, 오버랩은 최대에 수렴 — 균형이면 ×3, 병목이 율속. N은 보통 2~8.
• 이벤트는 GPU 타임라인에서 시간을 재고, 호스트를 막지 않고 스트림 간 의존성을 건다. 거짓 의존성은 nsys로 검출한다.

다음 장: 2.11 CUDA Graphs. 스트림+이벤트로 그린 의존 그래프를 반복 제출할 때 누적되는 CPU 런치 오버헤드를, 한 번 캡처해 단일 런치로 재생해 거의 0으로 만든다.