동일한 입력 데이터를 사용해 여러 번 연산을 수행했을 때 비트 단위까지 일치하는 결과를 도출한다면, 해당 연산은 결정론적(Deterministic)이라고 정의합니다. 언뜻 당연하고 단순해 보이는 속성일 수 있으나, 병렬 프로그래밍과 부동 소수점 연산이 개입되는 실제 현장에서는 이를 보장하기가 매우 까다롭습니다. 부동 소수점 덧셈과 곱셈은 엄밀히 따져 결합법칙이 성립하지 않기 때문입니다. 즉, 유한한 정밀도로 중간 결과를 저장할 때 발생하는 반올림 오차로 인해 (a + b) + c의 결과가 a + (b + c)와 일치하지 않을 수 있습니다.
NVIDIA CUDA Core Compute Libraries (CCCL) 3.1 버전부터는 속도 최적화 병렬 디바이스 알고리즘을 위한 저수준 CUDA 라이브러리인 CUB에 새로운 단일 단계가 추가되었습니다. 이 API는 실행 환경 설정을 지원하여 사용자가 알고리즘의 동작 방식을 정교하게 커스터마이징할 수 있도록 돕습니다. 개발자는 이 환경 설정을 활용해 reduce 알고리즘의 결정론적 속성을 직접 구성하게 됩니다. 참고로, 기존의 2단계 API는 실행 환경 설정을 지원하지 않으므로 오직 새로운 단일 단계 API를 통해서만 이 제어가 가능합니다.
다음은 CUB에서 결정론 수준을 지정하는 방법을 보여주는 코드 예시입니다. (전체 예제는 Compiler Explorer에서 확인하실 수 있습니다.)
auto input = thrust::device_vector<float>{0.0f, 1.0f, 2.0f, 3.0f};
auto output = thrust::device_vector<float>(1);
auto env = cuda::execution::require(cuda::execution::determinism::not_guaranteed); // can be not_guaranteed, run_to_run (default), or gpu_to_gpu
auto error = cub::DeviceReduce::Sum(input.begin(), output.begin(), input.size(), env);
if (error != cudaSuccess)
{
std::cerr << "cub::DeviceReduce::Sum failed with status: " << error << std::endl;
}
assert(output[0] == 6.0f);
우선 입력 및 출력 벡터를 지정하는 것부터 시작합니다. 그 다음 cuda::execution::require()를 사용하여 cuda::std::execution::env 객체를 생성하고, 결정론 수준을 not_guaranteed로 설정합니다.
reduce 알고리즘에서 사용할 수 있는 결정론 수준은 다음과 같이 세 가지입니다.
not_guaranteedrun_to_rungpu_to_gpu
결정론 미보장 (Determinism not guaranteed)
부동 소수점 리덕션 연산에서 결과값은 요소들이 결합되는 순서에 따라 달라질 수 있습니다. 만약 두 번의 실행에서 리덕션 연산자가 서로 다른 순서로 적용된다면, 최종 결과값에 미세한 차이가 발생할 수 있습니다. 많은 애플리케이션에서는 이러한 사소한 차이를 허용합니다. 엄격한 결정론 요구 사항을 완화함으로써, 리덕션 구현체는 연산 순서를 자유롭게 재배치할 수 있게 되며 이는 곧 실행 성능의 향상으로 이어집니다.
CUB에서 not_guaranteed 설정은 결정론 수준을 완화합니다. 이를 통해 원자적 연산(Atomic operations)을 사용하여 블록 수준의 부분 합계와 최종 리덕션 값을 모두 계산할 수 있습니다. 스레드 간에 순서 없이 실행되는 원자적 연산의 특성상 실행 시마다 연산 순서가 달라질 수 있지만, 이 덕분에 단 한 번의 커널 실행만으로 전체 리덕션 프로세스를 완료할 수 있습니다. 원자적 연산이 각 블록의 부분 합계를 결과값에 직접 통합하기 때문입니다.
결정론을 보장하지 않는 이 리덕션 변체는 일반적으로 실행 간 결정론을 보장하는 버전보다 빠릅니다. 특히 입력 배열의 크기가 작을 때 유리한데, 단일 커널로 리덕션을 수행함으로써 여러 번의 커널 호출로 인한 지연 시간을 줄이고, 불필요한 데이터 이동과 추가적인 동기화 과정을 최소화할 수 있기 때문입니다. 다만, 결정론적 동작의 부재로 인해 반복 실행 시 결과가 미세하게 다를 수 있다는 점이 트레이드오프입니다.
실행 간 결정론 (Run-to-run determinism)
비결정론적 리덕션이 잠재적인 성능 이득을 제공하는 반면, CUB는 실행 시마다 일관된 결과를 보장하는 모드도 제공합니다. 기본적으로 cub::DeviceReduce는 실행 간 결정론을 따르며, 이는 단일 단계 API에서 결정론 수준을 run_to_run으로 설정하는 것과 같습니다. 이 모드에서 동일한 입력, 커널 실행 구성 및 GPU를 사용한 여러 번의 호출은 동일한 출력을 생성합니다.
이러한 결정론은 실행 시마다 업데이트 순서가 달라질 수 있는 원자적 연산(Atomics)에 의존하는 대신 리덕션을 고정된 계층적 트리 구조로 구성함으로써 달성됩니다. 리덕션의 각 단계에서 요소들은 먼저 개별 스레드 내에서 결합됩니다. 중간 결과들은 셔플 지시어를 사용하여 워프 내 스레드 간에 리덕션된 후, 공유 메모리를 사용하는 블록 단위 리덕션이 뒤따릅니다. 마지막으로 두 번째 커널이 블록당 결과를 집계하여 최종 출력을 생성합니다. 이 시퀀스는 미리 결정되어 있으며 스레드 실행의 상대적 타이밍과 무관하기 때문에 동일한 입력, 커널 구성 및 GPU는 동일한 비트 결과를 생성합니다.
GPU 간 결정론 (GPU-to-GPU determinism)
최고 수준의 재현성이 필요한 애플리케이션을 위해 CUB는 동일한 입력에 대해 서로 다른 GPU에서도 동일한 결과를 보장하는 GPU 간 결정론을 제공합니다. 이 모드는 단일 단계 API에서 결정론 수준을 gpu_to_gpu로 설정하는 것과 대응됩니다.
이러한 수준의 결정론을 달성하기 위해 CUB는 재현 가능한 부동 소수점 누산기(RFA, Reproducible Floating-point Accumulator)를 사용합니다. 이는 NVIDIA GTC 2024 세션인 “HPC에 과학적 방법론 복원하기: 고성능 재현 가능 병렬 리덕션“을 기반으로 한 솔루션입니다. RFA는 서로 다른 지수를 가진 숫자를 더할 때 발생하는 부동 소수점 비결합성에 대응하기 위해 모든 입력값을 고정된 수의 지수 범위(기본값은 3개의 빈)로 그룹화합니다. 이처럼 고정되고 구조화된 누적 순서는 최종 결과가 GPU 아키텍처에 관계없이 독립적임을 보장합니다.
최종 결과의 정확도는 빈(bin)의 개수에 따라 달라집니다. 빈이 많을수록 정확도는 높아지지만 중간 합산 횟수가 늘어나 성능이 저하될 수 있습니다. 현재 구현체는 성능과 정확도의 균형을 맞춘 최적의 기본값인 3개의 빈을 사용합니다. 이 구성은 단순히 엄격한 결정론을 따를 뿐만 아니라, 병렬 리덕션에서 전통적으로 사용되는 표준 쌍별 합산(Pairwise summation)보다 더 타이트한 오차 범위를 제공하여 수치적으로 정확한 결과를 보장한다는 점에 주목할 가치가 있습니다.
결정론 수준에 따라 결과과 변화됩니다.
세 가지 결정론 수준은 여러 번의 실행 과정에서 발생하는 변동 폭에서 차이를 보입니다.
- 결정론 미보장(Not-guaranteed) 수준은 호출할 때마다 약간씩 다른 합산 값을 생성합니다.
- 실행 간 결정론(Run-to-run) 수준은 단일 GPU의 모든 호출에서 동일한 값을 보장하지만, 다른 GPU를 사용하면 결과가 달라질 수 있습니다.
- GPU 간 결정론(GPU-to-gpu) 수준은 리덕션을 실행하는 GPU에 관계없이 모든 호출에서 합산 값이 동일함을 보장합니다.
그림 1은 각 결정론 수준별 배열 합산 결과를 실행 횟수에 따라 녹색, 파란색, 빨간색 원으로 표시한 것입니다. 평평한 수평선은 리덕션이 동일한 결과를 생성하고 있음을 나타냅니다.
결정론 수준별 성능 비교
선택한 결정론(Determinism) 수준은 cub::DeviceReduce의 성능에 직접적인 영향을 미칩니다. 요구 사항이 가장 완화된 결정론 미보장(Not-guaranteed) 수준은 가장 높은 성능을 제공합니다. 기본 설정인 실행 간 결정론(Run-to-run)은 준수한 성능을 보여주지만, 결정론 미보장 수준보다는 약간 느립니다. 반면, 서로 다른 GPU 간에도 엄격한 재현성을 강제하는 GPU 간 결정론(GPU-to-GPU)은 성능이 크게 저하될 수 있으며, 대규모 문제 크기에서는 실행 시간이 20%에서 30%까지 증가할 수 있습니다.
그림 2는 NVIDIA H200 GPU에서 float32 및 float64 입력을 사용하여 각 결정론 요구 사항에 따른 성능을 비교한 결과입니다(수치가 낮을수록 성능이 좋음). 이 차트는 결정론 수준의 선택이 다양한 데이터 유형에 걸쳐 실행 시간에 어떤 영향을 미치는지 명확하게 보여줍니다.
결론
단일 단계(Single-phase) API와 명시적인 결정론 수준의 도입으로, CUB는 리덕션 알고리즘의 동작과 성능을 모두 제어할 수 있는 강화된 도구 상자를 제공하게 되었습니다. 사용자는 고성능과 유연성을 갖춘 미보장(Not-guaranteed) 모드부터 신뢰할 수 있는 기본 설정인 실행 간(Run-to-run) 결정론, 그리고 가장 엄격한 GPU 간(GPU-to-GPU) 재현성까지 자신의 필요에 가장 적합한 수준을 선택할 수 있습니다.
CUB의 결정론 제어는 리덕션에만 국한되지 않습니다. NVIDIA는 개발자들이 더 폭넓은 병렬 CUDA 프리미티브에서 재현성을 제어할 수 있도록 이러한 기능을 추가 알고리즘으로 확장할 계획입니다. 확장된 결정론 지원에 관한 최신 업데이트와 논의는 진행 중인 GitHub 이슈에서 확인하실 수 있으며, 로드맵을 확인하거나 결정론적 버전이 필요한 알고리즘에 대한 피드백을 공유하실 수 있습니다.
