GPU CUDA 기반 데이터 병렬처리 지원 파이썬 Numba 라이브러리 사용법 소개

이 글은 GPU CUDA 병렬처리를 지원하는 Numba 라이브러리를 간략히 소개한다. CUDA는 현재 딥러닝 기술의 기반처럼 사용되는 사실상 산업표준이다. 딥러닝은 모든 연산이 텐서 행렬 계산이므로, NVIDIA GPU에 내장된 수많은 계산 유닛(실수 계산에 특화된 CPU)들을 사용한다. 

CUDA 아키텍처(Multi-Process Service :: GPU Deployment and Management Documentation)

CUDA는 내장된 수많은 계산 유닛에 입력 데이터를 할당하고, 행렬연산을 하여, 출력된 데이터를 CPU 메모리가 접근할 수 있도록 데이터 고속 전송/교환하는 역할을 한다. 그러므로, 딥러닝 모델 학습 성능은 GPU CUDA 성능에 직접적 영향을 받는다. 이 글은 파이썬에서 CUDA를 이용해 수치해석 등 계산성능을 극대화할 수 있는 방법과 간단한 예제를 살펴본다.

• Github - CUDA programming example (github.com)

GPU CUDA 소개

CUDA는 게임 화면에 렌더링되는 3차원 이미지를 2차원 픽셀에 맵핑하기 위한 수많은 행렬처리를 실시간 병렬 처리할 수 있도록 개발되어 왔다. 이런 이유로, 행렬 고속 연산이 필요한 딥러닝 학습에 적극 사용된 것이다. 

CUDA 기반 실시간 텐서 행렬 연산 결과

CUDA는 오랫동안 개발자의 요구사항을 반영해 발전되어, 개발 플랫폼으로서 탄탄한 생태계를 구축했다.

NVIDIA 개발자 사이트 (NVIDIA Developer)

CUDA 기반 레이트레이싱 렌더링 결과 (Ray Tracey's blog: GPU path tracing tutorial 3: GPU)

사실, 많은 스타트업이 이런 기능을 지원하는 딥러닝용 AI 칩을 FPGA 기법 등을 이용해 개발, 홍보하고 있으나, 이런 개발자 지원도구와 플랫폼 생태계 없다면, 산업계에서는 의미가 없다고 볼 수 있다. 

Numba 소개

Numba(넘바)는 파이썬 기반 CUDA GPU 프로그래밍을 지원한다. 넘바는 컴파일 기술을 지원하여, CPU와 GPU모드에서 코딩되는 데이터구조, 함수호출을 추상화한다. 넘바는 NVIDIA의 CUDA 함수와 설정을 랩핑한 고수준의 함수 API를 제공한다. 이를 통해, 개발자가 CUDA의 세부 설정에 신경쓸 필요없이, 데이터 병렬 처리 개발에만 집중할 수 있다. 

개발환경

넘바 개발 환경은 다음과 같다. 

• NVIDIA Compute Capability 5.0 이상 CUDA 지원 GPU 장착 PC (2023.12 시점)
• NVIDIA CUDA 11.2 이상 (링크 참고)
• NVIDIA TX1, TX2, 자비애, 젯슨 나노 
• GTX 9, 10, 16 시리즈. RTX 20, 30, 40 시리즈. H100 시리즈

CONDA 환경의 경우, 다음과 같이 CUDA 툴킷을 자동 설치할 수 있다.

conda install cudatoolkit 

 

넘바는 cuda python을 이용해 NVIDIA GPU CUDA와 바인딩한다. 

conda install nvidia::cuda-python

설치는 다음과 같다.

conda install numba

CUDA 프로그래밍 개념

이 장은 CUDA 프로그래밍 개념을 간단히 짚고 넘어간다. 사실 넘바는 이 부분까지 어느정도 추상화해 놓았으므로, 이런 개념이 있구나 정도만 생각해도 된다. 

CUDA는 실시간 렌더링에 필요한 텍스처 공유 메모리 구조와 기존 CPU설계에 영향을 받아, 그리드 > 블럭 > 쓰레드로 아키텍처가 설계되었다. 쓰레드는 행렬 계산을 지원하는 프로세스라 생각하면 된다.  

CUDA 논리-물리 구조

각 쓰레드들은 행렬 병렬 계산할 수 있다. 다음 그림은 이를 보여준다.

CUDA 그리드 계층 구조 (Execution model of a CUDA program on NVidia's GPU: Hierarchy grid)

각 쓰레드에서 병렬 실행되는 함수를 커널이라 한다. 다음은 함수가 실행될 그리드, 블럭, 쓰레드가 명시된 CUDA 코드 예시를 보여준다. 하나의 블럭에서 256개 쓰레드를 사용한다고 가정하면, 다음과 같이 코딩된다. 

__global__

void add(int n, float *x, float *y)

{

  int index = threadIdx.x;

  int stride = blockDim.x;

  for (int i = index; i < n; i += stride)

      y[i] = x[i] + y[i];

}

CUDA를 사용한 커널함수 add 호출은 다음과 같다. <<<>>> 키워드를 통해 사용할 블록, 쓰레드 수를 지정하고 있다. 

add<<<1, 256>>>(N, x, y);

add함수가 호출되면, 다음과 같이 내부적으로 실행된다. 

병렬 데이터 처리를 위한 쓰레드 인덱스 계산 예시(An Even Easier Introduction to CUDA | NVIDIA Technical Blog)

이런 방식을 통해, 3차원, 4차원 등 고차원 대용량 행렬을 그리드, 블럭, 쓰레드로 나눠, 병렬처리할 수 있다. 모든 수치해석모델은 행렬 텐서로 계산되므로, 사실상, 수학적인 계산 모델은 CUDA에서 실행될 수 있다.

행렬(선형대수) 연산 예(선형대수)

넘바에서 CUDA 코딩 개념

넘바에서 간단한 커널 함수를 만들어보자. 배열을 입력받아, 배열값을 하나 증가하는 병렬 처리 함수를 만들어 본다. 

from numba import cuda

@cuda.jit

def increment_by_one(an_array):

    tx = cuda.threadIdx.x     # cuda 객체에 포함된 쓰레드 ID 획득

    ty = cuda.blockIdx.x      # 그리드내 블럭 ID 획득

    bw = cuda.blockDim.x   # 블럭 폭(쓰레드 갯수) 획득

    # Compute flattened index inside the array

    pos = tx + ty * bw        # 해당 스레드 접근 인덱스 계산

    if pos < an_array.size:    # 전달된 배열 경계 체크

        an_array[pos] += 1   # 배열값 하나씩 증가

이때 사용된 @는 cuda 데코레이터로서, 파이썬 코드를 CUDA에서 처리될 수 있는 코드로 자동 컴파일하는 역할을 한다(참고). 선언된 커널 함수는 GPU 다중쓰레드실행 아키텍처에 직접 맵핑된다. 함수 호출은 다음과 같다. 이때 커널함수는 비동기적으로 실행된다. 

threadsperblock = 32

blockspergrid = (an_array.size + (threadsperblock - 1)) // threadsperblock

increment_by_one[blockspergrid, threadsperblock](an_array)

커널에 전달된 배열값은 CPU메모리와 CUDA메모리간에 복사되어져야 한다. 이를 위해, NumPy로부터 CPU > CUDA 복사하는 다음 함수를 제공한다. 

cpu_mem = np.arrage(10)

cuda_mem = cuda.to_device(cpu_mem)

CUDA > CPU 복사 함수는 다음과 같다. 

cpu_mem = cuda_mem.copy_to_host()

다음은 커널 함수 내 코드 실행 시 참고사항이다. 커널 함수는 GPU 쓰레드 내에서 실행되므로, 제약사항이 많다. 

• 커널 함수 내에서 메모리를 동적 할당 가능(참고 - Memory management).
• 예외처리 지원 안됨. 일반적으로 Zero Divide 에러 발생 안됨 
• list, dic, set 사용 안됨
• print 사용 가능
• 재귀함수 지원
• abs(), bool, complex, enumerate(), float, int, len(), min(), max(), pow(), range(), round(), zip() 지원
• cmath, math 라이브러리 지원
• operator 연산자 add, eq, ge, sub, xor 등 함수 지원
• NumPy 일부 함수 지원(예. sin, cos, tan, deg2rad 등). 배열 생성 등 메모리 할당 관련 함수는 지원 안됨
• Fastmath, cuda.random, cuda.select_device 함수 등 지원
• 공유 메모리 락 등 처리 위한 atomic, thread, lock 함수 지원

커널 함수 디버깅을 위해서는 다음과 같은 pdb를 사용해야 한다. 

from pdb import set_trace; set_trace()

이외, 기존 쿠다 코드인 cu 루틴을 호출할 수 있는 기능을 지원한다. 

@cuda.jit(link=['functions.cu'])

def multiply_vectors(r, x, y):

    i = cuda.grid(1)

    if i < len(r):

        r[i] = mul(x[i], y[i])

병렬 처리 코딩해 보기

이제 모든 재료가 모였으니, 두 개 거대한 벡터 1차원 행렬을 만들어, 난수값을 할당하고, 이를 더하는 함수를 만들고, 이를 병렬처리해 보자. 다음을 코딩하고, 실행해 본다.

import numpy as np

from numba import cuda

from tqdm import tqdm

import time

@cuda.jit

def f(a, b, c):

    tid = cuda.grid(1)    # threadIdx.x + (blockIdx.x * blockDim.x)

    # if tid % 100 == 0:

    #     print('tid=', tid)

    size = len(c)

    if tid < size:

        c[tid] = a[tid] + b[tid]

N = 10000000

a = cuda.to_device(np.random.random(N))   # 천만개 난수 생성 및 a 배열 값 할당

b = cuda.to_device(np.random.random(N))   # b 배열 값 할당

c = cuda.device_array_like(a)                     # 계산된 값 전달받을 c 배열 생성

# CPU 모드 커널 함수 계산 성능 측정

start_time_cpu = time.time()               

f.forall(len(a))(a, b, c)

end_time_cpu = time.time()

cpu_execution_time = end_time_cpu - start_time_cpu

print(f'CPU len={len(a)}, ', c.copy_to_host())

# CUDA 모드 커널 함수 계산 성능 측정

start_time_cuda = time.time()

nthreads = 256  # Enough threads per block for several warps per block

nblocks = (len(a) // nthreads) + 1  # Enough blocks to cover the entire vector depending on its length

f[nblocks, nthreads](a, b, c)

end_time_cuda = time.time()

cuda_execution_time = end_time_cuda - start_time_cuda

print(f'CUDA len={len(a)}, ', c.copy_to_host())

print(f"CUDA threads: {nthreads}, blocks: {nblocks}")

print(f"CPU Execution Time: {cpu_execution_time:.6f} seconds")

print(f"CUDA Execution Time: {cuda_execution_time:.6f} seconds")

 

실행 결과는 다음과 같다. CUDA 병렬처리 계산속도가 CPU 모드 보다 4배 가까이 빠른 것을 볼 수 있다. 

마무리 

이 글에서 GPU CUDA 병렬처리를 지원하는 Numba 라이브러리를 간략히 소개하고 사용법을 정리해 보았다. Numba를 이용해, 3차원, 4차원 등 고차원 대용량 행렬을 그리드, 블럭, 쓰레드로 나눠, 병렬처리할 수 있다.  

모든 수치해석모델은 행렬 텐서로 계산되므로, 사실상, 수학적인 계산 모델은 CUDA에서 실행될 수 있다. 이를 통해, 행렬 텐서 계산량이 많은 딥러닝 뿐 아니라, 계산 기하학, 컴퓨터 그래픽스, 시뮬레이션 등을 효과적으로 계산할 수 있다. 

단, 병렬처리를 위해서는 공유되는 메모리(변수, 배열, 행렬 등)의 수정, 접근 사용 시 병렬성을 고려한 코딩 기법이 선행 되어야 한다(동시성, 원자성 데이터 Lock 처리 등. 참고 - CUDA Programming: A Developers Guide to Parallel Computing with GPUs). 아울러, CUDA 자체에서 지원되는 라이브러리 함수들을 제외하고는 커널 함수에서는 전혀 사용할 수 없으므로, 복잡한 라이브러리 함수 사용시 이를 CUDA 함수를 이용해 포팅하는 노력이 있어야 한다.

참고로, 다음 링크 확인해 보면, 좀 더 다양한 사용 예제를 확인할 수 있다.

Example (Laplace’s equation simulation example)

Numba Mission Statement

레퍼런스

• Applied Linear Algebra, Stanford
• NVIDIA/cuda-python: CUDA Python Low-level Bindings (github.com)
• Accelerating The Real-Time Ray Tracing Ecosystem: DXR For GeForce RTX and GeForce GTX | GeForce News | NVIDIA
• Multi-Process Service :: GPU Deployment and Management Documentation (nvidia.com)
• Examples — Numba documentation
• Introduction to Numba: CUDA Programming (nyu-cds.github.io)
• Microsoft, CPU programming
• CUDA Programming: A Developers Guide to Parallel Computing with GPUs