GPU의 폭발적인 성능 향상
(hazyresearch.stanford.edu)- AI 연산 비용이 커지는 상황에서 Hazy Research는 NVIDIA H100의 tensor core를 쉬지 않게 유지하는 것이 GPU 성능 최적화의 핵심이라고 정리함
- H100은 half-precision 행렬곱에서 989 TFLOPs를 내지만 일반 연산은 약 60 TFLOPs에 그쳐, tensor core가 멈추는 순간 활용률이 크게 떨어짐
- 최대 성능에 가까워지려면 WGMMA, shared memory 배치, 주소 생성, occupancy를 함께 다뤄야 하며
wgmma.mma_async없이는 마이크로벤치마크에서 피크의 약 63%에 머무름 - 공개된 CUDA 내장 DSL ThunderKittens는 tile·vector 추상화로 swizzling과 register layout 같은 복잡성을 감싸 FlashAttention 계열 커널 작성을 단순화함
- H100용 FlashAttention-2 forward 커널은 약 100줄로 작성되며 FlashAttention-2보다 약 30% 빠르고, Based linear attention 커널은 215 TFLOPs로 동작함
H100 성능을 좌우하는 조건
- AI는 많은 컴퓨트를 사용하며, Hazy Research는 최근 몇 년간 AI가 더 적은 컴퓨트를 쓰거나 주어진 컴퓨트에서 더 효율적으로 돌도록 하는 작업을 해옴
- 컴퓨트 절감 예시로 Based, Monarch Mixer, H3, Hyena, S4가 있음
- 효율적 실행 예시로 FlashAttention, FlashAttention-2, FlashFFTConv가 있음
- 실용적 목표는 GPU를 빠르게 만드는 데서 배운 내용을 정리하고, 빠른 커널 작성을 돕는 CUDA 내장 DSL ThunderKittens를 공개하는 것임
- 더 넓게는 하드웨어 이해가 AI 컴퓨트를 바라보는 방식을 어떻게 바꿨는지 다룸
NVIDIA H100의 구조와 병목
- H100 SXM GPU는 다음 구성을 기준으로 논의됨
- 80GB HBM3, 대역폭 3TB/s
- 50MB L2 캐시, 대역폭 12TB/s, GPU 전반에 25MB 섹션 2개로 나뉘고 crossbar로 연결됨
-
132개 SM
- 각 SM은 최대 227KB shared memory를 포함한 256KB L1 캐시를 가지며, 함께 약 33TB/s 대역폭을 가짐
- Hopper의 새 하드웨어인 Tensor Memory Accelerator(TMA) 가 비동기 주소 생성과 메모리 fetch를 담당함
- 각 SM은 4개 quadrant로 구성되고, 각 quadrant에는 warp scheduler, 512개 vector register, 행렬곱용 tensor core, 병렬 내장 명령들이 있음
- 모든 컴퓨트는 SM에서 일어나며, 대부분은 register에서 처리됨
- H100에서 성능을 내는 핵심은 tensor core를 계속 fed 상태로 유지하는 것임
- H100은 half-precision 행렬곱에서 989 TFLOPs, “그 외” 연산에서 약 60 TFLOPs를 제공함
- tensor core가 사용되는 사이클에는 최소 94%의 하드웨어 활용률에 도달함
- tensor core가 사용되지 않는 사이클에는 최대 6% 활용률에 머묾
WGMMA: 필요하지만 까다로운 명령
- H100에는 warp group matrix multiply accumulate 명령인
wgmma.mma_async가 있음- PTX에서는
wgmma.mma_async - SASS에서는
HGMMA/IGMMA/QGMMA/BGMMA
- PTX에서는
- 이전 GPU의
wmma.mma.sync,mma.sync는 32개 thread로 구성된 warp 하나가 tensor core에 데이터를 넣고 결과를 기다리는 동기 방식이었음 wgmma.mma_async는 연속된 128개 thread가 SM의 모든 quadrant에 걸쳐 협력적으로 동기화하고, shared memory에서 직접 비동기 행렬곱을 시작함- warp들은 행렬곱이 진행되는 동안 register로 다른 작업을 할 수 있음
- 결과는 원하는 시점에 기다릴 수 있음
- 마이크로벤치마크에서 H100의 전체 compute를 끌어내려면 이 명령들이 필요했음
- 사용하지 않으면 GPU가 피크 활용률의 약 63% 에서 머무르는 것으로 관찰됨
- tensor core가 로컬 리소스에서도 깊은 하드웨어 파이프라인을 요구하기 때문일 수 있음
- 가장 큰 난점은 memory layout의 복잡성임
- unswizzled shared memory layout은 coalescing이 매우 나빠 L2 대역폭을 많이 요구함
- swizzled layout은 문서가 잘못되어 있어 파악하는 데 시간이 걸렸음
- swizzled layout은 특정 행렬 shape에서만 동작하는 것으로 보이고,
wgmma.mma_async의 다른 기능과 잘 맞지 않음 - 하드웨어는 tensor core로 가는 중 sub-matrix transpose를 할 수 있지만, layout이 swizzled가 아닐 때만 가능함
- FlashAttention 같은 커널에서는 TMA와 L2 캐시가 충분히 빨라 이 문제를 어느 정도 숨길 수 있음
- 하드웨어를 완전히 쓰려면 memory request를 coalescing하고 bank conflict를 피해야 하므로 layout 제어가 중요함
Shared memory와 bank conflict
- Shared memory의 single-access latency는 약 30 cycles로 보이며, 이 시간 동안 SM의 tensor core는 거의 두 번의 32x32 정방 행렬곱을 수행할 수 있음
- FlashAttention 같은 이전 작업에서는 주로 HBM-SRAM 병목에 집중했고, 과거에는 이 병목이 실제로 중요했음
- HBM이 빨라지고 tensor core가 칩의 다른 부분보다 더 빠르게 커지면서, shared memory의 작은 latency도 제거하거나 숨겨야 하는 대상이 됨
- Shared memory는 32개 bank로 나뉘어 있어 조심하지 않으면 bank conflict가 발생함
- 같은 memory bank에 여러 다른 memory 조각을 동시에 요청하면 요청이 직렬화됨
- 경험상 커널이 불균형하게 느려질 수 있음
- WGMMA와 MMA 명령이 요구하는 register layout은 단순하게 쓰면 bank conflict를 겪을 수 있음
- 해결책은 여러 swizzling 패턴으로 shared memory를 재배치해 conflict를 피하는 것임
- 가능하면 register와 shared memory 사이 이동을 피하고, 필요할 때는 WGMMA와 TMA 같은 내장 하드웨어로 비동기 데이터 이동을 수행하는 편이 유리함
- 실제 warp를 사용한 동기 이동은 가장 일반적이지만 최악의 fallback에 가까움
주소 생성과 TMA
- H100은 tensor core와 memory가 모두 빨라서, fetch할 memory address를 생성하는 작업 자체가 칩 리소스의 상당 부분을 차지함
- 복잡한 interleaved pattern이나 swizzling pattern이 추가되면 더 두드러짐
- NVIDIA의 Tensor Memory Accelerator(TMA) 는 global/shared memory의 다차원 tensor layout을 지정하고, 해당 tensor의 subtile을 비동기로 fetch한 뒤 완료 시 barrier를 트리거할 수 있게 함
- TMA는 주소 생성 비용을 줄이고 pipeline 구성도 쉽게 만듦
- TMA는
wgmma.mma_async처럼 H100의 잠재력을 끌어내는 데 필수로 평가됨- 경험상 WGMMA보다 더 중요할 수도 있음
- register 리소스와 instruction dispatch를 절약함
- global memory에 비동기 reduction을 수행하는 기능도 있어 복잡한 backward kernel에서 유용함
- TMA 역시 swizzling mode를 이해하려면 일부 reverse engineering이 필요했지만, WGMMA보다는 덜 고통스러웠음
Occupancy가 숨겨주는 비용
- CUDA에서 occupancy는 같은 실행 하드웨어에 co-scheduled된 thread 수를 뜻함
- SM quadrant의 warp scheduler는 매 cycle마다 명령을 받을 준비가 된 warp에 instruction을 issue하려고 함
- H100은 이전 세대보다 occupancy에 덜 의존하는 면이 있음
- 비동기 기능 덕분에 단일 instruction stream도 memory fetch, matrix multiply, shared memory reduction, register math를 동시에 바쁘게 만들 수 있음
- 하지만 occupancy는 실수와 동기화 비용을 숨기는 데 매우 유용함
- 완벽히 설계된 pipeline은 추가 occupancy 없이도 빠를 수 있음
- 실제 관찰에서는 NVIDIA GPU가 occupancy를 염두에 두고 설계된 것으로 보였음
- synchronization과 실수 가능성이 많기 때문에 occupancy를 높이면 실현 하드웨어 활용률이 좋아지는 경우가 많았음
- H100에서는 occupancy가 유용한 수준이지만, A100과 RTX 4090에서는 각각 더 중요해졌다고 봄
- H100 대비 동기 instruction dispatch에 더 의존하기 때문일 가능성을 언급함
ThunderKittens: CUDA 안의 작은 DSL
- ThunderKittens는 H100에서 빠른 커널을 쉽게 작성하기 위해 만든 CUDA 내장 DSL임
- 초기에는 연구실 내부 사용을 위해 만들었고, 이후 공개됨
- 이름은 kittens가 귀엽고 코드에서
kittens::를 입력하게 하는 것이 재미있다고 생각해 붙임 - ThunderKittens는 단순함을 목표로 하며 네 가지 templated type을 제공함
- Register tiles: register file 위의 2D tensor
- Register vectors: register file 위의 1D tensor
- Shared tiles: shared memory 안의 2D tensor
- Shared vectors: shared memory 안의 1D tensor
- Tile은 height, width, layout으로 parameterized됨
- Register vector는 length와 layout으로 parameterized되고, shared vector는 length만 사용함
- shared vector는 일반적으로 bank conflict를 겪지 않음
- 제공 연산은 warp level 또는 협력 warp group level에서 tile·vector를 조작함
- initializer: shared vector를 zero로 만드는 작업 등
- unary op:
exp등 - binary op:
mul등 - row/column op:
row_sum등
- ThunderKittens는 CUDA 안에 내장되어 있어 Triton 같은 라이브러리와 달리 추상화가 “gracefully” 실패한다고 설명함
- 빠진 기능이 있으면 원하는 방식으로 확장할 수 있음
FlashAttention 예시와 성능
- ThunderKittens 예시로 RTX 4090용 간단한 forward FlashAttention 커널이 제시됨
- headdim=64만 다룸
n은 256의 배수여야 함- 약 60줄의 CUDA 코드로 작성됨
- 하드웨어 활용률은 75%
- 복잡성 대부분은 swizzling pattern이나 register layout이 아니라 알고리듬 자체에 있음
- H100용 FlashAttention-2 forward pass도 ThunderKittens로 작성됨
- TMA, WGMMA, swizzling mode, descriptor의 복잡성을 ThunderKittens가 감쌈
- 커널은 약 100줄
- H100에서 FlashAttention-2보다 약 30% 빠름
- ThunderKittens는 GPU에서 쓸 수 있는 “mini-pytorch”처럼 layout과 instruction을 감싸고 primitive를 제공함
- Based linear attention과 앞으로 공개될 다른 architecture용 kernel도 함께 공개됨
- Based linear attention kernel은 215 TFLOPs로 동작함
- 알고리듬 자체의 recompute를 고려하면 300 TFLOPs를 넘음
- Linear attention은 이론적으로 더 효율적이지만, 실제 하드웨어에서는 역사적으로 효율이 크게 낮았음
- 이 결과가 높은 throughput application 범위를 넓힐 수 있다고 봄
Tile 중심 사고
- ThunderKittens가 잘 작동한 이유는 모든 것을 하려 하지 않기 때문이라고 봄
- CUDA는 ThunderKittens보다 훨씬 표현력이 높음
- ThunderKittens는 작고 단순한 DSL임
- 핵심 추상화는 small tile이며, 이는 AI와 하드웨어가 향하는 방향과 맞는다고 봄
- ThunderKittens는 16보다 작은 차원을 지원하지 않음
- 하드웨어도 그런 작은 차원을 특별히 원하지 않는다고 봄
- “matrix multiply가 16x16보다 작다면 그것이 AI인지 확신할 수 있느냐”는 식으로 문제를 제기함
- CPU 시대의 32-bit word를 register로 보는 관점은 AI 하드웨어에는 맞지 않는다고 봄
- CUDA의 1024-bit vector register는 올바른 방향의 한 단계로 봄
- 여기서 register는 16x16 tile의 데이터임
- AI는 여전히 matrix multiply, reduction, reshape가 중심이므로 tile 추상화가 AI와 하드웨어 모두에 맞는다고 봄
- 앞으로는 AI 아이디어를 하드웨어에 잘 매핑되는 방식으로 재정렬해야 함
- recurrent state 크기는 SM에 들어갈 수 있을 만큼 커야 함
- compute density는 하드웨어가 요구하는 수준보다 낮아서는 안 됨
- 하드웨어에서 배운 내용을 AI 설계에 맞추는 것이 앞으로의 중요한 방향임
AMD 지원 계획
- ThunderKittens의 AMD hardware 지원이 곧 나올 예정임
댓글과 토론
Hacker News 의견들
-
"행렬 곱셈이 16x16보다 작다면, 그게 정말 AI인지 확신하나?"라는 질문이 흥미로움
AI 하드웨어의 요구사항이 점점 또렷해지고 있음. GPU는 원래 완전히 다른 용도로 설계됐지만 행렬 곱셈 하드웨어가 좋아서 AI에 쓰였고, "AI GPU"는 실제 GPU에 있는 일부 기능을 덜어낼 수 있음
숫자 표현도 16비트 부동소수점, 8비트, 2비트, 1비트처럼 더 짧아지는 흐름이 있고, 언젠가 적정 지점이 정해질 것임. 이 글은 16x16 타일을 선호하는 하드웨어가 꽤 타당하다는 걸 보여줌. 지금 누군가는 이미 VHDL로 이런 걸 작성 중이거나 곧 그럴 가능성이 큼
결국 "AI" 연산만 최대한 불필요한 하드웨어 짐 없이 수행하는 더 단순하고 덜 범용적이며 저렴한 장치가 나올 것 같음- GPU는 이미 가능한 한 군더더기 없는 AI 기계로 진화해 왔음. 적어도 Nervana가 설립된 2014년부터 GPU가 낡은 기술이라 AI에 부적합하다는 말이 있었지만, GPU가 이렇게 빠르게 AI 기계로 진화할 줄은 예상 못 했던 듯함
- Apple은 이미 몇 년 전부터 이런 방향으로 가고 있음. 다이 위의 NPU는 GPU나 CPU와 완전히 다름[1]
Nvidia도 아마 작업 중이겠지만, 게임/엔터테인먼트/암호화폐/AI를 한데 묶은 장치, 즉 비디오 카드 형태를 유지하는 편이 사업적으로는 더 나은 선택일 수 있음
[1] https://github.com/hollance/neural-engine/blob/master/docs/a... - "NVIDIA의 거짓말"이라는 대목이 경쟁의 깊이를 보여줌. 문서 오류가 완전히 우연일 리는 없고, 도표는 훔치거나 베끼기 쉬우니 Nvidia가 일부러 남겨둘 효용이 있었을 수 있음
Naveen Rao의 Nervana가 Nvidia 자체 드라이버보다 빠른 Nvidia Maxwell 드라이버를 만들던 때를 떠올리게 됨. 빠르게 성장하는 제품의 문서 실수가 전부 경쟁 대응책은 아니지만, 연구자들이 wgmma를 역공학하는 데 오래 걸렸고 H100을 둘러싼 미중 정치 상황까지 고려하면 Nvidia가 해자를 지키려고 예전 수법을 쓰는 것처럼 보임
그래서 H100의 특이성을 과하게 파고들기보다는, "AI가 원하는 하드웨어가 무엇인가"에는 상업적 상황도 포함된다고 봐야 함 - AMD는 이미 Versal 라인의 2세대까지 와 있음
https://www.amd.com/en/products/accelerators/alveo/v80.html
XDNA Architecture
https://www.amd.com/en/technologies/xdna.html - Google도 이런 장치를 거의 10년째 만들고 있지 않나?
-
"NVIDIA의 거짓말. 실제 128b swizzled wgmma 배치에 대해 엄청나게 오해를 부르는 표현이다. 이 도표 때문에 되돌릴 수 없는 인생 3주를 날렸으니 공개 망신을 준다"는 대목이 인상적임
AI의 진전 중 엄청난 부분이 행렬 곱셈 최적화 같은 공학에 있고, 그 공학의 상당 부분이 NVIDIA 칩 역공학이라는 사실에 놀랄 사람이 있을지 궁금함- 구조 자체는 큰 차이를 만들지 않음. 충분히 큰 자료로 충분히 큰 모델을 학습시키면 구조와 상관없이 비슷한 결과가 나오는 경향이 있음. 그래서 AI의 발전 대부분은 이제 행렬을 매우 빠르게 곱할 수 있게 된 덕분이라고 볼 수 있음
-
워프 스케줄러, 4개 사분면, 텐서 메모리 가속기, unswizzled wgmma 배치…
GPU 용어와 Star Trek식 테크노배블의 경계가 점점 더 흐려짐- 글을 읽으면서 어느 정도 알고는 있었지만, "텐서 가속기로 사분면을 워프 중"이라고 하면 정말 Star Trek 느낌이 남
다른 글들을 볼 때도 가끔 이런 생각이 들었음. 누군가가 여기 글 링크를 받아 읽으면 어떤 느낌일까 싶음. 워프 코어를 토론하는 Trek 덕후 행사장에 들어온 느낌일 듯함 - 이 말을 보고 한발 물러서서 용어들을 새 눈으로 보게 됐고, 정말 맞는 말이라 웃음이 났음
- 글을 읽으면서 어느 정도 알고는 있었지만, "텐서 가속기로 사분면을 워프 중"이라고 하면 정말 Star Trek 느낌이 남
-
AI 추론의 전력 소모를 줄이고 속도를 높이려면 아날로그 근사 회로로 전환하는 게 가장 좋을 것 같음
완벽한 부동소수점 곱셈과 덧셈이 필요한 게 아니라, 두 입력 전압을 받아 곱셈 결과에 충분히 가까운 출력 전압을 내는 장치가 필요할 뿐임- 이 방향으로 일하는 사람을 아는데, 큰 난관은 기존 칩 제조 기술로 아날로그 논리가 가능한 것을 만드는 방법, 안테나처럼 동작하지 않는 설계, 그리고 물리 칩마다 제조 공차가 달라서 실행할 모델을 칩별로 미세조정해야 할 가능성이라고 들었음
큰 장점은 float16을 16개 선으로 표현하는 대신 선 1개의 전압으로 그 숫자를 표현한다는 것임. 이론상 float32보다 훨씬 높은 정밀도도 가능할 수 있음. 또 값을 산술논리장치에 로드하지 않고 직접 연결할 수 있어서, 다이 면적과 전력 절감이 잠재적으로 여러 자릿수 규모가 될 수 있음 - 아날로그 회로가 실용적으로 유용해지려면 아직 멀었다고 보지만, 부정확성을 받아들일 수 있는 곳은 잡음 있는 디지털 회로일 수 있음
예를 들어 출력 비트 백만 개 중 하나가 뒤집히는 걸 받아들이고 성능/전력 비율을 개선하는 식임. 단일 무한대 값 하나가 전체를 망칠 수 있는 float32에서는 어렵겠지만, int8에서는 0을 원했는데 가끔 128이 나오는 정도는 견딜 수 있을 것 같음
[1] H100의 행렬 부동소수점 유닛이 실제로 IEEE 754를 준수하는지는 잘 모르겠음 - 한 걸음 더 나아가, 실제 생물학적 뇌가 작동하는 방식과 비슷하면서도 쉽게 생산 가능한 무언가가 필요하다고 봄
생물학적 신경망은 일반적인 인공신경망처럼 완전 연결에 가깝지 않고, 뉴런의 입출력 연결 계수는 10 미만이라 매우 국소적임. 생물학에는 우리가 아는 한 역전파도 없고, 대신 피드백과 순환이 있음
아직 모르는 중추신경계 기능에 필수적인 보조 세포나 과정이 있을 수도 있음. 고수준에서도 상당한 양의 "하드코딩된" 연결성이 있을 가능성이 있고, 이미 일부는 알려져 있음. 예컨대 귀의 청각 뉴런은 연결되어 있으며 소리의 위치를 찾기 위해 합성곱과 비슷한 일이 일어남. 이는 창발 현상이 아니라 훈련 없이도 가능한 기능임
생명은 수십억 년과 비슷한 수의 세대를 거치며 이를 찾아냈으니 놀랍지 않음. 이론적으로는 소프트웨어로도 가능하겠지만, 영장류/인간 뇌의 1조 개 이상 뉴런을 고려하면 오늘날의 천 코어급 기계로도 극도로 어려움. "클라우드"라고 해도 필요한 연결성과 지연시간을 충족하지 못할 것임
이런 접근으로 벌레나 곤충 정도를 성공적으로 모델링할 수 있다면 멋질 듯함 - 충분한 범위와 정밀도를 동시에 만족시키는 건 거의 불가능해 보임
- 솔직히 디버깅하기에는 악몽 같아 보임
- 이 방향으로 일하는 사람을 아는데, 큰 난관은 기존 칩 제조 기술로 아날로그 논리가 가능한 것을 만드는 방법, 안테나처럼 동작하지 않는 설계, 그리고 물리 칩마다 제조 공차가 달라서 실행할 모델을 칩별로 미세조정해야 할 가능성이라고 들었음
-
이 글을 보니 CS 149 병렬 프로그래밍 수업에서 느꼈던 즐거움이 다시 떠오름
- Kayvon과 Kunle는 대단함. 두 학기 전에 CS149 Parallel Programming을 들었고 정말 좋았음 :)
-
이 글의 문체가 정말 인상적이고, AMD MI300x에서 이걸 보는 게 기대됨. 내 장비에서 시간을 써보고 싶으면 알려달라
- AMD 제품으로 AI 작업을 많이 해봤는지 궁금함. RTX 4090에 2500달러 이상을 쓰고 싶지는 않아서, 실험용이나 입문용으로 RX 7900XTX를 고려 중임
실제로 얼마나 잘 작동할지, 아니면 조금 더 모아서 7900 XT 대신 XTX를 사는 게 나을지, VRAM이 줄어들면 실사용성이 얼마나 영향을 받을지도 궁금함 - 좋은 글은 명확하고 모호하지 않아야 함. 말로 할 때는 중간에 끊고 clarification을 요구할 수 있지만, 글은 메시지를 전달할 기회가 한 번뿐임
독자가 저자들이 무슨 말을 하려는지 알려고 knowyourmeme.com까지 찾아봐야 해서는 안 됨. 이 제목이 무슨 뜻인지조차 모르겠고, 그만큼 목표를 크게 벗어났다고 봄 - 정말? Wallstreetbets 시절 PTSD가 떠오름
- AMD 제품으로 AI 작업을 많이 해봤는지 궁금함. RTX 4090에 2500달러 이상을 쓰고 싶지는 않아서, 실험용이나 입문용으로 RX 7900XTX를 고려 중임
-
이런 글을 완전히 이해하려면 어디서부터 시작해야 할지, 어떤 로드맵을 따라가야 할지 궁금함
- GPU 프로그래밍을 배우기 좋은 강의가 있음. 4.0 강의쯤에서 필요한 기본기를 얻을 수 있음: https://youtube.com/playlist?list=PLzn6LN6WhlN06hIOA_ge6Srgd...
그리고 벡터-행렬 곱셈을 수행하는 CUDA 커널을 직접 작성해보는 게 좋음. pycuda를 쓰면 커널에 집중하고 나머지는 Python으로 작성할 수 있음. ChatGPT에게 4000개 원소 벡터와 4000x12000 행렬을 곱하는 구현을 직접 만들고 싶다고 말하고 전체 과정을 안내해달라고 하면 됨
GPU 임대는 Runpod가 좋고, 지금은 저가형 GPU부터 H100까지 있음. 처음에는 낮은 등급 GPU로 시작하면 됨 - 깊게 들어가려면 Spiral 행렬 곱셈 재생목록을 보면 좋을 듯함: https://www.youtube.com/playlist?list=PL04PGV4cTuIWT_NXvvZsn...
Spiral로 행렬 곱셈 커널을 구현하고 최적화하는 데 2개월을 썼음
- GPU 프로그래밍을 배우기 좋은 강의가 있음. 4.0 강의쯤에서 필요한 기본기를 얻을 수 있음: https://youtube.com/playlist?list=PLzn6LN6WhlN06hIOA_ge6Srgd...
-
GitHub README의 그래프(https://github.com/HazyResearch/ThunderKittens/blob/main/att...)가 너무 어지러움. 이런 물결 막대그래프가 합법이긴 한가? :P
- 동의함. 무슨 착시를 넣으려는 것 같음. 차라리 막대 없이 숫자만 보는 게 낫겠음
- matplotlib의 xkcd 테마처럼 보임[1]. 그래도 물결이 너무 과하다는 데 동의함
[1]: https://matplotlib.org/stable/gallery/showcase/xkcd.html#sph...
-
ThunderKittens라는 이름이 훌륭함. ThunderKittens가 순전파보다 한 자릿수 더 어려운 FlashAttention 역전파를 다루는 걸 보고 싶음
- 좋은 소식이 있음. TK에 최적화된 causal 및 non-causal 버전의 FlashAttention 역전파를 이미 포함했음. 한번 봐주면 좋겠음
causal: https://github.com/HazyResearch/ThunderKittens/blob/main/exa...
non-causal: https://github.com/HazyResearch/ThunderKittens/blob/main/exa...
- 좋은 소식이 있음. TK에 최적화된 causal 및 non-causal 버전의 FlashAttention 역전파를 이미 포함했음. 한번 봐주면 좋겠음
-
이런 연구는 이미 오늘날 NPU를 만드는 팀들이 해온 것 아닌가? 예를 들어 Groq 칩은 AI 전용 구조를 사용하기 때문에 지금 같은 성능을 낼 수 있음. 소비자 쪽에서는 Apple Silicon도 꽤 유능함
이 분야 사람은 아니지만, 비교적 느린 경로로 통신하는 범용 프로세서만으로는 한계가 있을 것 같음. 하드웨어 수준에서 설계를 다시 생각하고, 결국 소비자 시장에서 가격을 낮추는 쪽이 더 나은 장기 전략으로 보임- 소비자 쪽에서 Apple Silicon이 꽤 유능하다는 말은 확신이 안 듦. reddit의 localllama 서브레딧을 보면 유용한 속도를 내려고 애쓰는 CPU 사용자들이 좌절하는 글이 많음
몇백 달러로 Nvidia GPU를 사거나 900달러에 4050 6GB VRAM 게이밍 노트북을 살 수 있는데, CPU 기반 AI를 유능하다고 부르기는 어려움
직장에도 GPU가 없어서 CPU 기반으로 해봤지만, 작은 모델을 쓰고 기다리는 것 말고는 현실적이지 않았음. 결국 GPU 컴퓨터를 요청하게 됨
"기술적으로 가능하다"와 "실제로 쓰기 좋다"는 다름. Nvidia는 사용하기 정말 좋았고, CPU는 고통스럽고 답답했음
- 소비자 쪽에서 Apple Silicon이 꽤 유능하다는 말은 확신이 안 듦. reddit의 localllama 서브레딧을 보면 유용한 속도를 내려고 애쓰는 CPU 사용자들이 좌절하는 글이 많음