VectorWare ejecuta el SIMD portátil de Rust sin cambios en la GPU

⏱️ Lectura: 11 min

VectorWare, una startup dedicada a construir software nativo para GPU, logró que el SIMD en GPU corra con exactamente el mismo código Rust que ya usa la CPU. La empresa lo anunció en su blog oficial, mostrando una función que compila a una instrucción vectorial de x86-64 en la laptop y a una instrucción de warp en la tarjeta gráfica, sin cambiar una sola línea.

📑 En este artículo

1. TL;DR
2. Qué pasó
3. Contexto e historia
4. Detalles técnicos: así funciona SIMD en GPU
5. Cómo empezar a probarlo
6. Impacto y análisis
7. Qué sigue
8. Preguntas frecuentes
   1. ¿Qué es el SIMD portátil de Rust (core::simd)?
   2. ¿Qué diferencia hay entre SIMD y SIMT?
   3. ¿Necesito reescribir mi código para que corra en la GPU?
   4. ¿core::simd está disponible en Rust estable?
   5. ¿Qué alternativas existen hoy para programar GPU en Rust?
   6. ¿El compilador de VectorWare para GPU es de código abierto?
9. Referencias

El logro importa porque hasta ahora programar SIMD en GPU implicaba aprender un lenguaje de shaders distinto o escribir kernels de CUDA aparte. Con este anuncio, la abstracción Simd<T, N> de Rust deja de ser exclusiva de la CPU.

TL;DR

• VectorWare logró compilar el SIMD portátil de Rust (core::simd) para que también corra en la GPU.
• La misma función, sin cambios, compila a una instrucción de CPU y a una instrucción de warp en GPU.
• Un Simd de 32 elementos i16 ocupa las 32 lanes de un warp: vpaddw en CPU, add.s16 en GPU (PTX).
• El anuncio se apoya en trabajo previo de VectorWare, que ya mapeó std::thread a un warp de GPU.
• core::simd vive en core, no en std, y no depende del soporte de std que la empresa construyó antes.
• NVIDIA llama SIMT a su modelo, pero VectorWare sostiene que un warp es, en esencia, una unidad SIMD.
• La API sigue detrás del feature nightly portable_simd, sin fecha pública de estabilización en Rust.

Qué pasó

VectorWare publicó un post técnico donde demuestra que el tipo genérico Simd<T, N> de Rust, pensado originalmente para vectorizar código en la CPU, también compila hacia instrucciones de warp en la GPU. El ejemplo que muestran define una función que multiplica dos vectores de 32 elementos, aplica una máscara con una comparación y reduce el resultado a un escalar. Esa misma función, sin tocar una línea, corre en la CPU del autor y en una GPU NVIDIA.

La compañía ya había resuelto un problema parecido antes: mapear std::thread de Rust a un warp de la GPU, de forma que cada hilo lógico del programa ocupe un warp físico. Ese trabajo previo, según el blog de VectorWare, es la base sobre la que construyeron el soporte de SIMD en GPU.

Contexto e historia

Durante años, escribir SIMD en Rust significó usar los intrínsecos específicos de cada arquitectura en core::arch: funciones como _mm256_add_ps en x86-64 o vaddq_f32 en ARM. Cada arquitectura tenía su propio nombre de función, así que un programa multiplataforma necesitaba una implementación separada por cada una.

El proyecto portable-simd de Rust resolvió ese problema agregando una capa de abstracción: un tipo genérico Simd<T, N> que representa un vector de N elementos de tipo T. El programador escribe aritmética, comparaciones y reducciones una sola vez contra Simd, y el compilador decide qué instrucciones vectoriales usar según el objetivo de compilación.

VectorWare descubrió que la GPU es, ni más ni menos, otra pieza de hardware vectorial a la que core::simd puede apuntar. Como beneficio adicional, core::simd vive en core y no en std, así que ni siquiera depende del soporte de std para GPU que la empresa construyó en su trabajo anterior.

Detalles técnicos: así funciona SIMD en GPU

NVIDIA llama a su modelo de ejecución SIMT, o Single Instruction, Multiple Thread. Un warp emite una sola instrucción y cada uno de sus 32 lanes la ejecuta sobre su propio dato. Eso es, en esencia, SIMD: una instrucción que opera sobre muchos elementos a la vez. El direccionamiento por lane que agrega SIMT no cambia esa naturaleza. Un warp es una unidad vectorial ancha, y un vector Simd portátil encaja directamente sobre esa unidad.

Un Simd<i16, 32> le da un elemento i16 a cada uno de los 32 lanes del warp. Sumar dos vectores de ese tipo compila, en la CPU, a una sola instrucción vpaddw de AVX-512. En la GPU, la misma suma se traduce a add.s16 en ensamblador PTX. El código fuente de Rust es idéntico en los dos casos.

#![feature(portable_simd)]
use core::simd::Simd;

fn sumar_vectores(a: Simd<f32, 8>, b: Simd<f32, 8>) -> Simd<f32, 8> {
    a + b
}

fn main() {
    let a = Simd::<f32, 8>::splat(3.0);
    let b = Simd::from_array([1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0]);
    println!("{:?}", sumar_vectores(a, b));
}

Esta función corre hoy en cualquier CPU x86-64 o ARM con Rust nightly, y compila a una sola instrucción vectorial. Es el tipo de primitiva que VectorWare ahora también compila hacia la GPU.

El ejemplo real de VectorWare combina varias operaciones: multiplicación elemento a elemento, una comparación que genera una máscara booleana por lane, un select que filtra según esa máscara, y una reducción horizontal. Es el patrón detrás de una activación tipo ReLU en una capa de red neuronal:

use core::simd::cmp::SimdPartialOrd;
use core::simd::num::SimdFloat;
use core::simd::{Select, Simd};

fn producto_relu(pesos: Simd<f32, 32>, entradas: Simd<f32, 32>) -> f32 {
    let productos = pesos * entradas;
    let son_positivos = productos.simd_gt(Simd::splat(0.0));
    let filtrados = son_positivos.select(productos, Simd::splat(0.0));
    filtrados.reduce_sum()
}

producto_relu multiplica 32 pesos por 32 entradas en paralelo, descarta los productos negativos y suma lo que queda en un solo escalar. Según VectorWare, esta función compila y corre igual en CPU que en GPU, sin ninguna anotación especial.

💭 Clave: el punto de entrada sigue siendo un fn main normal, sin atributos de kernel ni anotaciones de GPU. El toolchain de VectorWare decide por detrás qué partes del programa compilar como kernel.

Así queda la jerarquía completa de paralelismo, con core::simd manejando el nivel más bajo en los dos mundos:

flowchart TD
    subgraph CPU["CPU"]
    T["thread"] --> L0["lane 0"]
    T --> L1["lane 1"]
    T --> L2["lane 2"]
    T --> LN["lane N"]
    end
    subgraph GPU["GPU"]
    W["warp"] --> G0["lane 0"]
    W --> G1["lane 1"]
    W --> G2["lane 2"]
    W --> GN["lane N"]
    end
    L0 -.-> G0
    L1 -.-> G1
    L2 -.-> G2
    LN -.-> GN

Cómo empezar a probarlo

El frontend de este trabajo, core::simd, es parte del Rust oficial y cualquiera puede probarlo hoy en la CPU sin esperar al toolchain de VectorWare. Los pasos son iguales en Windows, macOS y Linux porque dependen de rustup, no del sistema operativo:

# Windows, macOS y Linux (misma secuencia con rustup)
rustup toolchain install nightly
rustup override set nightly
cargo new simd-demo
cd simd-demo

Después, agregá #![feature(portable_simd)] como primera línea de src/main.rs, pegá el ejemplo de sumar_vectores de arriba y corré cargo run.

Para confirmar que el compilador generó una instrucción vectorial y no un bucle escalar, instalá cargo-show-asm e inspeccioná el ensamblador generado:

cargo install cargo-show-asm
cargo asm --release simd_demo::sumar_vectores

Si aparece una sola instrucción como vaddps (o addps sin AVX) en la salida, la vectorización funcionó. Ese mismo método, revisar el ensamblador generado, es el que usó VectorWare para mostrar que su compilador produce add.s16 en PTX del lado de la GPU.

⚠️ Ojo: el compilador que lleva este código a la GPU es propio de VectorWare y no forma parte del Rust oficial. Lo que se puede reproducir hoy con rustup es solo la mitad del anuncio: la ejecución en CPU.

Impacto y análisis

Rust ya tiene varios caminos para programar GPU: rust-gpu compila a SPIR-V para usar con Vulkan, wgpu corre shaders de cómputo escritos en WGSL, y crates como cudarc exponen bindings directos a CUDA. Lo distinto en el enfoque de VectorWare es que no hay un lenguaje de shader separado: la misma función Rust con core::simd apunta a los dos backends.

Opción	Cuándo usarla	Ventaja	Limitación
core::simd + VectorWare	Código que ya usa SIMD portátil y quiere correr también en GPU	Un solo código fuente para CPU y GPU	Toolchain propietario, no integrado en rustc oficial
rust-gpu	Shaders gráficos o de cómputo dentro de un pipeline Vulkan	Genera SPIR-V estándar	Requiere pensar en términos de shader, no de SIMD portátil
wgpu (WGSL)	Apps multiplataforma que ya usan wgpu para gráficos	Corre en web, desktop y móvil	El cómputo se escribe en un lenguaje aparte, WGSL
cudarc	Control fino sobre hardware NVIDIA específico	Acceso directo a la API de CUDA	Atado a NVIDIA, sin portabilidad a otros fabricantes

Una limitación real, no mencionada en el anuncio pero inherente al modelo: el tamaño de warp de 32 lanes es específico de NVIDIA. Un Simd<T, 32> pensado para ese número no aprovecha completo un wavefront de 64 lanes en hardware de otro fabricante, así que el tamaño del vector portátil sigue atando el código a una arquitectura concreta, aunque la sintaxis sea portátil.

Qué sigue

core::simd sigue detrás de un feature nightly, sin fecha pública de estabilización en el repositorio oficial de Rust. Eso significa que cualquier proyecto que lo adopte hoy, para CPU o para GPU, acepta el riesgo de que la API cambie antes de llegar a Rust estable. VectorWare no detalló en su post si planea abrir el código de su compilador para GPU o mantenerlo como producto propietario, ni qué arquitecturas además de NVIDIA piensa soportar.

📖 Resumen en Telegram: Ver resumen

Probalo vos: instalá Rust nightly, agregá #![feature(portable_simd)] a un crate nuevo y compará con cargo asm el ensamblador que genera core::simd en tu propia CPU.

Preguntas frecuentes

¿Qué es el SIMD portátil de Rust (core::simd)?

Es un tipo genérico, Simd<T, N>, que representa un vector de N elementos de tipo T. Reemplaza a los intrínsecos específicos de cada arquitectura, como _mm256_add_ps en x86-64, con una sola API que el compilador traduce según el hardware objetivo.

¿Qué diferencia hay entre SIMD y SIMT?

SIMD ejecuta una instrucción sobre varios datos dentro de un mismo hilo. SIMT, el modelo que usa NVIDIA en sus GPU, agrega direccionamiento por lane dentro de un warp, pero según el análisis de VectorWare sigue siendo SIMD en el fondo: un warp de 32 lanes es una unidad vectorial ancha.

¿Necesito reescribir mi código para que corra en la GPU?

Según el anuncio de VectorWare, no. La misma función que usa Simd<T, N> compila igual para CPU o para GPU; lo único que cambia es el backend del compilador que la procesa.

¿core::simd está disponible en Rust estable?

No. Requiere el feature nightly #![feature(portable_simd)] y no tiene fecha pública de estabilización en el repositorio oficial de Rust.

¿Qué alternativas existen hoy para programar GPU en Rust?

Entre las más usadas están rust-gpu (compila a SPIR-V), wgpu con shaders en WGSL, y cudarc con bindings directos a CUDA.

¿El compilador de VectorWare para GPU es de código abierto?

VectorWare no publicó el código de su toolchain de GPU junto con este anuncio. Solo core::simd, el frontend que sí corre en CPU hoy, es parte del proyecto oficial portable-simd de Rust.

Referencias

• VectorWare, “Rust SIMD on the GPU”: el anuncio original con el ejemplo de código y los diagramas de mapeo lane a lane.
• rust-lang/portable-simd en GitHub: repositorio oficial de la implementación de core::simd.
• NVIDIA, “Parallel Thread Execution ISA”: documentación oficial del modelo SIMT y el ensamblador PTX.
• Wikipedia, “Single instruction, multiple data”: contexto histórico sobre el modelo SIMD.

📱 ¿Te gusta este contenido? Únete a nuestro canal de Telegram @programacion donde publicamos a diario lo más relevante de tecnología, IA y desarrollo. Resúmenes rápidos, contenido fresco todos los días.

Imagen destacada: Foto de Anne Nygård en Unsplash