NVIDIA GPU에서 고성능 알고리즘을 구현할 때 CUB 라이브러리는 필수적이지만, 기존의 '두 단계(two-phase)' API는 메모리 크기 추정과 할당을 분리해야 해서 코드가 번잡해지곤 했습니다. 이번 포스팅에서는 CUDA 13.1부터 사용 가능해진 새로운 단일 호출(single-call) API가 어떻게 개발 경험을 개선하는지 살펴보겠습니다. 자세한 내용은 NVIDIA 개발자 블로그의 근거자료를 참고하세요.
기존 두 단계 API의 문제점
기존 방식은 먼저 필요한 임시 메모리 크기를 확인한 후, 메모리를 할당하고 알고리즘을 실행하는 두 번의 호출이 필요했습니다.
// 첫 번째 호출: 임시 저장소 크기 확인
cub::DeviceScan::ExclusiveSum(nullptr, temp_storage_bytes, d_input, d_output, num_items);
// 필요한 임시 저장소 할당
cudaMalloc(&d_temp_storage, temp_storage_bytes);
// 두 번째 호출: 실제 스캔 알고리즘 실행
cub::DeviceScan::ExclusiveSum(d_temp_storage, temp_storage_bytes, d_input, d_output, num_items);
이 방식은 유연성을 제공하지만, 반복적인 코드 작성과 어떤 인자가 호출 간에 변경 가능한지 명확하지 않다는 단점이 있었습니다.
새로운 단일 호출 API의 장점
새 API는 모든 과정을 한 줄로 단순화합니다. 내부적으로 메모리 할당을 처리하므로 성능 오버헤드는 없습니다.
| 기능 | 기존 두 단계 API | 새로운 단일 호출 API |
|---|---|---|
| 호출 횟수 | 2회 (크기 확인 + 실행) | 1회 |
| 명시적 할당 | 필요 (cudaMalloc) | 불필요 (내부 처리) |
| 코드 가독성 | 낮음 (보일러플레이트多) | 높음 |
| 성능 오버헤드 | 없음 | 없음 |
| 실행 환경 제어 | 제한적 | env 인자로 유연한 제어 가능 |
새 API는 env 인자를 통해 메모리 리소스나 CUDA 스트림을 전달할 수 있어, 기존의 유연성도 유지합니다.
실무 적용 팁 및 결론
cub::DeviceReduce::Sum, cub::DeviceScan::ExclusiveSum 등 주요 알고리즘들이 새 인터페이스를 지원합니다. env 인자를 활용하면 사용자 정의 메모리 풀이나 특정 스트림을 조합해 실행 환경을 세밀하게 제어할 수 있습니다. 성능 손실 없이 코드를 간결하게 만들 수 있는 이 새로운 표준을 적극 도입해 GPU의 계산 능력을 최대한 활용해 보세요. CUDA 13.1 이상 버전에서 바로 사용 가능합니다.