SCALE - Menjalankan CUDA di GPU AMD tanpa modifikasi

• SCALE adalah toolkit pemrograman GPGPU yang memungkinkan aplikasi CUDA dikompilasi secara native untuk GPU AMD
• Tidak perlu memodifikasi program CUDA maupun sistem build, dan dukungan untuk lebih banyak vendor GPU serta API CUDA sedang dalam pengembangan

Bagaimana cara kerjanya?

• SCALE memiliki beberapa inovasi utama dibandingkan solusi GPGPU lintas platform lainnya
  • Menerima program CUDA apa adanya. Tidak perlu di-porting ke bahasa lain. Ini juga berlaku bahkan jika program menggunakan inline PTX asm
  • Kompiler SCALE menerima opsi baris perintah dan dialek CUDA yang sama dengan nvcc, sehingga berfungsi sebagai pengganti langsung
  • "Menyamar" sebagai instalasi NVIDIA CUDA Toolkit sehingga alat dan skrip build yang ada tetap bekerja seperti biasa

Proyek apa saja yang telah diuji?

• SCALE divalidasi dengan mengompilasi proyek CUDA open source dan menjalankan pengujiannya
• Saat ini proyek open source berikut termasuk dalam pengujian otomatis nightly dan semuanya lulus sepenuhnya
  • NVIDIA Thrust, Blender Cycles, AMGX, llama-cpp, faiss, xgboost, GOMC, stdgpu, hashcat

GPU apa saja yang didukung?

• Target GPU berikut didukung dan termasuk dalam pengujian nightly
  • AMD gfx1030 (Navi 21, RDNA 2.0)
  • AMD gfx1100 (Navi 31, RDNA 3.0)
• Target GPU berikut telah melalui pengujian manual sementara dan "tampaknya berfungsi"
  • AMD gfx1010
  • AMD gfx1101
• Dukungan untuk GPU berikut sedang dikerjakan
  • AMD gfx900 (Vega 10, GCN 5.0)
• Jika Anda ingin dukungan cepat untuk arsitektur GPU AMD tertentu, silakan hubungi mereka

Komponen SCALE

• Kompiler kompatibel nvcc yang dapat mengompilasi CUDA dialek nvcc untuk GPU AMD, termasuk PTX asm
• Implementasi runtime CUDA dan API driver untuk GPU AMD
• Library wrapper open source yang menyediakan API "CUDA-X" dengan mendelegasikan ke library ROCm. Library seperti cuBLAS dan cuSOLVER ditangani dengan cara ini

Perbedaan SCALE dengan solusi lain

• Alih-alih menyediakan cara baru untuk menulis perangkat lunak GPGPU, SCALE memungkinkan program yang ditulis dalam bahasa CUDA yang digunakan luas untuk langsung dikompilasi bagi GPU AMD
• SCALE bertujuan untuk sepenuhnya kompatibel dengan NVIDIA CUDA. Mereka meyakini pengguna seharusnya dapat mendukung banyak vendor GPU tanpa harus memelihara banyak codebase atau berkompromi pada performa
• Bahasa SCALE adalah superset dari NVIDIA CUDA, yang menawarkan ekstensi bahasa opsional agar penulisan kode GPU menjadi lebih mudah dan efisien bagi pengguna yang ingin melampaui nvcc
• SCALE adalah pekerjaan yang masih berlangsung. Jika ada API yang belum tersedia dan menghalangi penggunaan, silakan hubungi mereka. Prioritas pengembangan akan disesuaikan

Ringkasan GN⁺

• SCALE adalah toolkit penting yang memungkinkan aplikasi CUDA dikompilasi secara native untuk GPU AMD
• Ada keuntungan besar bagi pengembang karena dapat dijalankan di GPU AMD tanpa perlu memodifikasi program CUDA yang sudah ada
• Karena menargetkan kompatibilitas penuh dengan NVIDIA CUDA, ini menguntungkan untuk mendukung banyak vendor GPU
• Ini adalah proyek yang masih berjalan, dan jika API yang dibutuhkan belum tersedia, Anda dapat menghubungi tim pengembang untuk menyesuaikan prioritas
• Proyek dengan fungsi serupa termasuk ROCm dan HIP

[GN⁺]

Komentar Hacker News

• Ada pendapat bahwa banyak orang berpikir AMD harus mendukung lapisan translasi, tetapi itu adalah ide yang buruk

  • CUDA tidak dirancang agar netral terhadap vendor, dan Nvidia dapat menciptakan kesulitan secara teknis maupun hukum
  • Misalnya, menjalankan cuDNN atau cuBLAS di atas ini bisa melanggar perjanjian lisensi
  • Library Nvidia ini akan menjadi bagian dari batas API yang harus diimplementasikan ulang dan didukung oleh AMD

• Ada pendapat bahwa mengejar kompatibilitas bug adalah hal yang bodoh

  • Pengguna CUDA yang penting bersifat open source
  • AMD dapat mengimplementasikan dukungan secara langsung di proyek upstream seperti pytorch atau llama.cpp
  • Jika dukungannya ada, komunitas dapat memeliharanya

• Ada pendapat bahwa sulit memahami bagaimana kode yang sangat bergantung pada hardware bisa "langsung berfungsi" di AMD

  • Sebagian besar kode CUDA yang serius memahami register file dan ukuran shared memory, instruksi wgmma, tata letak memori dan register tensor core yang optimal, instruksi tensor memory accelerator, dan sebagainya

• Ada pendapat bahwa jika ini benar maka mengesankan, tetapi proyek ini tidak open source dan kurang detail yang tepat tentang cara kerjanya

  • Tidak begitu paham mengapa belakangan ini orang mengharapkan proyek menjadi open source atau setidaknya source-available

• Ada pendapat bahwa penyebab utama valuasi Nvidia yang tinggi adalah karena AMD tidak berinvestasi untuk membuat GPU-nya berguna bagi ML

  • Mungkin AMD takut pada tindakan antimonopoli, atau ada sesuatu dalam pendekatan hardware-nya yang membatasi daya saing
  • Perusahaan ini tampaknya melewatkan miliaran dolar selama lonjakan permintaan GPU untuk penambangan kripto dan lonjakan permintaan akibat ledakan AI saat ini

• Ada pendapat bahwa AMD sudah melakukan terlalu banyak kesalahan sampai-sampai proyek seperti ini ingin diberi selamat

  • Terutama di Linux, sangat membuat frustrasi ketika fungsi laptop secara fisik ada tetapi tidak bisa digunakan

• Pernah bekerja di Spectral Compute beberapa tahun lalu

  • Tim teknisnya sangat cerdas dan sangat kompeten
  • Saat itu mereka bukan hanya menargetkan AMD, tetapi juga melampaui backend LLVM ptx dasar dan NVCC

• Pernah menulis sedikit CUDA

  • Penasaran seperti apa setup dasar untuk menulis kode bagi kartu AMD

• Ada pendapat bahwa proyek ini hebat

  • Menantikan AMD bisa bersaing langsung dengan Nvidia

• Bagus bahwa ada halaman mengenai keterbatasan saat ini, tetapi ada pendapat bahwa apa yang kebanyakan orang sebut sebagai "CUDA" hanyalah sebagian kecil dari kemampuan CUDA yang sebenarnya

  • Akan bagus jika ada tabel perbandingan untuk fitur lanjutan seperti warp shuffle, operasi atomik, DPX, TMA, MMA, dan lain-lain
  • Tabel yang memetakan instruksi PTX ke instruksi padanan RDNA atau daftar instruksi yang mengemulasikannya akan menjadi ideal

• Ada pendapat bahwa ini mungkin nyata karena secara teknis memungkinkan

  • Mem-parse inline PTX dan memetakannya ke AMDGPU akan menjadi pekerjaan yang sangat menyakitkan
  • Membuat source CUDA yang tidak menggunakan inline PTX menargetkan AMDGPU mirip dengan menggantinya ke HIP
  • Beberapa detail mungkin meragukan, misalnya model atomiknya mungkin tidak cocok atau Volta bisa memiliki model instruction pointer yang berbeda
  • Namun, ini bisa dilakukan dengan benar
  • AMD tidak akan melakukannya
  • CUDA umumnya bukan sesuatu yang sangat baik, dan tim legal akan menimbulkan masalah
  • Namun, pihak lain sangat mungkin bisa melakukannya