제2부 CUDA 프로그래밍

2.16 CUDA 라이브러리 II — Python (CuPy·Numba·Warp)

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

2.15에서 C++ 수치 라이브러리(Thrust·CUB·도메인)를 봤다. 이 장은 GPU 컴퓨팅의 Python 입구 — CuPy·Numba·cuda-python·RAPIDS·Warp — 를 구분하고, 같은 리덕션·정렬을 네 방식으로 작성하며 0-copy 연동을 본다. 이로써 제2부(CUDA)가 끝난다.

배우는 것:

• Python 진입점을 구분한다 — CuPy(NumPy 호환 배열), Numba @cuda.jit(파이썬으로 커널), cuda-python·cuda.core(저수준 바인딩), RAPIDS(cuDF·cuML), Warp(시뮬레이션 커널, 9.11 연결).
• 같은 작업을 직접 커널 ↔ Thrust ↔ CuPy ↔ Numba로 작성하며 추상화 사다리를 오르내린다.
• __cuda_array_interface__ 0-copy 연동과 전송·버전 함정.

선수 지식: 2.1~2.3(첫 커널·실행 구성), 2.8~2.9(워프 리덕션·아토믹), 1.10(메모리 바운드·융합), 2.12(컴파일·PTX). 빌드 환경은 0.3.

2.16.2 Python 진입 — 지도

GPU 컴퓨팅 입구가 Python으로 넓어진 이유는 분명하다 — 데이터과학·ML 생태계가 Python이고, 반복 실험에 빌드 사이클이 없는 편이 빠르다. 다만 진입점은 추상화 높이와 목적이 제각각이다.

이들을 묶는 접착제는 __cuda_array_interface__(CAI) 프로토콜이다 — CuPy 배열·Numba 디바이스 배열·PyTorch 텐서 등이 복사 없이 같은 디바이스 메모리를 가리키게 하는 규약(2.16.6). 덕분에 "CuPy로 데이터 만들고 Numba 커널로 처리하고 PyTorch로 넘긴다"가 0-copy로 가능하다.

2.16.3 CuPy — NumPy 호환 GPU 배열

CuPy는 NumPy(그리고 SciPy 상당 부분)의 드롭인 GPU 대체를 목표로 한다. import numpy as np를 import cupy as cp로 바꾸면 많은 코드가 그대로 GPU에서 실행된다. 내부적으로 cuBLAS·cuFFT·cuRAND·cuSPARSE·cuSOLVER·CUB/Thrust를 호출 — 2.15 라이브러리들을 파이썬으로 쓰는 가장 쉬운 길.

import cupy as cp
x = cp.random.rand(10_000_000, dtype=cp.float32)  # cuRAND가 디바이스에서 생성 (H2D 없음)
s = cp.sort(x)                                     # Thrust/CUB 정렬
total = cp.sum(x)                                  # CUB 리덕션
print(float(total), float(s[0]), float(s[-1]))     # float()이 D2H 1원소 (필요한 순간에만!)

pip install cupy-cuda12x   # CUDA 12.x면 그에 맞는 휠 (버전 정합 핵심 — 시점 의존)

커널 융합 — fuse·RawKernel. x * x 후 cp.sum은 융합 기회(1.10):

# (A) @cp.fuse: 여러 elementwise 연산을 한 커널로 융합
@cp.fuse()
def sq(a): return a * a
sumsq = cp.sum(sq(x))     # sq는 한 패스로 융합, 그 뒤 reduce

# (B) RawKernel: 실제 CUDA C 커널을 문자열로 넣어 컴파일 (2.1의 커널을 파이썬에서)
ker = cp.RawKernel(r'''
extern "C" __global__ void scale(float* y, const float* x, float a, int n){
    int i = blockIdx.x*blockDim.x + threadIdx.x;
    if (i < n) y[i] = a * x[i]; }''', 'scale')
y = cp.empty_like(x)
threads = 256; blocks = (x.size + threads - 1)//threads
ker((blocks,), (threads,), (y, x, cp.float32(2.0), cp.int32(x.size)))   # 2.3의 실행 구성; 스칼라는 폭 맞춰 전달

RawKernel은 C++ 커널을 그대로 파이썬에서 컴파일·실행하는 탈출구 — CuPy 추상화로 부족할 때 2.1~2.9 지식을 그대로.

2.16.4 Numba — @cuda.jit로 파이썬에서 커널

Numba는 파이썬 함수를 LLVM으로 JIT 컴파일해 PTX로 내린다(2.12의 흐름에 합류). @cuda.jit를 붙이면 파이썬 문법으로 CUDA 커널을 짠다 — threadIdx·blockIdx·공유 메모리·아토믹·워프 프리미티브까지 상당 부분 노출.

from numba import cuda
import numpy as np
@cuda.jit
def saxpy(y, x, a):
    i = cuda.grid(1)                       # blockIdx.x*blockDim.x + threadIdx.x 한 줄
    if i < x.size: y[i] = a * x[i] + y[i]
n = 1_000_000
x = np.ones(n, np.float32); y = np.zeros(n, np.float32)
d_x = cuda.to_device(x); d_y = cuda.to_device(y)   # 명시적 H2D
threads = 256; blocks = (n + threads - 1)//threads
saxpy[blocks, threads](d_y, d_x, np.float32(2.0))  # [blocks, threads] 실행 구성 (2.3)
y = d_y.copy_to_host()                              # 명시적 D2H

Numba는 공유 메모리(cuda.shared.array)·동기화(cuda.syncthreads())·아토믹(cuda.atomic.add)을 제공해 2.7·2.8의 패턴(타일링·블록 리덕션)을 파이썬으로 옮길 수 있다. 한계도 분명 — C++ 템플릿·CUB·라이브러리 연동·최신 텐서코어 경로는 C++가 앞선다. 빠른 프로토타이핑·교육·중간 규모 커스텀 커널에 강하다.

CuPy vs Numba 경계: 표준 배열 연산은 CuPy(라이브러리 호출), 고유 커스텀 커널 로직은 Numba @cuda.jit. 둘은 __cuda_array_interface__로 0-copy 연동(2.16.6).

2.16.5 cuda-python·RAPIDS·Warp

• cuda-python / cuda.core: CUDA Runtime·Driver·NVRTC API의 공식 파이썬 바인딩. CuPy·Numba가 추상화라면 이것은 밑단 C API를 거의 1:1로 — 스트림·이벤트·메모리·모듈 로드·커널 런치를 직접 호출. 더 파이썬다운 상위 cuda.core(API 진화 중). 라이브러리·프레임워크 저작자용, 대부분의 응용 개발자는 CuPy·Numba면 충분.
• RAPIDS — cuDF·cuML: 데이터과학 파이프라인 전체를 GPU에서. cuDF(pandas 호환 GPU 데이터프레임, Apache Arrow 기반), cuML(scikit-learn 호환 GPU ML), cuGraph·cuSpatial 등. 내부적으로 CUB·Thrust·cuSPARSE·cuSOLVER 위에 선다(2.15가 데이터프레임 한 줄 뒤에서 실행된다).
• Warp: 파이썬으로 미분 가능·고성능 시뮬레이션 커널. @wp.kernel로 커널 정의, 물리(충돌·강체·연체)·기하·미분(autodiff) 특화 자료형. Numba @cuda.jit와 결이 비슷하지만 시뮬레이션·로보틱스·미분 가능 물리에 맞춰짐.

import warp as wp
wp.init()
@wp.kernel
def integrate(pos: wp.array(dtype=wp.vec3), vel: wp.array(dtype=wp.vec3), dt: float):
    i = wp.tid()                       # 스레드 인덱스
    pos[i] = pos[i] + vel[i] * dt      # 단순 오일러 적분 (벡터 자료형 내장)
# wp.launch(integrate, dim=n, inputs=[pos, vel, 0.01])

Warp는 미분 가능성(역전파로 시뮬레이션 파라미터 학습) 덕에 Sim2Real·강화학습·로보틱스에서 쓰이며, 9.11(Isaac Sim)·9.13(Isaac Lab)로 이어진다.

2.16.6 워크스루 — 같은 작업을 네 방식 + 0-copy 연동

"배열의 합과 정렬"을 (A) 직접 커널 → (B) Thrust → (C) CuPy → (D) Numba로 작성하며 추상화 사다리를 오르내린다.

(A) 직접 커널 vs (B) Thrust (C++). 직접 커널(2.9)은 워프 셔플 + 블록 부분합 + 2차 합산까지 수십 줄. Thrust는 두 줄:

thrust::device_vector<int> d = h;
thrust::sort(d.begin(), d.end());
long long total = thrust::reduce(d.begin(), d.end(), 0LL);

교훈: 표준 리덕션·정렬에서 직접 커널은 거의 항상 손해. Thrust 내부는 CUB의 세대별 튜닝 커널이라 직접 작성한 첫 커널이 이기기 어렵다.

(C) CuPy + (D) Numba, 0-copy 연동.

import cupy as cp
from numba import cuda
x = cp.random.randint(0, 1000, size=1_000_000, dtype=cp.int32)   # (C) CuPy
s = cp.sort(x); total = int(cp.sum(x))     # int()에서만 D2H

@cuda.jit                                  # (D) Numba 커스텀 커널을 CuPy 배열 위에 0-copy
def add_const(a, c):
    i = cuda.grid(1)
    if i < a.size: a[i] += c
threads = 256; blocks = (x.size + threads - 1)//threads
add_const[blocks, threads](x, 7)   # x는 CuPy 배열 — 복사 없이 Numba 커널이 in-place 수정
print(int(x[0]))                    # +7 반영, D2H 없이 디바이스 위에서

핵심은 마지막 호출 — CuPy 배열을 Numba 커널에 그대로 넘겼다. __cuda_array_interface__가 양쪽이 같은 디바이스 포인터를 공유하게 해 H2D·D2H가 한 번도 일어나지 않았다. 이것이 "Python 진입점들을 0-copy로 잇는 접착제"의 실체다.

2.16.7 트레이드오프 — 의사결정 사다리

비NVIDIA 정직: SYCL·oneAPI(인텔, 이식성), Kokkos·RAJA(다중 백엔드 C++ 추상화), OpenMP target offload, 파이썬 쪽 PyTorch/JAX(딥러닝 텐서) 등. NVIDIA 라이브러리는 NVIDIA GPU에서 최고 성능·가장 성숙하나 이식성이 낮다. 멀티 벤더 이식성이 요구면 Kokkos·SYCL, NVIDIA 단일 타깃에서 성능 최우선이면 CUDA-X. — 견해.

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

• [ ] 같은 정렬·합을 C++(Thrust)와 Python(CuPy)으로 각각 5줄 안에 쓴다.
• [ ] "왜 직접 커널이 보통 Thrust보다 느린가"를 한 문장으로(세대 튜닝·알고리즘 분기·검증).
• [ ] CuPy 배열을 Numba 커널에 0-copy로 넘기는 조건(__cuda_array_interface__)을 설명한다.
• [ ] 결과를 매 원소 D2H로 끌지 않고, cupy 휠·Toolkit 버전을 맞추고, 핸들·플랜을 재사용한다.

2.16.9 흔한 함정

• 결과를 원소 단위 D2H로. Thrust d[i]·CuPy float(arr[i])를 루프에서 부르면 매번 PCIe 왕복. 디바이스에서 모은 뒤 한 번에.
• 버전 정합 무시. cupy-cuda11x를 CUDA 12 환경에, 또는 드라이버보다 새 Toolkit(0.3). cuSPARSE처럼 버전 간 시그니처가 바뀐 것은 사용 중 버전 문서로.
• CuPy/Numba/PyTorch 사이를 매번 .copy_to_host()로 왕복. __cuda_array_interface__로 0-copy 가능한데 호스트를 경유하면 대역폭을 버린다.
• Numba에서 C++ 수준 성능 기대. 강력하지만 최신 텐서코어·CUB·라이브러리 융합은 C++가 앞선다. 최상단 성능이면 C++/라이브러리로.
• 모든 것을 직접 커널로. 표준 연산은 라이브러리가 거의 항상 이긴다(2.15).

2.16.10 연습문제

연습 1 (직접 커널 → 라이브러리). 2.9의 워프 셔플 리덕션을 Thrust reduce 또는 CuPy sum으로 대체하고 코드 줄 수·시간을 비교하라.

연습 2 (0-copy 연동). CuPy로 난수 배열을 만들고 Numba @cuda.jit 커널로 in-place 변환한 뒤 다시 CuPy로 합을 구하라. 호스트 복사가 한 번도 안 일어남을 보장하는 메커니즘은?

연습 3 (진입점 선택). 각 작업에 무엇을 먼저: (a) 빠른 실험용 elementwise 변환 (b) 도메인 특화 충돌 검사 커널 (c) pandas ETL 가속 (d) 미분 가능 강체 시뮬.

연습 4 (CUB 빌딩블록). 2.7·2.9에서 직접 작성한 블록 리덕션을 cub::BlockReduce(C++)로 교체하라. TempStorage와 __syncthreads 처리가 어떻게 바뀌나?

2.16.11 요약 · 다음 장

• Python 진입은 목적별로 다르다: CuPy(NumPy 호환 배열·라이브러리 래퍼), Numba @cuda.jit(파이썬으로 커널), cuda-python·cuda.core(저수준), RAPIDS(데이터프레임·ML), Warp(시뮬레이션·9.11).
• __cuda_array_interface__ 가 이들을 0-copy로 잇는다.
• 실무 함정 셋: 결과를 원소 단위 D2H로 끌지 말 것, 핸들·플랜·워크스페이스 재사용, cupy 휠·Toolkit·드라이버 버전 맞추기.

제2부(CUDA 프로그래밍)가 끝났다. 다음 부 — 제3부 딥러닝 라이브러리(3.1~3.8). 3.1(cuBLAS·cuBLASLt)에서 이 장이 "위치만 표시한" GEMM 라이브러리를 본격적으로 파고, cuDNN·CUTLASS·어텐션·데이터로딩으로 이어진다. "라이브러리 vs 직접 커널" 판단과 핸들·워크스페이스·전송 함정은 그 라이브러리들에서 그대로 재등장한다.