NVIDIA CUDA-Oxide 0.2.006.06.2026 07:15

good-penguin.png

Опубликован выпуск 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 с генерацией host-setter«ов. На уровне PTX это соответствует .const и адресному пространству 4. Это важно для данных, которые должны быть доступны ядрам как константные значения на устройстве.

  • Больше device-математики. Через libdevice добавлены f32::max, f32::min, а также atan и atan2 для f32 и f64. Это расширяет набор реальных вычислительных ядер, которые можно выразить без ухода в C++/CUDA.

  • Исправления silent miscompile. Несколько случаев, где компилятор мог молча сгенерировать неверный PTX, теперь завершаются ошибкой сборки. Разработчики отдельно упоминают проблемы с wgmma, пропущенными деструкторами и «проглоченными» выражениями. Это важнее новых функций: для компилятора GPU-кода молчаливая неправильная генерация хуже обычного падения сборки.

  • Улучшение 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