CUTLASS

가장 단순한 layout — (8, 16) shape, (16, 1) stride 의 row-major 8×16 행렬. coordinate (i, j) 가 메모리 offset i·16 + j·1 = i·16 + j.

nested layout — ((4, 2), 16) shape 는 외부 4, 내부 2 의 두 단계 모드 + 마지막 모드 16. 같은 64 element 를 “4 그룹 × 2 element + 16” 으로 표현. 이게 hierarchical tiling 의 표현이 된다.

complex layout — column-major, transposed, swizzled (bank conflict 회피용 stride pattern) 모두 같은 Layout 추상으로 표현 가능.

// CUTLASS 3.x — CuTe Layout 의 예
using ALayout =
    cute::Layout<
      cute::Shape<_8, _16>,    // 8 × 16 tile
      cute::Stride<_16, _1>   // row-major
    >;

// nested — outer 4 그룹 × inner 8
using NestedLayout =
    cute::Layout<
      cute::Shape<cute::Shape<_4, _8>, _16>,
      cute::Stride<cute::Stride<_128, _16>, _1>
    >;

// 같은 데이터를 다양한 layout 으로 view
auto tA = cute::make_tensor(ptr, ALayout{});

예 — Ampere 의 m16n16k16.f16.f16.f32 MMA. 32 thread 가 한 warp 으로 일하는데:

• A (16×16): 8 element per thread, register 8 개에 분포.
• B (16×16): 8 element per thread, register 8 개.
• C (16×16, accumulator): fp32 4 element per thread.
• 각 thread 의 fragment 가 정확히 어디에 박혀야 하는지가 PTX 문서에 있음.

CUTLASS 는 이 매핑을 MMA_Atom 으로 추상화. 사용자는 using MMA = SM80_16x8x16_F32F16F16F32_TN; 같은 식으로 한 줄. 그 다음 tiled_mma 를 만들어 warp 안 thread 들에 자동 distribute.

// Ampere fp16 MMA atom
using MmaAtom = cute::MMA_Atom<
    cute::SM80_16x8x16_F32F16F16F32_TN
>;

// warp 마다 어떻게 tile 을 자르는지
using TiledMma = cute::TiledMMA<
    MmaAtom,
    cute::Layout<cute::Shape<_2, _2>>  // 2x2 warps
>;

// kernel 안에서 호출
auto tCgC = ...;   // thread 의 C fragment view
auto tArA = ...;   // A fragment
auto tBrB = ...;   // B fragment
cute::gemm(tiled_mma, tArA, tBrB, tCgC);

대표 epilogue 패턴.

• linear + bias + ReLU — Conv 또는 Linear 의 표준 후처리. cuBLAS 에는 일부만 있음, CUTLASS 는 자유.
• quantization — fp32 accumulator → int8 output. dequant scale 곱 + clamp + cast 가 한 자리에.
• FlashAttention 의 softmax + scaling — score 누산기 위에 online softmax 와 V 곱이 epilogue 안에 들어감 (사실은 main loop 도 일부 변형되지만 mental model 상).
• 2:4 sparsity dequant — sparse GEMM 의 결과 위에 metadata 기반 변환.

// epilogue functor 의 pattern
struct MyEpilogue {
    float alpha, beta, scale;

    template<class AccTensor, class CTensor>
    __device__ void
    operator()(AccTensor& acc, CTensor& bias) {
        // register 에서 처리 — HBM 안 거침
        cute::transform(acc, [&](float x){
            float y = alpha * x + beta * bias;
            return y > 0 ? y : 0;  // ReLU
        });
    }
};

이유 셋.

1. tensor core 의 layout 정확히 다뤄야 함 — FA 는 attention score 의 N×N tile 을 register 위에 누산하는데, 그 layout 이 MMA instruction 과 정확히 맞아야 함. CUTLASS 의 layout 추상이 이걸 표현 가능.
2. register accounting 이 critical — FA 의 m, l, O accumulator 가 register 에 박혀야 함 (§ L012 § 05 참조). CUTLASS 가 fragment 의 register 사용을 정확히 제어.
3. 하나의 kernel 안에 두 GEMM + softmax 를 fuse — Q×K^T 와 P×V 가 같은 kernel 안에서 register 위 누산기를 공유. Triton 으로도 가능하지만 CUTLASS 가 더 정밀한 layout 통제.

그래서 FA 의 algorithm 자체는 Triton 으로도 표현 가능하지만 — peak performance 의 마지막 5-10% 는 CUTLASS 위에서만 가능한 자리.

CUTLASS 로 내려가는 정당한 신호

• Triton 의 자동 layout 으로 못 잡는 bank conflict / register layout 문제 — NCU 의 instructions stalls 이 layout 관련.
• cuBLAS + 별도 kernel 로 짠 fusion 이 한 kernel 안에 합치면 1.3-2× 빨라질 잠재성이 보일 때.
• 새 hardware 의 새 instruction (TMA, WGMMA) 을 활용해야 할 때 — Triton 도 점차 따라잡지만 CUTLASS 가 항상 먼저.
• vendor library 의 90% 가 아니라 95% 가 필요한 production critical path.

안 가도 되는 자리

• memory-bound op (elementwise, normalize). CUTLASS 의 GEMM-grade 추상이 무의미.
• 일반적인 LLM training/inference 의 거의 모든 layer — Triton 또는 cuBLAS 으로 충분.
• fast prototyping. CUTLASS 는 컴파일 시간이 길고 디버깅이 어렵다.

Eric 의 직업적 입장 — “회사에서 CUTLASS 를 직접 짜는 사람은 보통 ML team 안에 1-2 명. 그 사람이 다른 ML 엔지니어와 협업해서 critical path 만 짠다.”