# GPU의 폭발적인 성능 향상

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

## Metadata

- GeekNews HTML: [https://news.hada.io/topic?id=14783](https://news.hada.io/topic?id=14783)
- GeekNews Markdown: [https://news.hada.io/topic/14783.md](https://news.hada.io/topic/14783.md)
- Type: GN+
- Author: [neo](https://news.hada.io/@neo)
- Published: 2024-05-13T10:18:48+09:00
- Updated: 2024-05-13T10:18:48+09:00
- Original source: [hazyresearch.stanford.edu](https://hazyresearch.stanford.edu/blog/2024-05-12-tk)
- Points: 1
- Comments: 1

## Topic Body

### H100 GPU의 특징

- 80GB HBM3 메모리와 3TB/s 대역폭 제공 (실제로는 조금 더 낮음)
- 50MB L2 캐시와 12TB/s 대역폭 제공. GPU에 두 개의 25MB 섹션으로 분할되어 크로스바로 연결됨 (크로스바는 성능 저하 요인)
- 132개의 Streaming Multiprocessor(SM)로 구성되며, 각 SM은 아래 구성을 가짐:
  - 256KB L1 캐시 내에 최대 227KB의 공유 메모리 제공 (합쳐서 약 33TB/s 대역폭)
  - 비동기 주소 생성 및 메모리 fetching이 가능한 Tensor Memory Accelerator(TMA) 제공. 온칩 메모리 네트워크 지원 등의 기능도 제공하지만 이번 포스트에서는 다루지 않음
  - 4개의 쿼드런트로 나뉘며, 각 쿼드런트는 warp scheduler, 512개 vector register(각각 32개의 4-byte 워드 포함), matrix multiply용 tensor core, sum/multiply 등의 병렬 연산을 지원하는 내장 instruction으로 구성됨

### H100 GPU 성능 최적화 팁

- Tensor Core를 최대한 활용하는 것이 중요. H100은 989 TFLOPs의 FP16 행렬곱 연산 성능을 가지므로, Tensor Core 활용도에 따라 전체 GPU 활용률이 크게 좌우됨
- 단, Tensor Core를 최대한 활용하기 위해서는 아래 사항들을 고려해야 함:
  - Warp Group Matrix Multiply Accumulate(WGMMA) 명령어 활용이 필수적이지만 사용이 까다로움
  - Shared Memory가 생각보다 빠르지 않고, 사용시 세심한 주의가 필요함
  - 주소 생성이 비용이 크므로 Tensor Memory Accelerator(TMA) 등을 활용해 최적화해야 함 
  - Occupancy를 높이는 것이 여전히 도움이 되며, register가 주요 리소스임
- 이런 특징들은 H100 뿐만 아니라 다른 GPU에도 어느정도 적용되지만, 특히 H100에서 Tensor Core 활용이 중요하고 까다로운 편

### ThunderKittens : H100에 최적화된 CUDA 임베디드 DSL

- NVIDIA H100 등 최신 GPU의 성능을 최대한 끌어내기 위해 개발된 CUDA 기반 임베디드 DSL
- 아래 4개의 타일 기반 템플릿 타입을 제공:
  - Register Tile (레지스터에 저장된 2D 텐서)
  - Register Vector (레지스터에 저장된 1D 텐서)  
  - Shared Tile (Shared Memory에 저장된 2D 텐서)
  - Shared Vector (Shared Memory에 저장된 1D 텐서)
- 또한 타일 조작을 위한 다양한 연산자 (exp, mul, sum 등)를 warp 또는 warp 그룹 레벨에서 제공
- 기존 Flash Attention, Flash Attention 2 커널을 ThunderKittens으로 구현하니 코드가 크게 간결해지고, H100에서 최대 30% 성능 향상
- Based Linear Attention 커널도 ThunderKittens으로 구현하여 215 TFLOPs 성능 달성 (알고리즘 특성 상 recompute 포함하면 300 TFLOPs 이상)

### 철학적 관점에서의 고찰

- ThunderKittens이 잘 작동하는 이유는 모든 것을 다 지원하려 하지 않고, 단순하면서도 GPU 아키텍처에 잘 들어맞는 타일 기반 추상화를 제공하기 때문
- 전통적인 32-bit 워드 대신 1024-bit 벡터 레지스터를 사용하는 것도 발전이지만, 16x16 타일을 레지스터의 단위로 보는 패러다임의 전환이 필요
- AI 워크로드가 결국 행렬곱, reduction, reshape 위주라는 점에서 타일 기반 접근이 타당하며, 하드웨어 관점에서도 systolic array 외에 작은 행렬곱을 지원하는 방향으로 진화할 것
- 더 나아가 하드웨어에 최적화된 형태로 AI 알고리즘을 설계하는 방식으로 사고를 전환할 필요가 있음. 예를 들어 RNN의 상태 크기를 SM에 들어갈 만큼으로 제한하고, 연산 밀도도 하드웨어가 요구하는 수준으로 맞추는 식의 접근이 필요

### GN⁺의 의견

- ThunderKittens은 CUDA에 익숙한 개발자들에게는 매력적인 선택지가 될 수 있음. 기존 커널 코드를 크게 건드리지 않고도 쉽게 성능 향상을 누릴 수 있기 때문
- 다만 초보자에게는 여전히 진입장벽이 높을 수 있음. 앞으로 더 다양한 예제 코드와 학습 자료가 뒷받침되어야 할 듯
- H100 뿐만 아니라 AMD 등 다른 GPU에도 ThunderKittens 지원을 확대한다는 계획이 흥미로움. 벤더 종속성을 낮추는데 기여할 수 있을 것으로 보임
- 궁극적으로는 AI 모델/알고리즘 자체를 하드웨어에 최적화된 형태로 설계하는 것이 매우 중요한 포인트임. 이를 위해서는 하드웨어 특성에 대한 깊은 이해가 선행되어야 하는데, ThunderKittens은 개발자들이 그런 인사이트를 얻는데 도움이 될 수 있을 것
- Hazy Research의 지속적인 연구개발과 오픈소스 기여가 CUDA 생태계 활성화에 큰 도움이 될 것으로 기대됨. 다만 장기적으로는 보다 추상화 레벨이 높은 프레임워크의 등장도 필요해 보임

## Comments



### Comment 25184

- Author: neo
- Created: 2024-05-13T10:19:29+09:00
- Points: 1

###### [Hacker News 의견](https://news.ycombinator.com/item?id=40337936) 
- AI 하드웨어의 요구 사항이 점점 명확해지고 있음. GPU는 원래 다른 목적으로 설계되었지만, 좋은 행렬 곱셈 하드웨어를 가지고 있어 AI에 사용되고 있음. "AI GPU"는 실제 GPU의 일부 기능을 제외할 수 있으며, 더 짧은 숫자(16비트, 8비트, 2비트, 1비트 부동 소수점)로 향하는 추세가 있음. 이 논문은 16x16 타일을 좋아하는 하드웨어가 많은 이점이 있음을 시사함.

- AI 전용 보조 프로세서(NPU)가 필요함. 특히 개발자, 전문가, 게이머 등을 위한 프로슈머급 데스크톱 시스템에서 필요함. GPU는 기업에서는 작동하지만 개인 컴퓨팅 측면에서 AI에 사용하기에는 불편함. 특히 VRAM 제한과 Vulkan 이외의 표준 개방형 API 부재가 문제임.

- AI 연구를 발전시키려면 신경과학, 심리학을 더 잘 연구해야 함. 또한 신경망의 그래프 토폴로지와 관련된 것들도 관련이 있을 수 있음. 

- CUTLASS를 사용자 친화적으로 만든 것인가?

- ThunderKittens 마스코트는 고양이/소니 아이보 분위기가 있음. AI로 잘 생성된 것으로 보임.

- 유니버설 베이직 컴퓨트(UBC)가 유니버설 베이직 인컴의 대체재로 고려된다면, 그것은 매우 디스토피아적인 미래가 될 것임. 엔비디아 같은 한 회사가 모든 컴퓨팅을 만든다고 상상해보라.