파이썬에서 NVIDIA CUDA Tile로 GPU 프로그래밍 간소화

NVIDIA CUDA 13.1은 GPU에 타일 기반 프로그래밍을 공식 도입하며, CUDA 탄생 이후 가장 큰 변화 중 하나를 만들어냈습니다. 타일 기반 커널을 사용하면 기존 단일 명령어 다중 스레드(SIMT) 방식보다 훨씬 높은 수준에서 알고리즘을 작성할 수 있고, 스레드 단위 작업 분할은 컴파일러와 런타임이 자동으로 처리합니다. 이 방식은 텐서 코어 같은 특수 하드웨어에 대한 복잡성을 감추고, 앞으로 등장할 GPU 아키텍처에서도 잘 동작하는 코드를 작성할 수 있도록 도와줍니다. 여기에 더해 NVIDIA cuTile Python이 새롭게 공개되면서, 이제 Python에서도 타일 기반 커널을 손쉽게 작성할 수 있게 되었습니다.

cuTile Python이란 무엇인가요?

cuTile Python은 CUDA Tile IR 사양을 기반으로 구축된 Python에서의 CUDA Tile 프로그래밍 모델입니다. 이를 통해 Python으로 타일 커널을 작성하고, 단일 명령어 다중 스레드(SIMT) 모델 대신 또는 그 외에 타일 기반 모델을 사용하여 GPU 커널을 구현할 수 있습니다.

기존 SIMT 프로그래밍에서는 GPU의 각 스레드가 수행할 작업을 명시적으로 지정해야 합니다. 이론적으로는 스레드가 서로 다른 코드 경로를 독립적으로 실행할 수 있지만, 실제로는 GPU 하드웨어 효율을 높이기 위해 대부분의 알고리즘이 각 스레드가 동일한 연산을 서로 다른 데이터 조각에 수행하는 방식으로 작성됩니다.

SIMT는 높은 유연성과 세밀한 제어를 제공하지만 최적의 성능을 위해 많은 수작업 튜닝이 필요할 수 있습니다. 반면 타일 기반 모델은 이러한 하드웨어 복잡성을 상당 부분 감춰줍니다. 개발자는 더 높은 수준의 알고리즘 설계에 집중할 수 있고, CUDA 컴파일러와 런타임이 타일 연산을 스레드 단위로 자동 분할해 GPU에서 실행합니다.

cuTile은 NVIDIA GPU를 위한 병렬 커널 작성을 위한 프로그래밍 모델입니다. 이 모델에서:

• 배열은 기본 데이터 구조입니다.
• 타일은 커널이 작동하는 배열의 하위 집합입니다.
• 커널은 블록에 의해 병렬로 실행되는 함수입니다.
• 블록은 GPU의 하위 집합이며, 타일에 대한 연산은 각 블록에서 병렬화됩니다.

cuTile은 블록 단위의 병렬 처리, 비동기 실행, 메모리 이동 등 GPU 프로그래밍에서 요구되는 저수준 작업들을 자동으로 처리합니다. 덕분에 텐서 코어, 공유 메모리, 텐서 메모리 가속기 같은 NVIDIA 하드웨어의 고급 기능도 별도의 코드 작성 없이 활용할 수 있습니다. 또한 cuTile은 다양한 NVIDIA GPU 아키텍처에서 높은 이식성을 제공해, 코드를 다시 손보지 않아도 최신 하드웨어 기능을 그대로 사용할 수 있습니다.

cuTile은 누구를 위한 것인가요?

cuTile은 범용 데이터 병렬 GPU 커널을 작성하는 개발자를 위한 모델이며 특히 AI/ML 워크로드에서 자주 등장하는 계산을 최적화하는 데 중점을 두고 개발되고 있습니다. 앞으로도 기능과 성능을 확장해 더 폭넓은 워크로드를 최적화할 수 있는 워크로드의 범위를 넓힐 예정입니다.

CUDA C++ 또는 CUDA Python이 지금까지 잘 작동해 왔는데 왜 cuTile을 사용하여 커널을 작성해야 하는지 궁금할 수 있습니다. CUDA 타일 모델을 설명하는 다른 게시물에서 이에 대해 더 자세히 이야기하고 있지만 핵심은, GPU 하드웨어가 더 복잡해짐에 따라, 개발자가 알고리즘에 더 집중하고 특정 하드웨어에 알고리즘을 매핑하는 데 덜 집중할 수 있도록 합리적인 수준의 추상화 계층을 제공하는 데 있습니다.

타일 프로그램은 향후 GPU 아키텍처와도 호환되며 텐서 코어를 효과적으로 활용할 수 있습니다 병렬 스레드 실행(PTX)이 GPU 프로그래밍을 위한 SIMT 모델의 기초가 되는 가상 명령어 세트 아키텍처(ISA)를 제공하는 것처럼, Tile IR은 타일 기반 프로그래밍을 위한 가상 ISA를 제공합니다. 이를 통해 더 높은 수준의 알고리즘 표현이 가능하며, 소프트웨어 및 하드웨어는 해당 표현을 투명하게 텐서 코어에 매핑하여 최고 성능을 제공합니다.

cuTile Python 예제

cuTile Python 코드는 어떤 모습일까요? CUDA C++를 경험해본 적 있다면 가장 기본적인 예제로 자주 등장하는 벡터 덧셈 커널을 기억하실 겁니다. 데이터가 호스트에서 디바이스로 이미 복사되었다고 가정하면, CUDA SIMT 방식의 벡터 덧셈 커널은 두 벡터의 각 요소를 더해 세 번째 벡터를 만드는 아래와 같은 형태를 띱니다. 가장 단순한 형태의 CUDA 커널 중 하나입니다.

__global__ void vecAdd(float* A, float* B, float* C, int vectorLength)
{
 /* calculate my thread index */
 int workIndex = threadIdx.x + blockIdx.x*blockDim.x;
 
 if(workIndex < vectorLength)
 {
  /* perform the vector addition */
  C[workIndex] = A[workIndex] + B[workIndex];
 }
}

이 커널에서는 스레드가 수행할 작업을 명시적으로 지정하고, 커널을 실행할 때 블록 수와 스레드 수를 프로그래머가 직접 결정해야 합니다.

이제 cuTile Python으로 작성된 동등한 코드를 살펴보겠습니다. 각 스레드가 수행하는 작업을 지정할 필요가 없습니다. 데이터를 타일로 분할하고 각 타일에 대한 수학적 연산을 지정하기만 하면 됩니다. 다른 모든 것은 자동으로 처리됩니다.

cuTile Python 커널은 다음과 같습니다:

import cuda.tile as ct
 
@ct.kernel
def vector_add(a, b, c, tile_size: ct.Constant[int]):
  # Get the 1D pid
  pid = ct.bid(0)
 
  # Load input tiles
  a_tile = ct.load(a, index=(pid,) , shape=(tile_size, ) )
  b_tile = ct.load(b, index=(pid,) , shape=(tile_size, ) )
 
  # Perform elementwise addition
  result = a_tile + b_tile
 
  # Store result
  ct.store(c, index=(pid, ), tile=result)

ct.bid(0)은 (이 경우) 0축을 따라 블록 ID를 얻는 함수입니다. 이는 SIMT 커널 작성자가 예를 들어 blockIdx.x와 threadIdx.x를 참조하는 방식과 동등합니다. ct.load()는 필수적인 인덱스와 모양을 가지고 디바이스 메모리에서 데이터 타일을 로드하는 함수입니다. 데이터가 타일로 로드되면, 이 타일은 계산에 사용될 수 있습니다. 모든 계산이 완료되면, ct.store()는 타일된 데이터를 GPU 디바이스 메모리에 다시 저장합니다.

종합하기

이제 이 vector_add 커널을 완전한 Python 스크립트를 사용하여 Python에서 호출하는 방법을 소개합니다. 아래 코드는 커널과 메인 함수를 포함한 전체 코드입니다.

"""
Example demonstrating simple vector addition.
Shows how to perform elementwise operations on vectors.
"""
 
from math import ceil
 
import cupy as cp
import numpy as np
import cuda.tile as ct
 
 
@ct.kernel
def vector_add(a, b, c, tile_size: ct.Constant[int]):
  # Get the 1D pid
  pid = ct.bid(0)
 
  # Load input tiles
  a_tile = ct.load(a, index=(pid,) , shape=(tile_size, ) )
  b_tile = ct.load(b, index=(pid,) , shape=(tile_size, ) )
 
  # Perform elementwise addition
  result = a_tile + b_tile
 
  # Store result
  ct.store(c, index=(pid, ), tile=result)
 
def test():
  # Create input data
  vector_size = 2**12
  tile_size = 2**4
  grid = (ceil(vector_size / tile_size),1,1)
 
  a = cp.random.uniform(-1, 1, vector_size)
  b = cp.random.uniform(-1, 1, vector_size)
  c = cp.zeros_like(a)
 
  # Launch kernel
  ct.launch(cp.cuda.get_current_stream(),
       grid, # 1D grid of processors
       vector_add,
       (a, b, c, tile_size))
 
  # Copy to host only to compare
  a_np = cp.asnumpy(a)
  b_np = cp.asnumpy(b)
  c_np = cp.asnumpy(c)
 
  # Verify results
  expected = a_np + b_np
  np.testing.assert_array_almost_equal(c_np, expected)
 
  print("✓ vector_add_example passed!")
 
if __name__ == "__main__":
  test()

cuTile Python과 CuPy를 포함하여 필요한 모든 소프트웨어를 이미 설치했다고 가정하면, 이 코드를 실행하는 것은 Python을 호출하는 것만큼 간단합니다.

$ python3 VectorAdd_quickstart.py
✓ vector_add_example passed!

축하합니다. 첫 번째 cuTile Python 프로그램을 실행하셨습니다!

개발자 도구

cuTile 커널은 SIMT 커널과 동일한 방식으로 NVIDIA Nsight Compute로 프로파일링할 수 있습니다.

$ ncu -o VecAddProfile --set detailed python3 VectorAdd_quickstart.py

프로필을 생성하고 Nsight Compute의 그래픽 버전으로 열었으면:

• vector_add 커널을 선택합니다.
• “Details” 탭을 선택합니다.
• “Tile Statistics” 보고서 섹션을 확장합니다.

그림 1과 유사한 이미지를 볼 수 있습니다.

타일 통계 보고서 섹션에는 지정된 타일 블록 수, 블록 크기(컴파일러가 선택), 그리고 기타 다양한 타일별 정보가 포함되어 있음을 확인하셔야 합니다.

소스 페이지는 또한 CUDA C 커널과 마찬가지로 소스 라인 수준에서 cuTile 커널 및 성능 메트릭을 지원합니다.

개발자는 cuTile을 어떻게 얻을 수 있나요?

cuTile Python 프로그램을 실행하려면 다음이 필요합니다:

• 컴퓨팅 능력 10.x 또는 12.x의 GPU (향후 CUDA 릴리스에서는 추가 GPU 아키텍처에 대한 지원을 추가할 예정입니다)
• NVIDIA 드라이버 R580 이상 (R590은 타일별 개발자 도구 지원에 필요합니다)
• CUDA Toolkit 13.1 이상
• Python 버전 3.10 이상
• cuTile Python 패키지: pip install cuda-tile

시작하기

더 자세히 알아보는 데 도움이 되는 몇 가지 비디오를 확인하세요:

• What is CUDA Tile?
• The Future is Tiled
• Deep Dive: How to Use cuTile Python

또한, cuTile Python 문서도 확인하세요.

이제 GitHub의 샘플 프로그램을 시도하고 cuTile Python으로 프로그래밍을 시작해보세요!