NVIDIA CUDA Tile C++, 드디어 C++ 개발자도 쓸 수 있다 GPU 커널 최적화의 새로운 패러다임

왜 지금 CUDA Tile C++인가?

GPU 프로그래밍을 해본 개발자라면 누구나 SIMT(Single Instruction, Multiple Threads) 모델에 익숙할 겁니다. 각 스레드가 수행할 작업을 일일이 지정하고, 블록과 그리드 크기를 수동으로 조절해야 했죠. 하지만 CUDA 13.3에서 공개된 **CUDA Tile C++**은 이러한 패러다임을 완전히 뒤집습니다.

타일(Tile) 기반 프로그래밍은 다차원 배열을 기본 데이터 단위로 삼고, 커널이 이 배열의 일부(타일) 단위로 연산을 수행하도록 추상화합니다. 개발자는 '스레드가 뭘 해야 하는지'가 아니라 '데이터를 어떻게 자르고 연산할지'에 집중하면 됩니다. 나머지(병렬화, 비동기 처리, 메모리 이동)는 CUDA Tile 컴파일러와 런타임이 알아서 처리합니다.

이미 Python에서는 CUDA 13.1부터 지원되었지만, C++ 생태계에서도 타일 기반 프로그래밍이 가능해졌다는 점이 이번 업데이트의 핵심입니다. 기존 대규모 C++ GPU 코드베이스를 유지하면서도, 최신 NVIDIA 하드웨어(Tensor Core, Shared Memory, TMA)의 성능을 직접 타겟팅하지 않고도 활용할 수 있게 되었습니다.

참고 자료: 이 글의 상세 내용은 NVIDIA 공식 블로그 포스트를 근거로 작성되었습니다.

핵심 개념: 벡터 덧셈으로 보는 CUDA Tile C++

가장 기본적인 GPU 커널인 벡터 덧셈(Vector Add)을 CUDA Tile C++로 작성해보겠습니다. 기존 SIMT 방식과의 차이를 바로 체감할 수 있습니다.

기존 SIMT 방식 (CUDA C++)

__global__ void vecAdd(float* A, float* B, float* C, int vectorLength) {
    /* 내 스레드 인덱스 계산 */
    int workIndex = threadIdx.x + blockIdx.x * blockDim.x;
    if (workIndex < vectorLength) {
        /* 벡터 덧셈 수행 */
        C[workIndex] = A[workIndex] + B[workIndex];
    }
}

CUDA Tile C++ 방식

#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;

    // 16바이트 정렬 보장 (성능 최적화)
    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;   // 블록 인덱스 획득

    // 입력 배열을 1024개 요소 단위 타일로 분할 후 로드
    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);

    // 요소별 덧셈
    auto oTile = aTile + bTile;

    // 결과 저장
    auto oView = ct::partition_view{
        ct::tensor_span{out, ct::extents{n}},
        ct::shape{1024_ic}
    };
    oView.store_masked(oTile, bx);
}

핵심 차이점:

• __tile_global__ 키워드로 타일 커널임을 명시
• ct::tensor_span으로 배열을 다차원 뷰로 래핑 (C++23 std::mdspan과 유사)
• ct::partition_view로 배열을 타일 크기(여기서는 1024)로 자름
• load_masked / store_masked로 경계 조건(배열 크기가 타일 크기로 나누어떨어지지 않을 때)을 자동 처리
• 스레드 인덱스 계산 코드가 완전히 사라짐

메인 함수에서 커널 실행

int main() {
    constexpr size_t N = 2ULL << 25;   // 67,108,864개 요소
    constexpr int TILE_SIZE = 1024;
    constexpr int BLOCKS = (N + TILE_SIZE - 1) / TILE_SIZE;

    // ... (호스트 배열 할당 및 초기화, 디바이스 메모리 할당 및 복사 생략) ...

    // 타일 커널 실행: 첫 번째 인수 = 타일 블록 수, 두 번째 인수 = 반드시 1
    vectorAdd<<<BLOCKS, 1>>>(d_a, d_b, d_out, N);

    // ... (결과 검증, 메모리 해제) ...
}

주의: 타일 커널 실행 시 <<<BLOCKS, 1>>>에서 두 번째 인자는 반드시 1이어야 합니다. 실제 스레드 수는 컴파일러가 결정합니다.

컴파일 명령어:

nvcc -std=c++20 --enable-tile -arch sm_120 -o vectorAdd vectorAdd.cu

실행 결과:

N: 67108864
Max error: 0.000000e+00

심화: 행렬 곱셈과 Tensor Core 활용

벡터 덧셈이 입문 예제였다면, 행렬 곱셈(Matrix Multiply)은 CUDA Tile C++의 진가를 보여줍니다. 특히 ct::mma 함수를 통해 Tensor Core를 자동 활용할 수 있다는 점이 압권입니다.

#include "cuda_tile.h"

// 8x24 * 24x16 = 8x16 행렬 곱셈 (K=24, 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);

    // 행렬 shape 정의
    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};

    // 파티션 뷰: a는 4x8 타일, b는 8x4 타일 -> c는 4x4 타일
    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}};

    // 누산기 타일 (4x4, 0으로 초기화)
    using f32x4x4 = ct::tile<float, ct::extents{4_ic, 4_ic}>>;
    auto accTile = ct::full<float>(0);

    auto [xBlock, yBlock, dummy] = ct::bid();

    // K dimension을 8씩 스트라이드하며 누적 곱셈
    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);  // Tensor Core 자동 활용!
    }

    cView.store_masked(accTile, xBlock, yBlock);
}

이 코드에서 ct::mma가 수행하는 행렬 곱셈-누산(Matrix Multiply-Accumulate) 연산은 NVIDIA Tensor Core의 핵심 기능입니다. 기존에는 wmma 네임스페이스나 PTX 어셈블리로 직접 제어해야 했지만, CUDA Tile C++에서는 단 한 줄의 코드로 자동 매핑됩니다.

Nsight Compute 프로파일링

CUDA Tile C++ 커널도 기존 SIMT 커널과 동일하게 ncu로 프로파일링할 수 있습니다.

ncu -o VecAddProfile --set detailed ./vectorAdd

Nsight Compute GUI에서 'Tile Statistics' 섹션을 열면 타일 블록 수, 컴파일러가 선택한 블록 크기 등 타일 특화 정보를 확인할 수 있습니다. 소스 레벨 성능 메트릭도 그대로 지원됩니다.

실무 적용 조언과 주의사항

한국 개발 생태계에서의 적용 맥락

국내 GPU 프로그래밍 환경은 아직까지 SIMT 모델에 깊이 뿌리박혀 있습니다. 특히 금융권 HPC, 자율주행 스타트업, AI 반도체 설계 회사에서 대규모 CUDA C++ 코드베이스를 운영 중인데, 이번 CUDA Tile C++ 도입 시 다음과 같은 점을 고려해야 합니다.

• 기존 코드와의 혼용 가능: __global__ SIMT 커널과 __tile_global__ 타일 커널을 같은 프로젝트에서 함께 사용할 수 있습니다. 점진적 마이그레이션이 가능합니다.
• 의존성 문제: cuda_tile.h 헤더는 CUDA 13.3 이상에서만 제공됩니다. 사내 빌드 시스템이 최신 CUDA 툴킷을 지원하는지 먼저 확인하세요.
• 교육 비용: 팀원들이 타일 추상화에 익숙해지는 데 시간이 필요합니다. 특히 '스레드 인덱스를 직접 계산하지 않는다'는 개념이 초기에는 혼란을 줄 수 있습니다.

이 기술의 한계

1. 하드웨어 의존성: CUDA Tile C++는 Compute Capability 8.0(Ampere) 이상 GPU에서만 동작합니다. 구형 GPU(V100, P100 등)에서는 사용할 수 없습니다.
2. 디버깅의 어려움: 컴파일러가 스레드 구성을 자동 결정하기 때문에, 성능 병목이 발생했을 때 원인 파악이 까다로울 수 있습니다. Nsight Compute의 Tile Statistics 섹션을 적극 활용하세요.
3. 모든 커널에 적합하지 않음: 타일 모델은 규칙적인 데이터 병렬 처리에 최적화되어 있습니다. 조건 분기가 많거나 비규칙적인 메모리 접근 패턴을 가진 커널은 여전히 SIMT 방식이 더 나을 수 있습니다.

다음 단계 학습 방향

• CUDA Tile C++ 공식 문서와 API 레퍼런스 매뉴얼을 정독하세요.
• NVIDIA Nsight Compute의 Tile Statistics 기능을 활용한 성능 분석 방법을 익혀보세요.
• 기존 SIMT 커널 중에서도 타일 모델로 쉽게 변환할 수 있는 패턴(행렬 연산, 컨볼루션, FFT 등)부터 적용해보는 것을 추천합니다.

함께 보면 좋은 글

• Google AI Edge Gallery 공개: 스마트폰에서 동작하는 온디바이스 함수 호출의 모든 것
• Astro, Cloudflare에 합류하다: 콘텐츠 중심 웹사이트의 미래가 바뀐다