참고: BASIC을 이용한 CUDA 타일 프로그래밍은 만우절 농담이지만, 실제로 작동하며 CUDA의 유연성을 보여주는 실제 프로젝트이기도 합니다.
CUDA 13.1에서 도입된 CUDA 타일은 세밀한 병렬 처리를 더욱 쉽고 유연하게 만들기 위해 설계된 차세대 타일 기반 GPU 프로그래밍 패러다임입니다. 이 모델의 핵심 강점 중 하나는 언어의 개방성입니다. 어떤 프로그래밍 언어라도 CUDA 타일을 대상으로 삼을 수 있어, 개발자들은 다양한 에코시스템에 타일 기반 GPU 가속을 도입할 수 있습니다.
오랫동안 소외되었던 언어에 대한 숙련된 개발자들의 압도적인 요청에 부응하여, NVIDIA는 GPU용 cuTile BASIC을 출시했습니다. 이제 BASIC 언어에서도 CUDA 타일 프로그래밍을 만나보실 수 있습니다.
cuTile BASIC이란 무엇인가요?
cuTile BASIC은 CUDA 타일 IR 사양을 기반으로 구축된, BASIC 언어로 표현된 CUDA 타일 프로그래밍 모델입니다. 멀티스레드 프로그래밍 개념이 생기기도 전에 탄생한 BASIC 언어의 특성상, 타일 기반 모델은 놀랍도록 자연스럽게 어우러집니다.
cuTile BASIC은 GPU의 강력한 성능과 BASIC 언어의 시대착오적인 매력 및 구문론적 단순함이 완벽하게 결합된 결과물입니다. 픽셀이 도드라지던 시절의 우아한 언어인 BASIC을 사용하여, 코드 라인 번호를 수동으로 매기는 즐거움과 상상할 수 없었던 빠른 실행 속도를 동시에 경험해 보세요!
cuTile BASIC은 누구를 위한 것인가요?
BASIC은 가장 오래된 프로그래밍 언어 중 하나로, 300바우드 모뎀의 핸드셰이킹 소리를 그리워하는 세대의 개발자들에게 경외의 대상입니다. 많은 개발자에게 BASIC은 컴퓨터 프로그래밍의 첫 입문서와 같았습니다.
이제 뇌리에 BASIC이 각인된 개발자들은 사상 처음으로 레거시 애플리케이션을 NVIDIA GPU 가속 컴퓨팅으로 이식할 수 있게 되었습니다. 이는 BASIC 언어로는 상상조차 할 수 없었던 성능과 기능을 열어주어, 여러분의 루나 랜더(Lunar Lander) 게임이 아르테미스 미션보다 더 빠르게 달 표면을 누빌 수 있게 해줍니다.
수학 시간 몰래 공학용 계산기로 게임을 코딩하던 영광의 시절을 세계에서 가장 강력한 GPU 위에서 다시 재현해 보세요.
설치하기
먼저 PIP를 사용하여 cuTile BASIC을 설치합니다.
pip install git+https://github.com/nvidia/cuda-tile.git@basic-experimentalcuTile BASIC을 실행하기 위한 전체 하드웨어 및 소프트웨어 요구 사항은 이 게시물 끝에 나열되어 있습니다. (64KB 이상의 RAM 권장)
cuTile BASIC 예제
CUDA C++를 배워보셨다면 정형화된 벡터 덧셈 커널을 접해보셨을 것입니다. CUDA C++에서의 벡터 덧셈 커널은 두 벡터를 가져와 요소별로 더해 세 번째 벡터를 생성하며, 대략 다음과 같은 모습입니다.
이는 작성할 수 있는 가장 간단한 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 BASIC으로 작성된 동일한 코드를 살펴보겠습니다. 여기서는 각 스레드가 무엇을 하는지 지정할 필요가 없습니다. 데이터를 타일로 나누고 이 타일들에 어떤 수학적 연산이 일어나야 하는지만 지정하면 됩니다. 그 외의 모든 것은 시스템이 우리 대신 처리해 줍니다.
cuTile BASIC 벡터 덧셈 커널은 아래와 같습니다.
10 REM Vector Add: C = A + B
20 INPUT N, A(), B()
30 DIM A(N), B(N), C(N)
40 TILE A(128), B(128), C(128)
50 LET C(BID) = A(BID) + B(BID)
60 OUTPUT C
70 END이 예제는 매우 기초적이며, 표준 BASIC에 다음과 같은 세 가지 추가 사항이 포함되어 있습니다.
- 배열을 인덱싱하면 배열의 하위 집합인 타일이 반환됩니다.
- BID는 타일 블록 인덱스를 지정하는 내장 변수입니다.
- TILE은 배열을 어떤 크기의 타일로 분할할지 지정합니다.
덧셈 연산 외에 다른 것은 지정할 필요가 없었다는 점에 주목해 주세요. 그 외의 모든 것은 cuTile BASIC에 의해 처리됩니다.
실행 방법 요약
이제 이 벡터 덧셈 커널을 BASIC에서 실행하는 방법을 보여드리겠습니다. 기본적인 워크플로는 먼저 BASIC 함수를 cubin 파일로 컴파일한 다음 GPU에서 실행하는 방식입니다. 간단히 설명하기 위해 지루한 파이썬 호스트 및 래퍼 코드는 생략했지만, GitHub 리포지토리에서 해당 내용을 확인할 수 있습니다.
적절한 버전의 CUDA 툴킷과 파이썬이 설치되어 있고 GitHub에서 cuTile BASIC 저장소를 다운로드했다면 다음 명령을 실행할 수 있습니다.
$ python examples/vector_add.py
[1/2] Compiling to cubin ...
Arrays: ['A', 'B', 'C'], tile_shapes={'A': [128], 'B': [128], 'C': [128]}, grid_size=8
[2/2] Launching kernel on GPU ...
Results (showing 5 samples of 1024):
C[ 0] = 0.0 (expected 0.0)
C[ 1] = 3.0 (expected 3.0)
C[ 511] = 1533.0 (expected 1533.0)
C[ 512] = 1536.0 (expected 1536.0)
C[1023] = 3069.0 (expected 3069.0)
VERIFICATION PASSED (max_diff=0.000000, 1024 elements)출력 결과가 동일하다면 축하합니다! 방금 여러분의 첫 번째 cuTile BASIC 프로그램을 실행하셨습니다. 어쩌면 BASIC으로 작성된 생애 첫 프로그램일지도 모르겠네요! 맥스 헤드룸(Max Headroom)이 자랑스러워할 것입니다.
BASIC 행렬 곱셈
BASIC은 아주 적은 줄의 코드로 일반적인 알고리즘을 표현할 수 있을 만큼 단순한 언어입니다. 아래에 표시된 BASIC 기반의 행렬 곱셈(GEMM) 커널을 살펴보십시오.
10 REM GEMM: C(M,N) = A(M,K) * B(K,N)
15 INPUT M, N, K, A(), B()
20 DIM A(M, K), B(K, N), C(M, N)
30 TILE A(128, 32), B(32, 128), C(128, 128), ACC(128, 128)
40 LET TILEM = INT(BID / INT(N / 128))
50 LET TILEN = BID MOD INT(N / 128)
60 LET ACC = 0.0
70 FOR KI = 0 TO INT(K / 32) - 1
80 LET ACC = MMA(A(TILEM, KI), B(KI, TILEN), ACC)
90 NEXT KI
100 LET C(TILEM, TILEN) = ACC
110 OUTPUT C
120 END이 커널에서는 표준 BASIC 구문 외에도, TILE이 A, B, C를 어떻게 타일링할지와 누산기 타일인 ACC의 크기를 지정합니다. MMA는 행렬 곱셈 및 누산을 위한 함수 호출입니다. 이 코드가 얼마나 단순한지 주목해 보십시오. 데이터를 타일로 나누는 방법을 지정하고 알고리즘을 상위 수준에서 정의하기만 하면, 보이지 않는 곳에서 CUDA Tile이 그 외의 모든 것을 처리합니다.
이 예제는 GitHub 저장소의 예제 폴더에서도 확인할 수 있습니다. 이를 실행하면 다음과 같은 출력이 생성됩니다.
$ python examples/gemm.py
[1/2] Compiling to cubin ...
M=512, N=512, K=512, tile_shapes={'A': [128, 32], 'B': [32, 128], 'C': [128, 128]}, grid_size=16
[2/2] Launching kernel on GPU ...
Results (showing 5 samples of 512x512 = 262144 elements):
C[0,0] = -0.1199 (expected -0.1199)
C[0,1] = -14.4456 (expected -14.4456)
C[256,0] = -15.8891 (expected -15.8891)
C[256,1] = -2.8646 (expected -2.8646)
C[511,511] = 11.4724 (expected 11.4724)
VERIFICATION PASSED (max_diff=0.000012, tol=0.005120)위에서 보여드린 것과 같은 행렬 곱셈은 거대 언어 모델과 같은 인공지능 도구의 핵심입니다. cuTile BASIC을 통해 개발자들은 이제 1MB의 시스템 메모리조차 상상하기 어려웠던 언어를 사용하여, 수조 개의 파라미터를 가진 모델이 활약하는 인공지능의 최전선을 탐험할 수 있습니다.
개발자가 cuTile을 얻는 방법
cuTile BASIC 프로그램을 실행하려면 다음 사항이 필요합니다:
- 컴퓨팅 성능 8.x, 10.x, 11.x 또는 12.x를 갖춘 GPU (향후 CUDA 릴리스에서 추가적인 GPU 아키텍처 지원이 추가될 예정입니다)
- NVIDIA 드라이버 R580 이상 (타일 전용 개발자 도구 지원을 위해서는 R590이 필요합니다)
- CUDA 툴킷 13.1 이상
- 파이썬 3.10 버전 이상
- cuTile BASIC 패키지
시작하기 모든 소프트웨어를 준비했다면 cuTile BASIC의 전체 문서를 확인하고 GitHub에 있는 모든 샘플 프로그램을 실행해 보며 지금 바로 cuTile BASIC 프로그래밍을 시작해 보십시오. 여러분의 현대적인 AI 또는 과학 계산 코드베이스를 역사적으로 중요한 언어로 포팅하면서도, 현존하는 가장 강력한 하드웨어에서 실행할 수 있는 기회를 즐기시기 바랍니다! 다만 여러분의 코모도 64에게는 이 사실을 알리지 마세요.
모든 언어에서의 CUDA Tile
BASIC이 고성능 병렬 컴퓨팅을 위해 개발자들이 가장 먼저 떠올리는 언어는 아닐 수 있지만, 이는 CUDA 소프트웨어 스택의 설계 덕분에 CUDA Tile이 거의 모든 프로그래밍 언어에서 사용될 수 있음을 보여주는 유익한 시연입니다. CUDA Tile IR 형식으로 컴파일함으로써, CUDA Tile은 거의 모든 언어에 적용될 수 있습니다… 심지어 BASIC에서도 말이죠!
편집자 주: 돌이켜보면, CUDA Tile 프로그래밍 모델에 대한 폭넓은 지원을 요청했던 개발자들은 아마 조금 더 구체적으로 요구했어야 했을지도 모릅니다. 2027년 4월 1일에 출시될 cuTile COBOL을 기대해 주세요.
