Outperforming cuBLAS on H100

감각이 헷갈리는 자리. cuBLAS 는 NVIDIA 가 직접 만든 GEMM 라이브러리다. 그런데 어떤 사이즈/조건에서 외부 사용자가 짠 커널이 더 빠를 수 있다. 이게 가능한 이유와 한계를 분명히 하는 게 이 강의의 가치.

강의 transcript 가 누락되었지만, 토픽이 가리키는 자리는 분명하다 — H100 의 새 hardware feature 들이 cuBLAS 에 다 반영되기까지 시간이 걸리고, 그 사이 특정 사이즈와 데이터 타입에 specialize 된 커널이 cuBLAS 의 general 한 path 를 이길 수 있다. 같은 결의 자료를 Modal blog 에서 동명 시리즈로 출간했다 — 본 노트는 Modal blog + Hopper 공식 문서를 베이스로 강의 토픽을 추정한다.

이 강의의 실용적 산출 — Hopper 의 모든 trick 이 한 자리에 모인 reference GEMM. L043 의 Turing 5단계 사다리를 출발점으로 하면, 이 강의가 그 사다리에 TMA, WGMMA, warp specialization, cluster, swizzle 까지 추가한 H100 버전.

cuBLAS 의 디자인 제약 — 모든 사이즈, 모든 dtype 조합, 모든 stride/layout, 모든 alpha/beta 조합에서 합리적으로 빠르게 동작해야 한다. 이 generality 가 비용으로 돌아오는 자리.

L043 의 Turing 사다리 — column-major → SMEM → register tile → double buffer — 가 Ampere(A100) 위에서 추가로 cp.async 를 만난다. Hopper(H100) 는 거기서 한 번 더, 5개의 새 도구가 추가된다.

H100 GEMM 이 cuBLAS 를 격파하려면 — 이 6개 도구를 한 커널 안에 모두 결합해야 한다. 하나라도 빠지면 격차가 작아진다. 강의의 본론은 그 결합의 모양.

L042 의 Mosaic GPU 에서도 본 핵심 패턴이 — Hopper GEMM 의 표준 구조다. 한 block 안의 warp group 들이 다른 역할을 맡는다.

SMEM 은 bank 32개로 분할. 한 word(4B) 가 한 bank 에 있고, warp 의 32 thread 가 같은 bank 의 다른 word 를 동시에 access 하면 bank conflict. 32-way conflict 면 32배 느려진다.

TMA 의 핵심 추상. copy 의 모든 결정 — global 의 base address, dimension, stride, swizzle, 데이터 타입, padding mode — 을 한 tensor map 객체에 한 번 묶는다. 그리고 이후엔 좌표만 주면 된다.

// host 측에서 descriptor 한 번 만들기
CUtensorMap tensor_map_a;
cuTensorMapEncodeTiled(
    &tensor_map_a,
    CU_TENSOR_MAP_DATA_TYPE_FLOAT16,         // dtype
    /*tensor_rank*/ 2,
    /*global_address*/ (void*)A_ptr,
    /*global_dim*/ {M, K},
    /*global_strides*/ {K * 2 /*bytes*/},
    /*box_dim*/ {BLOCK_M, BLOCK_K},          // SMEM 에 가져올 tile 크기
    /*element_strides*/ {1, 1},
    CU_TENSOR_MAP_INTERLEAVE_NONE,
    CU_TENSOR_MAP_SWIZZLE_128B,              // 128B swizzle 자동 적용
    CU_TENSOR_MAP_L2_PROMOTION_L2_64B,
    CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE
);

// device kernel 안에서 — 좌표만 넘기면 끝
asm volatile (
    "cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::complete_tx::bytes "
    "[%0], [%1, {%2, %3}], [%4];"
    :: "r"(smem_addr), "l"(&tensor_map_a),
       "r"(m_offset), "r"(k_offset),
       "r"(barrier_addr)
);

강의 결과 부분의 핵심 — 절대 throughput 보다 크기별 비율의 그래프가 중요하다. 어떤 (M, N, K) 에서 custom 이 cuBLAS 를 이기고, 어떤 자리에서는 지는지.

이 결과에서 cuBLAS 격파의 의미가 다시 분명해진다. 특정 사이즈에서 우월. 그리고 그 우월의 폭은 보통 5–15%. cuBLAS 가 이미 매우 좋다는 사실의 반증.

강의가 다뤘을 (또는 다뤘어야 할) 자리. custom GEMM 이 cuBLAS 를 한 사이즈에서 이긴다고 그대로 production 에 들어가지 않는다. 5가지 어려움.

그래서 production GEMM 은 fallback 으로 cuBLAS, hot path 만 custom 의 layered 접근. PyTorch 자체도 aten::matmul 안에서 사이즈별로 cuBLAS / cublasLt / Triton fused / custom 을 분기.

강의가 직접 CUDA C++ 로 풀었지만, NVIDIA 가 같은 일을 위한 라이브러리를 제공한다 — CUTLASS. 3.x 부터는 Hopper 의 모든 도구를 wrapper 로 노출.

손에 새기기 — 실습 시퀀스

1. cuBLAS baseline 측정 — H100 위에서 M=N=K ∈ {2048, 4096, 8192} 의 FP16 GEMM 을 cuBLAS 로 측정. throughput 의 peak 비율 기록.
2. Naive WGMMA 커널 — TMA 없이 직접 GMEM access 로 WGMMA 호출. 이 베이스라인이 얼마나 느린지 확인.
3. TMA 추가 — host 측에서 tensor map 만들고, kernel 안에서 cp.async.bulk.tensor PTX 호출. 베이스라인 대비 speedup.
4. warp specialization — producer 1 warp + consumer 2 WG 분리. setmaxnreg 로 register 분배.
5. SMEM swizzle 128B — TMA descriptor 와 WGMMA operand layout 매칭. NCU 의 SMEM bank conflict metric 으로 검증.
6. cluster + multicast — block cluster 활성화, A 또는 B 를 multicast. bandwidth 사용량 감소를 NCU 로 확인.
7. cuBLAS 와 비교 — 단계별로 cuBLAS 대비 % 그래프. 어느 단계에서 격차가 좁혀지는지 추적.

강의 transcript 가 누락된 강의 — 본 노트의 구체적 수치와 코드는 도메인 지식과 외부 자료에서 가져왔다. 영상 직접 확인이 필요한 사항들.

• 강의 발표가 보여준 정확한 사이즈와 격차 — 본 노트의 그래프는 도메인 지식 기반 추정. 실제 강의의 측정값은 영상 확인 필요.
• FP8 가 다뤄졌는가 — Hopper 의 새 dtype. cuBLAS 의 FP8 path 와 격차가 어떻게 다른지 — 영상의 specific 결과 확인.
• strategy: persistent vs split-K — 큰 K, 작은 M/N 의 사이즈에서 자주 쓰는 split-K 전략. 본 강의가 이 자리를 다뤘는지 미확인.
• cuBLAS internals — 강의가 cuBLAS 의 어느 path 와 비교했는지 (cuBLAS, cuBLASLt, cuBLAS-DX 의 차이). 영상 직접 확인.
• Hopper 의 FP32 accumulator strategy — FP16 input 에 FP32 accumulator 의 정확도/속도 trade-off 가 다뤄졌는가.
• Modal blog 와의 관계 — pranjalssh 가 Modal blog 에 거의 같은 토픽으로 시리즈 글을 게재. 강의가 그 글의 발표인지, 별도 자료인지 — Modal blog 직접 확인 필요.
• Blackwell 으로의 전환 — 강의 시점(2025-03) 기준. Blackwell 의 새 도구들 (TMEM, 5th gen TC) 과 본 강의의 결과가 어떻게 transferable 한지.