제2부 CUDA 프로그래밍

2.15 CUDA 라이브러리 I — Thrust·CUB·CCCL

지금까지 커널을 직접 작성하고, 타일링하고, 워프 프리미티브로 리덕션을 접고, Nsight로 프로파일링하는 법을 익혔다. 이 장은 정반대 질문을 던진다 — 언제 직접 커널을 쓰지 않을 것인가? 정렬·리덕션·FFT·희소·난수 같은 잘 알려진 문제는 NVIDIA가 수년간 갈아 넣은 라이브러리(CUDA-X)가 거의 항상 직접 작성한 첫 커널보다 빠르고 정확하다. 이 장은 C++ 수치 라이브러리, 다음 장(2.16)이 Python 진입점이다.

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

배우는 것:

CUDA-X 라이브러리 계층 지도를 그리고 "직접 커널 vs 라이브러리"를 판단한다.
Thrust(STL 유사 알고리즘)로 정렬·리덕션·스캔·변환을, CUB(블록·워프 빌딩블록)와의 역할 분담을 안다. CCCL 통합의 의미.
cuFFT·cuRAND·cuSPARSE·cuSOLVER가 각각 무엇을 푸는지.
핸들·워크스페이스·전송 함정.

선수 지식: 1.1(처리량 컴퓨팅), 1.10(메모리 바운드·융합), 2.1(첫 커널), 2.8~2.9(워프 리덕션·아토믹), 2.12(컴파일·PTX). "커널을 직접 작성할 수 있다"에서 출발해 그것을 언제 안 할지를 배운다. 딥러닝 전용 라이브러리(cuBLAS·cuDNN 등)는 제3부(3.1~).

2.15.2 왜 라이브러리인가 — 직접 커널의 한계

2.7에서 타일링 행렬곱을 직접 작성했을 때 cuBLAS 대비 한참 느렸을 것이다. 잘 알려진 문제일수록 라이브러리는 당신이 모르는 것까지 한다: 세대별 아키텍처 튜닝(SM 버전마다 다른 타일·언롤·텐서코어 경로), 수치 안정성(FFT 정밀도, 희소 솔버 피벗, 난수 통계 품질), 입력 크기·정렬에 따른 알고리즘 분기, 수년간의 버그 수정·회귀 테스트.

견해: 새 문제를 만나면 "이것은 이미 라이브러리에 있다"를 기본 가정으로 검색부터. 직접 커널은 프로파일링(2.14)으로 "라이브러리 경로가 병목이고 융합·특화로 이길 수 있다"를 확인한 뒤에 정당화된다.

2.15.3 CUDA-X 라이브러리 지도

가로축은 무슨 문제냐(수치·데이터 vs 딥러닝), 세로축은 추상화 높이(Python → 라이브러리 → 런타임 → 하드웨어). 이 장 초점은 왼쪽 아래(Thrust·CUB·cuFFT·cuRAND·cuSPARSE·cuSOLVER).

배포 형태 두 가지 (사용법에 직접 영향):

형태	예	링크	특징
헤더 온리(C++ 템플릿)	Thrust·CUB·libcudacxx(=CCCL)	별도 `-l` 불필요, `nvcc`로 컴파일	템플릿·functor로 융합·인라인. 컴파일 시간↑
사전 컴파일	cuFFT·cuRAND·cuSPARSE·cuSOLVER·cuBLAS	`-lcufft …` 링크	핸들·플랜 객체. 빠른 빌드

2.15.4 Thrust — STL 같은 병렬 알고리즘

Thrust는 C++ STL을 GPU로 옮긴 헤더 온리 템플릿 라이브러리다. thrust::sort·reduce·transform·inclusive_scan을 벡터에 적용하면 내부적으로 잘 튜닝된 커널이 실행된다. 핵심 셋:

컨테이너: device_vector<T>(디바이스), host_vector<T>(호스트). 대입(d = h)이 곧 cudaMemcpy.
알고리즘: 첫 인자에 실행 정책(thrust::device/thrust::host).
functor·iterator: 동작 커스터마이즈(thrust::plus<>·사용자 functor)와 가짜 시퀀스(counting_iterator·zip_iterator).

#include <cstdio>
#include <thrust/device_vector.h>
#include <thrust/sort.h>
#include <thrust/reduce.h>
int main() {
    thrust::host_vector<int>   h = {5,2,8,1,9,3,7,4,6,0};
    thrust::device_vector<int> d = h;            // H2D
    thrust::sort(d.begin(), d.end());            // GPU 정렬 (radix/merge 자동 선택)
    int total = thrust::reduce(d.begin(), d.end(), 0, thrust::plus<int>());  // GPU 리덕션
    printf("sum=%d, min=%d, max=%d\n", total, (int)d[0], (int)d[d.size()-1]); // d[i]는 D2H 1원소(느림)
    return 0;
}

nvcc -O3 -arch=sm_80 thrust_demo.cu -o thrust_demo   # Thrust는 Toolkit 동봉, 별도 -l 불필요
# 예상(개략): sum=45, min=0, max=9

functor로 융합 — transform_reduce. "각 원소를 제곱한 합"을 transform 후 reduce로 두 번 호출하면 중간 배열을 글로벌에 쓰고 다시 읽는다. transform_reduce는 한 패스로 융합한다:

struct square { __host__ __device__ float operator()(float x) const { return x*x; } };
float sumsq = thrust::transform_reduce(thrust::device, d.begin(), d.end(),
    square(), 0.0f, thrust::plus<float>());  // 중간 배열 없이 한 번에 → 대역폭 절약(1.10)

직접 커널을 작성하지 않고 메모리 왕복을 줄였다.

2.15.5 CUB — 블록·워프 빌딩블록과 CCCL

Thrust가 "디바이스 전체 알고리즘"이라면 CUB(CUDA UnBound)는 한 단계 아래 — 워프·블록·디바이스 수준 재사용 프리미티브다.

디바이스 수준: cub::DeviceReduce·DeviceScan·DeviceRadixSort·DeviceSelect. Thrust가 내부적으로 부르는 그 엔진. 알고리즘 선택·튜닝을 더 노출.
블록·워프 수준: cub::BlockReduce·BlockScan·WarpReduce — 내 커스텀 커널 안에서 블록 리덕션·스캔을 한 줄로. 2.9에서 직접 작성한 워프 셔플 리덕션을 검증 구현으로 대체.

#include <cub/cub.cuh>
template <int BLOCK>
__global__ void block_sum(const float* in, float* out, int n) {
    typedef cub::BlockReduce<float, BLOCK> BlockReduce;
    __shared__ typename BlockReduce::TempStorage temp;
    int i = blockIdx.x*BLOCK + threadIdx.x;
    float v = (i < n) ? in[i] : 0.0f;
    float block_total = BlockReduce(temp).Sum(v);        // 블록 내 합 (2.9 대체)
    if (threadIdx.x == 0) out[blockIdx.x] = block_total;
}

CUB 디바이스 API의 관용구는 두 번 호출 패턴(임시 저장소 크기 질의 → 할당 → 실행):

void* d_temp = nullptr; size_t temp_bytes = 0;
cub::DeviceReduce::Sum(d_temp, temp_bytes, d_in, d_out, n);  // 1차: temp_bytes만 채움
cudaMalloc(&d_temp, temp_bytes);
cub::DeviceReduce::Sum(d_temp, temp_bytes, d_in, d_out, n);  // 2차: 실제 실행
cudaFree(d_temp);

요지: Thrust는 "무엇을", CUB는 "어떻게"에 더 가깝다. 대부분 Thrust로 충분하고, 성능을 더 극대화하거나 커스텀 커널 내부 협력이 필요할 때 CUB로 내려간다.

CCCL — 셋을 하나로. 과거엔 Thrust·CUB·libcudacxx(cuda::std::·cuda::atomic)가 별도였다. 지금은 CCCL(CUDA Core Compute Libraries) 라는 하나의 묶음으로 통합돼 버전이 맞는다 — 상호 운용 보장, 버전 정합 문제↓(Toolkit 동봉 CCCL 기준). libcudacxx는 2.9에서 본 cuda::atomic의 출처다.

2.15.6 도메인 수치 라이브러리 — cuFFT·cuRAND·cuSPARSE·cuSOLVER

사전 컴파일·핸들 기반. "무엇을 푸는지"와 "공통 골격"을 잡으면 충분(세부 API는 방대하고 버전에 따라 변함 — A.3).

라이브러리	푸는 문제	핵심 객체	언제
cuFFT	1D·2D·3D 고속 푸리에 변환	`cufftHandle` 플랜	신호·영상, 스펙트럼, 주파수 영역 합성곱, 과학 시뮬
cuRAND	의사·준난수(균등·정규·포아송)	호스트(`curandGenerator`)·디바이스(`curandState`) API	몬테카를로, 초기화, 드롭아웃·증강
cuSPARSE	희소 행렬(SpMV·SpMM·변환)	행렬·벡터 디스크립터	희소 선형대수, 그래프, FEM, GNN 희소 곱
cuSOLVER	밀집·희소 솔버, 분해(LU·QR·Cholesky·SVD), 고유값	핸들 + 워크스페이스	최소제곱, 행렬 분해, 고유값, 과학 계산

공통 골격(이름만 다르고 흐름은 같다): ① 핸들/플랜 생성 → ② (필요 시) 워크스페이스 크기 질의 → cudaMalloc → ③ 디바이스 데이터 준비 → ④ 실행 → ⑤ 결과 회수 → ⑥ 정리(*Destroy, cudaFree).

핵심: 핸들·플랜·워크스페이스는 재사용하라. 매 호출마다 생성·파괴하면 그 오버헤드가 연산을 압도한다(반복 루프). cuRAND는 호스트 API(대량 생성)와 디바이스 API(커널 안 스레드별 curandState) 두 경로가 있어 데이터 흐름에 맞는 쪽을 고른다.

정직: 이 네 라이브러리의 시그니처·디스크립터 API는 CUDA 버전에 따라 적잖이 바뀌었다(특히 cuSPARSE generic API 전환). 구체 코드는 사용 중 Toolkit 버전 문서로(A.3).

2.15.7 핸들·워크스페이스·전송 함정

비효율 패턴 — 루프 안 매번 생성·파괴	좋은 패턴 — 한 번 만들고 재사용
`for step in range(N):`<br/> `create plan` (비쌈)<br/> `exec`<br/> `destroy plan` (비쌈)	`create handle/plan, alloc work`<br/>`for step in range(N):`<br/> `exec` (이것만 반복)<br/>`destroy handle/plan, free work`

2.5의 데이터 이동 교훈도 그대로 — 라이브러리 호출 자체보다 H2D·D2H 전송이 병목인 경우가 흔하다. 작은 연산을 위해 매번 호스트로 결과를 끌어오면(Thrust d[i] 단일 원소, cuRAND 후 즉시 D2H) PCIe 왕복이 GPU 연산을 압도한다. 데이터를 디바이스에 머무르게 하고 라이브러리·커스텀 커널을 디바이스 위에서 연쇄시키는 것이 정석.

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

[ ] "직접 커널 vs 라이브러리"를 판단하고, 직접 커널이 정당화되는 좁은 경우(없음/융합/오버헤드)를 제시한다.
[ ] Thrust(무엇을)·CUB(어떻게)·CCCL(통합)의 역할을 구분한다.
[ ] cuFFT·cuRAND·cuSPARSE·cuSOLVER가 각각 무엇을 푸는지, 핸들·워크스페이스 재사용 원칙을 안다.
[ ] 결과를 원소 단위 D2H로 끌지 않고 디바이스에 머무르게 한다.

2.15.9 흔한 함정

모든 것을 직접 커널로. 표준 연산은 라이브러리가 거의 항상 이긴다. 프로파일링(2.14)으로 라이브러리가 병목임을 확인하기 전엔 직접 커널을 정당화 말 것.
결과를 원소 단위 D2H로. Thrust d[i]·루프 안 D2H는 매번 PCIe 왕복. 디바이스에서 모은 뒤 한 번에.
루프 안 핸들·플랜·워크스페이스 생성·파괴. 한 번 만들고 재사용.
버전 정합 무시. cuSPARSE처럼 버전 간 시그니처가 바뀐 라이브러리는 사용 중 버전 문서로(A.3).
Thrust 컴파일 시간·템플릿 오류를 오판. 헤더 온리 템플릿의 대가. 오류는 보통 functor의 __host__ __device__ 누락·타입 불일치.

2.15.10 연습문제

연습 1 (직접 커널 → 라이브러리). 2.9에서 작성한 워프 셔플 리덕션을 Thrust reduce로 대체하라. 코드 줄 수와 (가능하면 측정한) 시간을 비교하라.

해설

직접 커널은 수십 줄(셔플+블록 부분합+2차 합산), Thrust는 한 줄. 성능은 보통 라이브러리가 같거나 더 빠르다(CUB 세대별 튜닝). 정확한 비는 GPU·크기 의존(측정값으로). 직접 커널이 이긴다면 그 차이가 융합·특화에서 오는지 의심.

연습 2 (융합). "각 원소를 제곱한 합"을 (a) transform 후 reduce, (b) transform_reduce로 작성하라. 왜 (b)가 대역폭에서 유리한가(1.10)?

정답

(a)는 중간 제곱 배열을 글로벌에 쓰고 다시 읽어 합 — 데이터 2회 왕복. (b)는 한 패스로 읽으며 제곱·누적해 글로벌 트래픽이 절반. 리덕션·elementwise는 메모리바운드라 트래픽이 곧 시간 → (b)가 빠르다.

연습 3 (라이브러리 선택). 각 작업에 무엇을 먼저: (a) 100만 키-값 정렬 (b) 신호 주파수 스펙트럼 (c) 도메인 특화 기하 충돌 (d) 희소 행렬-벡터 곱 (e) 빠른 실험용 elementwise 변환.

정답(예시)

(a) Thrust sort_by_key/CUB DeviceRadixSort. (b) cuFFT. (c) 직접 커널 또는 Warp/Numba(2.16) — 라이브러리에 없음. (d) cuSPARSE SpMV. (e) CuPy(또는 Numba, 2.16).

연습 4 (워크스페이스 재사용). cuFFT를 1000번 반복하는 루프에서 플랜을 안/밖에서 만드는 차이는? 어떤 프로파일링 신호(2.13)로 잡나?

정답

루프 안 생성은 매 반복 플랜 할당·튜닝·해제가 실제 FFT를 압도할 수 있다. 밖에서 한 번 만들면 1/1000로 분산. Nsight Systems 타임라인에서 FFT 커널 사이사이 플랜 생성·cudaMalloc·해제가 반복되면 신호(API 트레이스에 cufftPlan 호출이 루프마다).

2.15.11 요약 · 다음 장

표준 문제는 라이브러리 먼저. 정렬·리덕션·FFT·난수·희소·솔버는 CUDA-X가 세대 튜닝·검증까지 해 주며, 직접 작성한 첫 커널이 이기긴 어렵다. 직접 커널은 라이브러리에 없거나 융합·특화가 정당화될 때만.
Thrust=무엇을(STL 유사), CUB=어떻게(블록·워프 빌딩블록), CCCL=둘+libcudacxx 통합. Thrust로 뼈대, CUB로 커스텀 커널 내부 협력.
cuFFT·cuRAND·cuSPARSE·cuSOLVER는 핸들·워크스페이스를 재사용하고, 결과를 원소 단위 D2H로 끌지 말 것.

다음 장: 2.16 CUDA 라이브러리 II — Python. CuPy(NumPy 호환)·Numba(@cuda.jit)·cuda-python·RAPIDS·Warp 진입점과 __cuda_array_interface__ 0-copy 연동을 본다.