# AMD GPU가 ‘brrr’ 속도로 돌아가게 만드는 방법

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

## Metadata

- GeekNews HTML: [https://news.hada.io/topic?id=24390](https://news.hada.io/topic?id=24390)
- GeekNews Markdown: [https://news.hada.io/topic/24390.md](https://news.hada.io/topic/24390.md)
- Type: GN+
- Author: [neo](https://news.hada.io/@neo)
- Published: 2025-11-16T09:41:59+09:00
- Updated: 2025-11-16T09:41:59+09:00
- Original source: [hazyresearch.stanford.edu](https://hazyresearch.stanford.edu/blog/2025-11-09-amd-brr)
- Points: 1
- Comments: 1

## Topic Body

- **HipKittens**는 AMD GPU의 잠재 성능을 끌어내기 위해 설계된 **프로그래밍 프리미티브 모음**으로, 메모리 접근·스케줄링·캐시 재사용을 최적화함  
- AMD MI355X GPU는 **256개의 컴퓨트 유닛과 8개의 칩렛(XCD)** 구조를 가지며, **대형 레지스터 파일**과 **세밀한 매트릭스 코어 명령어**를 제공  
- NVIDIA와 달리 AMD는 **레지스터 재할당·비동기 행렬 명령어·mbarrier**가 없어, **wave specialization** 대신 **8-wave ping-pong**과 **4-wave interleave** 스케줄링이 효과적임  
- **HipKittens**는 **칩렛 인식(grid) 스케줄링**을 통해 L2 및 LLC 캐시 지역성을 개선, GEMM과 Attention 연산에서 **최대 대역폭 및 TFLOPS 향상**을 달성  
- 이 접근은 **AMD GPU 생태계의 소프트웨어 성숙도 부족**을 보완하며, **다양한 하드웨어 기반의 AI 컴퓨팅 확장성**을 높이는 기반 제공  

---

### AMD CDNA GPU 구조와 성능 특성
- AMD MI355X GPU는 **256개의 컴퓨트 유닛(CU)** 을 포함하며, 각 CU는 4개의 SIMD로 구성  
  - 한 SIMD는 64스레드로 구성된 **wave**를 실행하며, 이는 NVIDIA의 32스레드 warp와 대비됨  
- MI355X는 **B200 대비 70% 수준의 SRAM(165KB)** 을 가지며, **비동기 행렬 곱 명령어, 레지스터 재할당, 텐서 메모리 가속, mbarrier** 기능이 없음  
- 반면, **2배 큰 레지스터 파일**과 **60% 더 많은 프로세서 수(256 CU 대 160 SM)** 를 제공  
  - **작고 세밀한 매트릭스 코어 명령어**를 지원하며, **직접 글로벌→공유 메모리 로드(TMA 유사)** 기능 존재  
- AMD는 **8개의 칩렛(XCD)** 로 구성된 **칩렛 아키텍처**를 채택, 각 XCD는 독립 L2 캐시를 가지며, 상위에 **LLC 캐시**가 존재  
- 표에 따르면 MI355X는 **BF16 2.5 PFLOPs, MXFP8 5.0 PFLOPs, MXFP6 10.1 PFLOPs**의 연산 성능과 **288GB 메모리 용량, 8TB/s 대역폭**을 가짐  

### AMD용 커널 설계 과제
- **메모리 접근 최적화**: HIPCC 컴파일러 제약과 비공개 I/O 동작으로 인해 **데이터 배치와 스위즐(swizzle) 패턴** 설계가 중요  
- **프로세서 내 스케줄링**: AMD는 **공유 메모리 대신 레지스터 파일과 소형 매트릭스 명령어**를 활용해야 함  
- **프로세서 간 스케줄링**: 칩렛 기반 구조로 인해 **캐시 수준의 NUMA 효과**를 고려한 작업 분배 필요  

### HipKittens의 메모리 접근 패턴
- HipKittens(HK)는 **타일(tile)** 을 기본 데이터 단위로 사용하며, **PyTorch 유사 연산 함수**를 제공  
  - 타일은 **데이터 타입, 크기, 레이아웃**으로 정의되며, **C++ 템플릿 메타프로그래밍**으로 다양한 입력에 대응  
- **레지스터 스케줄링**: HIPCC는 특정 레지스터를 MFMA 입력으로 사용할 수 없어, HK는 **명시적 레지스터 고정 기능**을 제공  
  - 개발자가 직접 레지스터를 지정해 **최대 성능 커널 작성 가능**  
- **레지스터 레이아웃**: AMD는 데이터 타입·행렬 형태에 따라 레이아웃이 달라 **단일 스위즐 패턴 불가**  
  - 예시로 16×16 bf16 타일과 16×32 bf16 타일은 서로 다른 스위즐 패턴 필요  
- **명령어 페이즈 구조**: AMD의 공유 메모리 명령어는 **비연속적 페이즈 그룹**을 가지며, **내부 문서화 부족**  
  - HK는 이를 **리버스 엔지니어링한 솔버**를 제공  
- **주소 생성**: AMD는 **비동기 HBM→공유 메모리 로드**를 지원하며, **HBM 주소 스위즐**로 최적화 수행  

### 프로세서 내 스케줄링: Wave 패턴
- **Wave specialization**은 NVIDIA에서 효과적이지만, AMD는 **레지스터 재할당 부재**로 인해 성능 저하  
  - Producer wave가 불필요한 레지스터를 점유하고, Consumer wave는 부족한 레지스터로 인해 **스필(spill)** 발생  
- HK의 실험 결과, **wave specialization**은 AMD에서 **산술 강도 감소 및 메모리 병목**을 초래  
  - 예시: GEMM에서 HK 0/8 구성은 1605 TFLOPs, CUTLASS는 1570 TFLOPs  
- **대안 스케줄링 패턴**
  - **8-wave ping-pong**: 두 wave가 번갈아 **메모리/연산 클러스터**를 실행  
  - **4-wave interleave**: 한 wave가 **세밀하게 메모리와 연산을 교차 실행**  
  - 8-wave는 코드가 간결하고, 4-wave는 세밀하지만 코드가 길어짐  
  - GEMM과 Attention Forward에서 **8-wave가 SoTA 수준 성능** 달성  

### 프로세서 간 스케줄링: 칩렛 인식 접근
- AMD MI355X는 **8개의 XCD 칩렛**을 가지며, 각 칩렛은 **독립 L2 캐시**를 보유  
  - 스레드 블록이 **라운드로빈 방식**으로 칩렛에 배정되어, **그리드 순서가 캐시 재사용 효율**에 직접 영향  
- 단순 **row-major 배치**는 L2 캐시 재사용률이 낮아 **대역폭 손실** 발생  
  - 예: L2 55%, LLC 95%, 15.1 TB/s, 1113 TFLOPs  
- HK는 **칩렛 인식(grid) 스케줄링**을 도입, **L2·LLC 캐시 지역성**을 동시에 활용  
  - 스레드 블록을 **출력 행렬의 인접 영역 단위로 그룹화**하여 입력 데이터 재사용 극대화  

### 실제 커널 예시
- Attention Forward 및 BF16 GEMM 커널의 **핫 루프(hot loop)** 는 HK의 **8-wave ping-pong 스케줄**을 사용  
  - 각 루프는 **Compute–Memory 클러스터**를 번갈아 실행하며, **스케줄 배리어**로 동기화  
  - 코드 예시에서 **mma_AtB, load, exp2, col_sum** 등 HK 연산이 반복적으로 사용됨  

### 결론: Multi-silicon AI 시대의 AMD
- HipKittens는 AMD CDNA3·CDNA4에서 **경쟁력 있는 성능**을 달성  
  - 세 가지 핵심: **최적화된 메모리 접근**, **AMD 중심 wave 스케줄링**, **칩렛 인식 그리드 스케줄링**  
- HK 커널은 **AMD 기준 최고 성능**을 달성하며, **NVIDIA Blackwell 커널과도 경쟁 수준**  
- AI 컴퓨팅의 다양성을 위해 **AMD GPU 접근성 확대**가 필요하며, HipKittens는 이를 위한 **핵심 소프트웨어 기반** 제공  
- AMD의 **HIPCC 레지스터 스케줄링 개선**이 향후 중요한 발전 영역으로 지목됨

## Comments



### Comment 46366

- Author: neo
- Created: 2025-11-16T09:41:59+09:00
- Points: 1

###### [Hacker News 의견](https://news.ycombinator.com/item?id=45934416) 
- [HipKittens 관련 토론](https://news.ycombinator.com/item?id=45923188)을 참고하길 권함  
- 같은 연구를 다룬 [HipKittens: Fast and furious AMD kernels](https://news.ycombinator.com/item?id=45923188) 글도 있음. George Hotz와 AMD 직원들의 댓글이 달려 있음  
- 학계에서 이런 문제를 다루는 건 반갑지만, 이건 결국 **AMD 내부에서 해결해야 할 문제**라고 생각함  
  - 나는 하드웨어 회사는 하드웨어만 만드는 게 좋다고 봄. 그래야 **인센티브가 순수**하게 유지됨. 성능이 20% 떨어지더라도 그게 낫다고 생각함  
  - 완전히 동의함. AMD는 10년 전 이 문제를 미뤘고 이제서야 따라잡으려 함. 하드웨어는 훌륭하지만 **펌웨어 작성 능력 부족**으로 잠재력을 못 살리고 있음  
  - 하지만 이 연구팀은 Nvidia GPU용으로도 비슷한 소프트웨어를 만든 적이 있음. 뛰어난 연구자들이 자기 전문성을 발휘하는 것 같음  
  - 내가 알기로 AMD는 이미 여러 수준에서 이 문제를 다루고 있고, **tinycorp**와도 협력 중임  
- 글을 보면 AMD GPU의 **아키텍처적 복잡성** 때문에 최적화가 어렵다는 인상을 받음. 하지만 장기적으로는 AMD의 접근이 더 잘 확장될 수도 있음. Nvidia가 2개 칩렛을 쓰는 반면 AMD는 8개 칩렛 구조라 **메모리 지역성 문제**가 있음. 미래에는 칩렛 수가 더 늘어날 테니, 지금의 복잡성을 다루는 경험이 장기적으로 도움이 될 수도 있음  
  - AMD는 고성능을 위해 **warp specialization**이 필요 없어서 프로그래밍이 더 단순함  
- 많은 개발자들이 AMD GPU를 대중 개발자용으로 ‘go brrr’하게 만들려 했지만 실패했음. AMD가 왜 자체적으로 **소프트웨어 문제를 해결하지 않는지** 이해가 안 됨. 지금은 돈도 충분한데 개발자 고용을 안 하는 건 변명이 안 됨. 데이터센터용 GPU도 나쁘진 않지만, 개인이 ML·AI 실험을 할 때는 여전히 Nvidia가 훨씬 나음. 내 5년 된 RTX 3090이 지금까지 나온 AMD 소비자용 GPU보다 낫다고 느낌  
  - AMD 개발자 경험은 끔찍함. **드라이버 크래시 버그 리포트조차 받지 않음**  
  - 나는 최근 NVidia 5090에서 AMD R9700 32GB 두 장으로 추론 서버를 바꿨는데, 완전히 긍정적인 경험이었음. Fedora 커널에서 DKMS 설정 없이 바로 작동했고, ROCm으로 컨테이너 연결도 쉬웠음. Ollama와 Storyteller 설정만 바꾸면 끝이었음. **CUDA보다 훨씬 쾌적한 경험**이었음  
  - Nvidia는 심지어 **Unreal Engine 포크**까지 직접 유지함. AMD는 경쟁조차 안 되는 수준임  
  - Nvidia는 하드웨어 회사 중 유일하게 **소프트웨어 엔지니어에게 경쟁력 있는 보상**을 제공함. AMD는 여전히 소프트웨어를 ‘진짜 일’로 보지 않는 문화가 남아 있고, 이런 관성은 바꾸기 어려움  
- **Mojo**가 AMD GPU에서 개발자 경험(devX)을 개선하려는 아이디어를 갖고 있었는데, 그 진행 상황이 궁금함  
- AMD가 소프트웨어 개선에 **수십억 달러를 투자하지 않는 게 이해가 안 됨**. Nvidia는 세계에서 가장 가치 있는 회사이고, AMD는 유일한 경쟁자임  
  - AMD도 노력 중이지만, 매년 하드웨어를 갱신하는 조직 문화를 **소프트웨어 중심 문화로 전환**하는 건 어렵다고 봄. 소프트웨어는 하드웨어처럼 바로 수익을 내지 않기 때문에 경영진이 우선순위를 낮게 두는 경향이 있음. 또 외부 벤더들이 오픈소스로 코드를 제공하면서 단기적으로는 좋아 보여도 장기 품질에는 악영향을 줌. 하드웨어 트렌드를 한 번만 놓쳐도 경쟁사에 뒤처질 위험이 큼  
  - 여러 GPU 벤더에서 일해봤는데, Nvidia만이 소프트웨어를 **자산(asset)** 으로 보고 투자함. 다른 회사들은 비용으로만 여김  
- “go brr” 밈은 개인적으로 좋아하진 않지만, Stanford 같은 곳에서 쓰이는 걸 보니 재밌음  
  - 사실 이미 1년 전에 [ThunderKittens 발표](https://hazyresearch.stanford.edu/blog/2024-05-12-tk) 때 “go brr”를 썼음  
  - 이런 밈이 대학 공식 채널에 등장했다면, 이미 **유행이 끝났다는 신호**일 수도 있음  
- 프로젝트 자체는 훌륭하지만, 왜 AMD가 이런 걸 직접 하지 않는지 의문임. **성숙한 소프트웨어 스택의 중요성**을 AMD가 여전히 이해하지 못하는 것 같음. CUDA처럼 모든 카드에서 쓸 수 있는 통합 스택이 필요함. 한때 AMD가 언젠가 따라잡을 거라 믿었지만, 이제는 거의 포기한 상태임  
- 프로젝트는 좋지만, 글 자체는 **이상하게 작성된 느낌**이 듦  
  - 글이 너무 어색함. AI에 과도하게 의존했거나, AI 문체를 흉내 낸 것 같음. “part one을 확인하라”거나 “AMD GPU를 go brr 하게 만드는 법” 같은 문장이 반복됨. 기술적인 부분도 그래프로 설명해야 할 내용을 **100줄짜리 코드로 풀어쓴** 점이 특히 아쉬움