HipKittens: 빠르고 강력한 AMD 커널

 (hazyresearch.stanford.edu)
3P by GN⁺ | ★ favorite | 댓글 1개
  • 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 예고)

함께 보면 좋은 글 β

GN⁺   [-]
Hacker News 의견

  • 우리는 AMD와 계약을 맺고 Llama 405B를 MI350X에서 MLPerf용으로 학습 중임
    AMD의 상황이 확실히 좋아지고 있음. AMD GPU가 있다면 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 위에 구축된 프로젝트가 있는지 궁금함.
    이게 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가 항상 빠르다고 생각할 수 있지만, 메모리 용량 제한이 발목을 잡음.
      이 블로그 글이나 고급 Mac 결과를 보면 흥미로움
    • 나는 RTX7900XTX에서 Ollama로 Qwen3 Coder 30B를 돌리고 있음. 꽤 잘 작동함. 일부 부하는 시스템 메모리와 Ryzen 7 CPU로 넘어가는 듯함.
      Sonnet 4 API보다는 약간 느리지만, 로컬에서 이 정도 성능이면 충분히 만족스러움.
      Opencode + Ollama + Qwen3 Coder 조합은 ClaudeCode + Sonnet4의 훌륭한 대안임.
      단, AI가 전부 코딩을 대신해야 하는 상황이라면 다르게 느낄 수도 있음. 하지만 개인용 보조로는 완벽함

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

답변달기