# HipKittens: 빠르고 강력한 AMD 커널

> Clean Markdown view of GeekNews topic #24386. Use the original source for factual precision when an external source URL is present.

## Metadata

- GeekNews HTML: [https://news.hada.io/topic?id=24386](https://news.hada.io/topic?id=24386)
- GeekNews Markdown: [https://news.hada.io/topic/24386.md](https://news.hada.io/topic/24386.md)
- Type: GN+
- Author: [neo](https://news.hada.io/@neo)
- Published: 2025-11-16T08:33:08+09:00
- Updated: 2025-11-16T08:33:08+09:00
- Original source: [hazyresearch.stanford.edu](https://hazyresearch.stanford.edu/blog/2025-11-09-hk)
- Points: 3
- Comments: 1

## Topic Body

- **HipKittens**는 AMD GPU용 **고성능 커널과 C++ 기반 프로그래밍 프리미티브**를 제공해 AI 연산 효율을 높이는 프로젝트  
- 기존 AMD 생태계의 **AITER, PyTorch, Triton, TileLang, Composable Kernel** 등이 **성능 불안정과 미성숙한 지원**으로 한계 노출  
- HipKittens는 **타일 기반 추상화(tile abstraction)** 를 중심으로, **NVIDIA와 AMD 간 공통 인터페이스**를 유지하면서 하드웨어별 구현을 분리  
- 약 **500줄 이하의 코드로 작성된 커널**이 기존 AMD의 수작업 어셈블리 커널보다 빠른 성능을 달성  
- AI 연산을 **다중 실리콘 환경으로 확장**하기 위한 실질적 기반을 제시하며, **개방형 하드웨어 생태계**로의 전환 가능성 제시  

---

### AMD GPU 생태계의 현황과 문제점
- AI 연산은 지금까지 **단일 하드웨어 벤더 중심**으로 발전했으나, AMD GPU는 **최신 세대에서 최고 수준의 연산 성능과 메모리 대역폭**을 제공  
  - 예: AMD MI355X OAM은 **BF16 2.5 PFLOPs**, **MXFP8 5.0 PFLOPs**, **메모리 288GB**, **대역폭 8.0TB/s**  
- 그러나 **성숙한 소프트웨어 스택 부재**로 인해 실제 AI 워크플로우에서 이 성능이 활용되지 못함  
- AMD의 주요 소프트웨어 구성요소는 **AITER, PyTorch, Triton, Mojo, TileLang, Composable Kernel (CK)** 등으로 구성  
  - **AITER와 PyTorch**의 SDPA Llama GQA 역전파 커널은 각각 **30%, 24% 수준의 SoTA 대비 성능**만 달성  
  - **Mojo**는 뱅크 충돌(bank conflict)로 인해 **50% 수준의 성능**에 머무름  
  - **TileLang**은 CDNA3까지만 지원하며, **핵심 기능 부족 및 CK 의존성**으로 복잡성 증가  
  - **Triton**은 레지스터 수명 추적과 메모리 접근 최적화에 어려움 존재  
- 결과적으로, **최고 성능의 AMD 커널은 전문가가 어셈블리로 직접 작성해야 하는 상황**이며, 이는 확장성과 유지보수에 큰 제약  

### CUDA 중심 생태계와 비교
- AI 커뮤니티에서는 **CUDA 생태계의 진입장벽(CUDA moat)** 이 존재한다고 평가  
- 과거 NVIDIA 커널 개발도 **저수준 CUDA/CUTLASS 기반의 수년간 노력**이 필요했음  
- 최근에는 **DSL(도메인 특화 언어)** 과 **AI 보조 도구**의 발전으로 NVIDIA 커널 개발이 단순화  
  - 예: **ThunderKittens (TK)** , **CuTe DSL**, **TinyGrad “tinykittens”** , **TileLang**, **Gluon** 등  
- 이러한 흐름을 바탕으로, AMD에서도 **새로운 프로그래밍 프리미티브의 필요성**을 탐색  

### HipKittens의 설계 원리
- HipKittens는 **ThunderKittens(NVIDIA)** , **ThunderMittens(Apple Silicon)** 에 이어 **AMD용으로 개발된 버전**  
- **핵심 개념: 타일 기반 추상화(tile abstraction)**  
  1. **타일 추상화의 범용성** – NVIDIA에서 효과적이었던 타일 기반 연산 모델이 AMD에서도 자연스럽게 적용  
  2. **백엔드 구현의 아키텍처 특화** – 메모리 접근 패턴과 레지스터 스케줄링은 하드웨어별로 다르게 설계  
  3. **스케줄링 전략의 적응성** – AMD의 CDNA3·CDNA4에서는 wave 기반 스케줄링이 비효율적이지만, 타일 단위 스케줄링은 여전히 단순성과 성능을 유지  
- **인터페이스(타일 및 연산)** 와 **구현(하드웨어 매핑)** 을 분리함으로써, **다양한 GPU 아키텍처에 공통 적용 가능한 모델** 제시  

### 성능 결과
- **Attention Forward 커널**: 약 **500줄의 코드**로 작성되었으며, **AITER의 어셈블리 커널보다 우수한 성능** 달성  
  - 다양한 **헤드 차원(D)** 과 **시퀀스 길이(N)** 에서 **Causal/Non-Causal** 설정 모두에서 우위  
- **GEMM 커널**: **100줄 미만의 핵심 루프**로 구성되어, **BF16·FP8 연산 모두에서 최고 성능** 달성  
- **Attention Backward, Rotary, Fused Dropout-Residual-LayerNorm 커널**에서도 **기존 최고 성능 대비 향상** 확인  
- 모든 커널이 **가독성과 수정 용이성**을 유지하면서도 **수작업 어셈블리 수준의 성능** 확보  

### 다중 실리콘 AI로의 확장
- AI의 잠재력을 실현하려면 **다양하고 개방된 하드웨어 활용**이 필요  
- HipKittens는 AMD GPU를 **AI 개발자에게 실질적으로 접근 가능한 플랫폼**으로 만드는 것을 목표  
- AMD의 **오픈소스 소프트웨어 스택**과 결합해, **다중 실리콘 기반 AI 생태계**로의 전환 가능성 제시  
- 후속 글에서는 HipKittens의 **기술적 세부 구조**를 다룸 (part two 예고)

## Comments



### Comment 46363

- Author: neo
- Created: 2025-11-16T08:33:09+09:00
- Points: 1

###### [Hacker News 의견](https://news.ycombinator.com/item?id=45923188) 
- 우리는 AMD와 계약을 맺고 **Llama 405B**를 MI350X에서 MLPerf용으로 학습 중임  
  AMD의 상황이 확실히 좋아지고 있음. AMD GPU가 있다면 [pytorch.org](https://pytorch.org)에서 Linux+ROCm을 클릭해 PyTorch를 설치해볼 것을 추천함. 3년 전만 해도 절망적이었지만 지금은 대부분의 메인라인 기능이 잘 작동함. MI300X에서 nanochat을 돌렸는데 그냥 바로 됐음. MI350X도 마찬가지로 안정적임  
  물론 NVIDIA보다는 뒤처져 있고, **소프트웨어 생태계와 컴파일러, 드라이버**에 대한 투자가 많이 필요함. 하지만 2년 전의 절망적인 상황과 비교하면 지금은 희망이 보임.  
  AMD의 LLVM 백엔드가 부족한 부분을 보고 싶다면 HipKittens 코드를 CUDA Kittens와 비교해보면 좋음.  
  학습용으로는 NVIDIA와 Google이 1위, AMD가 2위, 그 외는 거의 없음. Intel과 Tenstorrent는 아직 멀었고, Huawei는 예제가 segfault로 죽었으며, Groq은 칩 판매를 포기했고, Cerebras는 구할 수조차 없음. Trainium은 인스턴스 하나 받는 데 5일이 걸려서 흥미를 잃었음
  - Tinygrad가 이런 **메모리 최적화나 warp 특화** 같은 기능을 표현하거나 탐색할 수 있으려면 얼마나 멀었는지 궁금함. 이런 기능을 tinygrad에 추가하는 게 얼마나 복잡한지도 알고 싶음
  - 소비자용 하드웨어(비-MI)에서도 ROCm + PyTorch를 돌리려면 **독점 커널 드라이버**가 필요한지 궁금함
  - “Cerebras는 어디서도 구할 수 없다”는 말이 오히려 **승리의 신호**처럼 들림
  - 나는 지난 2년간 AMD NeoCloud의 CEO로 일하고 있음. AMD가 이렇게 **턴어라운드**하는 걸 직접 보고 있어서 매우 반가움.  
    초기 설정은 여전히 약간 거칠지만 예전과 비교하면 훨씬 나아졌음. 예를 들어 한 달 전만 해도 nanochat이 잘 안 됐는데 지금은 잘 됨. 중요한 건 이제 사람들이 AMD 생태계에 관심을 가지고 있다는 점임.  
    AI는 **하드웨어 다양성**이 필요함. 모든 AI 하드웨어와 소프트웨어를 한 회사가 독점하는 건 주주에게는 좋을지 몰라도 기술 발전에는 해로움

- NVIDIA의 **기업가치**가 이해되지 않음. 지금은 Transformer 같은 소수의 알고리즘이 승리했고, 데이터가 더 중요해졌음. 과거처럼 다양한 HPC 코드가 필요하지 않기 때문에, 이제 경쟁사들은 좁은 알고리즘 집합만 최적화하면 됨.  
  vLLM과 Transformer를 효율적으로 돌릴 수 있다면 거대한 시장이 열림. 그렇다면 AMD나 Huawei도 쉽게 따라잡을 수 있을 것 같은데, NVIDIA의 **진짜 해자(moat)** 는 무엇인지 궁금함. InfiniBand만으로 충분한가?
  - 맞음, NVIDIA의 해자는 점점 약해지고 있음. MS, Google, Amazon, Apple 같은 **거대 기업들**이 모두 NVidia 종속을 피하려고 자체 칩을 개발 중임.  
    데이터센터용 GPU에서는 여전히 NVIDIA가 강하지만, Google은 TPU를 대규모로 사용하고 있고 OpenAI도 AMD 하드웨어를 주문함.  
    CUDA 생태계는 여전히 중요하지만, 모두가 거기서 벗어나려는 움직임이 활발함. AMD, Intel, Qualcomm 등도 이 경쟁에 뛰어들고 있음. HipKittens는 **중립적 소프트웨어 생태계**로 가는 중요한 발걸음처럼 보임
  - “소수의 알고리즘”을 구현하는 가장 쉬운 방법은 **범용 연산 하드웨어**를 만드는 것임. 하드웨어는 개발 주기가 길기 때문에 이런 접근이 안정적임
  - InfiniBand는 UEC로 대체되고 있음. **추론(inference)** 에는 InfiniBand가 필요하지 않음. 그래서 추론 시장에는 해자가 없음. AMD나 Google TPU를 쓰는 게 현명함
  - NVIDIA의 진짜 무기는 **방대한 CUDA 라이브러리 생태계**임. 거의 모든 분야에 대응하는 코드가 있음
  - Transformer는 하나의 기술이 아님. 구현 방식이 매우 다양함. 그래서 vLLM이나 TRL이 그렇게 간단하지 않음

- AMD가 이런 프로젝트에 자금을 지원할 법도 한데, 내 생각엔 또 **기회를 놓쳤을 것** 같음. GPU와 AI 관련해서는 늘 그런 식이었음
  - AMD는 제품을 내놓기 위한 최소한의 소프트웨어만 투자함. **성능 테스트 체계조차 부실**하고, 회귀 버그가 그대로 고객에게 전달됨.  
    HipKittens는 개선이지만, AMD 내부에는 커널 성능을 추적할 능력이 부족함. DevOps는 인도 TCS에 외주를 줬고, 그쪽은 상황을 잘 모름.  
    좋은 리더가 있는 팀은 자체적으로 **그림자 IT 팀**을 운영함. ROCm은 그런 팀이 없었고, 결국 대형 클라우드 고객들이 항의하면서 개선이 시작됨.  
    인재를 채용하더라도 **시장 이하의 연봉**을 제시함. Qualcomm이나 Walmart 수준으로 비교하기 때문임.  
    지난 4년간 **보너스가 전액 지급된 적이 없음**
  - “AMD는 기회를 놓치지 않고 놓친다”는 말이 딱 맞음. Instinct 하드웨어는 사실상 **NVidia와 경쟁 가능한 수준**이지만, 소프트웨어 지원이 끔찍했음.  
    과거 Radeon VII처럼 지원을 끊거나, 생태계를 자주 갈아엎어서 성숙하지 못했음. CUDA 호환성 부족과 투자 부재로 인해 대부분의 조직은 그냥 NVIDIA를 선택함.  
    그래도 최근에는 ROCm과 Instinct 생태계에 꾸준히 투자하고 있어서 조금씩 개선되고 있음. 다만 **네트워킹 분야**에서는 여전히 NVIDIA가 훨씬 앞서 있음
  - 성능 비교표를 보면 AMD가 지금쯤 NVIDIA 수준일 수도 있었음. 그런데 **소프트웨어 부재**로 실패함. 칩 설계가 더 어려운 일인데, 정작 소프트웨어를 이해하지 못한 게 문제임
  - 과거 ROCm에 최적화된 커널을 기여한 사람들이 있었는데, AMD가 **PR을 닫고 병합하지 않음**. 정말 이해할 수 없는 행동임
  - 지금은 자금이 지원되고 프로젝트가 **정상 작동 중**임

- [ThunderKittens](https://github.com/HazyResearch/ThunderKittens) 위에 구축된 프로젝트가 있는지 궁금함.  
  이게 HIP으로 포팅된 버전이라면, 단순한 CUDA 포트보다 훨씬 **실용적 가치**가 있을 것 같음

- “raw assembly vs cooked assembly”라는 표현이 흥미로움.  
  예전에는 CPU용 **어셈블리 코드를 직접 작성**하는 게 흔했음. GPU에서도 너무 두려워할 필요는 없음. 결국 컴파일러가 그 코드를 만들어내는 것임

- 수학적으로 보면 **추론(inference)** 은 기본적인 선형대수/BLAS 연산임.  
  dtype과 sparsity를 고려해 80%의 사용 사례를 커버하는 **간결한 추론 API**가 가능할지도 궁금함. CUDA처럼 복잡할 필요는 없을 듯함

- composable-kernel은 내 Gentoo 시스템에서 **OOM(메모리 부족)** 을 가장 자주 일으킨 소프트웨어임. Clang이 CK를 컴파일할 때 스레드당 2.5GB 이상을 사용함
  - Debian용 CK 패키지를 검토했는데, -j32로 빌드하다가 64GB 워크스테이션에서도 OOM이 났음. -j1로 돌리니 190시간 만에 성공함.  
    공식 빌드 서버에서는 아키텍처 수를 줄여야 할 듯함. 대부분의 의존 패키지는 헤더만 필요하다고 들음. 빌드 시간을 줄이기 위한 개선 작업이 진행 중임
  - 단일 커널을 위한 템플릿 인스턴스화에 10분 넘게 걸리는 건 **놀라운 수준**임.  
    `device_grouped_conv2d_fwd_xdl_ngchw_gkcyx_ngkhw_f16_instance`, 도대체 clang에게 무슨 짓을 하는 건지 모르겠음

- 이런 발전이 소비자용 AMD 칩에서도 **LLM을 잘 돌릴 수 있게** 해줄지 궁금함.  
  예를 들어 AMD CPU/iGPU와 RTX 5080이 함께 들어간 HP OMEN MAX 16 노트북을 고려 중인데, AMD 쪽이 RTX와 경쟁할 수 있을까?
  - dGPU가 항상 빠르다고 생각할 수 있지만, **메모리 용량 제한**이 발목을 잡음.  
    [이 블로그 글](https://www.ywian.com/blog/amd-ryzen-ai-max-plus-395-native-gpt-oss-120b-performance-explainer)이나 고급 Mac 결과를 보면 흥미로움
  - 나는 RTX7900XTX에서 Ollama로 Qwen3 Coder 30B를 돌리고 있음. 꽤 잘 작동함. 일부 부하는 시스템 메모리와 Ryzen 7 CPU로 넘어가는 듯함.  
    Sonnet 4 API보다는 약간 느리지만, **로컬에서 이 정도 성능**이면 충분히 만족스러움.  
    Opencode + Ollama + Qwen3 Coder 조합은 **ClaudeCode + Sonnet4**의 훌륭한 대안임.  
    단, AI가 전부 코딩을 대신해야 하는 상황이라면 다르게 느낄 수도 있음. 하지만 개인용 보조로는 완벽함

- AMD가 **B300**을 완전히 무시한 이유를 모르겠음