이제 개발자는 기존의 대규모 C++ GPU 코드베이스 안에서 NVIDIA CUDA Tile 프로그래밍을 활용해, 타일 기반 추상화로 고도로 최적화된 GPU 커널을 개발할 수 있습니다.
NVIDIA CUDA Tile은 NVIDIA CUDA 13.1과 함께 출시되며 GPU를 위한 타일 기반 프로그래밍을 선보였습니다. 최상위 언어 계층과, 어떤 고수준 프로그래밍 언어든 타깃으로 삼을 수 있는 또 하나의 중간 계층을 갖추도록 설계된 CUDA Tile은 텐서 코어, 공유 메모리, 텐서 메모리 액셀러레이터를 비롯한 NVIDIA 하드웨어의 고급 기능을 애플리케이션이 직접 겨냥하지 않고도 자동으로 활용합니다.
타일 기반 GPU 애플리케이션에서 가장 먼저 지원된 언어는 Python이었습니다. 새롭게 출시된 CUDA 13.3은 C++로 타일 커널을 작성하는 기능을 추가해, 개발자가 고도로 최적화된 GPU 커널을 구축할 수 있도록 합니다.
CUDA Tile C++란 무엇인가?
CUDA Tile C++는 CUDA Tile 프로그래밍 모델을 C++로 표현한 것으로, CUDA Tile IR 사양 위에 구축됩니다. 개발자는 C++로 타일 커널을 작성하고, SIMT(단일 명령 다중 스레드) 모델 대신, 또는 그와 더불어 타일 기반 모델로 GPU 커널을 표현할 수 있습니다.
복습하자면, 타일 모델은 다음과 같습니다.
- 다차원 배열이 기본 데이터 저장소입니다.
- 타일은 커널이 연산을 수행하는 배열의 일부입니다.
- 커널은 블록 단위로 병렬 실행되는 함수입니다.
- 블록은 GPU의 부분 집합이며, 타일에 대한 연산은 각 블록의 모든 스레드에 걸쳐 병렬화됩니다.
CUDA Tile C++는 블록 내부의 병렬성은 물론 비동기 처리, 메모리 이동, 그 밖의 GPU 프로그래밍 저수준 세부 사항을 자동으로 처리합니다. 또한 서로 다른 NVIDIA GPU 아키텍처 전반에 이식할 수 있어, 개발자는 코드를 다시 작성하지 않고도 최신 하드웨어 기능을 사용할 수 있습니다.
CUDA Tile C++ 벡터 덧셈 예제
SIMT용 CUDA C++에 익숙한 개발자라면 표준적인 벡터 덧셈 커널을 접해 봤을 것입니다. 데이터가 이미 GPU에 있다고 가정하면, 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];
}
}
이 커널에서는 각 스레드의 작업이 명시적으로 지정되며, 프로그래머는 이 커널을 실행할 때 실행할 블록과 스레드의 수를 지정합니다.
CUDA Tile C++로 작성한 동등한 코드를 보면, 각 스레드가 무엇을 하는지 지정할 필요가 없습니다. 데이터를 타일로 나누고 이 타일에 대한 수학적 연산만 지정하면, 나머지는 모두 처리됩니다.
CUDA Tile C++ 커널의 모습은 다음과 같습니다.
#include "cuda_tile.h"
__tile_global__ void vectorAdd(float* a, float* b, float* out, size_t n) {
/* set up the namespace */
namespace ct = cuda::tiles;
using namespace ct::literals;
/* attach shape to raw pointers */
auto aSpan = ct::tensor_span{a, ct::extents{n}};
auto bSpan = ct::tensor_span{b, ct::extents{n}};
auto oSpan = ct::tensor_span{out, ct::extents{n}};
/* partition each span into tiles of size 8 */
auto aView = ct::partition_view{aSpan, ct::shape{8_ic}};
auto bView = ct::partition_view{bSpan, ct::shape{8_ic}};
auto oView = ct::partition_view{oSpan, ct::shape{8_ic}};
/* load the a and b tiles from global memory */
int bx = ct::bid().x;
auto aTile = aView.load(bx); // load bx-th tile
auto bTile = bView.load(bx);
/* add the two tiles together, elementwise */
auto oTile = aTile + bTile;
/* store the result tile to the output partition. */
oView.store(oTile, bx);
}
단순한 vectorAdd 커널치고는 코드가 많아 보일 수 있지만, 놀라지 않아도 됩니다. 이처럼 다소 장황한 커널은 모든 단계를 순서대로 보여주기 위한 것입니다. 더 적은 코드로 같은 일을 하는 간소화된 버전은 다음과 같습니다.
- 첫 번째 차이는
__tile_global__을 사용해 이것이 타일 커널임을 컴파일러에 알린다는 점입니다. 배열 포인터와 배열 크기는 SIMT 커널에서와 마찬가지로 인자로 전달됩니다.
__tile_global__ void vectorAdd(float* a, float* b, float* out, std::size_t n) {
- 다음으로
cuda::tiles와ct::literals에 대한 네임스페이스를 설정합니다.
namespace ct = cuda::tiles;
using namespace ct::literals;
- 세 배열 각각에 대해
ct::tensor_span코드를 사용해 텐서 스팬(tensor span)을 만듭니다. 텐서 스팬은 본질적으로 메모리에 있는 다차원 배열을 가리키는 포인터로, C++23의std::mdspan과 유사합니다. 텐서 스팬은 배열의 형태(extents)와 배열 요소의 레이아웃(예: 행 우선 또는 열 우선)에 대한 정보를 함께 담습니다.ct::extents{}는 텐서 스팬에 배열의 차원이 무엇인지 알려줍니다. 1차원 배열은n을 사용합니다.
auto aSpan = ct::tensor_span{a, ct::extents{n}};
auto bSpan = ct::tensor_span{b, ct::extents{n}};
auto oSpan = ct::tensor_span{out, ct::extents{n}}
- 이제 텐서 스팬과 타일 형태로부터 파티션 뷰(partition view)를 만듭니다. 파티션 뷰는 텐서 스팬을 감싸는 래퍼로, 배열을 서로 겹치지 않는 고정 크기 파티션의 연속으로 표현합니다. 각 파티션의 크기는 shape 인자로 지정하며, 이는 반드시 컴파일 타임 인자여야 합니다. 이 예제에서
8_ic는ct::literals로 정의된 정수 상수입니다.ct::shape<8>{}와ct::shape{8_ic}는 이 맥락에서 동등합니다. 생성된 파티션 뷰는 본질적으로 원래 배열을 타일 크기인 8 단위의 청크로 나눈 것입니다.
auto aView = ct::partition_view{aSpan, ct::shape{8_ic}};
auto bView = ct::partition_view{bSpan, ct::shape{8_ic}};
auto oView = ct::partition_view{oSpan, ct::shape{8_ic}};
- X 차원의 블록 인덱스를
ct::bid().x로 얻어 입력 타일을 로드합니다. 다차원 블록을 다룬다면 Y와 Z 차원도 사용합니다. 그런 다음a와b타일을 로드합니다. 편의를 위해 auto를 사용하지만, 명시적으로 적으면aTile과bTile은ct::tile<float, ct::shape<8>>>타입입니다. 이들은 float 타입 요소를 갖는 크기 8의 1차원 타일입니다. 파티션 뷰를 사용하면 블록 인덱스를 손쉽게 전달할 수 있습니다. load 함수는 배열에서 올바른 청크를 자동으로 가져와 타일로 로드합니다.
int bx = ct::bid().x;
auto aTile = aView.load(bx);
auto bTile = bView.load(bx);
- 결과를 더하고 저장합니다. 이 한 줄의 코드는 입력 타일에 요소별 덧셈을 수행해 출력 타일에 저장합니다. 그 출력 타일을 X 차원의 동일한 블록 인덱스
bx로 인덱싱해oView파티션 뷰에 저장합니다.
/* add the two tiles together, elementwise. */
auto oTile = aTile + bTile;
/* store the result tile to the output partition. */
oView.store(oTile, bx);
완전한 벡터 덧셈 예제
다음 예제는 완전하고 실행 가능한 코드를 통해 이 벡터 덧셈 커널을 C++에서 어떻게 호출하는지 보여줍니다.
컴파일러가 최적화를 수행하도록 돕기 위해 몇 가지 유의할 점이 있습니다.
첫째, 최상의 성능을 위해 입력·출력 배열은 커널이 실행되는 동안 각각의 포인터로만 접근해야 합니다. 이 조건이 성립하면 배열에 대한 앨리어싱, 즉 다른 포인터나 심볼을 통한 접근이 발생하지 않습니다. 배열 포인터에 __restrict__ 데코레이터를 붙이면 이 사실을 컴파일러에 전달합니다.
기준 포인터가 16바이트 경계에 정렬된 배열을 사용하면 컴파일러가 더 효율적인 메모리 접근 패턴을 생성하는 데 도움이 됩니다. 각 커널 인자에 ct::assume_aligned<16>을 호출해 포인터가 정렬돼 있음을 컴파일러에 알립니다. 컴파일러가 이 정렬을 활용하도록 이 호출의 반환값을 사용합니다. cudaMalloc이나 유사한 CUDA API가 반환하는 포인터는 256바이트 정렬을 갖추므로 항상 이 조건을 충족합니다.
끝으로, 8보다 훨씬 큰 타일 크기를 사용합니다. 다음의 실행 가능한 코드에 이러한 조정을 적용하고, 타일 크기로 나누어떨어지지 않을 수 있는 데이터를 처리하는 load_masked와 store_masked의 사용을 추가합니다.
다음은 커널과 main 함수를 포함한 전체 코드입니다. 적용된 최적화와 줄어든 장황함에 주목하세요.
#include <cstdio>
#include <cstdlib>
#include "cuda_tile.h"
__tile_global__ void vectorAdd(float* __restrict__ a, float* __restrict__ b, float* __restrict__ out, size_t n) {
namespace ct = cuda::tiles;
using namespace ct::literals;
a = ct::assume_aligned(a, 16_ic);
b = ct::assume_aligned(b, 16_ic);
out = ct::assume_aligned(out, 16_ic);
int bx = ct::bid().x;
/* create partition views for the input tiles and load them */
auto aTile = ct::partition_view{ct::tensor_span{a, ct::extents{n}}, ct::shape{1024_ic}}.load_masked(bx);
auto bTile = ct::partition_view{ct::tensor_span{b, ct::extents{n}}, ct::shape{1024_ic}}.load_masked(bx);
/* add the two tiles together, elementwise. */
auto oTile = aTile + bTile;
/* create the partition view for the output tile and then store the output tile*/
auto oView = ct::partition_view{ct::tensor_span{out, ct::extents{n}}, ct::shape{1024_ic}};
oView.store_masked(oTile, bx);
}
/* define a macro to check for CUDA errors */
#define checkCudaError(X) do {\
auto ret = X;\
if (ret != cudaSuccess) {\
printf("\n error on line %d, CUDART error string : %s", __LINE__, cudaGetErrorString(ret));\
exit(1);\
}\
} while (0)
int main() {
constexpr size_t N = 2ULL << 25;
constexpr int TILE_SIZE = 1024;
constexpr int BLOCKS = (N + TILE_SIZE - 1) / TILE_SIZE;
/* declare and allocate the host arrays */
float* h_a = (float*)malloc(sizeof(float) * N);
float* h_b = (float*)malloc(sizeof(float) * N);
float* h_out = (float*)malloc(sizeof(float) * N);
/* initialize the host arrays */
for (size_t idx = 0; idx < N; ++idx) {
h_a[idx] = (float)rand() / RAND_MAX;
h_b[idx] = (float)rand() / RAND_MAX;
h_out[idx] = -1.0f;
}
/* declare the device arrays */
float* d_a{nullptr};
float* d_b{nullptr};
float* d_out{nullptr};
/* allocate the device arrays */
checkCudaError(cudaMalloc(&d_a, sizeof(float) * N));
checkCudaError(cudaMalloc(&d_b, sizeof(float) * N));
checkCudaError(cudaMalloc(&d_out, sizeof(float) * N));
/* copy the host arrays to the device arrays */
checkCudaError(cudaMemcpy(d_a, h_a, sizeof(float) * N, cudaMemcpyHostToDevice));
checkCudaError(cudaMemcpy(d_b, h_b, sizeof(float) * N, cudaMemcpyHostToDevice));
/* initialize the device output array to 0 */
checkCudaError(cudaMemset(d_out, -1, sizeof(float) * N));
/* launch the kernel */
vectorAdd<<<BLOCKS, 1>>>(d_a, d_b, d_out, N);
/* synchronize the device and check for errors */
checkCudaError(cudaDeviceSynchronize());
/* copy the device array out back to the host */
checkCudaError(cudaMemcpy(h_out, d_out, sizeof(float) * N, cudaMemcpyDeviceToHost));
/* compare the results to host results */
float max_err = 0.0f;
for (size_t idx = 0; idx < N; ++idx) {
float expected = h_a[idx] + h_b[idx];
max_err = fmaxf(max_err, fabsf(h_out[idx] - expected));
}
printf("N: %zu\n", N);
printf("Max error: %e\n", max_err);
checkCudaError(cudaFree(d_a));
checkCudaError(cudaFree(d_b));
checkCudaError(cudaFree(d_out));
free(h_a);
free(h_b);
free(h_out);
}
SIMT 커널 실행에 익숙하다면 과정은 비슷하지만 한 가지 구체적인 수정이 필요합니다. 이 커널은 다음과 같이 실행했습니다.
vectorAdd<<<BLOCKS, 1>>>(d_a, d_b, d_out, N);
타일 커널을 실행할 때 <<<>>>의 첫 번째 인자는 타일 블록의 수입니다(SIMT에서는 스레드 블록의 수에 해당합니다). 두 번째 인자는 반드시 1이어야 합니다. 커널 실행에 사용되는 스레드 수는 컴파일러가 결정하므로, 타일 커널을 실행할 때는 이 인자에 항상 1을 넣습니다.
NVIDIA Ampere 아키텍처 이상의 GPU에서 컴퓨트 능력 8.0으로 CUDA 13.3 이상을 실행하면, 이 명령들은 다음과 같은 출력을 만듭니다.
-arch sm_120 명령을 아키텍처에 맞게 조정하고, cuda_tile.h를 사용할 때는 -std=c++20을 포함하며, 타일 커널을 컴파일하려면 --enable-tile 옵션을 사용합니다.
$ nvcc -std=c++20 --enable-tile -arch sm_120 -o vectorAdd vectorAdd.cu
$ ./vectorAdd
N: 67108864
Max error: 0.000000e+00
이로써 첫 번째 CUDA Tile C++ 프로그램을 완성했습니다.
개발자 도구
타일 C++ 커널은 SIMT 커널과 동일한 방식으로 NVIDIA Nsight Compute로 프로파일링할 수 있습니다. 다음 명령은 Nsight Compute로 프로파일을 생성하는 방법을 보여줍니다.
$ ncu -o VecAddProfile --set detailed ./vectorAdd
그래픽 버전의 Nsight Compute로 프로파일을 생성하고 연 뒤에는 다음과 같이 진행합니다.
- 드롭다운 메뉴에서
vectorAdd커널을 선택합니다. - Details 탭을 선택합니다.
- Tile Statistics 보고 섹션을 펼칩니다.
그림 1은 Nsight Compute에서 생성한 프로파일을 보여줍니다.
Tile Statistics 보고 섹션에는 지정된 타일 블록 수, 컴파일러가 선택한 블록 크기, 그 밖의 타일 관련 정보가 포함된다는 점에 주목하세요.
소스 페이지 역시 CUDA C++ 커널과 마찬가지로 타일 커널과 소스 라인 수준의 성능 지표를 지원합니다.
행렬 곱셈
앞선 예제에서는 파티션 뷰의 로드·저장 세부 사항과 함께 vectorAdd를 살펴봤습니다. 이번 행렬 곱셈 예제는 매우 간단한 코드로 행렬 곱셈을 표현하는 방법을 보여줍니다.
이 커널은 MxK와 KxN 행렬을 곱해 MxN 행렬을 계산합니다. 이 커널에서 M=8, N=16이며, K는 8의 배수이기만 하면 가변일 수 있습니다. K=24로 설정합니다. 이처럼 매우 작은 크기는 개념 설명만을 위한 것입니다.
전체 커널과 함께 핵심 사항을 살펴봅니다.
#include "cuda_tile.h"
/* this kernel multiplies MxK and KxN matrices, where M=8 and N=16. K is variable but must be divisible by 8.*/
__tile_global__ void kernel(float* __restrict__ a, float* __restrict__ b, size_t length, float* __restrict__ c) {
namespace ct = cuda::tiles;
using namespace ct::literals;
a = ct::assume_aligned(a, 16_ic);
b = ct::assume_aligned(b, 16_ic);
c = ct::assume_aligned(c, 16_ic);
auto aShape = ct::extents{8_ic, length};
auto bShape = ct::extents{length, 16_ic};
auto cShape = ct::extents{8_ic, 16_ic};
auto aSpan = ct::tensor_span{a, aShape};
auto bSpan = ct::tensor_span{b, bShape};
auto cSpan = ct::tensor_span{c, cShape};
auto aView = ct::partition_view{aSpan, ct::shape{4_ic, 8_ic}};
auto bView = ct::partition_view{bSpan, ct::shape{8_ic, 4_ic}};
auto cView = ct::partition_view{cSpan, ct::shape{4_ic, 4_ic}};
using f32x4x4 = ct::tile<float, ct::shape<4, 4>>;
auto accTile = ct::full<f32x4x4>(0);
auto [xBlock, yBlock, dummy] = ct::bid();
for (auto idx : ct::irange(0, 1 + int(length - 1) / 8)) {
auto aTile = aView.load_masked(xBlock, idx);
auto bTile = bView.load_masked(idx, yBlock);
accTile = ct::mma(aTile, bTile, accTile);
}
cView.store_masked(accTile, xBlock, yBlock);
}
a, b, c행렬에 대해ct::extents객체로 extents를 만듭니다. 컴파일 타임 값이나 런타임 값을 사용할 수 있습니다. M=8, N=16이지만 K는 가변입니다. 이들은 다음 단계에서 텐서 스팬을 만드는 데 사용됩니다.
auto aShape = ct::extents{8_ic, length};
auto bShape = ct::extents{length, 16_ic};
auto cShape = ct::extents{8_ic, 16_ic};
- 텐서 스팬을 만듭니다. 이는
a,b,c의 정보를 담아 파티션 뷰를 생성합니다.
auto aSpan = ct::tensor_span{a, aShape};
auto bSpan = ct::tensor_span{b, bShape};
auto cSpan = ct::tensor_span{c, cShape};
a,b,c의 파티션 뷰를 만들되,a는 4×8로, 뷰b는 8×4로 분할합니다.a와b값으로 적절히 나누어떨어지기만 하면 조정할 수 있습니다. 이 차원들에 따라c뷰는 4×4로 결정됩니다.
auto aView = ct::partition_view{aSpan, ct::shape{4_ic, 8_ic}};
auto bView = ct::partition_view{bSpan, ct::shape{8_ic, 4_ic}};
auto cView = ct::partition_view{cSpan, ct::shape{4_ic, 4_ic}};
2차원 파티션은 2개 차원으로 인덱싱됩니다. a 행렬은 8×24이고, 파티션 뷰는 그림 2와 같이 4×8입니다.
aView와 bView의 파티션 뷰 크기는 행렬 곱셈 중 결과를 누적하는 데 사용하는 타일인 accTile의 형태도 결정합니다. 이 예제에서 accTile은 cView의 형태와 일치하는 4×4 타일입니다.
using f32x4x4 = ct::tile<float, ct::shape<4, 4>>;
auto accTile = ct::full<f32x4x4>(0);
ct::bid로 세 차원의 블록 인덱스를 얻어 루프를 실행합니다. 루프는 0부터 length / 8까지 반복하며, 이는 전체 K 차원을 8로 나눈 값에 해당합니다. 8로 나누는 것은aView와bView의 K 차원이 8인 것과 일치합니다. 루프 내부에서a와b의 타일을load_masked로 로드하고,ct::mma호출이 행렬 곱셈을 수행해 결과를accTile에 누적합니다.
auto [xBlock, yBlock, dummy] = ct::bid();
for (auto idx : ct::irange(0, int(length / 8))) {
auto aTile = aView.load_masked(xBlock, idx);
auto bTile = bView.load_masked(idx, yBlock);
accTile = ct::mma(aTile, bTile, accTile);
}
accTile의 값을c의 파티션 뷰인cView에 저장합니다. 이것으로 끝입니다. 커널 코드의 대부분은 데이터를 위한 뷰를 설정하고 데이터를 로드·저장하는 일에 쓰입니다. 커널의 연산 부분은 단순합니다.
cView.store_masked(accTile, xBlock, yBlock);
- 커널을 실행합니다.
cView의 차원 때문에dim3(2,4)를 사용합니다.
cView는 4×4이므로, 각 블록은 C 행렬의 4×4 청크를 계산합니다. C가 8×16이므로 cView 차원으로 C 행렬 차원을 나눕니다. 8/4=2이고 16/4=4이므로 dim3(2,4)로 커널을 실행합니다.
kernel<<<dim3(2, 4), 1>>>(d_a, d_b, K, d_c);
지금 CUDA Tile C++로 시작하기
CUDA Tile C++ 프로그램을 실행하려면 다음이 필요합니다.
- 컴퓨트 능력 8.x 이상의 GPU.
- NVIDIA Driver R580 이상. 타일 커널에 JIT 컴파일이 필요한 경우, NVIDIA 드라이버 버전은 코드 생성에 사용한 CUDA Toolkit과 연관된 버전과 같거나 더 새로워야 합니다. 예를 들어 CUDA Toolkit 13.3은 R610 이상의 드라이버를 요구합니다.
- CUDA Toolkit 13.3
이제 타일 기반 프로그래밍의 힘을 C++ 개발자도 사용할 수 있습니다. 지금 문서와 API 레퍼런스 매뉴얼, CUDA Toolkit 13.3을 확인해 타일 C++ 커널 작성을 시작하고, 가속 컴퓨팅의 새로운 표준을 경험해 보세요.
감사의 말
NVIDIA 기여자 Jaydeep Marathe와 Ezra Stein에게 감사드립니다.
