BarraCUDA - Trình biên dịch CUDA mã nguồn mở dành cho GPU AMD

• Trình biên dịch mã nguồn mở độc lập có thể chuyển trực tiếp mã nguồn CUDA C (.cu) thành mã máy cho AMD RDNA3(GFX11)
• Tạo nhị phân ELF .hsaco thông qua bộ phân tích từ vựng, parser và biểu diễn trung gian (BIR) riêng, không cần lớp LLVM hay HIP
• Được viết bằng hơn 15.000 dòng mã C99 và có thể build chỉ với một lệnh make
• Hỗ trợ các tính năng chính của CUDA như biến dựng sẵn cho luồng, bộ nhớ chia sẻ, phép toán nguyên tử, phép toán warp, cooperative groups
• Được phát hành theo giấy phép Apache 2.0, và trong tương lai hướng tới mở rộng sang các kiến trúc khác như Tenstorrent, Intel Arc, RISC-V

---

Tổng quan về BarraCUDA

• BarraCUDA là trình biên dịch CUDA dành cho GPU AMD, chuyển các tệp .cu thành mã máy GFX11
  • Kết quả đầu ra là nhị phân ELF .hsaco có thể chạy trên GPU AMD
  • Hoạt động hoàn toàn độc lập, không phụ thuộc LLVM
• Toàn bộ mã nguồn gồm khoảng 15.000 dòng viết bằng C99 và có thể build bằng một Makefile duy nhất
• Dự án được một nhà phát triển tại New Zealand phát triển cá nhân

Cách hoạt động

• Tệp .cu đầu vào được xử lý theo thứ tự tiền xử lý → phân tích từ vựng → parsing → phân tích ngữ nghĩa → tạo BIR → chọn lệnh → cấp phát thanh ghi → mã hóa nhị phân → xuất ELF
• BIR (BarraCUDA IR) là biểu diễn nội bộ dạng SSA, được thiết kế độc lập với kiến trúc
• Mọi mã hóa đều được kiểm chứng bằng llvm-objdump và đạt 0 lỗi giải mã

Tính năng được hỗ trợ

• Cú pháp cốt lõi của CUDA: __global__, __device__, __host__, threadIdx, blockIdx v.v.
• Tính năng CUDA: bộ nhớ __shared__, __syncthreads(), phép toán nguyên tử, shuffle/bỏ phiếu warp, kiểu vector, độ chính xác half, cooperative groups v.v.
• Tính năng trình biên dịch: bộ tiền xử lý C hoàn chỉnh, khôi phục lỗi, theo dõi vị trí mã nguồn, hỗ trợ truyền giá trị struct

Các mục chưa hỗ trợ

• Hiện vẫn chưa triển khai các mục như dùng unsigned đơn lẻ, toán tử gán kết hợp (+=, -= v.v.), const, bộ nhớ __constant__, mảng chia sẻ 2D, texture/surface, thực thi song song động
• Không hỗ trợ nhiều translation unit và không sinh mã host

Kiểm thử và lộ trình

• Được kiểm chứng với 14 tệp kiểm thử, hơn 35 kernel, và khoảng 27KB mã máy
• Mục tiêu ngắn hạn: hoàn thiện cú pháp và tăng khả năng tương thích với các tệp .cu dùng thực tế
• Mục tiêu trung hạn: tối ưu như lập lịch lệnh, cải thiện cấp phát thanh ghi, constant folding, loop-invariant code motion
• Mục tiêu dài hạn: bổ sung các backend mới như Tenstorrent, Intel Arc, RISC-V Vector Extension

Giấy phép

• Giấy phép Apache 2.0

[GN⁺]

Ý kiến trên Hacker News

• Khiếu hài hước kiểu New Zealand trong README của dự án khá ấn tượng
  Việc tự triển khai mã hóa instruction mà không phụ thuộc vào LLVM rất mới mẻ
  Điều này cho thấy cần lượng kiến thức mức thấp khổng lồ để bắt đầu một dự án như vậy
  Trong phe AMD, việc không hỗ trợ CUDA thường bị xem là cái cớ cho thế độc quyền của NVIDIA, nên những nỗ lực như thế này có vẻ sẽ giúp cân bằng thị trường

  • Chỉ ra rằng thứ được nói trong nguyên văn không phải LLM mà là LLVM
  • Tôi nghĩ dự án này chắc chắn cũng có đóng góp liên quan đến AI. Nhưng nếu dùng đúng cách thì AI có thể là một công cụ tuyệt vời
  • Nhắc đến “khiếu hài hước kiểu Oceania” làm tôi nhớ tới hoạt hình Beached Whale
  • Dòng chú thích trong mã nguồn “/* 80 keywords walk into a sorted bar */” rất dí dỏm (liên kết lexer.c)
  • Thấy cách diễn đạt “AI slope” nên đùa thêm rằng trong khoa học người ta gọi nó là gradient descent

• Tôi ngạc nhiên khi issue bên ngoài đầu tiên lại do geohot tạo (liên kết issue)
  Tôi muốn thấy những người như vậy hợp sức để phá vỡ thế độc quyền thị trường GPU của NVIDIA

• Chạy workload suy luận AI trên GPU NVIDIA có chi phí rất cao
  Những dự án như thế này rất quan trọng để các startup tạo ra những lựa chọn thay thế có thể chi trả được
  Tôi tò mò về hiệu năng ở các phép toán như conv2d hay attention

• “# It’s C99. It builds with gcc. There are no dependencies.”
  Sự đơn giản chỉ cần một dòng make thật sự là vẻ đẹp

  • Tôi rất thích cách tiếp cận đơn giản và trực quan như thế này

• Nhìn cách viết hoa trong tiêu đề bài đăng, giờ tôi mới hiểu vì sao GPU farm của công ty lại có tên là “barracuda”. Khá buồn cười

• Nếu các lập trình viên đầy nhiệt huyết làm được điều AMD không làm, thì đó sẽ là chuyện vừa buồn cười vừa đáng buồn

  • Tôi nghĩ AMD không hỗ trợ CUDA không phải vì không thể, mà là một lựa chọn chiến lược
    Hỗ trợ CUDA sẽ chỉ càng củng cố hệ sinh thái NVIDIA
    Nếu muốn một lựa chọn thay thế CUDA thì ZLUDA có thể thực dụng hơn
  • Mã nguồn mở không cần phải nhìn sắc mặt cổ đông nên thường tạo ra kết quả vượt thương mại
    Nhưng đáng tiếc là khi lớn lên thì cuối cùng lại bị các tập đoàn lớn mua lại rồi biến mất
    Như trường hợp của Linus và git, chỉ cần có ý chí và kiến thức là có thể phá được bức tường
  • Nhiều khi AMD không phải là “không làm được” mà là “không làm”
    Ví dụ như việc FSR4 không chính thức hỗ trợ các card cũ cũng là một trường hợp như vậy
  • Kết lại bằng câu đùa “chúng ta có HIP mà”

• Tôi tự hỏi liệu có thể hỗ trợ cả các kiến trúc AMD cũ hơn (GFX1010 v.v.) hay không

  • Khó nhưng tôi nghĩ hoàn toàn có thể. Trước đây ROCm rất tệ nên tôi phải tự vá, nhưng giờ đã khá hơn nhiều
  • Tôi cũng dùng RX5700m, nên đang tự mình làm phần hỗ trợ AMD đời cũ
    Tôi đang đọc tài liệu ISA và chỉnh lại mã hóa nhị phân
    Chỉ là đây là lĩnh vực LLM không làm tốt, nên phải tự hiểu và tự sửa

• CUDA hỗ trợ C++, nên tôi đã tự hỏi liệu đi theo hướng thuần C mà không có Clang/LLVM có quá hạn chế không

  • Thực ra hiện tại chỉ đang parse một phần các tính năng C++ mà CUDA dùng. Bản thân compiler được viết bằng C99
  • Nhìn vào test thì có vẻ một số tính năng C++ như template cũng xử lý được
  • Hỗ trợ AMD GX11 của LLVM rất hạn chế, nên không phụ thuộc vào nó sẽ tốt hơn
  • Lập trình viên thực thụ viết code không phụ thuộc AI
    Dạo này PR do AI tạo ra đang tràn ngập mã nguồn mở, nên việc dự án này tránh kiểu phụ thuộc đó khá ấn tượng

• Tôi nghĩ AMD nên tài trợ chính thức cho dự án như thế này
  Thế giới cần thoát khỏi thế độc quyền của NVIDIA

• Tôi đã tự hỏi liệu “OpenCL giờ coi như hết thời rồi sao?”
  Tôi không rõ, nhưng dù sao dự án này vẫn rất ấn tượng

  • CUDA đã giành được thị phần lớn hơn nhiều, còn OpenCL thì nền tảng yếu hơn
    Trước đây Apple cũng từng hỗ trợ nhưng giờ hình như không còn nữa
    Giống như mối quan hệ giữa Unix và Windows, về mặt kỹ thuật thì tốt nhưng thị trường lại nghiêng hẳn về một phía