▲

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

 (hazyresearch.stanford.edu)
1P by GN⁺ 5달전 | ★ favorite | 댓글 1개
  • 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 레지스터 스케줄링 개선이 향후 중요한 발전 영역으로 지목됨

함께 보면 좋은 글 β

▲
GN⁺ 5달전  [-]
Hacker News 의견
  • HipKittens 관련 토론을 참고하길 권함
  • 같은 연구를 다룬 HipKittens: Fast and furious AMD kernels 글도 있음. 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 발표 때 “go brr”를 썼음
    • 이런 밈이 대학 공식 채널에 등장했다면, 이미 유행이 끝났다는 신호일 수도 있음
  • 프로젝트 자체는 훌륭하지만, 왜 AMD가 이런 걸 직접 하지 않는지 의문임. 성숙한 소프트웨어 스택의 중요성을 AMD가 여전히 이해하지 못하는 것 같음. CUDA처럼 모든 카드에서 쓸 수 있는 통합 스택이 필요함. 한때 AMD가 언젠가 따라잡을 거라 믿었지만, 이제는 거의 포기한 상태임
  • 프로젝트는 좋지만, 글 자체는 이상하게 작성된 느낌이 듦
    • 글이 너무 어색함. AI에 과도하게 의존했거나, AI 문체를 흉내 낸 것 같음. “part one을 확인하라”거나 “AMD GPU를 go brr 하게 만드는 법” 같은 문장이 반복됨. 기술적인 부분도 그래프로 설명해야 할 내용을 100줄짜리 코드로 풀어쓴 점이 특히 아쉬움
답변달기