Skip to content
Search

SCALE

lastmod 2026-02-19created 2024-07-172 edits

Run CUDA unmodified on AMD GPUs

Features

  • SCALE is a GPGPU programming toolkit that lets CUDA applications be natively compiled for AMD GPUs
  • No need to modify CUDA programs or build systems, and support for more GPU vendors and CUDA APIs is in development

How does it work?

  • SCALE has several key innovations compared to other cross-platform GPGPU solutions
    • Accepts CUDA programs as-is. No need to port them to another language. This holds even when the program uses inline PTX asm
    • The SCALE compiler accepts the same command-line options and CUDA dialect as nvcc, working as a drop-in replacement
    • "Impersonates" an NVIDIA CUDA Toolkit installation so existing build tools and scripts work as-is

Which projects have been tested?

  • SCALE is validated by compiling open-source CUDA projects and running their tests
  • The following open-source projects are currently included in nightly automated testing and pass fully
    • NVIDIA Thrust, Blender Cycles, AMGX, llama-cpp, faiss, xgboost, GOMC, stdgpu, hashcat

Which GPUs are supported?

  • The following GPU targets are supported and included in Nightly testing
    • AMD gfx1030 (Navi 21, RDNA 2.0)
    • AMD gfx1100 (Navi 31, RDNA 3.0)
  • The following GPU targets have undergone ad hoc manual testing and "seem to work"
    • AMD gfx1010
    • AMD gfx1101
  • Work is underway to support the following GPU
    • AMD gfx900 (Vega 10, GCN 5.0)
  • Contact them if you want support for a specific AMD GPU architecture quickly

Components of SCALE

  • An nvcc-compatible compiler that can compile nvcc-dialect CUDA for AMD GPUs, including PTX asm
  • Implementations of the CUDA runtime and driver APIs for AMD GPUs
  • Open-source wrapper libraries that provide "CUDA-X" APIs by delegating to ROCm libraries. Libraries such as cuBLAS and cuSOLVER are handled this way

How SCALE differs from other solutions

  • Instead of offering a new way to write GPGPU software, SCALE lets programs written in the widely used CUDA language be compiled directly for AMD GPUs
  • SCALE aims to be fully compatible with NVIDIA CUDA. They believe users should be able to support multiple GPU vendors without maintaining multiple codebases or compromising performance
  • SCALE's language is a superset of NVIDIA CUDA, offering optional language extensions that make writing GPU code easier and more efficient for users who want to move away from nvcc
  • SCALE is a work in progress. If an API you need is missing, contact them. They will adjust development priorities

GN⁺'s summary

  • SCALE is an important toolkit that lets CUDA applications be natively compiled for AMD GPUs
  • Existing CUDA programs can run on AMD GPUs without modification, which is a big advantage for developers
  • Aims for full compatibility with NVIDIA CUDA, which helps in supporting multiple GPU vendors
  • A project in progress; if a needed API is missing, you can contact the development team to adjust priorities
  • Projects with similar functionality include ROCm and HIP

Hacker News comments

  • Run CUDA, unmodified, on AMD GPUs | Hacker News
  • Many people think AMD should support a translation layer, but one opinion is that this is a bad idea
    • CUDA was not designed to be vendor-neutral, and Nvidia can create technical and legal difficulties
    • For example, running cuDNN or cuBLAS on top of it may violate license agreements
    • These Nvidia libraries would become part of the API boundary that AMD would have to reimplement and support
  • One opinion is that chasing bug compatibility is foolish
    • The important CUDA users are open source
    • AMD can implement support directly in upstream projects like pytorch or llama.cpp
    • Once support exists, the community can maintain it
  • One opinion is that it's hard to understand how code that depends heavily on hardware could "just work" on AMD
    • Most serious CUDA code is aware of register file and shared memory sizes, wgmma instructions, optimal tensor core memory and register layouts, tensor memory accelerator instructions, etc.
  • One opinion is that it's impressive if true, but it is not open source and lacks precise details on how it works
    • Not sure why people these days expect projects to be open source or at least source-available
  • One opinion is that the main reason for Nvidia's high valuation is that AMD doesn't invest in making its GPUs useful for ML
    • AMD may fear antitrust action, or there may be something in its hardware approach that limits competitiveness
    • The company seems to have missed out on billions of dollars during the surge in GPU demand for cryptocurrency mining and the current AI boom
  • One opinion is that AMD has done so badly that they want to celebrate projects like this
    • Especially on Linux, it's very frustrating that laptop features physically exist but can't be used
  • Worked at Spectral Compute a few years ago
    • It was a very smart and capable technical team
    • At the time, it not only targeted AMD but also outperformed the stock LLVM ptx backend and NVCC
  • Have written a little CUDA
    • Wondering what the default setup is for writing code for AMD cards
  • One opinion is that this project is great
    • Looking forward to AMD competing directly with Nvidia
  • One opinion is that it's good there is a page on current limitations, but what most people describe as "CUDA" is a small part of actual CUDA functionality
    • A comparison table for advanced features such as warp shuffles, atomic operations, DPX, TMA, MMA, etc. would be nice
    • Ideally, a table mapping PTX instructions to their RDNA counterparts or to the list of instructions that emulate them
  • One opinion is that it could be real since it is technically possible
    • Parsing inline PTX and mapping it to AMDGPU would be a huge pain
    • Targeting AMDGPU from CUDA source that doesn't use inline PTX is similar to replacing it with HIP
    • Some details may be questionable, e.g. atomic models not matching or Volta having a different instruction pointer model
    • But it can be done correctly
    • AMD won't do it
    • CUDA is generally not a very good thing, and the legal team will cause problems
    • But others can certainly do it

See also

Favorite site

author Zer0 sitemap