VectorWare, которая строит первый GPU-native software company, объявила: Rust portable SIMD (core::simd) теперь работает на GPU. Раньше компания уже принесла потоки Rust на GPU, отображая каждый std::thread на warp, но параллельные lane'ы внутри warp'а не использовались. На CPU параллелизм внутри потока — это SIMD: одна инструкция обрабатывает несколько элементов. В Rust исторически SIMD писали через архитектурно-зависимые интринсики в core::arch, например _mm256_add_ps для x86-64 или vaddq_f32 для Arm. core::simd даёт абстракцию: тип Simd<T, N> — вектор из N элементов T, компилятор сам подставляет нужные инструкции.
В VectorWare поняли, что GPU — это просто ещё одно векторное устройство. NVIDIA называет свою модель SIMT: warp исполняет одну инструкцию сразу на 32 lanes. Это и есть SIMD, только lane'ы адресуются индивидуально. Поэтому Simd<i16, 32> даёт каждому из 32 lanes по одному элементу, а сложение двух таких векторов компилируется в одну warp-инструкцию. Получается полная иерархия: на CPU поток содержит SIMD lanes, на GPU наш std::thread — это warp, и в обоих случаях core::simd управляет lanes.
Код остаётся обычным Rust. Пример функции relu_dot — поэлементное умножение, сравнение с маской, select и горизонтальная редукция — работает и на CPU, и на GPU без изменений. Реализация прямая: поэлементные операции идут через обычные trait'ы, редукции через warp shuffle, перестановки тоже, маски через vote/ballot. Скалярные значения в коде считаются одинаково всеми lanes, поэтому просто реплицируются, как в CUDA.
Единственное несовпадение — ширина: на CPU Simd<T,N> позволяет любой N до 64, а GPU имеет фиксированные 32 или 64 lanes. Чтобы аккуратно размещать работу, в VectorWare закодировали специальный IR прямо в системе типов Rust: операции, shape выполнения и capacity — всё типизировано, многие некорректные программы просто не скомпилируются. IR без интерпретатора на GPU, каждая операция превращается в PTX. Для CPU есть reference interpreter, как Miri, для симуляции и дифференциального тестирования.
Плюсы: один и тот же код на CPU и GPU, обычные значения Simd<T,N> проходят через borrow checker, ничего GPU-специфичного не добавляется. Минусы: portable_simd пока unstable, векторы не кратные ширине warp'а работают медленнее, не все перестановки эффективны, а редукции становятся точками синхронизации. Пришлось менять компилятор для безопасности.
Дальше в планах — совместить SIMD, threads и async, опустить matrix-SIMD на тензорные ядра, авто-векторизация обычных циклов. Правда, текущие portable SIMD типы, возможно, не идеальны для общего векторного представления CPU и GPU.
Компания подчёркивает: будущие продукты будут поддерживать несколько языков, но Rust идеально подходит для GPU-native разработки.