CUTLASS CuTe 라이브러리와 메모리 주소 Layout (25.8.11)

CUTLASS: 엔비디아 GPU를 위한 고성능 행렬 연산 라이브러리

CUTLASS(CUDA Templates for Linear Algebra Subroutines)는 엔비디아(NVIDIA)에서 개발한 CUDA C++ 템플릿 라이브러리로, 고성능 행렬 곱셈(GEMM, General Matrix-Matrix Multiplication) 및 관련 연산을 구현하기 위한 추상화된 인터페이스를 제공합니다. 딥러닝과 고성능 컴퓨팅(HPC) 분야에서 핵심적인 역할을 하는 행렬 연산의 성능을 극대화하는 것을 목표로 합니다.
 
CUTLASS는 헤더 전용(header-only) 라이브러리로, 별도의 빌드 과정 없이 애플리케이션에 포함하여 사용할 수 있습니다. 이를 통해 개발자는 엔비디아 GPU의 아키텍처, 특히 텐서 코어(Tensor Core)의 성능을 최대한 활용하는 맞춤형 커널을 보다 쉽게 개발할 수 있습니다.
 

 

주요 특징 및 아키텍처

CUTLASS는 계층적 분해(hierarchical decomposition)와 데이터 이동 전략을 기반으로 설계되었습니다. 복잡한 행렬 연산을 더 작고 재사용 가능한 모듈식 소프트웨어 구성 요소로 분해하여, 다양한 수준의 병렬 처리 계층에 맞춰 최적화할 수 있도록 지원합니다. 주요 구성 요소는 다음과 같습니다.

• 스레드블록(Threadblock) 수준 연산: 대규모 행렬 곱셈을 처리하기 위한 기본 단위입니다.
• 워프(Warp) 수준 프리미티브: 스레드 블록 내에서 행렬 곱셈-누산(MMA, Multiply-Accumulate) 연산을 수행하는 기본 단위입니다.
• 에필로그(Epilogue): 행렬 곱셈 결과에 활성화 함수(Activation Function) 적용이나 다른 텐서와의 연산 등 후처리 작업을 효율적으로 수행하는 구성 요소입니다.
• 데이터 로딩/저장 유틸리티: GPU 메모리 간의 텐서 데이터를 효율적으로 이동시키는 기능을 제공합니다.

이러한 모듈식 접근 방식 덕분에 개발자는 타일링 크기, 데이터 타입, 알고리즘 정책 등을 C++ 템플릿 메타프로그래밍을 통해 유연하게 조정하여 특정 하드웨어와 문제에 최적화된 코드를 생성할 수 있습니다.
 
2025.06.26 - [TechStock&Review/SemiConduct] - NVIDIA 텐서 코어 진화 - Volta 에서 Blackwell 까지 (25.6.26)

NVIDIA 텐서 코어 진화 - Volta 에서 Blackwell 까지 (25.6.26)

 

다양한 데이터 타입 지원

CUTLASS의 큰 장점 중 하나는 폭넓은 데이터 타입 지원입니다. 이를 통해 혼합 정밀도(mixed-precision) 연산을 효과적으로 활용하여 성능과 정확도의 균형을 맞출 수 있습니다. 지원하는 주요 데이터 타입은 다음과 같습니다.

• 부동 소수점: FP64, FP32, TF32, FP16, BF16
• 8비트 부동 소수점: FP8 (e5m2, e4m3)
• 정수: 4비트 및 8비트 부호/비부호 정수
• 이진: 1비트 데이터 타입

CuTe: 차세대 프로그래밍 모델

CUTLASS 3.0부터는 CuTe(CUDA Template Engine)라는 새로운 상위 레벨 추상화가 도입되었습니다. CuTe는 텐서의 레이아웃과 스레드 계층을 정의하고 조작하기 위한 강력하고 표현력 있는 C++ 템플릿 라이브러리입니다. 이를 통해 개발자는 복잡한 텐서 연산을 더욱 직관적이고 유연하게 프로그래밍할 수 있으며, 특히 엔비디아의 최신 아키텍처인 호퍼(Hopper)의 텐서 메모리 가속기(TMA, Tensor Memory Accelerator)와 같은 새로운 하드웨어 기능을 효과적으로 활용할 수 있습니다.

CUTLASS vs. cuBLAS

cuBLAS는 엔비디아에서 제공하는 고도로 최적화된 BLAS(Basic Linear Algebra Subprograms) 라이브러리로, 별도의 커널 코드 작성 없이도 높은 성능을 제공합니다. 반면 CUTLASS는 개발자가 직접 커널 코드를 작성하고 특정 요구사항에 맞춰 세밀하게 튜닝할 수 있는 유연성을 제공합니다.
 

---

How do you read CUTLASS CuTe Layout? - CUTLASS CuTe Layout을 어떻게 읽나요?

https://x.com/SemiAnalysis_/status/1949862435160285298

X의 SemiAnalysis님(@SemiAnalysis_)

 
 
행렬 A의 TV 레이아웃 (이미지 왼쪽 아래)은 ((4, 8), (2, 2)) : (32, 1), (16, 8))입니다. 여기서 TV 레이아웃을 단계별로 유도해 보겠습니다.
 

Thread0 의 2x2 메트릭스 메모리 주소

 
TV (스레드-값) 레이아웃은 (스레드 인덱스, 값 인덱스)를 물리적 주소 인덱스로 매핑합니다. 스레드와 값 레이아웃을 다음과 같이 별도로 읽을 수 있습니다: T = (4, 8) : (32, 1) V = (2, 2) : (16, 8)
 
스레드 레이아웃은 스레드 인덱스를 물리적 주소 인덱스로 매핑합니다. (4, 8) : (32, 1)은 i를 0에서 4까지, j를 0에서 8까지 루프하며, i의 스트라이드 = 32, j의 스트라이드 = 1을 의미합니다. (i, j)를 컬럼 메이저 순서로 루프하여 스레드 레이아웃을 얻습니다. 따라서 시퀀스는 다음과 같이 시작합니다: 0, 32, 64, 96, 1, 33, 65, 97...
 
값 레이아웃은 값 인덱스를 물리적 주소 인덱스로 매핑합니다. 마찬가지로 (2, 2) : (16, 8)은 i를 0에서 2까지, j를 0에서 2까지 루프하며, i의 스트라이드 = 16, j의 스트라이드 = 8을 의미합니다. (i, j)를 컬럼 메이저 순서로 루프하여 값 레이아웃을 얻습니다. 따라서 시퀀스는 다음과 같습니다: 0, 16, 8, 24.
 
그 다음, 스레드 인덱스를 열에 배치하고 값 인덱스를 행에 배치합니다. 각 셀 (i, j)에 대해 스레드 i의 물리적 주소 인덱스를 베이스로 사용하고, 값 j의 물리적 주소 인덱스로 오프셋을 적용합니다. 예를 들어, 셀 (1, 1) = T(1) + V(1) = 32 + 16 = 48입니다. 결과는 다음과 같습니다:
 
이제 (스레드 인덱스, 값 인덱스)를 물리적 주소로 매핑했습니다. 물리적 배열은 컬럼 메이저입니다. 
예를 들어, 
(T0, V1)은 물리적 인덱스 16으로 매핑되며, 이는 행렬에서 (0, 1)이 됩니다. 
(T0, V3)은 물리적 인덱스 24로 매핑되며, 이는 행렬에서 (8, 1)이 됩니다.
 

 

---

배경 이해: 왜 이런 레이아웃이 필요한가?

• GPU 프로그래밍에서 행렬 연산(예: A * B = C)은 수천 개의 스레드가 병렬로 처리합니다. 각 스레드는 행렬의 일부 요소를 계산하지만, 메모리 접근 패턴이 비효율적이면 성능이 저하됩니다.
• CuTe 레이아웃은 텐서의 "모드"(차원)를 정의하여, 쓰레드가 어떤 데이터를 처리할지, 메모리에 어떻게 저장될지 매핑합니다. 이는 GPU의 공유 메모리(Shared Memory)나 레지스터를 효율적으로 사용하기 위함입니다.
• TV 레이아웃(Thread-Value Layout)은 스레드 인덱스와 값 인덱스를 결합하여 물리적 메모리 주소로 변환합니다. 이는 행렬 A의 데이터를 GPU 워프(Warp: 32개 스레드 그룹) 단위로 배분하는 데 유용합니다.
• 예시에서 총 요소 수: 32 스레드 (4x8) * 4 값 (2x2) = 128 요소. 이는 16x8 행렬(컬럼 메이저 저장)을 나타내며, GPU의 tensor core(텐서 코어) 연산에 최적화된 크기입니다.

단계별 의미 설명

1. TV 레이아웃의 기본 구조:

• ((4, 8), (2, 2)) : ((32, 1), (16, 8)) 형식은 두 부분으로 나뉩니다.
• Shape: ((4,8), (2,2)) – 스레드 차원(4x8=32 스레드), 값 차원(2x2=4 값).
• Stride: ((32,1), (16,8)) – 각 차원의 메모리 이동 거리(스트라이드).
• 이는 (스레드 인덱스, 값 인덱스) 쌍을 메모리 주소로 매핑합니다. 스레드는 GPU의 병렬 단위, 값은 각 스레드가 처리하는 데이터 조각입니다.

2. 스레드 레이아웃 (T = (4,8):(32,1)):

• Shape (4,8): 4행 x 8열의 2D 그리드로 스레드를 생각합니다. (총 32 스레드, GPU 워프 크기와 맞춤).
• Stride (32,1): 첫 번째 차원(행)의 이동 거리 32, 두 번째(열) 1.
• 컬럼 메이저 순서(열 우선): 메모리에서 열을 먼저 채웁니다.
  • 예:
  • (i=0, j=0) → 주소 0
  • (i=1, j=0) → 32
  • (i=2, j=0) → 64
  • (i=3, j=0) → 96
  • 다음 열 (i=0, j=1) → 1
  • 이는 스레드 0~31의 메모리 주소를 순차적으로 정의합니다. 메모리 뱅크 충돌(Bank Conflict)을 피하기 위해 스트라이드를 설계합니다.

3. 값 레이아웃 (V = (2,2):(16,8)):

• Shape (2,2): 각 스레드가 처리하는 2x2=4개의 값.
• Stride (16,8): 첫 번째 차원 이동 16, 두 번째 8.
• 컬럼 메이저:
  • (i=0, j=0) → 0
  • (i=1, j=0) → 16
  • (i=0, j=1) → 8
  • (i=1, j=1) → 24
• 이는 각 스레드가 로드하는 데이터의 오프셋을 정의합니다. 값이 메모리에서 산재되도록 하여, 병렬 로드 시 효율적입니다.

4. 결합: TV 매핑:

• 그리드 생성: 행=값 인덱스 (V0~V3), 열=스레드 인덱스 (T0~T31).
• 각 셀 = T 주소 (베이스) + V 주소 (오프셋). 예: T1 + V1 = 32 + 16 = 48.
• 결과 테이블(두 번째 이미지): 4x32 그리드의 주소 (0~127). 이는 전체 메모리 배열을 채웁니다.
• 물리적 배열은 컬럼 메이저: 주소 = 행 + 열 * 높이 (여기서 높이=16, 예: 주소 16 = 행 0, 열 1; 주소 24 = 행 8, 열 1).

 

더 넓은 의미와 응용

• 성능 최적화: GPU에서 행렬 곱셈 시, 이 레이아웃은 tensor core를 활용하여 FP16/INT8 같은 정밀도에서 고속 연산을 가능하게 합니다. 예: Hopper/Blackwell 아키텍처에서 GEMM 속도 향상.
• 왜 컬럼 메이저?: CUDA는 Fortran 스타일(컬럼 메이저)을 선호하며, 벡터화된 로드(LDG)에 적합합니다.
• 실제 사용: CUTLASS로 커스텀 커널 작성 시, CuTe를 사용해 레이아웃을 정의하면 코드가 간결해지고, 자동 최적화됩니다. 딥러닝 프레임워크(TensorFlow, PyTorch) 내부에서 유사한 기법이 적용됩니다.
• 잠재적 어려움: 초보자는 스트라이드 계산이 복잡할 수 있으나, 이 스레드는 단계별로 풀어 설명하여 이해를 돕습니다. 실제 코드에서 cute::Layout<> 템플릿으로 구현됩니다

---

CUTLASS CuTe 사용 코드 예시

CUTLASS의 CuTe(CuTe Tensor)는 C++ CUDA 템플릿을 사용해 텐서 레이아웃을 정의하고 조작하는 추상화 도구입니다. 이전 설명에서 다룬 TV 레이아웃((스레드, 값) 매핑)은 CuTe의 Layout 개념을 기반으로 하며, GPU 스레드와 메모리 접근을 최적화합니다. 아래는 CuTe를 사용하는 기본적인 코드 예시입니다. 이는 NVIDIA 공식 문서와 튜토리얼을 기반으로 하며, 실제로 CUTLASS 라이브러리를 설치하고 CUDA 컴파일러(nvcc)로 빌드해야 실행할 수 있습니다. (CUTLASS GitHub: https://github.com/NVIDIA/cutlass 에서 다운로드 가능)
 
CuTe 사용을 위해 #include <cute/layout.hpp>와 using namespace cute;를 추가합니다. 예시는 간단한 레이아웃 정의부터 시작해, 이전 TV 레이아웃과 유사한 매핑, 그리고 행렬 전치(transpose) 예시로 확장합니다.
 

1. 기본 Layout 정의와 사용 예시

CuTe의 Layout은 Shape(모양)과 Stride(이동 거리)를 사용해 정의됩니다. 이전 설명의 ((4,8),(2,2)):( (32,1),(16,8) ) TV 레이아웃을 CuTe로 표현하면 다음과 같습니다. 이 코드는 레이아웃을 생성하고, 좌표를 인덱스로 매핑하는 방법을 보여줍니다.
 

#include <cute/layout.hpp>
#include <cstdio>

using namespace cute;

int main() {
    // TV Layout 예시: Shape ((4,8),(2,2)), Stride ((32,1),(16,8))
    // 스레드 부분: (4,8):(32,1) - 컬럼 메이저
    // 값 부분: (2,2):(16,8) - 컬럼 메이저
    auto thread_shape = make_shape(Int<4>{}, Int<8>{});
    auto thread_stride = make_stride(Int<32>{}, Int<1>{});
    auto value_shape = make_shape(Int<2>{}, Int<2>{});
    auto value_stride = make_stride(Int<16>{}, Int<8>{});

    auto tv_shape = make_shape(thread_shape, value_shape);
    auto tv_stride = make_stride(thread_stride, value_stride);
    auto tv_layout = make_layout(tv_shape, tv_stride);

    // 레이아웃 출력 (print_layout 사용, CUTLASS에 포함)
    print_layout(tv_layout);
    printf("\n");

    // 예: (스레드 인덱스 1, 값 인덱스 1) -> 물리적 주소 계산
    // 이전 설명처럼 T(1) = 32, V(1) = 16, 총 48
    int thread_idx = 1;  // (i=1, j=0) in thread shape
    int value_idx = 1;   // (i=1, j=0) in value shape
    auto coord = make_coord(make_coord(1, 0), make_coord(1, 0));  // 다차원 좌표
    printf("TV Layout at coord ((1,0),(1,0)): %d\n", tv_layout(coord));

    // 전체 크기: size = 32 스레드 * 4 값 = 128
    printf("Size of layout: %d\n", size(tv_layout));

    return 0;
}

 

• 설명: make_shape와 make_stride로 계층적 레이아웃을 만듭니다. make_layout으로 결합합니다. tv_layout(coord)는 좌표를 인덱스로 매핑합니다. 이는 이전 X 포스트의 TV 매핑을 직접 구현한 것입니다. 컴파일: nvcc -o example example.cu -I/path/to/cutlass/include. 출력 예: coord ((1,0),(1,0))에서 48 (32 + 16).

 

2. Cosize 개념 예시 (메모리 범위 계산)

CuTe에서 cosize(layout)는 레이아웃이 커버하는 최대 메모리 범위(가능한 출력 값의 크기)를 계산합니다. 이전 TV 레이아웃의 경우 cosize는 128 (0~127)이 됩니다. 아래는 간단한 3x5 레이아웃 예시입니다.
 

#include <cstdio>
#include <cute/layout.hpp>

using namespace cute;

int main(int argc, char** argv) {
    Layout layout = make_layout(
        make_shape(Int<3>{}, Int<5>{}),
        make_stride(Int<1>{}, Int<4>{})
    );
    printf("layout: ");
    print_layout(layout);
    printf("\n");

    // domain: 입력 가능한 인덱스 [0 ~ size-1]
    printf("%17s", "domain: [");
    for (int i = 0; i < size(layout); i++) {
        if (i == 0) { printf("%d", i); } else { printf(", %2d", i); }
    }
    printf("]\n");

    // range: 실제 출력 값
    printf("%17s", "range: [");
    for (int i = 0; i < size(layout); i++) {
        if (i == 0) { printf("%d", layout(i)); } else { printf(", %2d", layout(i)); }
    }
    printf("]\n");

    // codomain: 가능한 출력 범위 [min ~ max]
    printf("%17s", "codomain: [");
    for (int i = layout(0); i < (layout(size(layout) - 1) + 1); i++) {
        if (i == 0) { printf("%d", i); } else { printf(", %2d", i); }
    }
    printf("]\n\n");

    // size와 cosize
    printf("%14s: %d\n", "size(layout)", (int)size(layout));
    printf("%14s: %d\n", "cosize(layout)", (int)cosize(layout));
    printf("%14s: %d\n", "cosize equiv:", (int)layout(size(layout) - 1) + 1);

    return 0;
}

 

• 설명: 이 코드는 레이아웃의 domain(입력 범위), range(출력 값), codomain(가능 출력 범위)을 보여줍니다. cosize는 메모리 할당 시 유용하며, 이전 TV 예시에서 물리적 배열 크기(128)를 계산하는 데 사용될 수 있습니다. 출력 예: size=15, cosize=17 (스트라이드 때문에 범위가 0~16).

 

3. 실제 응용: 행렬 전치(Matrix Transpose) 예시

CuTe를 사용해 GPU에서 행렬을 전치하는 예시입니다. 이는 GEMM 같은 연산에서 메모리 접근을 최적화합니다. 글로벌 메모리(GMEM)에서 공유 메모리(SMEM)를 통해 전치합니다. 이전 TV 레이아웃처럼 타일(tile)로 나누어 스레드에 배분합니다.
 

#include <cute/tensor.hpp>
#include <cute/layout.hpp>
#include <cuda_runtime.h>

using namespace cute;

__global__ void transpose_kernel(float* d_D, float* d_S, int M, int N) {
    extern __shared__ float smem[];

    // 레이아웃 정의: M x N 행렬, row-major
    auto tensor_shape = make_shape(M, N);
    auto gmemLayoutS = make_layout(tensor_shape, GenRowMajor{});
    auto gmemLayoutD = make_layout(make_shape(N, M), GenRowMajor{});
    Tensor tensor_S = make_tensor(make_gmem_ptr(d_S), gmemLayoutS);
    Tensor tensor_D = make_tensor(make_gmem_ptr(d_D), gmemLayoutD);

    // 타일 크기: 32x32 (TV 레이아웃처럼 블록 분할)
    using b = Int<32>;
    auto block_shape = make_shape(b{}, b{});
    Tensor tiled_S = tiled_divide(tensor_S, block_shape);
    Tensor tiled_D = tiled_divide(tensor_D, block_shape);  // 전치된 뷰

    // 블록 타일 선택
    Tensor tile_S = tiled_S(make_coord(_, _), blockIdx.x, blockIdx.y);
    Tensor tile_D = tiled_D(make_coord(_, _), blockIdx.y, blockIdx.x);  // 전치 인덱스

    // 스레드 레이아웃: 8x32 (총 256 스레드, 워프 최적화)
    auto thr_layout = make_layout(make_shape(Int<8>{}, Int<32>{}), GenRowMajor{});
    Tensor thr_tile_S = local_partition(tile_S, thr_layout, threadIdx.x);
    Tensor thr_tile_D = local_partition(tile_D, thr_layout, threadIdx.x);

    // SMEM 레이아웃: row-major와 col-major 뷰
    auto smemLayout = make_layout(block_shape, GenRowMajor{});
    auto smemLayoutT = make_layout(block_shape, GenColMajor{});
    Tensor sS = make_tensor(make_smem_ptr(smem), smemLayout);
    Tensor sD = make_tensor(make_smem_ptr(smem), smemLayoutT);

    // GMEM -> SMEM 복사 (coalesced load)
    Tensor thr_sS = local_partition(sS, thr_layout, threadIdx.x);
    copy(thr_tile_S, thr_sS);
    __syncthreads();

    // SMEM -> GMEM 복사 (전치된 뷰로 store)
    Tensor thr_sD = local_partition(sD, thr_layout, threadIdx.x);
    copy(thr_sD, thr_tile_D);
}

int main() {
    int M = 2048, N = 2048;
    float *h_S, *h_D, *d_S, *d_D;
    // 호스트/디바이스 메모리 할당 및 초기화 (생략)
    dim3 grid((M + 31) / 32, (N + 31) / 32);
    dim3 block(256);  // 8x32 스레드
    size_t smem_size = 32 * 32 * sizeof(float);
    transpose_kernel<<<grid, block, smem_size>>>(d_D, d_S, M, N);
    // 결과 확인 (생략)
    return 0;
}

 

• 설명: Tensor는 포인터와 레이아웃을 결합합니다. tiled_divide로 타일을 나누고, local_partition으로 스레드에 배분합니다 (이전 TV처럼). SMEM을 사용해 메모리 접근을 최적화합니다. 이는 uncoalesced 접근을 피하며, 성능이 크게 향상됩니다 (튜토리얼에서 0.79 ms, 10,900 GB/s). 이전 TV 레이아웃을 GEMM에 적용할 때 유사하게 사용됩니다.

 
CUTLASS CuTe 예시의 특징:

• 복잡성 및 유연성: C++ 템플릿 메타프로그래밍을 사용하여 커널의 동작(타일 크기, 데이터 로딩 방식, MMA 연산 종류 등)을 매우 세밀하게 정의합니다.
• "화이트박스" 접근: 개발자가 GPU 아키텍처를 이해하고 커널의 내부 동작을 직접 제어할 수 있습니다.
• 최고 성능 지향: 표준 라이브러리로 해결할 수 없는 특수한 연산(예: Fused-Attention)을 구현하거나, 특정 하드웨어(예: Hopper 아키텍처의 Tensor Memory Accelerator)의 성능을 극한까지 활용하고자 할 때 사용됩니다

---

cuBLAS 사용 예시: 간결하고 직관적인 API 호출

cuBLAS는 이미 고도로 최적화된 함수를 호출하는 방식이라 코드가 매우 간결합니다. 개발자는 메모리 할당과 데이터 전송 후, 원하는 연산에 맞는 cuBLAS 함수(cublasSgemm 등)를 호출하기만 하면 됩니다.

#include <iostream>
#include <vector>
#include <cublas_v2.h>
#include <cuda_runtime.h>

// CUDA 오류 체크를 위한 헬퍼 함수
void checkCuda(cudaError_t result) {
    if (result != cudaSuccess) {
        std::cerr << "CUDA Error: " << cudaGetErrorString(result) << std::endl;
        exit(-1);
    }
}

int main() {
    // 행렬 크기 정의 (M, N, K)
    int M = 256;
    int N = 512;
    int K = 128;

    // 호스트(CPU) 메모리에 데이터 생성
    std::vector<float> h_A(M * K);
    std::vector<float> h_B(K * N);
    std::vector<float> h_C(M * N, 0.0f);

    // 행렬 초기화 (예: 1.0f, 2.0f로 채우기)
    for (int i = 0; i < M * K; ++i) h_A[i] = 1.0f;
    for (int i = 0; i < K * N; ++i) h_B[i] = 2.0f;

    // 디바이스(GPU) 메모리 할당
    float *d_A, *d_B, *d_C;
    checkCuda(cudaMalloc((void**)&d_A, M * K * sizeof(float)));
    checkCuda(cudaMalloc((void**)&d_B, K * N * sizeof(float)));
    checkCuda(cudaMalloc((void**)&d_C, M * N * sizeof(float)));

    // 호스트 -> 디바이스 데이터 복사
    checkCuda(cudaMemcpy(d_A, h_A.data(), M * K * sizeof(float), cudaMemcpyHostToDevice));
    checkCuda(cudaMemcpy(d_B, h_B.data(), K * N * sizeof(float), cudaMemcpyHostToDevice));

    // cuBLAS 핸들 생성
    cublasHandle_t handle;
    cublasCreate(&handle);

    // 알파, 베타 값 설정 (C = alpha * A * B + beta * C)
    float alpha = 1.0f;
    float beta = 0.0f;

    // **[핵심] cuBLAS GEMM 함수 호출**
    cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, M, N, K, &alpha, d_A, M, d_B, K, &beta, d_C, M);

    // 디바이스 -> 호스트 결과 복사
    checkCuda(cudaMemcpy(h_C.data(), d_C, M * N * sizeof(float), cudaMemcpyDeviceToHost));

    // 핸들 파괴 및 메모리 해제
    cublasDestroy(handle);
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    // 결과 확인 (예시)
    std::cout << "cuBLAS execution finished. First element of C: " << h_C[0] << std::endl;

    return 0;
}

 
 
cuBLAS 예시의 특징:

• 단순함: cublasCreate, cublasSgemm, cublasDestroy 등 몇 개의 API 호출로 끝납니다.
• 고수준 추상화: GPU의 내부 아키텍처나 스레드 관리에 대해 전혀 신경 쓸 필요가 없습니다.
• "블랙박스" 형태: 엔비디아가 제공하는 최적의 구현을 그대로 사용합니다.

---

결론: 언제 무엇을 써야 할까? 

• cuBLAS: 대부분의 표준적인 딥러닝 및 HPC 애플리케이션에 적합합니다. 사용이 간편하고 이미 높은 성능이 보장되므로 빠른 개발과 안정적인 성능이 필요할 때 최선의 선택입니다.
• CUTLASS CuTe: 최고 수준의 성능 최적화가 필요하거나, cuBLAS가 지원하지 않는 새로운 종류의 연산자를 직접 개발해야 할 때 사용합니다. GPU 아키텍처에 대한 깊은 이해가 필요하며, 개발 시간과 복잡성이 훨씬 높습니다.

 

Work Cite

1. https://docs.nvidia.com/cutlass/media/docs/cpp/cute/01_layout.html
2. https://docs.nvidia.com/cutlass/media/docs/cpp/cute/00_quickstart.html
3. https://github.com/NVIDIA/cutlass
4. https://medium.com/%2540ngocson2vn/cutlass-cute-layout-cosize-4ca79f44d42c
5. https://github.com/NVIDIA/cutlass/blob/main/media/docs/cpp/terminology.md
6. https://www.reddit.com/r/Compilers/comments/1lidz5w/nvidia_cutlass_cute_dsl_for_tensor_layout_algebra/
7. https://stackoverflow.com/questions/79687584/distinction-cute-and-nvidia-cutlass
8. https://research.colfax-intl.com/tutorial-matrix-transpose-in-cutlass/
9. https://www.youtube.com/watch?v=PWWOGrLZtZg
10. https://docs.nvidia.com/cutlass/overview.html