NVIDIA CUDA Core Compute Libraries (CCCL)는 C++와 Python을 사용하는 CUDA 개발자를 위해 편리하고 효율적인 추상화를 제공합니다. 주요 기능은 다음과 같습니다.
- 병렬 알고리즘 – 정렬(sort), 스캔(scan), 리듀스(reduce) 등 호스트에서 실행되는 알고리즘으로, 일반적인 연산에 커스텀 커널을 직접 작성할 필요를 없애줍니다.
- 협력 알고리즘 – 블록 단위 또는 워프 단위 리덕션·스캔 등 디바이스 측 알고리즘으로, 커스텀 커널 개발을 단순화합니다.
- 언어 관용적 CUDA 추상화 – 메모리 할당, 리소스 관리, 하드웨어 기능을 포함한 CUDA 고유 연산을 위한 핵심 추상화입니다.
이 글에서는 CCCL에 새로 추가된 기능 그룹을 소개합니다. 이 기능 그룹은 CUDA 프로그래밍 모델의 핵심 개념에 현대화된 C++ 추상화를 제공하여, CUDA C++ 개발을 보다 안전하고 편리하게 만듭니다.
CCCL 런타임이란 무엇인가?
NVIDIA CCCL 런타임은 CUDA 13.2부터 제공되는 새로운 관용적 C++ API 모음으로, 스트림 관리, 메모리 할당, 커널 실행 등 핵심 CUDA 기능을 구현합니다.
기존의 NVIDIA CUDA 런타임은 원래 CUDA 드라이버 API 위에 편의 계층으로 개발되었습니다. 새로운 CCCL 런타임은 같은 목표를 지향하는 대안으로, 현대 C++에 맞게 업데이트된 설계를 채택하였습니다. 아래 그림 1은 앞서 언급한 세 가지 CUDA API 인터페이스 간의 관계를 보여줍니다.
CCCL 런타임은 <cuda/buffer>, <cuda/stream>, <cuda/launch> 등 CCCL 내의 헤더 모음입니다. 기존 CUDA 런타임 API가 C 소스 호환성 제약 아래에서는 구현할 수 없었던 더욱 편리하고 견고한 추상화를 제공하기 위해 현대 C++ 기능을 적극 활용합니다.
또한 20년에 걸친 CUDA 발전 과정에서 얻은 교훈을 API 설계에 반영할 기회로도 삼았습니다. 이 모든 변화에도 불구하고, CCCL 런타임은 기존 CUDA 런타임 API를 사용하는 주변 코드를 전면 재작성하지 않고도 점진적으로 도입할 수 있는 호환성 헬퍼를 제공합니다.
CUDA 프로그램이 복잡해지면서 여러 라이브러리가 디바이스, 스트림, 메모리를 공유하게 되면, 의존 관계를 명시적으로 드러내고 깔끔하게 조합되는 API에 대한 필요성이 커집니다. CCCL 런타임은 바로 그러한 역할을 담당하도록 설계되었습니다.
코드 살펴보기
새로운 CCCL 런타임 API로 구현한 전통적인 vectorAdd 예제를 소개합니다. CUDA 개발 경험이 있으시다면 전반적인 구조가 익숙하실 겁니다. 처음부터 모든 코드를 파악하려고 하기보다는, 기존 방식과의 차이점을 중심으로 가볍게 살펴보세요. 이어지는 본문에서 이 예제를 한 단계씩 분석하며 CCCL 런타임의 개념과 주요 설계 방향을 자세히 설명할 예정입니다.
#include <cuda/buffer>
#include <cuda/devices>
#include <cuda/launch>
#include <cuda/memory_pool>
#include <cuda/std/span>
#include <cuda/stream>
struct kernel {
template <typename Config>
__device__ void operator()(Config config,
cuda::std::span<const int> A,
cuda::std::span<const int> B,
cuda::std::span<int> C) {
auto tid = cuda::gpu_thread.rank(cuda::grid, config);
if (tid < A.size())
C[tid] = A[tid] + B[tid];
}
};
int main() {
// 1. Devices and streams
cuda::device_ref device = cuda::devices[0];
cuda::stream stream{device};
// 2. Memory allocation
auto pool = cuda::device_default_memory_pool(device);
int num_elements = 1000;
auto A = cuda::make_buffer<int>(stream, pool, num_elements, 1);
auto B = cuda::make_buffer<int>(stream, pool, num_elements, 2);
auto C = cuda::make_buffer<int>(stream, pool, num_elements, cuda::no_init);
// 3. Kernel launch
constexpr int threads_per_block = 256;
auto config = cuda::distribute<threads_per_block>(num_elements);
cuda::launch(stream, config, kernel{}, A, B, C);
// Make the CPU thread wait for the GPU work to finish.
stream.sync();
return 0;
}이 예제는 다음 세 가지 주요 섹션으로 나눌 수 있습니다.
1. 디바이스와 스트림
다음 코드는 CUDA 런타임 API로 스트림을 생성하는 방법을 보여줍니다.
cudaStream_t stream;
cudaStreamCreate(&stream); // associated with whichever device happens to be "current"이 코드는 스트림을 생성하지만, 스트림은 cudaStreamCreate가 호출될 때 현재 활성화된 디바이스와 연결됩니다. 이 호출만 보아서는 스트림이 어떤 디바이스와 연결되어 있는지 알 수 없습니다.
이를 CCCL 런타임 API를 사용한 방식과 비교해 보겠습니다.
cuda::device_ref device = cuda::devices[0];
cuda::stream stream{device};위 코드는 특정 디바이스에서 스트림을 생성하는 방법을 보여줍니다. 첫 번째 줄은 CCCL 런타임의 핵심 설계 원칙을 보여줍니다. CCCL 런타임은 원시 식별자 대신 전용 타입을 사용합니다. 디바이스는 단순 정수가 아닌 device_ref이고, 스트림은 불투명 포인터가 아닌 객체입니다. API 전반에 걸친 강타입(strong typing)은 런타임에 오류를 추적하는 대신 컴파일 타임에 실수를 잡아냅니다.
두 번째 줄은 또 다른 원칙인 의존 관계의 명시화를 보여줍니다. CCCL 런타임과 CUDA 런타임 API 모두에서 스트림은 디바이스와 연결됩니다. 차이는 그 방법에 있습니다. 여기서는 cuda::stream 생성자가 디바이스를 명시적 인수로 받는 반면, CUDA 런타임 API에서는 스트림이 생성될 때 활성화된 디바이스와 연결됩니다.
명시적 의존 관계는 로컬 추론을 가능하게 합니다. 전역 상태를 추적하지 않고도 함수를 읽는 것만으로 동작을 이해할 수 있습니다. 또한 조합성(composability)을 향상시킵니다. 여러 라이브러리를 함께 사용할 때, 서로 간섭을 피하기 위해 암묵적 상태를 호출 전후로 저장하고 복원할 필요가 없습니다.
이와 관련된 결과로, CCCL 런타임은 기본 스트림(default stream)을 노출하지 않습니다. 기본 스트림의 의미를 관리하려면 현재 디바이스를 추적해야 하는데, 이는 바로 우리가 벗어나고자 하는 암묵적 상태의 전형입니다. CUDA 런타임 API의 기본 스트림은 여전히 CCCL 런타임 타입으로 래핑할 수 있지만, 그 사용은 권장되지 않습니다. 기본 스트림과 관련된 모든 처리는 CUDA 런타임 API를 통해 직접 수행해야 합니다. API에 기본 스트림이 없으므로 “블로킹 스트림”의 개념도 적용되지 않으며, 모든 CCCL 런타임 스트림은 논블로킹(non-blocking)으로 생성됩니다.
리소스 소유권: 소유 타입과 참조 타입
std::string과 std::string_view의 예시를 따라, CCCL 런타임의 많은 CUDA 객체는 소유 타입과 _ref 접미사가 붙은 비소유 타입, 두 가지 타입을 갖습니다. cuda::stream은 내부의 cudaStream_t 핸들을 소유하며 소멸자에서 이를 해제합니다. cuda::stream_ref는 핸들의 수명을 관리하지 않고 보유하며, 간편하게 복사할 수 있습니다.
_ref 타입은 기존 코드와의 조합성을 위해 필수적입니다. 스트림 핸들의 수명이 다른 곳에서 관리되는 경우, cudaStream_t는 암묵적으로 cuda::stream_ref로 변환되며, 원시 핸들은 .get()으로 가져올 수 있습니다. 소유권을 이전하려면 cuda::stream::from_native_handle로 원시 핸들을 소유 타입으로 래핑하고, .release()로 소유권을 반환합니다.
void stream_type_example(cudaStream_t handle) {
cuda::stream_ref non_owning{handle};
assert(handle == non_owning.get());
cuda::stream owning = cuda::stream::from_native_handle(handle);
assert(handle == owning.get());
assert(handle == owning.release());
}이벤트, 메모리 풀 등 다른 CUDA 객체에도 동일한 패턴이 적용됩니다. cuda::device_ref에는 소유 타입이 없는데, 소유할 디바이스 상태가 없기 때문입니다.
2. 메모리 할당
auto pool = cuda::device_default_memory_pool(device);
auto A = cuda::make_buffer<int>(stream, pool, num_elements, 1);
auto B = cuda::make_buffer<int>(stream, pool, num_elements, 2);
auto C = cuda::make_buffer<int>(stream, pool, num_elements, cuda::no_init);다음 섹션에서는 디바이스 메모리를 비동기적으로 할당하고 초기화하는 방법을 살펴봅니다. 여기서 또 다른 설계 원칙을 확인할 수 있습니다. API는 기본적으로 비동기(asynchronous)입니다. 동기와 비동기 변형을 이름으로 구분하는 대신, CCCL 런타임은 간단한 규칙을 사용합니다. API의 첫 번째 인수가 스트림이라면, 해당 API는 스트림 순서대로 동작합니다. CUDA 런타임 API에서 두 가지 변형이 모두 존재하는 API에 대해 동기식 대응 함수는 제공하지 않을 계획입니다.
메모리 할당은 실제로 이 원칙이 가장 중요하게 적용되는 부분입니다. 메모리 풀을 통한 스트림 순서 메모리 관리는 CUDA 11.2부터 제공되어 왔으며(여기에서 설명), CUDA 13.0은 이를 관리 메모리와 호스트 메모리까지 확장하였습니다. 대부분의 경우 최대 성능을 달성하려면 메모리 풀링과 동기화 지점 최소화가 필수적이며, 스트림 순서 메모리 관리는 나머지 비동기 프로그래밍 모델과 자연스럽게 조합됩니다. 이러한 지침을 전달하기 위해 CCCL 런타임은 메모리 풀과 스트림 순서 할당을 기본으로 채택하였습니다. 최신 메모리 풀 타입이 아직 지원되지 않는 이전 CUDA 버전과 플랫폼에서는 스트림 순서 비의존 할당을 폴백(fallback)으로 제공하지만, 풀 지원이 보편화되면 제거할 계획입니다.
위 코드에서는 먼저 특정 디바이스의 기본 메모리 풀을 가져오는데, 이때 cudaMallocAsync의 암묵적 디바이스 선택에 의존하는 대신 명시적 인수로 전달합니다. 예제에서는 가능하면 기본 풀을 사용하는 것이 권장되지만, CCCL 런타임은 다른 풀 설정이 필요할 때 별도의 풀 객체를 생성하는 것도 지원합니다.
이 풀 참조를 사용해 새로운 cuda::make_buffer로 버퍼 세 개를 생성합니다. 스트림 순서 동작임을 나타내기 위해 스트림을 첫 번째 인수로 받습니다. 각 버퍼는 해당 스트림에 세 가지 작업을 제출합니다. 지정된 풀에서의 할당, 초기화, 그리고 버퍼가 스코프를 벗어날 때의 해제입니다.
커널이 덮어쓸 버퍼 C처럼 cuda::no_init으로 명시적으로 선택하지 않는 한, 초기화는 필수입니다. 초기화되지 않은 디바이스 메모리는 진단하기 어려운 버그의 흔한 원인이므로, 묵시적 기본값으로 두는 대신 명시적으로 선택하도록 설계하였습니다. 입력 버퍼 A와 B의 모든 요소는 각각 1과 2로 초기화됩니다. 버퍼는 다른 버퍼나 범위(range)에서 초기화하는 등 추가적인 초기화 모드도 지원합니다.
버퍼 수명과 해제
make_buffer에 전달된 스트림은 버퍼 내부에 저장되며, 버퍼가 소멸될 때 해제에 사용됩니다. 따라서 연산이 해제와 올바르게 순서를 지키도록, 버퍼에는 일반적으로 해당 버퍼의 사용에 대응하는 스트림을 보유해야 합니다. .set_stream()으로 나중에 스트림을 변경하거나 .destroy()로 특정 스트림에서 소멸을 수동으로 트리거하는 것도 가능하지만, 기본 동작은 일반적인 경우에 올바르게 작동하도록 설계되어 있습니다.
{
auto pool = cuda::device_default_memory_pool(device);
// Equivalent to cudaMallocFromPoolAsync on the stream, possibly along with initialization pushed into the stream as well. Saves the stream for future deallocation
auto buffer = cuda::make_buffer(allocation_stream, pool, /*... */);
// buffer usage...
}
// Closing bracket will call cudaFreeAsync on allocation_stream, there is also buffer.destroy(which_stream) to keep the behavior explicit3. 커널 실행
struct kernel {
template <typename Config>
__device__ void operator()(Config config,
cuda::std::span<const int> A,
cuda::std::span<const int> B,
cuda::std::span<int> C) {
auto tid = cuda::gpu_thread.rank(cuda::grid, config);
if (tid < A.size())
C[tid] = A[tid] + B[tid];
}
};
// ...
constexpr int threads_per_block = 256;
auto config = cuda::distribute<threads_per_block>(num_elements);
cuda::launch(stream, config, kernel{}, A, B, C);마지막 섹션에서는 cuda::launch로 GPU에서 커널을 설정하고 실행하는 방법을 살펴봅니다.
cuda::launch는 세 그룹의 인수를 받습니다.
- 실행할 스트림
- 스레드 계층(블록 및 그리드 크기)과 기타 실행 옵션을 인코딩하는 설정(configuration) 객체. 여기서
cuda::distribute는threads_per_block크기의 블록으로 구성된 적어도num_elements개의 스레드를 실행하는 설정을 생성합니다. 이는 많은 CUDA 개발자에게 익숙한(N + block_size - 1) / block_size패턴을 대체합니다. - 커널과 그 인수
컴파일 타임 설정 흐름
cuda::launch에서 가장 새로운 측면은 컴파일 타임 정보를 타입 시스템을 통해 호스트 실행 지점에서 디바이스 코드로 전달하는 방식입니다. 예를 들어, 블록 크기가 cuda::distribute의 템플릿 인수로 제공되는 점에 주목하십시오. 이는 블록 크기가 설정 객체의 타입에 인코딩됨을 의미합니다.
커널이 첫 번째 인수로 해당 설정을 받으면, cuda::launch가 이를 자동으로 전달합니다. 커널 내부에서는 그리드 내 호출 스레드의 순위를 계산할 때 이 정적 정보를 활용할 수 있습니다.
auto tid = cuda::gpu_thread.rank(cuda::grid, config);블록 크기를 컴파일 타임에 알 수 있으므로, 순위 계산에서는 x 차원만 사용하고 런타임 블록 크기 쿼리를 완전히 생략할 수 있습니다. 이는 단순한 예시이지만 메커니즘은 일반화됩니다. CCCL 문서에서는 설정에 내장된 정보를 활용하여 디바이스 코드를 특수화하는 추가 사례를 확인할 수 있습니다. 때로는 커널 구현이 그리드 및/또는 블록의 정확한 형태에 대한 가정을 두기도 합니다. 설정 객체 내 컴파일 타임 정보를 활용하면 커널 작성자가 커널과 호출 지점 간의 정렬을 검증하는 체크를 구현할 수 있습니다.
template <typename Config>
__global__ void kernel(Config conf) {
// Make sure the block is one dimensional with 256 threads
static_assert(cuda::gpu_thread.static_dims(cuda::block, conf).x == 256);
static_assert(cuda::gpu_thread.static_dims(cuda::block, conf).y == 1);
static_assert(cuda::gpu_thread.static_dims(cuda::block, conf).z == 1);
}커널 펑터
커널이 __global__ 함수가 아니라 __device__ 연산자(operator())를 가진 구조체(struct)라는 점을 알아채셨을 것입니다. cuda::launch는 기존의 __global__ 함수도 지원하지만, 이와 함께 __device__가 지정된 호출 연산자 타입인 커널 펑터를 도입했습니다. 실질적인 장점은 템플릿 인수가 자동으로 추론된다는 것입니다. 기존 __global__ 함수를 cuda::launch와 사용할 때는 명시적인 인스턴스화가 필요했던 반면, 커널 펑터를 사용하면 이 과정이 자동으로 이루어집니다.
template <typename T>
__global__ void kernel_function(T input) {
// body ...
}
struct kernel_functor {
template <typename T>
__device__ void operator()(T input) {
// body ...
}
};
// explicit template instantiation is required with a __global__ function
cuda::launch(stream, config, kernel_function<int>, 42);
// deduction from arguments for a functor with __device__ call operator
cuda::launch(stream, config, kernel_functor{}, 42);이것이 컴파일 타임 설정 흐름을 가능하게 하는 원리입니다. config 템플릿 매개변수가 cuda::launch에서 전달된 설정 객체로부터 추론됩니다. 커널 펑터는 디바이스 람다도 포함하며, CCCL 문서에 설명된 추가 기능도 갖고 있습니다.
자동 인수 변환
cuda::buffer는 내부 할당을 소유하지만, CUDA 커널은 간단히 복사 가능한(trivially copyable) 인수만 받을 수 있습니다. 버퍼를 cuda::launch에 전달하면 자동으로 cuda::std::span으로 변환됩니다. 스팬(span)을 직접 구성하거나 원시 포인터를 추출할 필요가 없습니다. 커널 시그니처는 디바이스 측에서 데이터가 실제로 사용되는 방식을 반영합니다.
다음 단계
본 포스트에서는 명시적 의존성, 강력한 타입 시스템, 기본 비동기 방식의 API, 그리고 기존 CUDA 코드와의 깔끔한 상호 운용성 등 CCCL 런타임의 핵심 개념들을 살펴보았습니다. 하지만 하나의 예제를 살펴보는 것만으로는 전체를 다 파악하기 어렵습니다.
CCCL 문서에는 추가적인 버퍼 초기화 모드, 이벤트 관리, 데이터 이동, 그리고 동적 공유 메모리 및 기타 론치 속성과 같은 고급 커널 론치 기능을 포함하여 각 API에 대한 더 자세한 내용이 다루어져 있습니다. CCCL 런타임 API는 CUDA 툴킷 13.2 이상 버전과 함께 제공되는 CCCL 3.2 이상 버전에서 사용할 수 있습니다. API별 자세한 가용 여부는 CCCL 문서를 참조해 주시기 바랍니다. 직접 사용해 보시고 많은 피드백을 남겨주시면 감사하겠습니다.
