SCALE - Ejecutar CUDA en GPUs de AMD sin modificaciones

• SCALE es un toolkit de programación GPGPU que permite compilar de forma nativa aplicaciones CUDA para GPUs de AMD
• No es necesario modificar los programas CUDA ni el sistema de compilación, y se está desarrollando soporte para más fabricantes de GPU y para la API de CUDA

¿Cómo funciona?

• SCALE tiene varias innovaciones clave en comparación con otras soluciones GPGPU multiplataforma
  • Acepta programas CUDA tal como están. No hace falta portarlos a otro lenguaje. Esto también aplica cuando el programa usa asm de PTX en línea
  • El compilador de SCALE acepta las mismas opciones de línea de comandos y el mismo dialecto de CUDA que nvcc, por lo que funciona como un reemplazo directo
  • "Simula" una instalación de NVIDIA CUDA Toolkit para que las herramientas y scripts de compilación existentes sigan funcionando sin cambios

¿Qué proyectos se han probado?

• SCALE se validó compilando proyectos CUDA de código abierto y ejecutando sus pruebas
• Actualmente, los siguientes proyectos de código abierto están incluidos en las pruebas automáticas nocturnas y las pasan por completo
  • NVIDIA Thrust, Blender Cycles, AMGX, llama-cpp, faiss, xgboost, GOMC, stdgpu, hashcat

¿Qué GPUs son compatibles?

• Se admiten los siguientes objetivos de GPU y están incluidos en las pruebas Nightly
  • AMD gfx1030 (Navi 21, RDNA 2.0)
  • AMD gfx1100 (Navi 31, RDNA 3.0)
• Los siguientes objetivos de GPU han pasado pruebas manuales ocasionales y "parecen funcionar"
  • AMD gfx1010
  • AMD gfx1101
• Se está trabajando en compatibilidad para las siguientes GPUs
  • AMD gfx900 (Vega 10, GCN 5.0)
• Si necesitas compatibilidad rápida para una arquitectura específica de GPU AMD, ponte en contacto con ellos

Componentes de SCALE

• Un compilador compatible con nvcc capaz de compilar CUDA en dialecto nvcc para GPUs de AMD, incluyendo PTX asm
• Implementación del runtime de CUDA y de la API del driver para GPUs de AMD
• Bibliotecas wrapper de código abierto que delegan en bibliotecas de ROCm para ofrecer la API "CUDA-X". Bibliotecas como cuBLAS y cuSOLVER se manejan de esta forma

Diferencias entre SCALE y otras soluciones

• En lugar de ofrecer una nueva forma de escribir software GPGPU, SCALE permite compilar directamente para GPUs de AMD programas escritos en el ampliamente usado lenguaje CUDA
• SCALE apunta a ser totalmente compatible con NVIDIA CUDA. Considera que los usuarios deberían poder dar soporte a varios fabricantes de GPU sin mantener múltiples bases de código ni sacrificar rendimiento
• El lenguaje de SCALE es un superconjunto de NVIDIA CUDA, y ofrece extensiones opcionales del lenguaje que hacen más fácil y eficiente escribir código GPU para quienes quieren ir más allá de nvcc
• SCALE es un trabajo en progreso. Si falta alguna API que impida su uso, piden que te pongas en contacto para ajustar las prioridades de desarrollo

Resumen de GN⁺

• SCALE es un toolkit importante que permite compilar de forma nativa aplicaciones CUDA para GPUs de AMD
• Poder ejecutar programas CUDA existentes en GPUs de AMD sin necesidad de modificarlos representa una gran ventaja para los desarrolladores
• Al apuntar a una compatibilidad total con NVIDIA CUDA, resulta favorable para dar soporte a múltiples fabricantes de GPU
• Es un proyecto en desarrollo, y si falta alguna API necesaria se puede contactar al equipo para ajustar prioridades
• Un proyecto con funciones similares es ROCm con HIP

[GN⁺]

Comentarios en Hacker News

• Algunas personas creen que AMD debería soportar una capa de traducción, pero hay quien opina que es una mala idea

  • CUDA no fue diseñado para ser neutral entre proveedores, y Nvidia podría crear dificultades técnicas y legales
  • Por ejemplo, ejecutar cuDNN o cuBLAS encima de esto podría violar los acuerdos de licencia
  • Estas bibliotecas de Nvidia pasarían a formar parte del límite de API que AMD tendría que reimplementar y dar soporte

• Hay quien opina que buscar compatibilidad con bugs es una tontería

  • Los usuarios importantes de CUDA son de código abierto
  • AMD podría implementar soporte directamente en proyectos upstream como pytorch o llama.cpp
  • Si existe soporte, la comunidad puede encargarse del mantenimiento

• Hay quien dice no entender cómo código muy dependiente del hardware podría simplemente "funcionar" en AMD

  • La mayoría del código CUDA serio tiene en cuenta el tamaño del archivo de registros y de la memoria compartida, instrucciones wgmma, disposiciones óptimas de memoria y registros para tensor cores, instrucciones del acelerador de memoria de tensores, etc.

• Si es cierto, es impresionante, pero hay quien señala que no es de código abierto y faltan detalles precisos sobre cómo funciona

  • No queda claro por qué hoy en día se espera que un proyecto sea open source o al menos source-available

• Hay quien opina que la principal razón de la alta valoración de Nvidia es que AMD no invierte en hacer que sus GPU sean útiles para ML

  • Puede que AMD tema medidas antimonopolio, o que haya algo en su enfoque de hardware que limite su competitividad
  • La empresa parece haber dejado pasar miles de millones de dólares tanto durante el auge de demanda de GPU para minería de criptomonedas como en el actual boom de la IA

• Hay quien quiere felicitar proyectos como este porque AMD lo ha hecho demasiado mal

  • Especialmente en Linux, es muy frustrante que funciones de una laptop existan físicamente pero no se puedan usar

• Alguien trabajó en Spectral Compute hace unos años

  • Era un equipo técnico muy inteligente y capaz
  • En ese momento no solo apuntaban a AMD, sino que también superaban al backend base de LLVM ptx y a NVCC

• Alguien ha escrito un poco de CUDA

  • Tiene curiosidad por saber cuál es la configuración básica para escribir código para tarjetas AMD

• Hay quien opina que este proyecto es excelente

  • Espera que AMD compita directamente con Nvidia

• Es bueno que haya una página sobre las limitaciones actuales, pero hay quien comenta que lo que la mayoría describe como "CUDA" es solo una pequeña parte de la funcionalidad real de CUDA

  • Sería bueno tener una tabla comparativa de funciones avanzadas como warp shuffle, operaciones atómicas, DPX, TMA, MMA, etc.
  • Una tabla que mapee instrucciones PTX a instrucciones equivalentes en RDNA, o a la lista de instrucciones que las emulan, sería ideal

• Hay quien opina que podría ser real porque es técnicamente posible

  • Parsear PTX inline y mapearlo a AMDGPU sería un dolor enorme
  • Hacer que código fuente CUDA sin PTX inline funcione apuntando a AMDGPU sería parecido a reemplazarlo con HIP
  • Algunos detalles pueden ser dudosos, por ejemplo, que el modelo atómico no coincida o que Volta tenga un modelo distinto de puntero de instrucción
  • Pero podría hacerse correctamente
  • AMD no lo hará
  • CUDA en general no es algo particularmente bueno y el equipo legal causaría problemas
  • Pero otras personas sí podrían hacerlo