Cách khiến GPU AMD chạy với tốc độ ‘brrr’

 (hazyresearch.stanford.edu)
1 điểm bởi GN⁺ 2025-11-16 | 1 bình luận | Chia sẻ qua WhatsApp
  • HipKittens là bộ primitive lập trình được thiết kế để khai thác hiệu năng tiềm năng của GPU AMD, tối ưu hóa truy cập bộ nhớ, lập lịch và tái sử dụng cache
  • GPU AMD MI355X có cấu trúc 256 compute unit và 8 chiplet (XCD), đồng thời cung cấp register file lớn và các lệnh matrix core chi tiết hơn
  • Khác với NVIDIA, AMD không có phân bổ lại register, lệnh ma trận bất đồng bộ, mbarrier, nên thay vì wave specialization, cách lập lịch 8-wave ping-pong và 4-wave interleave hiệu quả hơn
  • HipKittens cải thiện tính cục bộ của cache L2 và LLC thông qua lập lịch grid nhận biết chiplet, đạt được mức tăng băng thông và TFLOPS tối đa trong các phép toán GEMM và Attention
  • Cách tiếp cận này bù đắp cho độ trưởng thành phần mềm còn thiếu của hệ sinh thái GPU AMD, đồng thời tạo nền tảng để nâng cao khả năng mở rộng AI computing trên phần cứng đa dạng

Kiến trúc và đặc tính hiệu năng của GPU AMD CDNA

  • GPU AMD MI355X gồm 256 compute unit (CU), mỗi CU được cấu thành từ 4 SIMD
    • Mỗi SIMD thực thi một wave gồm 64 luồng, đối lập với warp 32 luồng của NVIDIA
  • MI355X có SRAM chỉ ở mức khoảng 70% của B200 (165KB) và không có các tính năng lệnh nhân ma trận bất đồng bộ, phân bổ lại register, tăng tốc tensor memory, mbarrier
  • Ngược lại, nó cung cấp register file lớn gấp 2 lần và nhiều bộ xử lý hơn 60% (256 CU so với 160 SM)
    • Hỗ trợ các lệnh matrix core nhỏ và chi tiết, cùng khả năng nạp trực tiếp từ global vào shared memory (tương tự TMA)
  • AMD áp dụng kiến trúc chiplet gồm 8 chiplet (XCD), mỗi XCD có cache L2 riêng biệt, phía trên là cache LLC
  • Theo bảng so sánh, MI355X đạt BF16 2.5 PFLOPs, MXFP8 5.0 PFLOPs, MXFP6 10.1 PFLOPs, cùng 288GB bộ nhớ và băng thông 8TB/s

Thách thức khi thiết kế kernel cho AMD

  • Tối ưu truy cập bộ nhớ: do các ràng buộc của compiler HIPCC và hành vi I/O không được công khai, việc thiết kế bố trí dữ liệu và mẫu swizzle là rất quan trọng
  • Lập lịch trong bộ xử lý: AMD cần tận dụng register file và các lệnh ma trận nhỏ thay vì dựa vào shared memory
  • Lập lịch giữa các bộ xử lý: do cấu trúc dựa trên chiplet, cần phân phối công việc có tính đến hiệu ứng NUMA ở cấp cache

Mẫu truy cập bộ nhớ của HipKittens

  • HipKittens (HK) dùng tile làm đơn vị dữ liệu cơ bản và cung cấp các hàm toán tử tương tự PyTorch
    • Tile được định nghĩa bởi kiểu dữ liệu, kích thước, layout và hỗ trợ nhiều đầu vào khác nhau thông qua C++ template metaprogramming
  • Lập lịch register: vì HIPCC không thể luôn dùng các register cụ thể làm đầu vào MFMA, HK cung cấp khả năng cố định register một cách tường minh
    • Nhờ đó, lập trình viên có thể tự chỉ định register để viết kernel đạt hiệu năng tối đa
  • Layout register: trên AMD, layout thay đổi tùy kiểu dữ liệu và dạng ma trận nên không thể dùng một mẫu swizzle duy nhất
    • Ví dụ, tile bf16 16×16 và tile bf16 16×32 cần các mẫu swizzle khác nhau
  • Cấu trúc phase của lệnh: các lệnh shared memory của AMD có nhóm phase không liên tục và tài liệu nội bộ còn thiếu
    • HK cung cấp solver được reverse engineering cho phần này
  • Sinh địa chỉ: AMD hỗ trợ nạp bất đồng bộ từ HBM vào shared memory, và tối ưu thông qua swizzle địa chỉ HBM

Lập lịch trong bộ xử lý: mẫu Wave

  • Wave specialization hiệu quả trên NVIDIA, nhưng trên AMD lại làm giảm hiệu năng do thiếu khả năng phân bổ lại register
    • Wave producer chiếm các register không cần thiết, còn wave consumer bị thiếu register nên phát sinh spill
  • Kết quả thử nghiệm của HK cho thấy wave specialization trên AMD gây giảm cường độ tính toán và tạo nút thắt bộ nhớ
    • Ví dụ: trong GEMM, cấu hình HK 0/8 đạt 1605 TFLOPs, CUTLASS đạt 1570 TFLOPs
  • Các mẫu lập lịch thay thế
    • 8-wave ping-pong: hai wave luân phiên chạy cụm bộ nhớ/tính toán
    • 4-wave interleave: một wave đan xen chi tiết giữa bộ nhớ và tính toán
    • 8-wave cho mã gọn hơn, còn 4-wave chi tiết hơn nhưng mã dài hơn
    • Trong GEMM và Attention Forward, 8-wave đạt hiệu năng ở mức SoTA

Lập lịch giữa các bộ xử lý: cách tiếp cận nhận biết chiplet

  • AMD MI355X có 8 chiplet XCD, mỗi chiplet sở hữu cache L2 độc lập
    • Các thread block được gán vào chiplet theo kiểu round-robin, nên thứ tự grid ảnh hưởng trực tiếp đến hiệu quả tái sử dụng cache
  • Cách bố trí row-major đơn giản có tỷ lệ tái sử dụng cache L2 thấp, gây mất băng thông
    • Ví dụ: L2 55%, LLC 95%, 15.1 TB/s, 1113 TFLOPs
  • HK đưa vào lập lịch grid nhận biết chiplet, tận dụng đồng thời tính cục bộ của cache L2 và LLC
    • Các thread block được nhóm theo các vùng lân cận của ma trận đầu ra để tối đa hóa việc tái sử dụng dữ liệu đầu vào

Ví dụ kernel thực tế

  • Hot loop của các kernel Attention Forward và BF16 GEMM sử dụng lịch 8-wave ping-pong của HK
    • Mỗi vòng lặp luân phiên thực thi cụm Compute–Memory và đồng bộ bằng schedule barrier
    • Trong ví dụ mã, các phép toán HK như mma_AtB, load, exp2, col_sum được dùng lặp lại

Kết luận: AMD trong kỷ nguyên AI đa silicon

  • HipKittens đạt được hiệu năng cạnh tranh trên AMD CDNA3 và CDNA4
    • Ba điểm cốt lõi gồm: truy cập bộ nhớ tối ưu, lập lịch wave tập trung cho AMD, và lập lịch grid nhận biết chiplet
  • Các kernel HK đạt hiệu năng tốt nhất theo chuẩn AMD và có mức cạnh tranh với cả kernel NVIDIA Blackwell
  • Để tăng tính đa dạng cho AI computing, cần mở rộng khả năng tiếp cận GPU AMD, và HipKittens cung cấp nền tảng phần mềm cốt lõi cho mục tiêu đó
  • Việc cải thiện lập lịch register trong HIPCC được xem là hướng phát triển quan trọng trong tương lai

Bài viết liên quan

1 bình luận

 
GN⁺ 2025-11-16
Ý kiến Hacker News
  • Khuyên nên tham khảo thảo luận liên quan đến HipKittens
  • Cũng có bài viết về cùng nghiên cứu là HipKittens: Fast and furious AMD kernels. Có bình luận của George Hotz và nhân viên AMD
  • Thật đáng mừng khi giới học thuật xử lý kiểu vấn đề này, nhưng rốt cuộc tôi nghĩ đây là vấn đề AMD phải tự giải quyết nội bộ
    • Tôi cho rằng các công ty phần cứng nên chỉ làm phần cứng. Như vậy các động lực mới được giữ thuần khiết. Dù hiệu năng có giảm 20% thì tôi vẫn nghĩ như thế tốt hơn
    • Hoàn toàn đồng ý. AMD đã trì hoãn vấn đề này từ 10 năm trước và giờ mới cố bắt kịp. Phần cứng rất tốt nhưng năng lực viết firmware còn thiếu nên không khai thác được hết tiềm năng
    • Nhưng nhóm nghiên cứu này trước đây cũng từng làm phần mềm tương tự cho GPU Nvidia. Có vẻ đây là những nhà nghiên cứu giỏi đang phát huy đúng chuyên môn của mình
    • Theo tôi biết thì AMD đã xử lý vấn đề này ở nhiều cấp độ và cũng đang hợp tác với tinycorp
  • Đọc bài viết tạo cảm giác việc tối ưu khó khăn vì độ phức tạp về mặt kiến trúc của GPU AMD. Tuy vậy, về dài hạn thì cách tiếp cận của AMD có thể mở rộng tốt hơn. Trong khi Nvidia dùng 2 chiplet thì AMD có cấu trúc 8 chiplet nên phát sinh vấn đề locality bộ nhớ. Tương lai số chiplet sẽ còn tăng, nên kinh nghiệm xử lý sự phức tạp hiện tại có thể sẽ hữu ích về lâu dài
    • AMD không cần warp specialization để đạt hiệu năng cao, nên việc lập trình đơn giản hơn
  • Nhiều lập trình viên đã cố làm cho GPU AMD thật sự ‘go brrr’ với các nhà phát triển phổ thông nhưng đều thất bại. Tôi không hiểu vì sao AMD không tự giải quyết vấn đề phần mềm. Giờ họ cũng có đủ tiền rồi, không thể lấy cớ không tuyển lập trình viên được nữa. GPU cho datacenter của họ không tệ, nhưng khi cá nhân muốn thử nghiệm ML·AI thì Nvidia vẫn tốt hơn rất nhiều. Tôi có cảm giác chiếc RTX 3090 đã 5 năm tuổi của mình vẫn tốt hơn mọi GPU tiêu dùng AMD từng ra mắt đến nay
    • Trải nghiệm lập trình viên với AMD thật kinh khủng. Họ thậm chí còn không tiếp nhận báo cáo lỗi crash driver
    • Gần đây tôi chuyển máy chủ suy luận từ NVidia 5090 sang hai chiếc AMD R9700 32GB, và trải nghiệm hoàn toàn tích cực. Trên kernel Fedora nó chạy ngay không cần cấu hình DKMS, và việc kết nối container với ROCm cũng rất dễ. Chỉ cần đổi cấu hình Ollama và Storyteller là xong. Trải nghiệm dễ chịu hơn CUDA rất nhiều
    • Nvidia thậm chí còn tự duy trì cả fork của Unreal Engine. AMD còn chưa ở cùng đẳng cấp để cạnh tranh
    • Nvidia là công ty phần cứng duy nhất còn trả đãi ngộ đủ sức cạnh tranh cho kỹ sư phần mềm. Ở AMD vẫn còn văn hóa không xem phần mềm là ‘công việc thực sự’, và kiểu quán tính đó rất khó thay đổi
  • Mojo từng có ý tưởng cải thiện trải nghiệm lập trình viên (devX) trên GPU AMD, nên tôi tò mò không biết tiến độ đến đâu rồi
  • Tôi không hiểu vì sao AMD không đầu tư hàng tỷ USD để cải thiện phần mềm. Nvidia là công ty giá trị nhất thế giới, còn AMD là đối thủ duy nhất
    • AMD cũng đang cố gắng, nhưng tôi nghĩ việc chuyển một văn hóa tổ chức hằng năm làm mới phần cứng sang văn hóa lấy phần mềm làm trung tâm là rất khó. Phần mềm không tạo doanh thu ngay như phần cứng nên ban điều hành thường có xu hướng ưu tiên thấp hơn. Ngoài ra, việc các vendor bên ngoài đóng góp mã nguồn mở có thể trông tốt trong ngắn hạn nhưng lại ảnh hưởng xấu đến chất lượng dài hạn. Chỉ cần lỡ một xu hướng phần cứng thôi là đã có nguy cơ tụt lại sau đối thủ
    • Tôi từng làm ở nhiều vendor GPU, và chỉ Nvidia mới xem phần mềm là tài sản (asset) để đầu tư. Các công ty khác chỉ coi đó là chi phí
  • Cá nhân tôi không thích meme “go brr”, nhưng thấy nó xuất hiện ở nơi như Stanford thì cũng khá thú vị
    • Thực ra họ đã dùng “go brr” từ 1 năm trước trong bài giới thiệu ThunderKittens
    • Nếu mấy meme kiểu này đã xuất hiện trên kênh chính thức của trường đại học, thì có khi đó là dấu hiệu trào lưu đã hết hot
  • Bản thân dự án thì rất tốt, nhưng tôi thắc mắc vì sao AMD không tự làm việc này. Có vẻ AMD vẫn chưa hiểu tầm quan trọng của một software stack trưởng thành. Họ cần một stack hợp nhất dùng được trên mọi card như CUDA. Trước đây tôi từng tin rằng AMD rồi sẽ bắt kịp, nhưng giờ thì gần như đã bỏ hy vọng
  • Dự án thì tốt nhưng bản thân bài viết lại tạo cảm giác được viết khá kỳ quặc
    • Bài viết quá gượng gạo. Có vẻ tác giả phụ thuộc quá nhiều vào AI, hoặc cố bắt chước văn phong AI. Những câu như “hãy xem part one” hay “cách làm AMD GPU go brr” cứ lặp đi lặp lại. Điều đáng tiếc nhất là các phần kỹ thuật lẽ ra nên giải thích bằng biểu đồ thì lại được diễn đạt bằng 100 dòng mã