SCALE - Chạy CUDA trên GPU AMD mà không cần chỉnh sửa

• SCALE là một bộ công cụ lập trình GPGPU cho phép biên dịch native các ứng dụng CUDA cho GPU AMD
• Không cần chỉnh sửa chương trình CUDA hay hệ thống build, và hiện đang phát triển thêm hỗ trợ cho nhiều nhà cung cấp GPU hơn cùng với API CUDA

Hoạt động như thế nào?

• So với các giải pháp GPGPU đa nền tảng khác, SCALE có một số đổi mới chính
  • Chấp nhận nguyên trạng các chương trình CUDA. Không cần port sang ngôn ngữ khác. Điều này áp dụng ngay cả khi chương trình sử dụng PTX asm nội tuyến
  • Trình biên dịch SCALE chấp nhận cùng các tùy chọn dòng lệnh và phương ngữ CUDA như nvcc, nên có thể hoạt động như một bản thay thế drop-in
  • "Giả làm" cài đặt NVIDIA CUDA Toolkit để các công cụ và script build hiện có vẫn hoạt động nguyên vẹn

Những dự án nào đã được kiểm thử?

• SCALE được xác thực bằng cách biên dịch các dự án CUDA mã nguồn mở và chạy các bài kiểm thử
• Hiện tại, các dự án mã nguồn mở sau đây được đưa vào kiểm thử tự động hằng đêm và đều vượt qua hoàn toàn
  • NVIDIA Thrust, Blender Cycles, AMGX, llama-cpp, faiss, xgboost, GOMC, stdgpu, hashcat

Những GPU nào được hỗ trợ?

• Các mục tiêu GPU sau được hỗ trợ và được đưa vào kiểm thử Nightly
  • AMD gfx1030 (Navi 21, RDNA 2.0)
  • AMD gfx1100 (Navi 31, RDNA 3.0)
• Các mục tiêu GPU sau đã trải qua kiểm thử thủ công tạm thời và "có vẻ hoạt động"
  • AMD gfx1010
  • AMD gfx1101
• Đang thực hiện hỗ trợ cho các GPU sau
  • AMD gfx900 (Vega 10, GCN 5.0)
• Nếu muốn đẩy nhanh hỗ trợ cho một kiến trúc GPU AMD cụ thể, hãy liên hệ

Các thành phần của SCALE

• Trình biên dịch tương thích nvcc có thể biên dịch CUDA theo phương ngữ nvcc cho GPU AMD, bao gồm cả PTX asm
• Triển khai API runtime và driver CUDA cho GPU AMD
• Thư viện wrapper mã nguồn mở cung cấp API "CUDA-X" bằng cách ủy quyền cho các thư viện ROCm. Các thư viện như cuBLAS và cuSOLVER được xử lý theo cách này

Điểm khác biệt giữa SCALE và các giải pháp khác

• Thay vì cung cấp một cách mới để viết phần mềm GPGPU, SCALE cho phép biên dịch trực tiếp các chương trình viết bằng ngôn ngữ CUDA phổ biến sang GPU AMD
• SCALE hướng tới khả năng tương thích hoàn toàn với NVIDIA CUDA. Họ tin rằng người dùng nên có thể hỗ trợ nhiều nhà cung cấp GPU mà không phải duy trì nhiều codebase hoặc đánh đổi hiệu năng
• Ngôn ngữ của SCALE là một siêu tập của NVIDIA CUDA, cung cấp các phần mở rộng ngôn ngữ tùy chọn giúp việc viết mã GPU dễ dàng và hiệu quả hơn cho những người dùng muốn vượt ra ngoài nvcc
• SCALE là một dự án đang trong quá trình phát triển. Nếu thiếu API nào cản trở việc sử dụng, hãy liên hệ. Họ sẽ điều chỉnh ưu tiên phát triển

Tổng hợp của GN⁺

• SCALE là một bộ công cụ quan trọng cho phép biên dịch native các ứng dụng CUDA cho GPU AMD
• Có thể chạy trên GPU AMD mà không cần chỉnh sửa các chương trình CUDA hiện có, mang lại lợi ích lớn cho nhà phát triển
• Hướng tới khả năng tương thích hoàn toàn với NVIDIA CUDA, thuận lợi cho việc hỗ trợ nhiều nhà cung cấp GPU
• Đây là một dự án đang được phát triển; nếu thiếu API cần thiết, có thể liên hệ với nhóm phát triển để điều chỉnh mức độ ưu tiên
• Các dự án có chức năng tương tự bao gồm ROCm và HIP

[GN⁺]

Ý kiến Hacker News

• Có ý kiến cho rằng nhiều người nghĩ AMD nên hỗ trợ một lớp dịch, nhưng đây là một ý tưởng tồi

  • CUDA không được thiết kế theo hướng trung lập với nhà cung cấp, và Nvidia có thể tạo ra khó khăn cả về kỹ thuật lẫn pháp lý
  • Ví dụ, việc chạy cuDNN hay cuBLAS trên lớp này có thể vi phạm thỏa thuận cấp phép
  • Các thư viện Nvidia này sẽ trở thành một phần của ranh giới API mà AMD phải tái triển khai và hỗ trợ

• Có ý kiến cho rằng theo đuổi khả năng tương thích cả với lỗi là điều ngớ ngẩn

  • Những người dùng CUDA quan trọng đều là mã nguồn mở
  • AMD có thể triển khai hỗ trợ trực tiếp trong các dự án upstream như pytorch hoặc llama.cpp
  • Nếu có hỗ trợ, cộng đồng có thể tự duy trì nó

• Có ý kiến cho rằng không thể hiểu nổi làm sao mã phụ thuộc rất nhiều vào phần cứng lại có thể "chạy được ngay" trên AMD

  • Phần lớn mã CUDA nghiêm túc đều nhận thức được register file và kích thước shared memory, lệnh wgmma, bố cục bộ nhớ và thanh ghi tối ưu cho tensor core, các lệnh Tensor Memory Accelerator, v.v.

• Nếu đúng thì rất ấn tượng, nhưng có ý kiến cho rằng dự án này không phải mã nguồn mở và cũng thiếu các chi tiết chính xác về cách nó hoạt động

  • Không rõ vì sao ngày nay người ta lại kỳ vọng các dự án phải là mã nguồn mở hoặc ít nhất là source-available

• Có ý kiến cho rằng nguyên nhân chính khiến Nvidia được định giá cao là vì AMD không đầu tư để biến GPU thành thứ hữu ích cho ML

  • Có thể AMD sợ các biện pháp chống độc quyền, hoặc có điều gì đó trong cách tiếp cận phần cứng đang hạn chế khả năng cạnh tranh của họ
  • Có vẻ như công ty đã bỏ lỡ hàng tỷ USD trong đợt bùng nổ nhu cầu GPU cho đào tiền mã hóa và cả đợt bùng nổ nhu cầu AI hiện nay

• Có ý kiến cho rằng AMD làm quá tệ đến mức người ta muốn chúc mừng những dự án như thế này

  • Đặc biệt trên Linux, việc phần cứng của laptop về mặt vật lý có tồn tại nhưng lại không dùng được thật sự rất bực bội

• Có người từng làm việc ở Spectral Compute vài năm trước

  • Đó là một đội ngũ kỹ thuật rất thông minh và có năng lực
  • Khi đó họ không chỉ nhắm tới AMD mà còn vượt qua cả backend LLVM ptx cơ bản và NVCC

• Có người từng viết một chút CUDA

  • Họ tò mò thiết lập cơ bản để viết mã cho card AMD là gì

• Có ý kiến cho rằng dự án này rất tuyệt

  • Họ mong chờ AMD cạnh tranh trực tiếp với Nvidia

• Có ý kiến cho rằng việc có một trang về các giới hạn hiện tại là tốt, nhưng thứ mà đa số mọi người gọi là "CUDA" thực ra chỉ là một phần nhỏ trong toàn bộ tính năng CUDA

  • Sẽ rất hay nếu có bảng so sánh cho các tính năng nâng cao như warp shuffle, atomic operations, DPX, TMA, MMA, v.v.
  • Lý tưởng nhất là có một bảng ánh xạ các lệnh PTX sang các lệnh tương ứng của RDNA hoặc danh sách các lệnh dùng để giả lập chúng

• Có ý kiến cho rằng vì về mặt kỹ thuật là khả thi nên chuyện này có thể là thật

  • Việc phân tích inline PTX và ánh xạ nó sang AMDGPU sẽ là một nỗi đau rất lớn
  • Việc nhắm AMDGPU từ mã nguồn CUDA không dùng inline PTX khá giống với việc thay thế bằng HIP
  • Một số chi tiết có thể đáng nghi, chẳng hạn như mô hình atomic không khớp hoặc Volta có thể có mô hình instruction pointer khác
  • Tuy nhiên, điều đó vẫn có thể được làm đúng
  • AMD sẽ không làm việc này
  • CUDA nói chung không phải thứ gì quá tốt đẹp và đội ngũ pháp lý sẽ gây rắc rối
  • Tuy vậy, những người khác hoàn toàn có thể làm được