# decuda Architecture
## Overview
decuda is a CUDA-to-target source translator. It does not compile or run
CUDA code; it performs mechanical syntax rewrites and flags everything
that needs human judgement.
## Pipeline
```
input (.cu/.cuh)
│
▼
┌─────────────┐
│ preprocess │ kernel<<<g,b>>>(args) → __decuda_launch(g,b,args)
│ │ skips comments/strings; captures launch sites
└──────┬──────┘
│
▼
┌─────────────┐
│ parser │ tree-sitter-cpp validation (AST unused)
│ │ regex sweep over ORIGINAL source → IR nodes
└──────┬──────┘
│
▼
┌─────────────┐
│ ir.rs │ TranslationUnit { source, path, nodes: Vec<IrNode> }
│ │ nodes carry byte spans into `source`
└──────┬──────┘
│
▼
┌─────────────┐
│ emit.rs │ byte-precise span substitution
│ │ cumulative byte-shift tracking
│ │ overlapping edits: earlier wins
└──────┬──────┘
│
▼
┌─────────────┐
│ targets/ │ per-backend banner + emit dispatch
│ hip.rs │
│ sycl.rs │
│ rust.rs │
│ opencl.rs │
└──────┬──────┘
│
▼
┌─────────────┐
│ migrate │ walk input → parse → emit → write → report
│ walker │ collect .cu/.cuh, map output paths
│ report │ JSON + human-readable summary
└─────────────┘
```
## Module Responsibilities
| `main.rs` | CLI binary entry point; tracing init |
| `cli.rs` | clap CLI surface; `Target` enum; subcommand dispatch |
| `preprocess.rs` | Launch-syntax rewrite; comment/string skipping; balanced-paren arg splitting |
| `parser.rs` | tree-sitter-cpp validation; regex IR harvest over original source |
| `ir.rs` | `TranslationUnit`, `IrNode` enum, `CudaQualifier`, `BuiltinKind`, span helpers |
| `cuda_db.rs` | `once_cell::Lazy` map of CUDA API → per-target mappings |
| `targets/mod.rs` | `TargetBackend` trait; `for_target()` dispatch |
| `targets/{hip,sycl,rust,opencl}.rs` | Per-backend banner headers |
| `emit.rs` | Span-substitution rewriter; byte-shift tracking; overlap detection |
| `migrate.rs` | Orchestrator: walk → parse → emit → write → report |
| `walker.rs` | `walkdir`-based `.cu`/`.cuh` collection; output path mapping |
| `report.rs` | `MigrationReport` JSON + human-readable stdout summary |
## Key Design Decisions
### 1. Preprocess for parseability, harvest from original
The kernel launch syntax `<<<...>>>` is not valid C++. We preprocess it
into `__decuda_launch(...)` so tree-sitter-cpp can parse the source.
However, all IR harvesting (qualifiers, builtins, runtime calls, etc.)
happens against the **original** source, not the preprocessed text. This
keeps byte spans correct.
### 2. AST unused, regex used
tree-sitter-cpp does not understand CUDA-specific identifiers. We use
the parser only for validation (does it parse?). The IR is harvested by
targeted regex sweeps over the linear source.
### 3. Byte-precise span rewriting
Each IR node carries a `(start, end)` byte span into the original
source. The rewriter applies substitutions in source order, tracking a
cumulative byte shift so each edit lands at its original semantic
position. Overlapping edits are detected; the earlier-starting edit
wins.
### 4. Mechanical, not semantic
decuda rewrites syntax and surfaces the rest as warnings. It does not
attempt semantic translation (e.g., `__shared__` → SYCL accessor is
flagged, not auto-generated). This keeps the tool honest about what it
can and cannot do.
### 5. CUDA API database is a one-liner per entry
Adding a new API mapping is a single `insert()` call in `cuda_db.rs`.
The four optional target strings (HIP, SYCL, Rust, OpenCL) control
per-backend rewriting. `None` means "no mapping; preserve verbatim and
warn."
## Data Flow
```
MigrateOptions
│
├── walker::collect_cuda_files → Vec<PathBuf>
│
└── for each file:
parser::translate_path → TranslationUnit
├── preprocess::preprocess → PreprocessedSource
│ └── launches: Vec<LaunchSite>
└── regex harvest → Vec<IrNode>
for each target:
targets::for_target(t) → Box<dyn TargetBackend>
backend.emit(&unit) → String
└── emit::emit(t, backend, unit)
└── apply_replacements(t, unit)
└── replacement_for(t, node) per node
walker::map_output_path → PathBuf
fs::write(output, text)
report.record_warning(...)
```