《GPU Mode》 L097 2025 · 후반 · MLSys 2026 High priority transcript · failed · HipKittens repo + ThunderKittens 자료 보강

HipKittens — ThunderKittens 정신을 AMD GPU 위에

Stanford HazyResearch 의 ThunderKittens (TK) 가 NVIDIA GPU 위에서 “tile primitive” 만으로 fast kernel 을 짜게 했다면, HipKittens 는 같은 정신을 AMD CDNA3/CDNA4 위에 옮긴 작업. RDNA/CDNA 의 wavefront 는 NVIDIA 의 warp 와 다르고, MFMA 는 WGMMA 와 다르다 — 그 차이를 어디까지 같은 추상으로, 어디부터 새 primitive 가 필요한지가 강의의 본질. MLSys 2026 에 발표된 디자인을 William Hu 가 GPU Mode 청중에게 옮긴 강의. transcript 가 실패해서 본 노트는 HazyResearch/HipKittens repo 와 일반 ROCm/CDNA 도메인 지식 기반.

HipKittens ThunderKittens CDNA3 / CDNA4 MI300X MFMA wavefront AITER ROCm

Speaker

William Hu

Stanford · HazyResearch · Simran Arora 와 공저 · ThunderKittens 의 AMD port

강의 번호

L097

스피커

William Hu

학습 우선순위

High · 정독

발표지

MLSys 2026 채택

하위 목차 · 12개 섹션↓ 클릭해서 이동

01강의가 풀려는 문제why this lecture exists 02AMD GPU 위 ThunderKittens 의 정신port the philosophy 03HIP 와 CK 통합ROCm stack 04RDNA / CDNA 차이arch divergence 05example — attentionfwd/bwd 06채택 사례 — AITER backendin the wild 07ROCm stack 안 위치stack positioning 08ThunderKittens 와의 차이divergence 09한계 — RDNA 미지원 · 학습 곡선limitations 10기억할 메모와 실습key takeaways 11다른 강의로 이어지는 길connections 12열린 질문open questions

§ 01강의가 풀려는 문제· why this lecture exists

“NVIDIA 위에서 작동하는 작은 추상이 AMD 위에서도 통하는가”

강의의 핵심 질문 — repo README 가 직접 인용. “우리는 ThunderKittens 의 primitive 가 AMD 위에서도 충분한지, 아니면 완전히 새로운 primitive 가 필요한지를 탐구한다.”

강의의 frame

“좋은 추상은 hardware 의 차이를 가린다 — 그러나 너무 가리면 성능을 못 낸다. ThunderKittens 의 ‘tile’ 추상이 NVIDIA 와 AMD 의 차이를 얼마까지 가리는가, 그리고 어디부터 따로 길을 가야 하는가가 HipKittens 의 본문.”

그리고 이 작업이 의미 있는 이유 — AMD MI300X / MI355X 가 LLM 학습/추론 시장에서 점유율을 늘리고 있지만, kernel 라이브러리 생태계는 NVIDIA 대비 한참 빈약. ThunderKittens 정신을 AMD 에 옮기는 것 자체가 ROCm 생태계의 한 큰 부재를 채운다.

강의의 끝에 손에 잡혀야 할 건 4개의 primitive 카테고리(tile / bulk / async memory / overlap pattern)와 1개의 비교(CDNA 의 MFMA vs Hopper 의 WGMMA). 그리고 “언제 같은 추상으로 통하고 언제 새 길이 필요한가” 의 판단.

“ThunderKittens 가 NVIDIA 위 좋은 추상이라면 — 그 추상이 hardware-specific 인지 아닌지를 AMD port 가 증명한다.”학습 노트

§ 02AMD GPU 위 ThunderKittens 의 정신· port the philosophy

tile 단위 primitive — “coalesced 메모리 + tensor core 친화 size + Python-같은 ergonomic”

ThunderKittens 의 핵심 아이디어를 빠르게 — 그리고 그게 HipKittens 에 어떻게 그대로 / 어떻게 다르게 옮겨졌는지.

tile operations

tensor core 단위에 맞는 사이즈의 small tile (16×16, 32×32 등)
register file 위에 거주 — register tile
coalesced 메모리 access 가 default
address 계산이 minimal — compile-time 결정

bulk compute functions

tile 위 high-level operations — matmul, sum, exp, layernorm
Python-inspired API 외피, 안은 inline assembly + HIP/CUDA
한 줄로 큰 일을 표현. mul(out, a, b) 같은 형태.

async memory ops

direct global → shared 의 async load
token 또는 explicit fence 로 완료 대기
load latency 를 다음 compute 로 hide
NVIDIA 는 cp.async / TMA, AMD 는 buffer load

compute/memory overlap

AMD 의 새 패턴 — 8-wave ping-pong, 4-wave interleave
NVIDIA 의 producer-consumer (warp specialization) 의 AMD 등가
각 wavefront 그룹이 phase 다른 일

이 4가지 카테고리의 공통점 — tile 이 1차 시민이라는 점. shared memory 와 register file 위에 작은 tile 을 두고, 그걸 hardware-friendly 한 단위로 굴린다. 같은 정신이 NVIDIA 와 AMD 에 모두.

차이점이 등장하는 자리 — overlap pattern

NVIDIA Hopper 의 warp specialization (block 의 warp 를 task 별로 분할) 이 AMD 위에서는 — wavefront 의 size 가 다르고, async 의 mechanism 도 다름. HipKittens 는 “8-wave ping-pong” / “4-wave interleave” 라는 새 이름의 패턴을 도입. tk 의 producer-consumer 와 같은 의도지만 hardware 에 맞게 재구성.

§ 03HIP 와 CK 통합· ROCm stack

HipKittens 는 ROCm 의 어디에 들어가는가

AMD 의 GPU 컴퓨팅 stack 의 layered 구조. HipKittens 가 그 안에서 차지하는 위치.

L7 · PyTorch PyTorch ROCm build — model 코드. tensor.cuda() 가 사실상 HIP API 호출.

L6 · AITER backend AMD 의 inference engine. HipKittens 가 AITER 의 backend 로 통합 — repo 가 명시.

L5 · HipKittens tile primitive · bulk ops · async memory · overlap pattern. kernel 작성의 ergonomic layer.

L4 · CK (Composable Kernel) AMD 의 hand-tuned kernel 라이브러리 — CUTLASS 등가물. HipKittens 가 일부 자리에서 CK 와 결합 또는 비교.

L3 · HIP CUDA-like API. 같은 model — kernel<<<...>>> launch · shared memory · barrier. 코드 90% 가 CUDA 와 동일하게 보임.

L2 · ROCm runtime scheduler · driver · memory manager. NVIDIA 의 CUDA runtime 등가.

L1 · GPU hardware CDNA3 (MI300X / MI325X) · CDNA4 (MI350X / MI355X). 각 architecture 별 MFMA 지원 차이.

HIP 의 전제

HIP 는 CUDA 와 거의 source-compatible 한 API. __global__, __shared__, blockIdx 모두 같음. HipKittens 의 코드는 ThunderKittens 와 90% 가 동일해 보인다 — hardware-specific 한 자리만 다름. 이게 port 의 비용을 크게 줄인다.

§ 04RDNA / CDNA 차이· arch divergence

같은 AMD 라도 CDNA 와 RDNA 는 다른 GPU — HipKittens 는 CDNA 위주

AMD GPU 의 두 큰 갈래. compute (CDNA) 와 graphics (RDNA) 가 hardware 가 다르다. HipKittens 는 CDNA 에 집중.

차원

CDNA3/4 (MI300X 등)

RDNA3/4 (Radeon)

사용 자리

data center · LLM 학습/추론

consumer GPU · 게임 + 일부 AI

wavefront size

64 thread (전통 AMD)

32 thread (Wave32 모드)

tensor core

MFMA — 큰 matrix ops, fp16/bf16/fp8

WMMA — 더 작은 matrix, 일부 정밀도

메모리

HBM3/3e — 192GB+, ~5TB/s

GDDR6/7 — 더 작음, 더 느림

CK 지원

우선 지원

제한적

HipKittens 지원

CDNA3 / CDNA4 우선

현재 미지원

대표 칩

MI300X · MI325X · MI350X · MI355X

RX 7900 XT · RX 9070 등

이 비교가 보여주는 사실 — CDNA 와 RDNA 는 wavefront size 부터 다르다. wavefront 64 vs 32 의 차이는 작아 보이지만, kernel primitive 의 거의 모든 자리에 영향. HipKittens 가 CDNA 만 우선 지원하는 결정은 자연스럽다.

MFMA — CDNA 의 핵심

NVIDIA 의 WGMMA 와 등가. fp16/bf16/fp8/int8 의 큰 matrix 한 번에. MFMA instruction 은 wavefront 단위로 issue — 한 wavefront (64 thread) 가 한 MFMA 를 같이 실행. NVIDIA 의 32-thread warp 와 다른 unit. HipKittens 의 tile size · register layout 결정이 이 사실 위에서 다르다.

§ 05example — attention· fwd/bwd

flash attention (forward + backward) 가 CDNA 위에서 도는 풍경

repo 의 가장 큰 example — flash attention 의 forward · backward · multi-head · grouped-query 모두 구현. 이 example 이 HipKittens 의 가치를 가장 명확히 보여준다.

FIG · CDNA3 위 8-wave ping-pong 패턴 (attention)load / compute 의 wavefront 분리

wave 0–3
load group

load K[0]
load V[0]
load K[1]
load V[1]
load K[2]

wave 4–7
compute group

QK^T[0]

PV[0] + softmax

QK^T[1]

PV[1] + softmax

tensor core
MFMA

MFMA chain — busy throughout

load group (wave 0–3) 이 다음 KV tile 을 미리 가져오는 동안 compute group (wave 4–7) 이 이전 tile 위에서 attention 을 계산. tensor core 가 내내 busy. NVIDIA 의 producer-consumer pattern 의 AMD 등가 — “ping-pong” 은 wavefront 사이의 그 핸드오버를 가리킨다.

fwd attention: causal · GQA · head dim 64/128. flash attention 2 의 알고리즘을 CDNA 위에서.
bwd attention: gradient 계산. fwd 보다 메모리 access 패턴이 다양해 더 까다로움.
rotary embedding: 별도 kernel 로 fused.
fused layer norm: 모델 forward 의 표준 자리.
BF16 GEMM: 가장 baseline 한 op — cuBLAS 등가의 hipBLAS 와 비교.

왜 attention 이 첫 example 인가

LLM 추론/학습의 hot path. attention kernel 한 개의 성능이 throughput 의 30–50% 를 결정. AMD 위에서 ROCm 의 native kernel 보다 빠른 attention 이 있다면 — AMD GPU 가 LLM workload 에서 NVIDIA 와 경쟁 가능한 자리에 더 가까워진다. HipKittens 의 의미가 이 한 자리에서 가장 크다.

§ 06채택 사례 — AITER backend· in the wild

AMD 의 inference engine 안에 흡수 — repo 가 명시

repo README 가 자랑하는 사실 — “integrated as an AITER backend”. AITER 는 AMD 의 자체 inference engine. HipKittens 의 kernel 이 그 안에 backend 로 들어갔다는 건 academic project 의 production 채택이라는 의미.

AITER 안 backend 로 통합: AMD 가 직접 채택. HazyResearch 의 academic 작업이 vendor 의 production stack 에.
MLSys 2026 채택: 학술 지면에서도 인정. ThunderKittens 의 정신이 hardware-portable 함의 증거.
오픈소스 채택: PyTorch ROCm 사용자가 직접 갖다 쓸 수 있음. vLLM 의 AMD backend 와 결합 가능성.
학습 자리 부족 채움: AMD 의 ROCm 생태계에 “학생이 따라 짤 수 있는 reference kernel” 가 부족했음. HipKittens 가 그 자리.

“ThunderKittens 가 academic 에서 industry 로 가는 사례를 만들었다면, HipKittens 는 그 사례를 multi-vendor 로 확장한다.”학습 노트

§ 07ROCm stack 안 위치· stack positioning

CK 와 같은 자리 — 그러나 다른 ergonomic

AMD 의 kernel 라이브러리 풍경. 어떤 자리에서 HipKittens 가 무엇과 경쟁/협력하는지.

라이브러리

scope · ergonomic

HipKittens 와의 관계

CK (Composable Kernel)

AMD 의 hand-tuned kernel — CUTLASS 등가

같은 자리 · 다른 추상 (CK 는 lower-level template)

hipBLAS / hipBLASLt

cuBLAS 의 AMD 등가 · GEMM 위주

HipKittens 가 attention 같은 fused 자리 보강

hipDNN

cuDNN 등가 — convolution / RNN

scope 다름 (HipKittens 는 LLM 위주)

Triton (AMD backend)

Triton 의 AMD GPU lowering — DSL

같은 ergonomic 영역 · DSL vs C++ 차이

AITER

AMD 의 inference engine

HipKittens 가 backend 로 통합

MIGraphX

AMD 의 graph compiler · ONNX 위주

scope 다름 (HipKittens 는 hand-written kernel)

HipKittens 의 차별 — “C++ template 으로 짜는 kernel 의 ergonomic 을 Python-같이”. CK 가 bare-metal C++ template 의 무거움이라면, HipKittens 는 그 위 layer 의 사용자 친화 wrapper. 학생이 일주일에 attention 을 짤 수 있게가 디자인 목표.

§ 08ThunderKittens 와의 차이· divergence

같은 정신 · 다른 hardware-specific 디테일

차원

ThunderKittens (NV)

HipKittens (AMD)

target arch

Hopper / Blackwell · Ampere

CDNA3 / CDNA4

execution unit

warp = 32 thread

wavefront = 64 thread

tensor core ISA

WGMMA · WMMA

MFMA

async mem

cp.async · TMA

buffer load · global async

overlap pattern

warp specialization · producer-consumer

8-wave ping-pong · 4-wave interleave

tile size

16×16, 32×32 (warp friendly)

16×16, 32×32 (wavefront friendly · MFMA shape)

CUDA register file

CDNA register file (조금 큼)

동일한 부분

tile API · bulk ops · Python-style ergonomic · attention example 의 알고리즘

차이의 분포

tile API 와 bulk ops 의 표면 — 거의 동일. 사용자 코드는 90% 같은 형태. 차이는 — overlap pattern 의 “8-wave ping-pong” 같은 새 이름, MFMA instruction 의 직접 호출, async load 의 mechanism. “정신은 같고 hardware-near 자리만 다르다” 의 깨끗한 예시.

§ 09한계 — RDNA 미지원 · 학습 곡선· limitations

현재의 한계와 그 이유

RDNA 미지원 — wavefront size 32 와 64 의 차이로 같은 코드가 안 도는 영역이 큼. consumer GPU 사용자가 직접 시도하기 어렵다.
CDNA3/CDNA4 우선 — 이전 세대 (MI100, MI200) 에서는 일부 instruction 미동작.
학습 자료 부족 — ThunderKittens 의 풍부한 tutorial 대비 HipKittens 의 자료가 적다.
학습 곡선 — wavefront 64 의 mental model, MFMA shape constraint, overlap pattern 의 새 패턴 — NVIDIA 경험자도 다시 학습 필요.
kernel coverage — attention · GEMM · layer norm · rotary 까지. 더 다양한 op 는 미래.
portability — 같은 HipKittens 코드가 NVIDIA 위에서 안 돔. ThunderKittens 와 코드 공유 어려움.

언제 HipKittens 를 쓸 가치

1. AMD MI300X / MI355X 위에서 LLM workload — 가장 자연스러운 사용처. 2. ROCm 사용자 + custom kernel 필요 — CK template 직접 짜기보다 ergonomic. 3. attention 같은 hot path 의 직접 튜닝. AMD 가 아니거나 단순 LLM 추론이면 vLLM-AMD 가 충분.

§ 10기억할 메모와 실습· key takeaways

다시 열었을 때 5분 안에 잡혀야 할 것

tile primitive

bulk ops

tile 위 high-level op — matmul · sum · exp · layernorm. inline assembly 외피 위 Python-같은 API.

async memory

global → shared 의 비차단 transfer. NV 는 cp.async/TMA, AMD 는 buffer load.

8-wave ping-pong

AMD 위 producer-consumer 의 새 이름. wave 0–3 load, wave 4–7 compute.

CDNA vs RDNA

wavefront 64 vs 32, HBM vs GDDR. HipKittens 는 CDNA 위주.

MFMA = WGMMA 등가

CDNA 의 큰 matrix instruction. wavefront 단위로 issue.

AITER backend 통합

academic 작업의 vendor production 채택. ROCm 생태계의 의미 있는 진전.

ThunderKittens 와 90% 동일

tile API · bulk ops 표면이 같음. 차이는 hardware-near 자리만.

실습 시퀀스

HipKittens build — repo clone, ROCm 환경에서 build. MI300X cloud instance (RunPod / TensorWave 등) 에서 시도.
BF16 GEMM 비교 — HipKittens GEMM 과 hipBLASLt 의 throughput 비교. 같은 size 에서.
flash attention 비교 — HipKittens attention 과 ROCm 의 native flash attention 의 throughput · accuracy.
NVIDIA 와 비교 — H100 위 ThunderKittens 의 attention 과 MI300X 위 HipKittens 의 attention 을 같은 input size 로. peak vs achieved.
코드 diff — TK 의 attention 코드와 HK 의 attention 코드를 line by line. 어디가 같고 어디가 다른지.

YouTubeL097 강의 영상

RepoHazyResearch/HipKittens

SisterThunderKittens (NVIDIA)

PaperMLSys 2026 채택 — “HipKittens: Tiles on AMD” (정확 제목 확인 필요)

ReferenceROCm/composable_kernel · ROCm/AITER

§ 11다른 강의로 이어지는 길· connections

kernel / multi-vendor 시리즈 안에서

L001

How to profile CUDA kernels in PyTorch

profiling 의 기본 — HipKittens 도 NCU 등가의 rocprof 로 측정

L025

Speaking Composable Kernel (CK)

AMD CK 의 deep dive — § 07 의 baseline

L096

TLX

NVIDIA Hopper 의 hardware-near layer — HipKittens 의 자매

L018

Fusing Kernels

attention 같은 fused kernel 의 디자인

L094

tvm-ffi

multi-framework / multi-vendor kernel 배포의 layer

§ 12열린 질문· open questions

다음에 다시 들었을 때 직접 검증해야 할 것들

transcript 가 실패해서 본 노트의 비교 표 · 패턴 이름 (8-wave ping-pong 등) 은 repo README + 일반 CDNA 도메인 지식 결합. 강의에서 다른 framing 가능성.
정확한 throughput 비교 숫자 — vs CK · vs hipBLASLt · vs ThunderKittens 의 정량 비교가 강의에 등장했을 가능성.
RDNA 지원 로드맵 — wavefront 32 모드 지원이 향후 계획에 있는지.
새 primitive 의 디테일 — “8-wave ping-pong” 외에 ThunderKittens 에 없는 어떤 새 primitive 가 도입됐는지의 정확한 목록.
MFMA 의 정확한 size 매핑 — TK 의 16×16 / 32×32 와 HK 의 MFMA 16×16×16 / 32×32×8 사이 어떤 mapping 이 자연스러운지.
fp8 지원 — CDNA4 (MI350X) 의 fp8 MFMA 지원이 HipKittens 에 어떻게 들어왔는지.
academic vs production 의 gap — AITER backend 통합 시 어떤 부분을 추가로 다듬어야 했는지.

검증 메모

본 노트의 stack 위치 · primitive 분류 · attention pipeline 시각화는 HazyResearch/HipKittens 의 README 와 일반 ROCm/CDNA 도메인 지식의 결합이다. 강의에서 보여준 정확한 기술 디테일 (instruction 이름, register layout, 정확한 throughput 수치) 은 영상 직접 시청 후 보강 필요. 특히 “8-wave ping-pong” / “4-wave interleave” 의 정확한 작동은 repo 의 attention 코드 직접 참조.

← Lecture 096TLX