ROUGE-V

A compiler that translates CUDA PTX into native code for RISC-V.

English · Русский · 中文

nvcc does not produce machine code for a GPU. It produces PTX — a textual assembly with a published specification. Anything that can read PTX can compile it for a different target.

ROUGE-V does that: PTX → LLVM IR → native code, for RISC-V and its vector extension rv64gcv. It is a compiler, not an emulator — the translated kernel is compiled, not interpreted instruction by instruction.

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

Requires CMake ≥ 3.20, a C++17 compiler, and clang. nvcc is optional — it enables the vendor-PTX test.

No build needed — two binaries with no dependencies beyond system 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

Run it on an NVIDIA GPU (sm_75+, driver only, no toolkit needed at run time):

rouge-run --target nvptx64-nvidia-cuda kernel.ptx --emit-ptx kernel_gpu.ptx
# then fill the 3 EDIT points in software/rouge-compiler/examples/run_on_gpu.c
# (PTX path, kernel name, buffers/launch) and build it against libcuda

(Arch Linux: packaging/PKGBUILD for a manual build. packaging/install.sh does the above.)

The question a user has is "will you translate my kernel?". One command answers it:

./build/rouge-compiler/rouge-run kernel.ptx

input: kernel.ptx
kernels found: 1
  @gemm_tile  params=3  instructions=67
  shared memory: 1024 bytes in 2 variable(s)
    shA          @   0     512 bytes  (.b16 x 256)
    shB          @ 512     512 bytes  (.b16 x 256)
translation: ok (10869 bytes of LLVM IR)

OK

Outside the supported subset you get a refusal, not a wrong answer:

rouge-run: REFUSED: unsupported PTX op 'mma.sync.aligned.m16n8k16...' @ 4
This instruction is outside the supported subset. Nothing was emitted;
nothing was approximated.

Emit a native object, or inspect the IR directly:

./build/rouge-compiler/rouge-run kernel.ptx -o kernel.o --print-ir
./build/rouge-compiler/ptx2ir kernel.ptx kernel.ll

Target RISC-V instead of the host:

clang --target=riscv64-unknown-elf -march=rv64gcv -c kernel.ll -o kernel-rv64.o

The triple itself is chosen on the translator side with --target, which is just a prefix of ptx2ir's arguments — the rest of the command is unchanged. The same PTX therefore builds for four platforms:

./build/rouge-compiler/ptx2ir --target amdgcn-amd-amdhsa kernel.ptx kernel-amd.ll
clang --target=amdgcn-amd-amdhsa -mcpu=gfx1100 -c kernel-amd.ll -o kernel-amd.o

All seven canonical kernels in software/rouge-compiler/tests/kernels/ — vadd, block_reduce, atomic_reduce, fp16_reduce, gemm_tile, shfl_reduce, atom_cas — build for all four triples.

The NVIDIA kernels do not just compile: ctest compiler_nvptx_gpu_exec* loads the generated PTX through the CUDA Driver API (driver-side JIT) on a real GeForce MX450 (sm_75) and checks every output bit-exactly against a host reference — all seven canonical kernels, two datasets for fp16_reduce/gemm_tile, plus a two-block grid proving the multi-block path. compiler_nvcc_gpu_exec additionally runs genuine vendor output (nvcc real_nvcc.cu) end to end. On machines without an sm_75+ NVIDIA GPU the tests report SKIP, never failure.

The AMD row stays a compile-time claim: no AMD device is involved anywhere in the project. Neither row says anything about performance — only bit-identity against the host reference was measured. The target changes the address spaces in the emitted IR, which is what lets those backends emit global_load/ global_store instead of scalar accesses — see docs/06-compiler-architecture.md.

Where the PTX comes from:

nvcc -arch=sm_75 -ptx kernel.cu -o kernel.ptx   # build from CUDA source
cuobjdump -ptx ./my_app > kernel.ptx             # extract from a fatbin

QUICKSTART.md covers the rest, including the minimal host driver needed to launch a translated kernel.

$ ctest --test-dir build
...
 7/21 aot_native_block_reduce ..............   Passed
 8/21 aot_native_atomic_reduce .............   Passed
13/21 compiler_nvcc_ptx ....................   Passed
19/21 compiler_nvptx_gpu_exec ..............   Passed
21/21 compiler_nvcc_gpu_exec ...............   Passed
100% tests passed out of 21

Both execution paths are built on the same parser. That is deliberate: one grammar, one set of rules for registers, parameters, .shared layout and labels, so the paths cannot drift apart. The suite is built around the requirement that they agree bit-for-bit on every kernel.

Tensor instructions (mma), warp f16x2 vector forms, and cvt.sat rounding modes are not implemented yet and are reported as errors.

docs/06-compiler-architecture.md — launch descriptor, shared-memory layout, bar.sync model, atomics, FP16/BF16 lowering, and the planned SIMT → RVV path.

If you maintain a project that uses ROUGE-V — a fork, a binding, a backend — please note in your README that it is not built by the ROUGE-V team and is not affiliated with it.

ROUGE-V contains no NVIDIA source code, headers, or binaries, and performs no reverse engineering of them. The PTX it compiles is produced by the user's own nvcc installation, under the user's NVIDIA license, from a published specification. The CUDA API headers under software/rouge-cuda/include/ are written from scratch against public API documentation. The clean-room rule is enforced in CONTRIBUTING.md.

ROUGE-V is not affiliated with, endorsed by, or supported by NVIDIA. CUDA and other trademarks are the property of their respective owners and are used here only to describe compatibility.

Read CONTRIBUTING.md first — it covers the build, the test workflow, and the clean-room requirement that applies to all contributions. Issues and pull requests are open.

• LLVM / MLIR — the IR, backend, and pass infrastructure this is written against
• RISC-V International — the rv64gcv ISA and Vector Extension specifications
• NVIDIA — the published PTX ISA and CUDA API documentation, which define the input format

Apache 2.0 WITH LLVM-exception — the same license as LLVM, so this code can live inside LLVM and MLIR without friction.