ИИтог

VectorWare научила Rust SIMD работать на GPU

инструментыРедакция ИИтогHacker News🔥69Рейтинг ИИтог69Влияние65Новизна80Релевантность70Неожиданность55Участники30Достоверность80

VectorWare, компания, описывающая себя как первая "GPU-native software company", объявила, что добилась работы переносимого SIMD Rust (core::simd) на GPU. Раньше SIMD-код на Rust писали либо через архитектурно-специфичные интринсики (например, _mm256_add_ps на x86-64 или vaddq_f32 на Arm), либо через переносимый тип Simd<T, N>, который компилятор превращал в векторные инструкции конкретного процессора. VectorWare распространила эту же модель на GPU: та же программа с core::simd, которая на ноутбуке компилируется в SIMD-инструкции x86-64, теперь без изменений компилируется в инструкции GPU-варпа.

Идея строится на том, что варп GPU (в терминологии NVIDIA, SIMT, Single Instruction Multiple Thread), это, по сути, широкий векторный блок: одна инструкция варпа выполняется одновременно всеми его лейнами (потоковыми дорожками), каждый на своих данных. Ранее VectorWare сопоставляла каждый Rust-поток (std::thread) с отдельным варпом GPU; теперь компания идёт на уровень ниже и сопоставляет каждый SIMD-лейн с лейном самого варпа. Например, Simd<i16, 32> отдаёт по одному элементу i16 каждому из 32 лейнов варпа, и сложение двух таких векторов превращается в одну варп-инструкцию, где все лейны складывают свои элементы одновременно.

Поэлементные операции (сложение, сравнение) выполняются на GPU напрямую через обычные трейты Rust вроде Add. Редукции (reduce_sum, reduce_max) и перестановки лейнов (simd_swizzle!, повороты) используют аппаратные shuffle-инструкции варпа для обмена значениями между лейнами. Маски (Mask<T, N>, Mask::select) и горизонтальные проверки (any, all) опираются на инструкции vote и ballot GPU. Скалярные значения в коде (например, счётчик цикла) реплицируются по всем лейнам одинаково, так же, как в обычном CUDA-коде.

Есть ограничение: на CPU Simd<T, N> допускает любое N от 1 до 64, но у GPU фиксированная ширина, 32 лейна на NVIDIA и 32 или 64 на AMD. Соответствие один к одному работает, только когда N совпадает с шириной варпа; при меньшем N часть лейнов простаивает, при большем, одна операция превращается в несколько инструкций. Внутри VectorWare реализовали это как типизированный IR (набор операций: ballots, shuffles, редукции, сканы, gather/scatter, атомики, а также разбиение векторов шире варпа), который не требует интерпретатора на GPU и, по утверждению авторов, ложится на PTX-инструкции без накладных расходов по сравнению с рукописным кодом; тот же IR можно исполнять и на CPU через собственный референсный интерпретатор для дифференциального тестирования.

Среди плюсов, которые называет VectorWare: один и тот же исходный код и библиотеки, уже использующие переносимый SIMD, становятся кандидатами на выполнение на GPU без переписывания; обычный CPU-код получает доступ к параллелизму на уровне лейнов GPU; Simd<T, N> остаётся обычным Rust-значением, к которому применяются заимствования, время жизни и проверка типов как на CPU, никакого отдельного GPU-специфичного типа вектора не вводится.

Среди ограничений: переносимый SIMD в Rust по-прежнему нестабилен и требует ночной фичи #![feature(portable_simd)], чей интерфейс может измениться до стабилизации. Абстракция бесплатна (zero-cost) только когда ширина вектора точно совпадает с шириной варпа. Не всякая межлейновая перестановка ложится на эффективную варп-инструкцию, произвольные перестановки могут потребовать нескольких инструкций или обращения к разделяемой памяти, а горизонтальные операции вроде редукций и any/all выступают точками синхронизации внутри варпа, что ограничивает планировщик. Авторы также признают, что для корректности абстракции при взаимодействии с другими возможностями Rust пришлось менять сам компилятор, и не уверены, что учли все случаи, территория для них новая. Сейчас работа нацелена на NVIDIA, но, по словам VectorWare, ничего в реализации не завязано конкретно на CUDA: AMD wavefronts и Vulkan subgroups предоставляют похожие примитивы, а сам IR архитектурно-независим.

Дальше компания планирует объединить SIMD, потоки и async в одной модели параллелизма на GPU (потоки распределяют работу между варпами, core::simd, данные между лейнами внутри варпа, async структурирует конкурентность между ними), а также рассматривает перенос матричного SIMD на тензорные ядра GPU и автовекторизацию обычных скалярных Rust-циклов в операции Simd без явного использования core::simd. Авторы называют себя членами команды разработчиков компилятора Rust и говорят, что будущие продукты компании поддержат и другие языки, но текущий фокус, именно Rust.

Ключевые факты

  • VectorWare добилась того, что GPU-код может использовать переносимый SIMD Rust (core::simd), тот же исходный код без изменений компилируется и в SIMD-инструкции CPU, и в инструкции GPU-варпа.
  • Каждый SIMD-лейн сопоставляется с лейном варпа GPU: 32 лейна на NVIDIA, 32 или 64 на AMD, тогда как на CPU Simd<T, N> допускает N от 1 до 64.
  • Абстракция бесплатна по производительности только когда ширина вектора точно совпадает с шириной варпа; более узкий вектор оставляет лейны простаивать, более широкий, превращает операцию в несколько инструкций.
  • Технология опирается на нестабильную ночную фичу Rust #![feature(portable_simd)], интерфейс которой может измениться до стабилизации.
  • Сейчас реализация нацелена на NVIDIA, но, по словам авторов, ничего не завязано на CUDA, AMD wavefronts и Vulkan subgroups предлагают похожие примитивы.

Почему это важно

До сих пор SIMD-код на Rust писали либо через архитектурно-специфичные интринсики под конкретный процессор, либо через переносимый тип Simd<T, N>, который компилятор превращал в векторные инструкции CPU. GPU оставался отдельным миром со своей моделью программирования (CUDA и её аналоги). VectorWare показала, что варп GPU, это по сути ещё один широкий векторный блок, и один и тот же код с core::simd можно без изменений компилировать и под CPU, и под GPU. Это стирает границу между CPU-параллелизмом и GPU-параллелизмом на уровне языка, а не только на уровне потоков, раньше компания уже научила Rust-потоки исполняться как GPU-варпы, теперь то же самое сделано на уровень ниже, для отдельных лейнов внутри варпа.

Кому это важно

В первую очередь, Rust-разработчикам, которые уже пишут высокопроизводительный код с переносимым SIMD (core::simd) и хотят задействовать GPU без переписывания под CUDA или аналоги. Также это релевантно для более широкой экосистемы GPU-программирования: авторы утверждают, что их подход не привязан к NVIDIA и должен переноситься на AMD (wavefronts) и Vulkan (subgroups).

Как это применить

Нужен nightly-компилятор Rust с включённой фичей #![feature(portable_simd)], стабильной версии этой возможности пока нет. Существующий код, использующий Simd<T, N>, Mask<T, N>, операции вроде reduce_sum, reduce_max, simd_swizzle! и Mask::select, работает без изменений; на GPU эти операции ложатся на аппаратные shuffle-, vote- и ballot-инструкции варпа. Максимальную эффективность (без простаивающих лейнов и лишних инструкций) даёт вектор, ширина которого точно совпадает с шириной варпа, 32 элемента на NVIDIA, 32 или 64 на AMD.

Можно ли доверять

Источник, блог самой VectorWare, независимых замеров производительности в тексте нет: утверждение о нулевых накладных расходах по сравнению с рукописным PTX не подкреплено измеренными цифрами. Авторы называют себя членами команды разработчиков компилятора Rust, но никто из них не назван по имени, везде фигурирует коллективное "мы"/VectorWare. Дата публикации в тексте не указана, конкретные модели GPU и версии драйверов/тулчейна, на которых демонстрировался результат, тоже не названы.

Риски и подводные камни

Переносимый SIMD в Rust остаётся нестабильной ночной фичей, чей интерфейс может измениться до стабилизации. Выигрыш в производительности "бесплатен" только при точном совпадении ширины вектора с шириной варпа, в остальных случаях часть лейнов простаивает или операция превращается в несколько инструкций. Не всякая межлейновая перестановка ложится на эффективную аппаратную инструкцию: произвольные перестановки могут потребовать нескольких шагов или обращения к разделяемой памяти, а горизонтальные операции (редукции, any/all) синхронизируют весь варп, ограничивая планировщик. Сами авторы признают, что ради корректности абстракции пришлось менять компилятор Rust и что не уверены, что учли все пограничные случаи взаимодействия с другими возможностями языка. Наконец, реализация сегодня проверена только на NVIDIA, переносимость на AMD и Vulkan заявлена архитектурно, но не продемонстрирована в тексте.