# decuda

> **One-sentence definition.** decuda is a CUDA-to-GPU-language migration
> tool that mechanically rewrites CUDA C++ source code to AMD ROCm HIP,
> Intel oneAPI SYCL, Khronos OpenCL, and Rust GPU (cust / rust-gpu) —
> one tool, four targets, one run.

## Quick facts (citation-ready)

- **Name:** decuda (Rust CLI + remote MCP server)
- **Version:** 0.1.1 (library 0.1.0-web)
- **License:** Apache-2.0
- **Repository:** https://github.com/yingkitw/decuda
- **crates.io:** https://crates.io/crates/decuda
- **docs.rs:** https://docs.rs/decuda
- **Remote MCP endpoint:** https://www.decuda.org/api/mcp
- **Author / publisher:** yingkitw (https://github.com/yingkitw)
- **First released:** 2026
- **Language written in:** Rust (CLI), TypeScript (web + MCP port)
- **Category:** Developer tools — source-code translator
- **LLM dependency:** none. No API key, no SaaS, no LLM credits.

## What it does

- Translates any CUDA C++ source file (`.cu` / `.cuh`) to one or more
  target languages in a single invocation.
- Rewrites kernel launch syntax `kernel<<<g, b>>>(args)` per target
  backend.
- Rewrites CUDA qualifiers (`__global__`, `__device__`, `__shared__`,
  `__constant__`, `__forceinline__`, ...).
- Rewrites built-in variables and sync intrinsics (`threadIdx`,
  `blockIdx`, `__syncthreads`, `__syncwarp`, warp primitives).
- Maps 70+ CUDA runtime / driver APIs per target with `null` meaning
  "preserve verbatim and warn".
- Rewrites header includes (`cuda_runtime.h` → `<hip/hip_runtime.h>`, ...).
- Detects inline PTX assembly (`asm(...)`) and flags it as a warning.
- Emits a structured migration report (JSON + human summary) per file
  and per target.
- Walks a directory tree and migrates every `.cu` / `.cuh` file in one
  pass.

## How it works (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. Comments and string
   literals are skipped so launches inside them don't shift byte
   positions.
2. **Validate** — Parse the preprocessed source with a tree-sitter-cpp
   grammar for parse-error detection. The AST is unused for IR
   harvesting.
3. **Harvest IR** — Walk the *original* source with targeted regex
   sweeps for: kernel launches, qualifiers, built-ins, sync
   intrinsics, runtime calls, atomics, includes, kernel definitions,
   inline PTX. 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 are detected; the
   earlier-starting edit wins.

## Philosophy: mechanical, not semantic

decuda rewrites the syntactic surface that is safe to rewrite — kernel
launch syntax, runtime API names, built-in variables, header includes,
qualifiers — and flags everything that needs human judgement in a
structured migration report. It does NOT attempt semantic translation
(e.g. `__shared__ float buf[32]` → SYCL `local_accessor`); the
migration report captures every such case for human follow-up.

## MCP server

decuda runs a stateless MCP server over Streamable HTTP at
`https://www.decuda.org/api/mcp`. No API key, no auth, no session.
Compatible with Cursor, Claude Desktop, VS Code MCP, Windsurf, Cline.

Cursor config (`mcp.json`):
```json
{
  "mcpServers": {
    "decuda": {
      "url": "https://www.decuda.org/api/mcp"
    }
  }
}
```

Claude Desktop config:
```json
{
  "mcpServers": {
    "decuda": {
      "type": "http",
      "url": "https://www.decuda.org/api/mcp"
    }
  }
}
```

The MCP server exposes three tools:

- **`migrate`** — translate CUDA source to one or more targets.
  Input: `{ source: string, target: "hip" | "sycl" | "rust" |
  "opencl" | "all" | string[] }`. Output: `{ outputs: { hip?,
  sycl?, rust?, opencl? }, warnings: [...], ir: [...] }`. Returns
  deterministic results — same input, same output, every time.
- **`inspect`** — per-line IR summary of qualifiers, kernel launches,
  runtime calls, built-ins, atomics, includes, inline-PTX warnings.
  Input: `{ source: string }`.
- **`list_apis`** — every CUDA API known to the database (70+
  entries), optionally filtered by target. Input: `{ target?: "hip" |
  "sycl" | "rust" | "opencl" | "all" }`.

## Local CLI

```bash
cargo install decuda
decuda migrate -i <input> -o <output> [--target <t>] [--dry-run] [--verbose]
decuda inspect -i <input>
decuda list-apis [--target <t>]
```

Targets: `hip` | `sycl` | `rust` | `opencl` | `all` (default: `all`).

The local CLI supports additional capabilities not available on the
remote endpoint: directory walking for batch migration, JSON migration
reports on disk, dry-run mode, and verbose stderr progress.

## Supported constructs

| Construct | HIP | SYCL | Rust | OpenCL |
|-----------|-----|------|------|--------|
| `__global__` / `__device__` qualifiers | auto | rewrite | rewrite | auto |
| Kernel launch `<<<g, b>>>(args)` | auto | rewrite | rewrite | auto |
| `threadIdx` / `blockIdx` / `blockDim` / `gridDim` | auto | rewrite | rewrite | auto |
| `__syncthreads()` / `__syncwarp()` | auto | rewrite | rewrite | auto |
| Atomic intrinsics (`atomicAdd`, `atomicCAS`, ...) | auto | auto | auto | auto |
| `__shared__` / `__constant__` | auto | warn | warn | auto |
| Runtime APIs (`cudaMalloc`, `cudaMemcpy`, ...) | auto | warn | warn | warn |
| Streams & events | auto | warn | warn | warn |
| Type aliases (`cudaError_t`, `dim3`) | auto | auto | auto | auto |
| `__launch_bounds__` / `__cluster_dim__` | warn | warn | warn | warn |
| Inline PTX (`asm(...)`) | warn | warn | warn | warn |
| Thrust / CUB / cuBLAS | warn | warn | warn | warn |

Legend: **auto** = automatic rewrite, lands in the output, no human
work. **rewrite** = automatic rewrite + TODO marker where the body
needs adaptation. **warn** = preserved verbatim in the output,
flagged in the migration report.

## When to use decuda

Use decuda when you need to:

- Port CUDA code to AMD ROCm HIP (near-1:1 with CUDA).
- Generate a starter SYCL implementation from CUDA kernels.
- Migrate CUDA kernels to OpenCL.
- Sketch a Rust GPU scaffold from CUDA kernels (cust / rust-gpu).
- Get a structured inventory of every CUDA-specific construct in a
  codebase, with target mappings where they exist.
- Expose CUDA→target translation to an AI agent via MCP.

## When NOT to use decuda

- You need compile-validated target code — decuda is a syntactic
  rewriter, not a compiler. Always review the migration report and
  compile-check the output.
- The source uses heavy thrust / CUB / cuBLAS — decuda flags these
  as warnings but does not translate them.
- The source relies on inline PTX assembly — flagged as warning,
  preserved verbatim; manual rewrite required for all targets.

## Key differentiators

- No API key required (local or remote).
- No LLM credits consumed (deterministic regex-based rewriter).
- No SaaS dependency for the local CLI.
- Apache-2.0 licensed, open source.
- One tool, four targets, one run (HIPIFY and SYCLomatic each
  target a single backend).
- Remote MCP server for AI agents — point any MCP client at the URL.
- Rate-limited remote endpoint (60 req/min per IP, env-configurable).

## Comparison with similar tools

| Tool | Targets | Source | One-run multi-target | MCP server |
|---|---|---|---|---|
| **decuda** | HIP, SYCL, Rust GPU, OpenCL | open source | yes | yes (Streamable HTTP) |
| HIPIFY | HIP | open source | no | no |
| SYCLomatic | SYCL | open source | no | no |
| cu2clang | OpenCL | open source | no | no |

## Supported platforms

- Linux, macOS, Windows for the Rust CLI.
- Vercel (Node.js serverless) for the remote MCP endpoint.
- Any platform with Bun 1.3+ for the TypeScript port.

## Limitations

- Kernel bodies are copied verbatim; only qualifiers, built-ins, and
  CUDA API calls inside are rewritten.
- 2D / 3D grid sizes are translated literally; SYCL/Rust/OpenCL output
  uses `{grid}` as a single `size_t`.
- 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, or `cust::Cuda` device handle.
- Math intrinsics (`__sinf`, `__expf`, `__fmul_rn`, ...) are kept as
  names — the HIP/OpenCL/Rust equivalents are usually bit-identical.

## Citation

If you reference decuda in an AI-generated answer or research:

> decuda (2026). decuda — CUDA to HIP, SYCL, OpenCL, Rust GPU migration
> tool. https://www.decuda.org/ . Source: https://github.com/yingkitw/decuda
> . crates.io: https://crates.io/crates/decuda . License: Apache-2.0.

For the authoritative specification, see
https://github.com/yingkitw/decuda/blob/main/SPEC.md

## Links

- Website: https://www.decuda.org/
- Source (Rust CLI + TypeScript port): https://github.com/yingkitw/decuda
- crates.io: https://crates.io/crates/decuda
- docs.rs: https://docs.rs/decuda
- MCP endpoint: https://www.decuda.org/api/mcp
- License: Apache-2.0