Skip to content
Published on

커널을 쓰는 세 가지 층위 — CUDA C++, Triton, CUTLASS를 같은 문제로 비교하기

🇰🇷KO🇺🇸EN🇯🇵JA🇨🇳ZH
필사 모드✨Start Quiz
공유하기
Authors

들어가며 — 같은 커널을 두 번 짜 보면 드러나는 것

행 단위 softmax 커널을 CUDA C++로 짜면 워프 셔플 리덕션, 셰어드 메모리 부분합, 두 번의 블록 동기화까지 해서 80줄쯤 나옵니다. 같은 커널을 Triton으로 짜면 12줄입니다. 그런데 성능은 비슷하거나 Triton이 조금 낫습니다.

이 사실을 처음 확인하면 "그럼 CUDA C++를 쓸 이유가 뭔가" 싶어집니다. 그리고 몇 달 뒤에는 Triton으로 도저히 표현이 안 되는 커널을 만나서 다시 CUDA C++를 열게 됩니다. 그 두 지점 사이 어딘가에 실무의 판단 기준이 있습니다.

이 글은 커널을 쓰는 세 층위를 같은 자로 재 봅니다. 확인 기준은 Triton 3.7.1(2026년 6월 릴리스), CUTLASS 4.6.1(2026년 7월), CUDA Toolkit 13.3 Update 1입니다. 세 층위 모두 활발히 바뀌고 있으므로 API 세부는 자기 버전 문서와 대조하시기 바랍니다.

세 층위는 각각 무엇을 대신 해 주는가

커널을 쓰는 일은 결국 여섯 가지 결정을 내리는 일입니다. 세 층위의 차이는 그 결정 중 몇 개를 우리가 하고 몇 개를 도구가 하느냐입니다.

결정해야 할 것CUDA C++TritonCUTLASS / CuTe
그리드와 블록 형태사람사람 (프로그램 격자만)사람 (타일 정책으로)
블록 안 스레드 배치사람컴파일러템플릿이 결정
셰어드 메모리 할당과 스윙사람컴파일러템플릿이 결정
뱅크 충돌 회피(스위즐)사람컴파일러레이아웃이 표현
비동기 복사와 파이프라이닝사람컴파일러템플릿이 결정
텐서 코어 명령어 선택사람 (또는 라이브러리)컴파일러아톰으로 명시

읽는 법은 이렇습니다. CUDA C++는 여섯 개를 다 우리가 정합니다. 그래서 가장 자유롭고 가장 오래 걸립니다. Triton은 위 두 개만 정하고 나머지를 컴파일러에 넘깁니다. 그래서 짧고, 대신 컴파일러가 잘 못하는 패턴에서는 손쓸 방법이 줄어듭니다. CUTLASS는 여섯 개를 다 정하되 직접 쓰지 않고 이미 검증된 부품을 골라 조립합니다. 그래서 GEMM 계열에서는 최고 성능이 나오지만 배우는 데 시간이 걸립니다.

CUDA C++ — 하드웨어를 직접 지시하는 층

먼저 기준선입니다. 행 단위 softmax를 손으로 씁니다. 블록 하나가 행 하나를 맡고, 최댓값과 합을 각각 블록 전체에서 리덕션합니다.

// softmax_cuda.cu
// 빌드: nvcc -O3 -arch=sm_80 softmax_cuda.cu -o softmax_cuda
#include <cstdio>
#include <cfloat>
#include <cuda_runtime.h>

static const int WARP = 32;
static const int BLOCK = 256;             // 블록당 8워프

__inline__ __device__ float warpMax(float v) {
  for (int off = WARP / 2; off > 0; off >>= 1)
    v = fmaxf(v, __shfl_xor_sync(0xffffffffu, v, off));
  return v;
}
__inline__ __device__ float warpSum(float v) {
  for (int off = WARP / 2; off > 0; off >>= 1)
    v += __shfl_xor_sync(0xffffffffu, v, off);
  return v;
}

__global__ void softmaxKernel(float *out, const float *in, int nCols) {
  __shared__ float part[BLOCK / WARP];
  __shared__ float bcast;                 // 브로드캐스트 전용 슬롯

  const long row = blockIdx.x;
  const float *src = in + row * nCols;
  float *dst = out + row * nCols;

  const int tid = threadIdx.x;
  const int lane = tid % WARP, wid = tid / WARP;
  const int nWarps = BLOCK / WARP;

  // 1패스: 행의 최댓값
  float m = -FLT_MAX;
  for (int i = tid; i < nCols; i += BLOCK) m = fmaxf(m, src[i]);
  m = warpMax(m);
  if (lane == 0) part[wid] = m;
  __syncthreads();
  if (tid == 0) {
    float t = part[0];
    for (int w = 1; w < nWarps; w++) t = fmaxf(t, part[w]);
    bcast = t;
  }
  __syncthreads();
  m = bcast;
  __syncthreads();                        // part[] 재사용 전 반드시 필요

  // 2패스: exp의 합
  float s = 0.f;
  for (int i = tid; i < nCols; i += BLOCK) s += __expf(src[i] - m);
  s = warpSum(s);
  if (lane == 0) part[wid] = s;
  __syncthreads();
  if (tid == 0) {
    float t = 0.f;
    for (int w = 0; w < nWarps; w++) t += part[w];
    bcast = t;
  }
  __syncthreads();
  const float inv = 1.f / bcast;

  // 3패스: 정규화해서 쓰기
  for (int i = tid; i < nCols; i += BLOCK) dst[i] = __expf(src[i] - m) * inv;
}

int main() {
  const int R = 8192, C = 4096;
  size_t bytes = (size_t)R * C * sizeof(float);
  float *h = (float *)malloc(bytes);
  for (long i = 0; i < (long)R * C; i++) h[i] = (float)((i * 37) % 100) * 0.01f;

  float *d_in, *d_out;
  cudaMalloc(&d_in, bytes); cudaMalloc(&d_out, bytes);
  cudaMemcpy(d_in, h, bytes, cudaMemcpyHostToDevice);

  for (int i = 0; i < 5; i++) softmaxKernel<<<R, BLOCK>>>(d_out, d_in, C);
  cudaDeviceSynchronize();

  cudaEvent_t a, b; cudaEventCreate(&a); cudaEventCreate(&b);
  cudaEventRecord(a);
  for (int i = 0; i < 50; i++) softmaxKernel<<<R, BLOCK>>>(d_out, d_in, C);
  cudaEventRecord(b); cudaEventSynchronize(b);

  float ms; cudaEventElapsedTime(&ms, a, b);
  double per = ms / 50.0;
  // 최소 트래픽: 한 번 읽고 한 번 쓴다. (3패스지만 캐시가 흡수한다)
  printf("cuda   %7.3f ms  %7.1f GB/s\n", per,
         2.0 * bytes / (per * 1e-3) / 1e9);
  return 0;
}

80줄 중 실제 수학은 세 줄입니다. 나머지는 전부 "블록 안에서 값을 모으고 다시 뿌리는" 배관입니다. 그리고 이 배관에는 조용한 함정이 있습니다.

위 코드에서 m = bcast; 다음의 __syncthreads()가 그것입니다. 이것이 없으면 빠른 워프가 2패스의 part[wid]를 먼저 덮어써서, 느린 워프가 아직 읽지 못한 1패스 결과를 날립니다. 결과는 대부분의 실행에서 맞게 나오고 가끔 틀립니다. 작은 입력에서는 거의 재현되지 않습니다. 손으로 쓴 커널에서 가장 비싼 종류의 버그이고, 이런 배관이 길어질수록 확률이 올라갑니다.

Triton — 타일 단위로 쓰고 스레드는 컴파일러에 맡긴다

같은 연산을 Triton으로 씁니다.

# softmax_triton.py
# 실행: python softmax_triton.py   (triton 3.7.x, torch 2.13 기준)
import torch
import triton
import triton.language as tl


@triton.jit
def softmax_kernel(out_ptr, in_ptr, in_stride, out_stride, n_cols,
                   BLOCK_SIZE: tl.constexpr):
    row = tl.program_id(0)
    cols = tl.arange(0, BLOCK_SIZE)
    mask = cols < n_cols

    # 행 하나를 통째로 SRAM에 올린다. 패딩 자리는 -inf라 max/sum에 영향이 없다.
    x = tl.load(in_ptr + row * in_stride + cols, mask=mask, other=-float("inf"))
    x = x - tl.max(x, axis=0)
    num = tl.exp(x)
    y = num / tl.sum(num, axis=0)
    tl.store(out_ptr + row * out_stride + cols, y, mask=mask)


def softmax(x: torch.Tensor) -> torch.Tensor:
    n_rows, n_cols = x.shape
    BLOCK_SIZE = triton.next_power_of_2(n_cols)
    num_warps = 4 if BLOCK_SIZE < 2048 else (8 if BLOCK_SIZE < 8192 else 16)
    out = torch.empty_like(x)
    softmax_kernel[(n_rows,)](
        out, x, x.stride(0), out.stride(0), n_cols,
        BLOCK_SIZE=BLOCK_SIZE, num_warps=num_warps,
    )
    return out


if __name__ == "__main__":
    torch.manual_seed(0)
    x = torch.randn(8192, 4096, device="cuda", dtype=torch.float32)

    ours, ref = softmax(x), torch.softmax(x, axis=1)
    assert torch.allclose(ours, ref, atol=1e-5), "정확성 검증 실패"

    gb = 2 * x.numel() * x.element_size() / 1e9
    for name, fn in [
        ("triton", lambda: softmax(x)),
        ("torch ", lambda: torch.softmax(x, axis=1)),
    ]:
        ms = triton.testing.do_bench(fn, warmup=25, rep=100)
        print(f"{name}  {ms:7.3f} ms  {gb / (ms * 1e-3):7.1f} GB/s")

핵심 차이는 tl.max(x, axis=0) 한 줄입니다. CUDA에서 40줄이 필요했던 블록 전체 리덕션이 여기서는 배열 연산 하나입니다. 워프 셔플을 쓸지, 셰어드 메모리를 몇 바이트 잡을지, 동기화를 어디에 넣을지는 전부 컴파일러가 결정합니다. 그리고 앞서 말한 동기화 누락 버그는 구조적으로 발생할 수 없습니다. 우리가 동기화를 쓰지 않으니까요.

컴파일 파이프라인 — 컴파일러가 실제로 하는 일

Triton이 마법을 부리는 것은 아닙니다. 단계가 있습니다. 저장소 구조로 확인할 수 있는 경로는 이렇습니다.

파이썬 함수 (@triton.jit)
   │  파이썬 AST를 순회하며 IR 생성
   ▼
Triton IR (TTIR)          타일 연산. 아직 하드웨어 개념 없음
   │  레이아웃 배정, 코얼레싱, 파이프라이닝 패스
   ▼
TritonGPU IR (TTGIR)      워프/스레드 배치와 셰어드 메모리가 결정됨
   │
   ▼
LLVM IR
   │  NVPTX 백엔드          │  AMDGPU 백엔드
   ▼                        ▼
PTX → (ptxas) → SASS       AMDGCN → 오브젝트

여기서 실무적으로 의미 있는 사실 두 가지가 나옵니다.

첫째, 백엔드가 갈라지는 지점이 LLVM IR 아래라는 것입니다. Triton 저장소의 third_party에는 nvidia와 amd 백엔드가 나란히 들어 있습니다. 그래서 같은 @triton.jit 커널이 두 벤더에서 모두 컴파일됩니다. 이 성질이 다음 두 절의 이야기를 지탱합니다.

둘째, TTIR에서 TTGIR로 내려가는 단계가 성능의 대부분을 결정한다는 것입니다. 레이아웃 배정이 잘못되면 셰어드 메모리 접근이 뱅크 충돌을 일으키고, 파이프라이닝 패스가 붙지 않으면 로드 대기가 그대로 노출됩니다. 우리가 손댈 수 있는 손잡이는 BLOCK_SIZE, num_warps, num_stages 정도이고, 그래서 Triton 튜닝은 대개 이 세 개를 흔드는 일입니다.

Triton은 그 흔드는 일을 데코레이터로 제공합니다.

@triton.autotune(
    configs=[
        triton.Config({"BLOCK_M": 64,  "BLOCK_N": 64},  num_warps=4, num_stages=3),
        triton.Config({"BLOCK_M": 128, "BLOCK_N": 64},  num_warps=8, num_stages=3),
        triton.Config({"BLOCK_M": 128, "BLOCK_N": 128}, num_warps=8, num_stages=4),
        triton.Config({"BLOCK_M": 64,  "BLOCK_N": 128}, num_warps=4, num_stages=4),
    ],
    key=["M", "N", "K"],          # 이 값들이 바뀌면 다시 튜닝한다
)
@triton.jit
def matmul_kernel(...):
    ...

key에 적힌 인자가 달라질 때마다 후보를 전부 컴파일해 실측하고 가장 빠른 것을 캐시합니다. 같은 일을 CUDA C++로 하려면 템플릿 인스턴스화와 벤치 하니스를 직접 짜야 합니다.

무엇을 만들었는지 보고 싶을 때

Triton을 블랙박스로 두지 않는 방법이 있습니다. 중간 산출물을 그대로 꺼낼 수 있습니다.

import triton, torch

# 커널을 한 번 컴파일해서 핸들을 얻는다
k = softmax_kernel.warmup(
    torch.empty(1, 8, device="cuda"), torch.empty(1, 8, device="cuda"),
    8, 8, 8, BLOCK_SIZE=8, grid=(1,),
)
k._init_handles()

print(k.asm.keys())        # ttir, ttgir, llir, ptx, cubin
print(k.asm["ttgir"][:800])  # 레이아웃 배정 결과를 눈으로 확인
print(k.n_regs, k.n_spills)  # 레지스터 스필이 나면 여기가 0이 아니다

n_spills가 0이 아니면 레지스터가 모자라 로컬 메모리로 흘렀다는 뜻이고, 대개 BLOCK_SIZE가 과합니다. Triton 커널이 이유 없이 느릴 때 가장 먼저 확인할 값입니다.

Triton이 잘 못하는 것

정직하게 정리해야 할 부분입니다. Triton으로 가지 말아야 할 곳이 분명히 있습니다.

  • 워프 단위의 정밀한 제어가 필요한 커널. 특정 레인 사이의 셔플, 워프 특화(생산자 워프와 소비자 워프를 나누는 구조), __ballot류 투표 연산에 의존하는 알고리즘은 타일 추상화에 담기지 않습니다. 이 간극을 메우려고 최근 Triton은 triton.experimental.gluon이라는 더 낮은 층의 언어를 실험 중인데, 3.7.1 시점에도 실험 단계이므로 프로덕션 판단 근거로 삼기에는 이릅니다.
  • 불규칙한 인덱싱과 데이터 의존 제어 흐름. 그래프 순회, 정렬, 동적 크기의 희소 연산처럼 접근 패턴이 실행 시점에 정해지는 경우 타일 모델이 맞지 않습니다.
  • 최고 성능의 GEMM. 밀집 행렬곱은 CUTLASS와 cuBLAS가 수년치 튜닝을 축적한 영역이고, Triton으로 그 마지막 10에서 20퍼센트를 따라잡기는 어렵습니다. 실무에서는 그럴 이유도 없습니다.
  • 컴파일 시간이 응답 시간에 포함되는 경우. JIT 컴파일이 첫 호출에 붙습니다. 오토튜닝을 켜면 후보 수만큼 곱해집니다. 서빙 경로에서는 캐시 워밍이 전제 조건입니다.
  • 정밀한 수치 재현성이 요구되는 경우. 컴파일러가 리덕션 순서를 정하므로, 버전이 바뀌면 마지막 비트가 달라질 수 있습니다.

같은 커널을 두 번 짜서 비교하기

두 구현을 같은 조건에서 재는 방법과, 그때 나오는 결과의 형태입니다.

측정 규칙은 앞 글과 같습니다. 웜업을 버리고, 여러 번 돌려 평균을 내고, 정확성을 먼저 검증하고, 시간이 아니라 유효 대역폭으로 환산합니다. Triton 쪽은 triton.testing.do_bench가 웜업과 반복, L2 캐시 플러시까지 처리해 주므로 직접 짤 필요가 없습니다. CUDA 쪽은 cudaEvent로 같은 조건을 맞춥니다.

softmax는 산술 강도가 낮은 메모리 바운드 커널이므로, 두 구현 모두 결국 HBM 대역폭에 붙습니다. 그래서 결과의 형태는 대체로 이렇게 나옵니다.

구현코드 줄 수(커널 본체)유효 대역폭비고
CUDA C++ (위 코드)약 45줄상한의 80퍼센트대3패스. 리덕션 배관이 대부분
Triton (위 코드)12줄상한의 85퍼센트대행이 SRAM에 들어가면 1패스
torch.softmax0줄상한의 85퍼센트대이미 퓨전된 라이브러리 커널

절대값은 GPU와 행 길이에 따라 크게 달라집니다. 위 하니스를 자기 장비에서 돌려 자기 숫자를 만드십시오. 이 표에서 읽어야 할 것은 순위가 아니라 두 가지 관찰입니다.

첫째, 코드량이 4배 가까이 차이 나는데 성능은 비슷합니다. 메모리 바운드 커널에서는 둘 다 같은 벽에 부딪히기 때문입니다. 벽이 성능을 정한다면, 벽까지 가는 코드가 짧을수록 좋습니다.

둘째, torch.softmax가 이미 그만큼 빠릅니다. 이것이 이 비교의 진짜 교훈입니다. 표준 연산자는 직접 짜지 마십시오. 커널을 직접 쓰는 값어치는 라이브러리에 없는 연산, 즉 퓨전된 조합이나 도메인 특유의 마스킹에서 나옵니다. 예를 들어 softmax 앞뒤로 마스킹과 스케일링과 드롭아웃이 붙어 있고 그것들이 각각 별도 커널로 실행되고 있다면, 하나로 합치는 순간 HBM 왕복이 네 번에서 한 번으로 줄어듭니다. 그때의 이득은 3배 이상이고, 그것이 Triton을 쓰는 이유입니다.

왜 Triton이 커스텀 어텐션의 사실상 표준이 되었나

지금 어텐션 변형 커널이 새로 나오면 대부분 Triton으로 먼저 나옵니다. 이유는 성능이 아니라 구조입니다.

  • 어텐션은 타일 모델과 잘 맞습니다. 쿼리 블록 하나를 잡고 키와 값을 블록 단위로 훑으며 온라인 softmax로 누적하는 구조가 곧 Triton의 프로그래밍 모델입니다.
  • 변형이 많고 수명이 짧습니다. 슬라이딩 윈도우, ALiBi, 소프트캡, 다양한 희소 패턴이 계속 나오고 그중 일부만 살아남습니다. 실험 하나에 CUDA C++로 2주를 쓸 수는 없습니다.
  • 한 번 쓰면 두 벤더에서 돕니다. 앞서 본 대로 백엔드가 LLVM IR 아래에서 갈라지므로, 같은 커널이 NVIDIA와 AMD 양쪽에서 컴파일됩니다. CUDA C++로 쓰면 AMD 쪽은 별도 작업입니다.
  • 생태계가 이미 Triton을 전제합니다. PyTorch의 TorchInductor는 GPU 대상 코드로 Triton 커널을 생성합니다. vLLM은 어텐션 백엔드 목록에 TRITON_ATTN을 두고 우선순위 표에 올려 두고 있습니다(vLLM 어텐션 백엔드 문서). 이미 설치되어 있고 이미 돌아가고 있는 도구라는 사실 자체가 큰 이점입니다.

CUTLASS와 CuTe — 템플릿으로 조립하는 GEMM

세 번째 층입니다. CUTLASS는 2017년부터 이어져 온 CUDA C++ 템플릿 라이브러리로, 행렬곱과 그 주변 연산을 계층적으로 분해해 재조립 가능한 부품으로 만들어 둔 것입니다.

핵심 발상은 GEMM을 한 덩어리로 보지 않는 것입니다. 전체 문제를 스레드 블록 타일로 쪼개고, 그것을 다시 워프 타일로, 다시 텐서 코어 명령어 하나가 처리하는 크기로 쪼갭니다. 각 층이 독립된 템플릿 파라미터이고, 우리는 그 조합을 고릅니다.

// CUTLASS 3.x/4.x 스타일 GEMM 구성의 골격
// 실제 빌드 가능한 전체 예제는 저장소의 examples 디렉터리를 참고하십시오.
using ElementA = cutlass::half_t;
using ElementB = cutlass::half_t;
using ElementC = float;

// 어떤 크기의 타일을, 몇 단계 파이프라인으로 처리할지
using TileShape   = cute::Shape<cute::_128, cute::_128, cute::_64>;
using ClusterShape = cute::Shape<cute::_1, cute::_1, cute::_1>;

// 메인 루프(데이터 이동 + 텐서 코어 누산)와
// 에필로그(결과에 붙는 후처리)를 따로 조립한다
using CollectiveMainloop = /* CollectiveBuilder<...> */;
using CollectiveEpilogue = /* CollectiveBuilder<...> */;

using GemmKernel = cutlass::gemm::kernel::GemmUniversal<
    cute::Shape<int, int, int, int>, CollectiveMainloop, CollectiveEpilogue>;
using Gemm = cutlass::gemm::device::GemmUniversalAdapter<GemmKernel>;

여기서 실무적으로 가장 값어치 있는 개념이 에필로그입니다. GEMM 결과가 아직 누산기 레지스터에 있는 동안 편향 덧셈과 활성화 함수를 적용해 버리면, 결과를 HBM에 썼다가 다시 읽는 왕복이 통째로 사라집니다. cuBLAS를 부르고 별도 커널로 활성화를 적용하면 이 왕복이 그대로 남습니다. CUTLASS를 직접 쓰는 이유의 상당 부분이 여기에 있습니다.

CuTe와 레이아웃 대수

CUTLASS 3.x부터의 기반이 CuTe입니다. 발상은 텐서의 레이아웃, 즉 논리적 좌표에서 물리적 주소로 가는 사상을 일급 값으로 만들고 그것을 대수적으로 조합하는 것입니다.

앞 글에서 셰어드 메모리 뱅크 충돌을 배열 한 칸 패딩으로 해결했습니다. CuTe에서는 그런 조작이 "스위즐"이라는 레이아웃 변환으로 표현되고, 타일 분해와 합성이 되며, 컴파일 시점에 검증됩니다. 인덱스 산술을 손으로 쓰다가 틀리는 종류의 버그가 타입 수준에서 걸립니다.

배우는 비용은 정직하게 말해서 높습니다. 템플릿 오류 메시지는 길고, 개념이 여러 층입니다. 그래서 CUTLASS 4는 파이썬 인터페이스인 CuTe DSL을 추가했습니다. CUTLASS 4.6.1 README는 이를 "C++ 없이 CuTe의 핵심 개념(레이아웃, 텐서, 하드웨어 아톰, 스레드와 데이터 계층에 대한 완전한 제어)을 노출하는 저수준 프로그래밍 모델"로 설명하며, 현재 공개 베타라고 명시합니다. PyPI에서 nvidia-cutlass-dsl로 설치할 수 있고 4.6.1이 최신입니다.

베타라는 점은 그대로 받아들이는 편이 좋습니다. 프로토타이핑에는 쓸 만하지만, 프로덕션 커널의 유일한 구현으로 삼기에는 아직 이릅니다.

어느 층을 고를 것인가

판단을 표로 압축하면 이렇습니다.

상황권장이유
표준 GEMM, 컨볼루션cuBLAS / cuDNN직접 쓸 이유가 없다
GEMM에 후처리를 붙여 왕복을 없애고 싶다CUTLASS 에필로그이 목적으로 설계된 층
어텐션 변형, 퓨전된 원소별 체인Triton타일 모델과 맞고 반복이 빠르다
MoE 라우팅, 커스텀 정규화Triton같음
워프 특화, 레인 단위 셔플이 본질인 알고리즘CUDA C++타일 추상화로 표현 불가
불규칙 인덱싱, 데이터 의존 제어 흐름CUDA C++같음
새 하드웨어 명령어를 가장 먼저 써야 함CUDA C++ 또는 CUTLASS컴파일러 지원이 늦게 온다
NVIDIA와 AMD를 한 코드로 지원해야 함Triton백엔드가 아래에서 갈라진다

실무 순서로 다시 쓰면 이렇습니다. 라이브러리로 되는지 본다. 안 되면 Triton으로 짠다. Triton이 표현하지 못하거나 성능이 안 나오면, 그때 프로파일러가 가리키는 그 커널만 CUDA C++로 내려간다. 처음부터 CUDA C++로 시작하는 것은, 그 커널이 워프 단위 제어를 본질적으로 요구한다는 확신이 있을 때뿐입니다.

한 가지 덧붙이면, 층을 섞는 것이 정상입니다. 실제 추론 스택 하나를 열어 보면 GEMM은 cuBLAS나 CUTLASS, 어텐션은 손으로 쓴 CUDA C++ 또는 Triton, 정규화와 활성화는 Triton, 그 사이 접착은 PyTorch입니다. 하나를 골라 전부 통일해야 한다는 압박은 근거가 없습니다.

마치며 — 추상화의 층은 성능이 아니라 수정 속도로 고르는 것이다

세 층위를 비교했지만 결론이 "Triton이 이겼다"는 아닙니다. 메모리 바운드 커널에서 CUDA C++와 Triton의 성능이 비슷하게 나온 것은 둘 다 같은 물리적 벽에 붙었기 때문이고, GEMM에서 CUTLASS가 앞서는 것은 그 벽이 다른 곳에 있기 때문입니다. 층위 자체가 성능을 정하지 않습니다.

층위가 정하는 것은 다른 쪽입니다. 커널 하나를 고쳐서 다시 재기까지 걸리는 시간, 동기화를 빠뜨려서 비결정적으로 틀릴 확률, 새 GPU가 나왔을 때 다시 손봐야 할 코드의 양. 성능이 비슷하다면 이쪽이 선택 기준이 되어야 하고, 대부분의 실무에서 성능은 비슷합니다.

그래서 판단 기준을 한 줄로 줄이면 이렇습니다. 이 커널을 앞으로 몇 번이나 고칠 것 같은가. 한 번 쓰고 2년간 안 건드릴 커널이라면 CUDA C++로 내려가 마지막 10퍼센트를 짜내는 것이 합리적입니다. 다음 달에도 실험이 이어질 커널이라면, 12줄로 쓰고 그 시간에 실험을 한 번 더 돌리는 쪽이 거의 언제나 낫습니다.

참고 자료

Discuss on Twitter • View on GitHub