Build a Prod Ready CUDA library

Oscar 의 출발점은 단순하다 — 회사에서 그는 CUDA 를 아는 유일한 엔지니어였고, 동료 비전·미디어 엔지니어들에게 CUDA 를 가르치려고 9년 가까운 노력을 쏟았는데 성과는 미미했다. 그래서 방향을 뒤집었다 — “이 사람들이 CUDA 를 모르더라도, 내가 만든 라이브러리를 부르기만 하면 좋은 성능을 얻게 한다.”

강의의 두 사례는 둘 다 같은 메시지를 다른 도메인에서 보여준다.

• Case A (컴퓨터 비전 파이프라인) — image-resize, color-convert, normalize, crop 같은 op 들이 OpenCV 위에서 따로 돌면 GPU 위에서도 따로 돌 텐데, 그것들을 한 커널 안에 fuse 한 라이브러리(FK / FastKAT) 를 회사 동료들에게 제공한다.
• Case B (실시간 미디어 그래프) — 비디오 프로세싱 그래프에서 producer 와 consumer 가 서로 기다리지 않게 만들기 위해 streams + delay buffer + graph synchronization 패턴을 깐다.

그래서 이 강의의 끝에 손에 잡혀야 하는 건 단순한 “fast kernel” 가 아니라 — 4가지 라이브러리 패턴이다: vertical fusion, horizontal fusion, producer-consumer streams, delay-buffered graph iteration.

Oscar 가 강의 초반에 못 박는 진단 — 회사에서 새 GPU 코드 컨설팅 문의가 들어왔을 때 가장 먼저 들여다보는 건 커널이 아니라 호스트 코드다. 왜냐하면 대부분의 application 에서 GPU 는 잘 도는데 — 호스트가 GPU 를 제대로 먹이지 못하기 때문에 비어 있다.

Oscar 는 이 진단을 바탕으로 — 라이브러리를 설계할 때 “호스트가 단순한 호출 한 줄로 GPU 위에 큰 일감을 던지게 하는 것” 이 가장 중요한 디자인 결정이라고 본다. 그래서 다음 섹션의 추상화 층 분리가 의미를 갖는다.

Oscar 의 핵심 디자인 원칙 — 라이브러리는 같은 코드를 세 가지 시각으로 동시에 만족시켜야 한다. 각 층은 서로 다른 추상화 수준으로 같은 기능을 노출한다.

그리고 이 디자인의 실용적인 장점 — 새로운 동료가 회사에 들어왔을 때, 그 사람이 어떤 layer 에서 일을 시작할지를 명시적으로 정할 수 있다. 모두를 닌자로 만들 필요가 없다.

Memory-bound op 들의 가장 흔한 낭비 — resize 가 결과를 HBM 에 쓰고, 다음 op 인 normalize 가 그걸 다시 읽고, 또 쓰고, 다음 crop 이 또 읽는다. 같은 픽셀이 HBM 사이를 여러 번 왕복한다. vertical fusion 은 이 chain 을 한 커널로 합쳐서 — 픽셀이 한 번 SRAM 에 올라오면 그 위에서 모든 op 가 도는 구조다.

// FK 스타일 fusion (의사코드)
auto pipeline = fk::compose(
    fk::read(input),
    fk::resize(out_w, out_h),
    fk::normalize(mean, std),
    fk::crop(box),
    fk::write(output)
);

// 한 번의 launch 로 SRAM 위에서 모두 합성
pipeline.launch(stream);

Vertical fusion 이 “한 픽셀이 여러 stage 를 통과한다” 면, horizontal fusion 은 “한 입력에서 여러 출력이 동시에 나간다”. 비전 파이프라인에서 흔하다 — 같은 frame 에서 작은 thumbnail, 큰 inference 입력, 시각화용 overlay 가 동시에 필요한 경우.

이 패턴이 가장 많이 만나는 자리는 — multi-resolution / multi-output 모델 추론 파이프라인. 같은 이미지에서 detector 와 classifier 가 다른 해상도를 요구하는데, 각 모델이 자기 전처리를 따로 부르면 H2D copy 가 두 배가 된다.

Vertical/horizontal fusion 으로 단일 iteration 안의 HBM 왕복을 줄였다면, 다음 도구는 iteration 사이의 latency를 줄이는 것. CUDA streams 와 producer-consumer 패턴이 표준 도구다.

// 두 stream — pinned host buffer + async copy
cudaStream_t copy_s, compute_s;
cudaStreamCreate(&copy_s);
cudaStreamCreate(&compute_s);

// host pinned buffer 풀
float* h[N];
for(int i=0; i<N; ++i)
    cudaMallocHost(&h[i], bytes);

for(int i=0; i<num_iters; ++i){
    // 다음 입력을 copy stream 으로 미리
    cudaMemcpyAsync(d_in[i+1], h[i+1], bytes,
                    cudaMemcpyHostToDevice, copy_s);
    // 현재 입력으로 compute stream 위 kernel
    kernel<<<..., compute_s>>>(d_in[i], d_out[i]);
}

Oscar 가 강의 후반부에 가장 길게 깐 패턴 — delay buffer. iteration 사이에 작업이 의존성을 가질 때, 매번 sync 하면 파이프라인이 멈춘다. 그런데 어떤 종류의 데이터는 몇 step 늦어도 결과가 같다 — 그 자리에 delay buffer 를 박아 sync 자체를 우회한다.

강의의 첫 사례. Oscar 의 회사에서는 OpenCV 의 GPU 모듈을 쓴다 — 그런데 OpenCV 는 op 단위로 host API 를 노출하기 때문에 여러 op 를 묶어 부르면 자연스럽게 vertical fusion 기회를 놓친다. FK (FastKAT) 가 그 위에 얹히는 layer.

이 사례에서 반복되는 디자인 결정 — “user 가 손을 한 번도 안 대게 한다”. FK 가 도입되기 전과 후의 user 코드 한 줄도 안 바뀐다. 라이브러리가 위에서 자동으로 fusion graph 를 만든다. 새 user 가 회사에 들어와도 OpenCV 만 알면 된다.

강의에서 짧지만 인상적으로 다룬 부분 — Oscar 가 회사의 일부 product 를 Windows 환경에 배포하면서 만난 시스템 차이. 같은 GPU, 같은 코드인데 Windows 에서 더 느리고 더 비결정적이다.

이 강의의 가치는 단일 커널 최적화가 아니라 — 여러 커널을 묶어 production 라이브러리로 만드는 패턴 카탈로그다. 다음 라이브러리 디자인을 시작할 때 빠르게 복원해야 하는 핵심.

손에 새기기 — 실습 시퀀스

1. vertical fusion baseline — torchvision 의 Resize → Normalize → CenterCrop 전처리 chain 을 (1) 단계별 따로 호출, (2) torch.compile 로 compose, (3) 직접 Triton 으로 fuse. nsys 로 커널 수와 H2D 트래픽 비교.
2. two-stream copy/compute — pinned host buffer pool 을 만들고 cudaMemcpyAsync + 두 stream 으로 image preprocessing 파이프라인을 깐다. 단일 stream 대비 throughput 측정.
3. CUDA Graphs capture — 같은 iteration 을 한 번 capture 해서 graph instance 로 재사용. launch overhead 가 N→1 로 줄어드는 시점을 nsys 로 확인.
4. delay buffer toy — frame n 의 monitoring 통계가 frame n+5 의 control loop 에 들어가는 toy 그래프. host sync 없이 pointer swap 만으로 도는 형태로 구현.
5. WDDM vs Linux 비교 — 같은 코드를 듀얼부트나 두 환경에서 돌려 launch overhead variance 측정. variance 가 큰 자리가 어디에 박히는지 확인.
6. FK 스타일 fusion library — C++ template + lambda 로 작은 영상 변환 op 3개를 fuse 하는 toy 라이브러리 작성. user-level API 한 줄, internal 에서 한 커널로 launch.

L010 의 “라이브러리 디자인” 시각이 다른 강의의 단일 도구들과 어떻게 맞물리는지.

이 강의는 산업 사례 위주라서 — 본문 안의 구체적 수치와 코드 디테일이 슬라이드/repo 에 흩어져 있다. 다음 회독 때 채울 자리들.

• FK 라이브러리의 정확한 API 형태 — 본 노트는 fk::compose(...) 같은 의사코드 수준으로만 적었다. morousg 의 실제 repo 확인 필요. 확인 필요.
• vertical fusion 의 실제 측정 수치 — § 08 의 “3-4× speedup” 은 패턴 재구성이며 강의 슬라이드의 정확한 숫자와 다를 수 있다. 확인 필요.
• delay buffer 의 정확한 응용 도메인 — 강의에서 영상처리 파이프라인의 어떤 specific feedback loop 에 적용했는지 — 슬라이드 직접 확인.
• CUDA Graphs 와 stream capture 의 trade-off — graph capture 모드(global/threadlocal) 별로 호환 가능한 launch 패턴이 다르다. 별도 강의에서 본격적으로 다뤄질 가치.
• Windows WDDM vs MCDM — Windows 11 의 새 driver 모델(MCDM)이 compute workload 에서 어떻게 변화하는지 — 강의 시점(2024 March)이후 갱신된 정보 추적 필요.
• fusion 의 한계 — compute-bound 일 때 — 강의는 memory-bound op chain 위주. compute-bound 가 섞이면 fusion 의 이득이 사라지는 자리가 있는데, 그 경계는 어디인가.