TL;DR: Проблема переносимости CUDA для меня оказалась не столько в самих арифметических операциях, сколько в огромном количестве инфраструктуры вокруг них. Поэтому вместо попытки повторять CUDA runtime я решил перехватить программу раньше — на уровне PTX. Я написал ROUGE‑V — открытый AOT‑компилятор, разработанный по принципу clean‑room: берёт PTX, сгенерированный официальным nvcc 13.4, транслирует его в LLVM IR, дальше собирает обычный clang. Исполняется на x86-64 и на живой GeForce MX450 (побитово с хост‑эталоном); AMD и RISC‑V — пока только сборка, без вранья про «любое железо».

Мой GitHub: github.com/MrModelOS/ROUGE‑V (автор: MrModelOS).

Покрытие: ctest — 21/21; 8 ядер (7 канонических + 1 настоящее от nvcc) исполнены на GeForce MX450.

1. Зачем я сделал свой транслятор, когда есть другие проекты?

Один из очевидных способов бороться с vendor lock‑in NVIDIA — реимплементировать runtime или перехватывать API на лету. Это вечная гонка за чужим интерфейсом: вендор поменял поведение — ты догоняешь.

Я пошёл по пути Clean‑Room разработки.

Компилятор nvcc генерирует не машинный код, а PTX — текстовый промежуточный ассемблер с публично опубликованной спецификацией (выхлоп нашего nvcc 13.4 — PTX ISA 9.4). Я выбрал его как точку перехвата:

  1. Беру стандартный .cu‑код и компилирую его через nvcc в .ptx.

  2. Мой транслятор ptx2ir парсит PTX и строит эквивалентное представление в LLVM IR.

  3. Стандартный clang / LLVM собирает конечный бинарник под целевую архитектуру.

Никаких проприетарных бинарников и заголовков NVIDIA я не копирую и не реверсирую — правило clean‑room закреплено в CONTRIBUTING.md.

2. Архитектура транслятора и 4 шага до запуска

Пайплайн у меня выглядит так:

CUDA / .cu
    │
    ▼
nvcc --ptx
    │
    ▼
PTX
    │
    ▼
ptx2ir
    │
    ▼
LLVM IR
    │
    ▼
LLVM / clang
    │
    ▼
native code

NVPTX → GeForce MX450 → bit-exact check

Честная легенда: AMD (amdgcn) и RISC‑V (rv64gcv) — кросс‑компиляция, исполнения на железе у меня не было. Исполняется — хост и NVIDIA.

На практике весь процесс занимает 4 команды (проверено на живом стенде):

# 1. Скомпилировал .cu в PTX вендорным nvcc 13.4
nvcc -arch=sm_75 -ptx sq.cu -o sq.ptx
# 2. Транслировал PTX в LLVM IR своим ptx2ir
ptx2ir --target nvptx64-nvidia-cuda sq.ptx sq.ll
# 3. Собрал настоящим бэкендом clang в PTX ядра
clang --target=nvptx64-nvidia-cuda -march=sm_75 -S sq.ll -o sq_gpu.ptx
# 4. Запустил на видеокарте через CUDA Driver API
./gosq
# Output: OK: 1024 квадрата посчитаны на GeForce MX450

Что внутри sq_gpu.ptx — настоящее ядро, а не заглушка:

.visible .entry _Z2sqPKfPfj(
    mul.lo.s32  %r7, %r6, %r5;
    ld.global.b32   %r12, [%rd9];
    mul.rn.f32  %r14, %r13, %r13;
    st.global.b32   [%rd14], %r15;

3. Грабли: как я ловил тихую порчу данных

ручной PTX → всё зелёное → настоящий nvcc → парсер ломается    → чиню грамматику → реальная GPU → вылазят address-space баги    → исправляю pipeline

Написать парсер для add/mul — просто. Сложности начались на выхлопе свежего nvcc.

Камень 1: однострочные сигнатуры. Современный nvcc пишет .visible .entry _Z...(…) в одну строку, а мой парсер ждал скобку на отдельной строке — и тело ядра оставалось пустым. Плюс 37 мелких пробелов в грамматике: and/or/xor, сдвиги, selp, div/rem, mad.lo.s32, bfi/bfe, f16x2, варп‑shfl. Базовая арифметика, которую компилятор генерирует в каждом ядре.

Камень 2: инверсия предикатов. Запись @!%p1 bra LABEL означает «перейти, если предикат ложен» — флаг ! обязан явно инвертироваться в IR. У меня это закрыто на уровне чтения предикатных операндов: ведущий ! превращается в xor i1 …, true, и то же правило действует в интерпретаторе. Тихой порчи здесь нет именно потому, что инверсия — часть контракта чтения операнда, а не «надежда, что бэкенд угадает».

Камень 3: setp и div. setp.lt.f32 у меня сравнивал биты целочисленным icmp вместо fcmp, а div.s32 эмитил mul. Оба не падали — молча считали неправильно. Нашлись только сверкой с настоящим входом от nvcc.

Мораль, которую я вынес: тест на собственных данных доказывает, что код работает на твоих данных. Нужен чужой вход — настоящий PTX из CI (compiler_nvcc_ptx).

А потом видеокарта нашла ещё два бага, невидимых на хосте: вывод матрицы писался в shared вместо global (метка регистра переживала перезапись — теперь пространство берётся из суффикса инструкции), а адрес device‑глобала читался как загрузка из него. На плоской хост‑памяти оба молчали.

4. Мои результаты и тесты

ctest — 21/21, ноль ворнингов. Из них:

  • 7 канонических ядер (vadd, block_reduce, atomic_reduce, fp16_reduce, gemm_tile, shfl_reduce, atom_cas) исполнены на GeForce MX450 и сошлись побитово с хост‑эталоном;

  • настоящее ядро от nvcc (манглированное имя, device-глобал, syncthreads) — насквозь nvcc → ptx2ir → GPU. Честная оговорка: в его исходнике гонка (редукция без syncthreads), вендорная сборка гонится так же — эталон принимает доказанный интервал, а не одно число;

  • shfl.sync.bfly проверен отдельным NVPTX‑тестом (shfl_reduce) на реальной GPU — через настоящий NVVM‑интринсик; остальные warp‑коллективы покрытыми тестами не считаю: vote / activemask / предикатный shfl — явный отказ, а не тихий фолбэк;

  • без карты sm_75+ GPU‑тесты говорят честный SKIP, а не «успех».

Производительности я не мерял — только бит‑идентичность. Кто обещает скорость без замеров, тот врёт; я не обещаю.

5. Что дальше?

  • SIMT → RVV векторизация: сворачивать 32 потока в векторные линии (MLIR‑контур уже анализирует паттерны, полноценной векторизации пока нет);

  • mma и тензорные инструкции — пока осознанный отказ: «примерно правильно» опаснее, чем никак;

  • больше настоящих CUDA‑бинарников на входе вместо синтетики;

  • публичные бенчмарки — но сначала повторяемый стенд.

Где пощупать

Без сборки, два бинарника (кроме системного libc зависимостей нет):

curl -fsSL https://github.com/MrModelOS/ROUGE-V/releases/download/v0.1/rouge-v-0.1-linux-x86_64.tar.gz | tar xz -C ~/.local
export PATH="$HOME/.local/bin:$PATH"
ptx2ir kernel.ptx kernel.ll

Из исходников:

git clone https://github.com/MrModelOS/ROUGE-V.git && cd ROUGE-V
cmake -B build -G Ninja -DCMAKE_BUILD_TYPE=Release && cmake --build build
ctest --test-dir build --output-on-failure

Архитектура — docs/06-compiler-architecture.md. Лицензия — Apache 2.0 WITH LLVM‑exception.

Про юридическую сторону — прямо

  • Я не копирую и не реверс‑инжинирю проприетарные бинарники и заголовки NVIDIA.

  • PTX, который я перевожу, выпускает nvcc пользователя — из его собственной лицензии, по открытой спецификации PTX ISA.

  • ROUGE‑V не связан с NVIDIA и не одобрен ими. Торговые марки — только чтобы описать совместимость.

P. S.

Если возьмёте PTX от реального nvcc и мой транслятор скажет отказ — расскажите, что именно: это самый полезный баг‑репорт для проекта.