decuda — CUDA to HIP / SYCL / OpenCL / Rust GPU Migration Tool

Migrate CUDA C++ source code to HIP, SYCL, OpenCL, or Rust GPU automatically. decuda is a command-line tool that performs the mechanical rewrites — kernel launch syntax, runtime API names, built-in variables, header includes, qualifiers — and flags everything that needs human judgement in a structured migration report.

CUDA is the dominant GPU programming model, but vendor lock-in to NVIDIA hardware limits portability. decuda helps you port CUDA kernels to:

• AMD ROCm HIP — near source-compatible, minimal changes needed
• Intel oneAPI SYCL — C++ cross-platform GPU programming
• Khronos OpenCL — open standard for heterogeneous computing
• Rust GPU (cust / rust-gpu) — Rust-native GPU compute

It is a starting point for migration, not a finished translator. The philosophy is mechanical, not semantic: do the syntactic rewrites that are safe, and surface everything else as warnings so a human engineer can make the right call.

Table of Contents

• Supported CUDA Constructs
• Quick Start
• Usage
• How It Works
• Examples
• CUDA API Database
• Limitations
• Comparison with Other Tools
• License

Supported CUDA Constructs

auto = automatic rewrite, warn = flagged in the migration report, preserved as-is in output.

Quick Start

cargo install decuda

# Migrate a single CUDA file to all four backends.
decuda migrate -i src/kernels.cu -o out/

# Migrate to HIP only.
decuda migrate -i src/kernels.cu -o out/ --target hip

# Migrate an entire directory of .cu / .cuh files.
decuda migrate -i src/ -o out/ --target all --verbose

Usage

migrate — translate CUDA to target GPU language

decuda migrate -i <input> -o <output> [options]

inspect — show the IR extracted from a CUDA file

decuda inspect -i src/kernels.cu

list-apis — list all CUDA APIs in the database

decuda list-apis --target hip

Output layout

Each target writes its output under <output>/<target>/, preserving the directory layout of the input:

A migration-report.json.<timestamp> is written next to the output directories, capturing per-file warnings (unsupported APIs, mismatched qualifiers, …). A human-readable summary is printed to stdout.

How It Works

decuda processes CUDA source in four stages:

1. Pre-process — Replace CUDA's kernel<<<g,b>>>(args) launch syntax with an ordinary function call (__decuda_launch(...)) so the source is parseable by a stock C++ grammar (tree-sitter-cpp). Comments and string literals are skipped so a launch inside a comment never shifts byte positions.

2. Validate — Parse the preprocessed source with tree-sitter-cpp. Parse errors are surfaced but the AST is unused (tree-sitter doesn't understand CUDA-specific identifiers).

3. Harvest IR — Walk the original source with targeted regex sweeps for: kernel launches, function/storage qualifiers (__global__, __shared__, ...), built-in variables (threadIdx, ...), synchronization intrinsics (__syncthreads), CUDA runtime API calls (cudaMalloc, cudaStreamCreate, ...), atomic intrinsics, and header includes. Each match becomes an IR node with a byte span.

4. Emit — Compute per-node replacements for the target backend and apply them as byte-precise span substitutions in source order, tracking the cumulative byte shift so each edit lands at its original semantic position. Overlapping edits (e.g. a cuda_runtime runtime call landing inside an #include <cuda_runtime.h> line) are detected and the earlier one wins.

Examples

The examples/cu/ directory contains 14 CUDA input files covering basic through advanced constructs:

Regenerate all outputs:

cargo run -- migrate -i examples/cu -o examples/out --target all

See examples/README.md for the full layout.

CUDA API Database

The CUDA-API database is in src/cuda_db.rs. Adding a new API entry is a one-liner:

insert(&mut m, "cudaFoo", api("cuda_runtime.h",
    Some("hipFoo"),     // HIP
    Some("syclFoo"),    // SYCL
    Some("rustFoo"),    // Rust GPU
    Some("clFoo"),      // OpenCL
    "note string"));

The four Option<&str> are per-target mappings. None means "no mapping; preserve verbatim and warn." Run decuda list-apis to see all entries.

Limitations

• Kernel bodies are copied verbatim — qualifiers, built-in variables, and cudaXxx calls inside are rewritten, but algorithmic structure stays as-is. SYCL, Rust, and OpenCL kernels typically need a manual rewrite of the body (e.g. converting __shared__ float buf[32] to a SYCL local_accessor).
• 2D/3D grid sizes — dim3 grid/block sizes are translated literally; SYCL/Rust/OpenCL output uses {grid} as a single size_t. The user must translate dim3 to sycl::range<3> / (grid_x, grid_y, grid_z) etc.
• Host code is not rewritten end-to-end — cuda_runtime.h is mapped to the closest target header, but you'll still need to set up the SYCL queue, OpenCL context+queue, cust::Cuda device handle, etc.
• Math intrinsics (__sinf, __expf, __fmul_rn, ...) are kept as names — the HIP/OpenCL/Rust equivalents are usually bit-identical.

Comparison with Other Tools

decuda's unique advantage: one tool, four targets, one run. HIPIFY only targets HIP, SYCLomatic only targets SYCL, and cu2clang only targets OpenCL. decuda emits all four in a single invocation with a unified migration report.

License

Apache-2.0.