# decuda — CUDA to HIP / SYCL / OpenCL / Rust GPU Migration Tool
[](https://crates.io/crates/decuda)
[](https://docs.rs/decuda)
[](https://github.com/yingkitw/decuda)
[](https://github.com/yingkitw/decuda/actions)
> **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](#supported-cuda-constructs)
- [Quick Start](#quick-start)
- [Usage](#usage)
- [How It Works](#how-it-works)
- [Examples](#examples)
- [CUDA API Database](#cuda-api-database)
- [Limitations](#limitations)
- [Comparison with Other Tools](#comparison-with-other-tools)
- [License](#license)
## Supported CUDA Constructs
| `__global__` / `__device__` qualifiers | auto | auto | auto | auto |
| Kernel launch `kernel<<<g,b>>>(...)` | auto | auto | auto | auto |
| `threadIdx` / `blockIdx` / `blockDim` | auto | auto | auto | auto |
| `__syncthreads()` / `__syncwarp(...)` | auto | auto | auto | auto |
| `atomicAdd` / `atomicCAS` / ... | auto | auto | auto | auto |
| `__shared__` / `__constant__` | auto | auto | auto | auto |
| `cudaMalloc` / `cudaMemcpy` / ... | auto | warn | warn | warn |
| `cudaStreamCreate` / events | auto | warn | warn | warn |
| `cudaError_t` / `dim3` aliases | auto | auto | auto | auto |
| `__launch_bounds__` | warn | warn | warn | warn |
| Inline PTX / texture references | warn | warn | warn | warn |
| Thrust / CUB / cuBLAS wrappers | warn | warn | warn | warn |
**auto** = automatic rewrite, **warn** = flagged in the migration report,
preserved as-is in output.
## Quick Start
```bash
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
```bash
decuda migrate -i <input> -o <output> [options]
```
| `-i, --input` | Input `.cu`/`.cuh` file or directory |
| `-o, --output` | Output directory; each target writes into `<output>/<target>/` |
| `-t, --target` | `hip`, `sycl`, `rust`, `opencl`, or `all` (default: `all`) |
| `--dry-run` | Parse and plan only; do not write output files |
| `-v, --verbose` | Emit progress to stderr |
| `--filter <substr>` | Only process files whose path contains `<substr>` |
### `inspect` — show the IR extracted from a CUDA file
```bash
decuda inspect -i src/kernels.cu
```
### `list-apis` — list all CUDA APIs in the database
```bash
decuda list-apis --target hip
```
### Output layout
Each target writes its output under `<output>/<target>/`, preserving the
directory layout of the input:
| `hip` | `.hip.cpp` | `<output>/hip/` |
| `sycl` | `.sycl.cpp` | `<output>/sycl/` |
| `rust` | `.rs` | `<output>/rust/` |
| `opencl` | `.cl` | `<output>/opencl/` |
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:
| `saxpy.cu` | kernels, launches, malloc/free, shared memory |
| `histogram.cu` | atomics, shared bins, grid-stride, warp intrinsics |
| `transpose.cu` | 2D dim3 grid/block, shared-memory tile |
| `stream_pipeline.cu` | streams, events, async memcpy, smem+stream launches |
| `device_helpers.cu` | `__device__` helpers, inline hints, `__constant__`, `__launch_bounds__` |
| `reduction.cu` | warp-shuffle reduction, `__shfl_sync`, tree reduction, `atomicAdd` |
| `stencil_3d.cu` | 3D 7-point stencil, shared-memory halo, 3D dim3 |
| `device_management.cu` | multi-GPU, error handling, pinned memory |
| `managed_memory.cu` | `cudaMallocManaged`, `__managed__`, unified memory |
| `warp_primitives.cu` | `__shfl_sync`, `__ballot_sync`, `__any_sync`, `__all_sync` |
| `rich.cu` | atomics, shared mem, warp intrinsics, 2D launches, constant mem |
| `launch_in_comment.cu` | launches inside comments/strings (preprocessor test) |
| `headers_only.cu` | every include-replacement path |
| `empty.cu` | no CUDA constructs (banner + verbatim source) |
Regenerate all outputs:
```bash
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:
```rust
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
| **HIP target** | yes | yes | — | — |
| **SYCL target** | yes | — | yes | — |
| **OpenCL target** | yes | — | — | yes |
| **Rust GPU target** | yes | — | — | — |
| **Multi-target in one run** | yes | no | no | no |
| **Migration report** | JSON + stdout | YAML | YAML | — |
| **Dry-run mode** | yes | yes | yes | — |
| **Directory walker** | yes | yes | yes | no |
| **Open source** | yes | yes | yes | yes |
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.