- HipKittens là một dự án cung cấp các kernel hiệu năng cao và các primitive lập trình dựa trên C++ cho GPU AMD nhằm nâng cao hiệu quả tính toán AI
- Trong hệ sinh thái AMD hiện có, AITER, PyTorch, Triton, TileLang, Composable Kernel đều bộc lộ giới hạn do hiệu năng thiếu ổn định và mức hỗ trợ còn non trẻ
- HipKittens lấy trừu tượng hóa dựa trên tile (tile abstraction) làm trung tâm, duy trì giao diện chung giữa NVIDIA và AMD đồng thời tách riêng phần triển khai theo từng phần cứng
- Các kernel được viết với chưa tới khoảng 500 dòng mã đạt hiệu năng nhanh hơn các kernel assembly thủ công hiện có của AMD
- Dự án đưa ra một nền tảng thực tiễn để mở rộng tác vụ AI sang môi trường đa silicon, đồng thời gợi mở khả năng chuyển dịch sang hệ sinh thái phần cứng mở
Hiện trạng và vấn đề của hệ sinh thái GPU AMD
- Tính toán AI cho đến nay chủ yếu phát triển xoay quanh một nhà cung cấp phần cứng duy nhất, nhưng GPU AMD ở thế hệ mới nhất lại cung cấp hiệu năng tính toán và băng thông bộ nhớ ở mức hàng đầu
- Ví dụ: AMD MI355X OAM đạt BF16 2.5 PFLOPs, MXFP8 5.0 PFLOPs, bộ nhớ 288GB, băng thông 8.0TB/s
- Tuy nhiên, do thiếu một software stack đủ trưởng thành, hiệu năng này không được khai thác đầy đủ trong các workflow AI thực tế
- Các thành phần phần mềm chính của AMD gồm AITER, PyTorch, Triton, Mojo, TileLang, Composable Kernel (CK)
- Các kernel backpropagation SDPA Llama GQA của AITER và PyTorch lần lượt chỉ đạt 30% và 24% hiệu năng so với SoTA
- Mojo bị giới hạn ở mức hiệu năng khoảng 50% do xung đột bank (bank conflict)
- TileLang chỉ hỗ trợ đến CDNA3, đồng thời thiếu tính năng cốt lõi và phụ thuộc vào CK, làm tăng độ phức tạp
- Triton gặp khó khăn trong việc theo dõi vòng đời thanh ghi và tối ưu hóa truy cập bộ nhớ
- Kết quả là, các kernel AMD có hiệu năng cao nhất hiện vẫn phải do chuyên gia tự viết bằng assembly, điều này tạo ra rào cản lớn về khả năng mở rộng và bảo trì
So sánh với hệ sinh thái xoay quanh CUDA
- Trong cộng đồng AI, người ta đánh giá có tồn tại rào hào CUDA (CUDA moat)
- Trước đây, phát triển kernel trên NVIDIA cũng đòi hỏi nhiều năm nỗ lực dựa trên CUDA/CUTLASS cấp thấp
- Gần đây, nhờ sự phát triển của DSL (ngôn ngữ đặc thù miền) và các công cụ hỗ trợ bởi AI, việc phát triển kernel cho NVIDIA đã trở nên đơn giản hơn
- Ví dụ: ThunderKittens (TK), CuTe DSL, TinyGrad “tinykittens”, TileLang, Gluon
- Từ xu hướng này, nhu cầu về các primitive lập trình mới cho AMD cũng được đặt ra
Nguyên lý thiết kế của HipKittens
- HipKittens là phiên bản được phát triển cho AMD, tiếp nối ThunderKittens (NVIDIA) và ThunderMittens (Apple Silicon)
- Khái niệm cốt lõi: trừu tượng hóa dựa trên tile (tile abstraction)
- Tính phổ quát của trừu tượng hóa tile – mô hình tính toán dựa trên tile vốn hiệu quả trên NVIDIA cũng được áp dụng tự nhiên trên AMD
- Triển khai backend đặc thù theo kiến trúc – mẫu truy cập bộ nhớ và lịch biểu thanh ghi được thiết kế khác nhau tùy phần cứng
- Tính thích nghi của chiến lược scheduling – trên CDNA3·CDNA4 của AMD, scheduling theo wave kém hiệu quả, nhưng scheduling theo đơn vị tile vẫn giữ được sự đơn giản và hiệu năng
- Bằng cách tách giao diện (tile và phép toán) khỏi triển khai (ánh xạ phần cứng), dự án đưa ra một mô hình có thể áp dụng chung cho nhiều kiến trúc GPU khác nhau
Kết quả hiệu năng
- Kernel Attention Forward: được viết với khoảng 500 dòng mã và đạt hiệu năng vượt kernel assembly của AITER
- Vượt trội trong cả cấu hình Causal/Non-Causal ở nhiều chiều head (D) và độ dài chuỗi (N) khác nhau
- Kernel GEMM: gồm vòng lặp cốt lõi dưới 100 dòng, đạt hiệu năng hàng đầu ở cả phép toán BF16 và FP8
- Các kernel Attention Backward, Rotary, Fused Dropout-Residual-LayerNorm cũng được xác nhận cải thiện so với mức hiệu năng tốt nhất trước đó
- Tất cả các kernel đều duy trì tính dễ đọc và dễ chỉnh sửa, đồng thời vẫn đạt hiệu năng ở cấp độ assembly thủ công
Mở rộng sang AI đa silicon
- Để hiện thực hóa tiềm năng của AI, cần khai thác phần cứng đa dạng và cởi mở hơn
- HipKittens đặt mục tiêu biến GPU AMD thành một nền tảng thực sự dễ tiếp cận đối với các nhà phát triển AI
- Khi kết hợp với software stack mã nguồn mở của AMD, dự án cho thấy khả năng chuyển dịch sang hệ sinh thái AI dựa trên đa silicon
- Bài viết tiếp theo sẽ đi sâu vào cấu trúc kỹ thuật chi tiết của HipKittens (báo trước phần hai)
1 bình luận
Ý kiến trên Hacker News
Chúng tôi đã ký hợp đồng với AMD và đang huấn luyện Llama 405B trên MI350X cho MLPerf
Tình hình của AMD rõ ràng đang tốt lên. Nếu bạn có GPU AMD, tôi khuyên nên vào pytorch.org rồi chọn Linux+ROCm để cài PyTorch. Cách đây 3 năm còn rất tuyệt vọng, nhưng giờ hầu hết các tính năng chính đều hoạt động tốt. Tôi đã chạy nanochat trên MI300X và nó chạy ngay lập tức. MI350X cũng ổn định tương tự
Tất nhiên vẫn còn thua NVIDIA, và vẫn cần đầu tư nhiều vào hệ sinh thái phần mềm, trình biên dịch và driver. Nhưng nếu so với tình cảnh tuyệt vọng của 2 năm trước thì giờ đã thấy hy vọng.
Nếu muốn thấy những chỗ backend LLVM của AMD còn thiếu, hãy so sánh mã HipKittens với CUDA Kittens.
Với tác vụ huấn luyện, NVIDIA và Google đứng đầu, AMD đứng thứ hai, còn lại thì gần như không đáng kể. Intel và Tenstorrent còn rất xa, Huawei thì ví dụ chạy bị chết vì segfault, Groq đã từ bỏ việc bán chip, Cerebras thì không thể mua ở đâu được. Trainium mất 5 ngày mới xin được một instance nên tôi cũng hết hứng thú
Thiết lập ban đầu vẫn hơi thô ráp, nhưng đã tốt hơn rất nhiều so với trước đây. Ví dụ chỉ mới một tháng trước thôi nanochat còn chạy chưa ổn, còn bây giờ thì đã chạy tốt. Điều quan trọng là giờ mọi người đã quan tâm đến hệ sinh thái AMD.
AI cần sự đa dạng phần cứng. Việc một công ty độc chiếm toàn bộ phần cứng và phần mềm AI có thể tốt cho cổ đông, nhưng lại có hại cho tiến bộ công nghệ
Tôi không hiểu giá trị doanh nghiệp của NVIDIA. Hiện tại chỉ có một số ít thuật toán như Transformer đang thắng thế, và dữ liệu đã trở nên quan trọng hơn. Không còn cần nhiều loại mã HPC đa dạng như trước, nên giờ đối thủ chỉ cần tối ưu cho một tập thuật toán hẹp hơn là đủ.
Nếu có thể chạy vLLM và Transformer hiệu quả thì một thị trường khổng lồ sẽ mở ra. Nếu vậy thì có vẻ AMD hay Huawei cũng sẽ dễ bắt kịp, nên tôi tò mò hào lũy (moat) thực sự của NVIDIA là gì. Chỉ InfiniBand thôi đã đủ chưa?
Trong GPU cho trung tâm dữ liệu, NVIDIA vẫn mạnh, nhưng Google đang dùng TPU ở quy mô lớn và OpenAI cũng đã đặt hàng phần cứng AMD.
Hệ sinh thái CUDA vẫn rất quan trọng, nhưng ai cũng đang tích cực tìm cách thoát khỏi nó. AMD, Intel, Qualcomm và các hãng khác cũng đang lao vào cuộc cạnh tranh này. HipKittens trông giống như một bước tiến quan trọng hướng đến hệ sinh thái phần mềm trung lập
AMD có vẻ cũng nên tài trợ cho những dự án như thế này, nhưng tôi nghĩ có lẽ họ lại bỏ lỡ cơ hội nữa rồi. Với GPU và AI thì họ lúc nào cũng như vậy
HipKittens là một cải thiện, nhưng bên trong AMD vẫn thiếu năng lực theo dõi hiệu năng kernel. DevOps được thuê ngoài cho TCS ở Ấn Độ, và bên đó cũng không thực sự nắm tình hình.
Những đội có lãnh đạo giỏi thì tự vận hành đội IT trong bóng tối của riêng mình. ROCm thì không có đội như vậy, và cuối cùng chỉ bắt đầu cải thiện khi các khách hàng cloud lớn lên tiếng phản đối.
Ngay cả khi tuyển được nhân tài, họ vẫn đưa ra mức lương dưới thị trường. Vì họ so với mặt bằng kiểu Qualcomm hay Walmart.
Trong 4 năm qua, thưởng chưa từng được trả đủ 100%
Trước đây họ từng cắt hỗ trợ như với Radeon VII, hoặc liên tục đập đi xây lại hệ sinh thái nên mãi không trưởng thành nổi. Vì thiếu tương thích CUDA và thiếu đầu tư, phần lớn tổ chức đơn giản là chọn NVIDIA.
Dù vậy, gần đây họ đã liên tục đầu tư vào ROCm và hệ sinh thái Instinct nên đang cải thiện dần. Tuy nhiên ở mảng mạng kết nối thì NVIDIA vẫn dẫn trước rất xa
Tôi tò mò có dự án nào được xây dựng trên ThunderKittens hay chưa.
Nếu đây là phiên bản được port sang HIP thì có vẻ nó sẽ có giá trị thực tiễn lớn hơn nhiều so với một bản port CUDA đơn thuần
Cụm từ “raw assembly vs cooked assembly” nghe rất thú vị.
Ngày xưa việc tự viết mã assembly cho CPU là chuyện rất phổ biến. Với GPU cũng không cần phải sợ hãi quá mức. Rốt cuộc thì chính trình biên dịch cũng là thứ tạo ra đoạn mã đó
Xét về mặt toán học thì suy luận (inference) chỉ là các phép toán đại số tuyến tính/BLAS cơ bản.
Tôi tự hỏi liệu có thể có một API suy luận gọn nhẹ bao phủ 80% trường hợp sử dụng, có tính đến dtype và sparsity hay không. Có lẽ không cần phải phức tạp như CUDA
composable-kernel là phần mềm khiến hệ thống Gentoo của tôi bị OOM (hết bộ nhớ) thường xuyên nhất. Khi Clang biên dịch CK, nó dùng hơn 2.5GB cho mỗi luồng
Có lẽ trên máy build chính thức phải giảm số lượng kiến trúc xuống. Tôi nghe nói phần lớn các gói phụ thuộc chỉ cần header. Hiện cũng đang có các nỗ lực cải thiện để giảm thời gian build
device_grouped_conv2d_fwd_xdl_ngchw_gkcyx_ngkhw_f16_instance, tôi thật sự không hiểu họ đang làm gì với clang nữaTôi tò mò liệu những tiến bộ như thế này có giúp chạy LLM tốt trên chip AMD tiêu dùng hay không.
Ví dụ tôi đang cân nhắc laptop HP OMEN MAX 16 có AMD CPU/iGPU và RTX 5080, liệu phía AMD có thể cạnh tranh với RTX được không?
Bài blog này hoặc kết quả từ các máy Mac cao cấp khá thú vị
Nó hơi chậm hơn API Sonnet 4 một chút, nhưng với hiệu năng như vậy ở máy local thì tôi hoàn toàn hài lòng.
Tổ hợp Opencode + Ollama + Qwen3 Coder là một lựa chọn thay thế tuyệt vời cho ClaudeCode + Sonnet4.
Dĩ nhiên nếu bạn cần AI làm thay toàn bộ việc lập trình thì có thể sẽ thấy khác. Nhưng làm trợ lý cá nhân thì hoàn hảo
Tôi không hiểu vì sao AMD lại phớt lờ hoàn toàn B300