Опубликован выпуск CUDA-Oxide 0.2.0 — экспериментального компилятора NVIDIA Labs, позволяющего писать код под CUDA-ядра на чистом Rust и компилировать их напрямую в PTX. Проект использует собственный backend для rustc, поддерживает модель SIMT и собирается через команду cargo oxide, при этом host-код и device-код могут находиться в одном дереве исходников. Релиз состоялся 5 июня 2026 года и назван первым «community release»: после открытия ветки 0.1.0 в проект приняли 37 pull request от 23 участников.
Главное изменение CUDA-Oxide 0.2.0 — переход к самодостаточному исполняемому файлу. Сгенерированные GPU-артефакты — PTX, NVVM-IR, LTOIR и cubin — теперь могут встраиваться прямо в host-бинарник через новый формат oxide-artifacts. Благодаря этому Rust-программа с CUDA-ядрами больше не обязана таскать за собой отдельные .ptx-файлы, а загрузка ядер происходит из самого запущенного executable.
Основные изменения:
Встраивание GPU-ядер в host-бинарник. Новый механизм #[cuda_module] генерирует типизированный интерфейс запуска ядер и прячет низкоуровневую загрузку модулей. Для разработчика это выглядит ближе к обычному Rust API: срезы и буферы сопоставляются с CUDA-буферами, а запуск получает более строгую типизацию.
Поддержка CUDA constant memory. Добавлены #[constant] и Constant<T> с генерацией host-setter’ов. На уровне PTX это соответствует .const и адресному пространству 4. Это важно для данных, которые должны быть доступны ядрам как константные значения на устройстве.
Больше device-математики. Через libdevice добавлены f32::max, f32::min, а также atan и atan2 для f32 и f64. Это расширяет набор реальных вычислительных ядер, которые можно выразить без ухода в C++/CUDA.
Исправления silent miscompile. Несколько скрытых ошибок компиляции (silent miscompiles) теперь приводят к сбою на этапе сборки вместо генерации некорректного PTX (пропуск wgmma accumulate, исключённые деструкторы, поглощённые операторы).
Улучшение host runtime. Для передачи данных на устройство добавлен контракт безопасности DeviceCopy, появились pinned host buffers с stream-ordered transfers, проброс асинхронных ошибок и низкоуровневый выход к raw CUmodule для интеграции с CUDA Driver API. Также исправлено несовпадение аллокаторов DeviceBuffer, добавлена совместимость с CUDA 12.8/13 и aarch64.
Интеграция с CUDA Tile. В релиз добавлен пример взаимодействия cuTile ↔ SIMT: Tile-ядро и SIMT PTX-ядро CUDA-Oxide могут работать на одном CUDA stream с общими device-тензорами. Это показывает место CUDA-Oxide в более широкой Rust/CUDA-стратегии NVIDIA: проект не заменяет Tile DSL, а закрывает сторону явного SIMT-программирования на Rust.
Сближение с upstream Pliron/LLVM. Внутренний dialect-llvm заменён на внешний pliron-llvm из сообщества pliron-org, а локальная часть переименована в llvm-export. Тем самым проект сокращает собственный форк инфраструктуры и ближе следует за upstream-разработкой.
Упрощение сборки и входа в проект. cargo oxide теперь учитывает RUSTFLAGS, CUDA_HOME, host-компилятор nvcc, умеет находить llc из Rust toolchain, автоматически определяет целевой GPU и сбрасывает устаревший cache backend’а после обновления. Добавлены devcontainer и Nix flake для воспроизводимой среды разработки.
CUDA-Oxide пока остаётся ранней alpha-версией: проект распространяется через git, ещё не опубликован на crates.io и поддерживает только Linux. В README также прямо предупреждается, что API может ломаться, а баги и недоработанные возможности ожидаемы.
С лицензированием есть нюанс: основная часть workspace распространяется под Apache License 2.0, но crate cuda-bindings использует NVIDIA Software License. Поэтому CUDA-Oxide нельзя описывать как полностью однородный Apache-проект; это открытая разработка NVIDIA Labs вокруг CUDA-экосистемы с отдельной лицензией для низкоуровневых биндингов к CUDA.
Источник: linux.org.ru
