oneapi-agent-toolkit

Evidence-driven MCP toolkit that makes coding agents fluent in oneAPI/SYCL.

SYCLomatic migrates most of a CUDA code base to SYCL. What remains is the last mile: code it could not translate, output full of dpct:: helpers, and no proof the port still behaves like the original. Coding agents can do this work, but left alone they drift back to CUDA idioms and report success without checking.

oneapi-agent-toolkit is an open-source MCP server and agent skill that gives Claude Code, Codex, Cursor and other MCP-capable agents (or the bundled standalone loop) five tools: migrate (SYCLomatic), compile (icpx), scan_helpers, verify and benchmark. Inspired by the evidence-first design of the REA project, every tool returns structured evidence and its limitations, never a bare claim. A port counts as done only when behavioural verification passes and no dpct helpers remain, both on the same source.

We ran 120 ports of four workloads (an audio waveshaper, a 64-tap FIR filter, SAXPY and HeCBench softmax) on two Claude models, and graded every final port with a hidden test. With the toolkit, 40 of 40 ports were correct and helper-free. Without its checks, 23 of 40 were, and one unchecked softmax port compiled cleanly yet got 332,679 of 524,288 outputs wrong. Giving the agent a shell to test its own code removed wrong output but not unfinished migration: only 14 of 40 ports were helper-free. SYCLomatic shows the same failure mode: its softmax output compiled with zero warnings while expf had silently become a reduced-precision builtin.

Then we closed the loop from CUDA to GPU with no human edits. On five HeCBench programs the agent had never seen, Claude Sonnet produced 15 of 15 correct, helper-free ports with the toolkit and 6 of 15 without (Claude Haiku: 15 of 15 and 12 of 15). A one-command GPU test, which rents a machine, runs the checks and always deletes it, then ran the agent's unedited ports on NVIDIA GPUs and an AMD Instinct MI350X: all 10 matched both the C++ reference and the original CUDA built for that GPU.

Each result above links to the committed output behind it: traces, final sources and tool results for every run. A tool result carries a status, the exact commands that ran, parsed diagnostics, numbers, and an explicit list of limitations. "It works" means verify returned ok on the current code. "It's portable" means scan_helpers found zero helpers. Nothing is reported that was not measured.

Without the toolkit's checks, every step the agent saw said OK and it declared itself done; the hidden check found 332,679 of 524,288 outputs wrong. Click for the step-through replay (all replays).

Check it yourself

To re-run it: the Quickstart installs the toolchain; scripts/ablation.py repeats the with/without comparison with your own model key, and oneapi-agent remote-test <task.json> --provider runpod repeats a GPU check with your own GPU-cloud key.

The last-mile problem, observed

These results come from running this repo's pipeline (examples/run_pipeline.py) on its own examples. Full output is in examples/RESULTS.md. Toolchain: icpx 2026.1.1, c2s/SYCLomatic 21.0.0 (dpct version string), OpenCL CPU device.

For cuBLAS, c2s 21.0.0 emits the oneMKL product API (oneapi::mkl::blas), not the UXL Foundation's open oneMath; scan_helpers flags that as an onemkl_product_api advisory, and the idiomatic port uses oneapi::math::blas.

For softmax, c2s produced code that compiles cleanly with zero migration warnings, but it translated expf to the reduced-precision sycl::native::exp, giving up to 3,473 ULP of error. Only behavioural verification caught it. Swapping in sycl::exp brings it back to 9.3e-8. Whether the wrong results exceed the tolerance depends on the CPU's instruction set: they fail on our AVX-512 Xeon (22,156 wrong, up to 4.1e-6) but stay just inside 1e-6 (6.9e-7) on AVX2 machines such as GitHub's CI runners. The silent expf → native::exp substitution happens everywhere, and CI asserts it on every push. scan_helpers now flags sycl::native:: math as an advisory. That is the evidence-first loop in practice.

Related MCP server: gwaya

A real model closing the gap

oneapi-agent loop with claude-sonnet-5-5 driving the real tools (c2s 21.0.0, icpx 2026.1.1, OpenCL CPU device). The loop declares success only when compile, verify and a zero-helper scan all hold on the same source tree. Traces and reports: docs/real-runs/.

All four end with 0 dpct helpers and a passing verify. That is one run per example, so treat iteration counts as anecdotes, not a distribution.

Checked on NVIDIA and AMD GPUs

oneapi-agent remote-test rents a GPU (Runpod, Vast.ai, Vultr or DigitalOcean), builds the C++ reference and the original CUDA there (nvcc, or hipify-perl + hipcc on AMD), verifies the SYCL port against both, benchmarks it, and always deletes the machine and confirms it is gone. Spending is capped per run. How it works and what broke on real clouds: docs/gpu-testing.md.

The agent's own ports, unedited (docs/agent-ports/). Max abs error vs the C++ reference / vs the original CUDA on the same GPU:

The hand-written reference ports (docs/gpu-testing.md): softmax passes on an RTX 2000 Ada (Runpod), an A16 (Vultr), an RTX A4000 (Vast.ai) and an AMD MI350X (DigitalOcean), and all six HeCBench programs pass on the MI350X and on NVIDIA GPUs. Making the softmax port use an in-order queue, as CUDA's default stream is, cut its per-launch cost on the MI350X from 114.6 to 24.4 us (the original CUDA, hipified: 8.9 us).

The first manual run, on an NVIDIA RTX 2000 Ada with scripts/gpu-run.sh (docs/gpu-run/rtx2000ada/), checked the four original examples:

Kernel timings come from single small launches and are not a performance claim. Intel GPUs have not been run yet: Intel's cloud has no public API, so that path is a bundle run by hand.

Does the toolkit actually help?

Same tasks, same prompts. Five runs per example (20 per cell), on two models, in three arms:

• toolkit: the normal loop.

• baseline: no scan_helpers, verify or benchmark; the run ends when the agent says it is done.

• baseline + shell: the baseline plus a bash shell, so the agent can build and run its own tests.

A hidden check the agent never sees graded every final port the same way: compile, verify against the reference, count dpct helpers. Details and all 120 runs: docs/ablation/round2/. The first, smaller round is in docs/ablation/.

Cells: ports that are correct and have 0 dpct helpers.

• Wrong output: without any checks, agents shipped wrong output in 6 of 40 runs, including a softmax port that compiled cleanly but got 332,679 of 524,288 outputs wrong. With a shell, none did. Testing catches bugs.

• Unfinished conversion: with a shell, agents still left SYCLomatic's dpct helpers in 26 of 40 runs, though the prompt asked for zero. They stopped once the code worked. The helper scan is what finishes the port.

• Cost: the toolkit arm costs about the same as the baseline: $1.37 vs $1.23 for 20 Sonnet runs.

5 runs per cell, two models from one vendor, CPU device only: indicative, not a benchmark.

Five more programs, the agent's ports run on GPUs

The same comparison on five HeCBench programs the agent had never seen (stencil1d, matrix_rotate, atomic_reduction, scan, adam), 3 runs per cell:

Nearly every miss was dpct helpers left in while the agent reported success; one Haiku baseline port hard-coded a GPU selector and crashed on the CPU grader. The Sonnet toolkit ports then passed unedited on NVIDIA and AMD GPUs (above). Details and all 90 runs: docs/agent-ports/.

Architecture

flowchart LR
    subgraph Agents
        CC[Claude Code / Codex / Cursor<br/>+ skills/oneapi-port/SKILL.md]
        LOOP[oneapi-agent loop<br/>standalone, Anthropic SDK]
    end
    CC <-- MCP stdio --> SRV[oneapi-agent serve<br/>MCP server]
    LOOP --> LIB
    SRV --> LIB[oneapi_agent.tools<br/>ToolResult: status · summary · evidence · limitations]
    LIB --> MIG[migrate] --> C2S[(SYCLomatic c2s)]
    LIB --> CMP[compile] --> ICPX[(icpx -fsycl)]
    LIB --> SCAN[scan_helpers] --> SRC[(migrated sources)]
    LIB --> VER[verify] --> RUN[(reference vs candidate<br/>array stats / audio null test)]
    LIB --> BEN[benchmark] --> DEV[(sycl-ls devices<br/>ONEAPI_DEVICE_SELECTOR)]
    LOOP --> TRACE[(runs/&lt;ts&gt;/trace.jsonl<br/>report.md · best/)]

Statuses are ok | fail | error | toolchain_missing | timeout. A missing binary is never faked: the result is toolchain_missing, with the exact command that could not run.

Quickstart

git clone https://github.com/checocadena/oneapi-agent-toolkit && cd oneapi-agent-toolkit
pip install -e ".[all]"                  # Python 3.11+; not on PyPI yet
oneapi-agent --summary doctor            # what toolchain was found

# Optional, real toolchain (see docs/install.md); this is the route verified here:
scripts/install-toolchain.sh /opt && source scripts/oneapi-env.sh

# Use the tools directly
oneapi-agent migrate examples/saxpy/saxpy.cu -o /tmp/saxpy_mig
oneapi-agent compile /tmp/saxpy_mig/saxpy.dp.cpp -o /tmp/saxpy_sycl
oneapi-agent scan-helpers /tmp/saxpy_mig
g++ -O2 examples/saxpy/saxpy_ref.cpp -o /tmp/saxpy_ref
oneapi-agent verify --ref "/tmp/saxpy_ref {output}" --cand "/tmp/saxpy_sycl {output}" --atol 1e-6
oneapi-agent benchmark "/tmp/saxpy_sycl /tmp/out.bin" --runs 5

# Register the MCP server (and the oneapi-port skill) with your agents, in one command
oneapi-agent setup --dry-run             # show what would change
oneapi-agent setup                       # every detected agent: claude, codex, cursor, gemini
oneapi-agent setup --agent claude        # or pick

# Verify on this machine's GPU (NVIDIA, AMD or Intel), or bundle the test for another machine
oneapi-agent gpu-test examples/saxpy/task.json
oneapi-agent bundle examples/saxpy/task.json   # -> gputest-saxpy.sh; `bash` it on any GPU box, then `oneapi-agent ingest`

# Standalone agent loop
oneapi-agent loop --dry-run --mock-tools                 # no toolchain, no API key
oneapi-agent loop --dry-run --config examples/saxpy/task.json   # scripted model, real tools
ONEAPI_AGENT_API_KEY=... oneapi-agent loop --config examples/fir/task.json   # real model

The loop's model comes from ONEAPI_AGENT_MODEL (default claude-sonnet-5-5).

• The API key comes from ONEAPI_AGENT_API_KEY if set, otherwise ANTHROPIC_API_KEY. Use ONEAPI_AGENT_API_KEY inside Claude Code sessions, which blank ANTHROPIC_API_KEY; with it the loop always talks to https://api.anthropic.com (override with ONEAPI_AGENT_BASE_URL), never an inherited ANTHROPIC_BASE_URL.

• Optional: ONEAPI_AGENT_EFFORT sets the effort level.

• ONEAPI_AGENT_FALLBACKS=0 disables the server-side refusal fallback, which is on by default.

The loop declares success only when compile, verify and a zero-helper scan all hold for the same current source tree. The model cannot pick verify's commands or tolerances.

Example run output

This is real output from oneapi-agent loop --dry-run --config examples/saxpy/task.json. The model is scripted, so no LLM was involved, but every tool result comes from a real tool run. Full files: docs/example-run/report.md and trace.jsonl.

| it | tool         | status | summary |
|----|--------------|--------|---------|
| 1  | migrate      | ok     | migrated 1 CUDA file(s) -> 1 file(s); 1 DPCT warning(s) (DPCT1009x1) |
| 3  | compile      | ok     | compiled 1 source(s) -> …/work/build/candidate (0 warning(s)) |
| 4  | scan_helpers | fail   | 11 dpct helper use(s) remain in 1 file(s) (events_timing=3, macro=2, device_queue_management=2, memory_management=2) |
| 5  | write_file   | ok     | wrote migrated/saxpy.dp.cpp (1386 chars) |
| 5  | compile      | fail   | compile failed with 15 error(s) at …/saxpy.dp.cpp:15: use of undeclared identifier 'sycl' |
| 6  | write_file   | ok     | wrote migrated/saxpy.dp.cpp (1411 chars) |
| 6  | compile      | ok     | compiled 1 source(s) -> …/work/build/candidate (0 warning(s)) |
| 6  | scan_helpers | ok     | no dpct helpers in 1 file(s) |
| 7  | verify       | ok     | PASS: 1/1 case(s) match (worst max abs err 4.77e-07, atol 1e-06) |
| 7  | benchmark    | ok     | 5 run(s) each; Intel(R) Xeon(R) Processor @ 2.80GHz: median 217.41 ms, p95 239.39 ms |

What failed and what fixed it
- it 5: `compile:saxpy.dp.cpp:15:use of undeclared identifier 'sycl'` → fixed by iteration 6 after editing saxpy.dp.cpp

Benchmark numbers are whole-process wall time on a 4-vCPU cloud VM's OpenCL CPU device. They include SYCL start-up and JIT compilation. See examples/README.md before quoting them.

Status

Roadmap

• Ablation on models from other vendors; harder, multi-file tasks.

• cuFFT → oneMath, and a multi-file HeCBench app.

• Intel GPU validation (Arc/Max through Level Zero).

• AOT builds (-fsycl-targets=) and in-process kernel timing by default, so benchmarks exclude JIT.

• verify input generators (random arrays, edge cases such as denormals, NaN and huge values) and multi-output programs.

• More scan rules: sycl::ext::oneapi::experimental usage, implicit double on devices without fp64, USM malloc_shared in hot loops.

• Multi-file CUDA projects using compile_commands.json (c2s -p).

Credits

• REA (MIT) inspired the design. Its evidence-first approach, where tools return structured evidence rather than claims, is the principle this toolkit is built on.

• SYCLomatic (Intel; Apache-2.0 with LLVM exceptions) does the bulk CUDA→SYCL migration that this toolkit drives and post-processes.

• HeCBench (Zheming Jin et al.; BSD-3-Clause) is the source of the softmax example. See examples/softmax_hecbench/LICENSE.HeCBench.

• Intel oneAPI DPC++/C++ Compiler (icpx).

License

MIT; see LICENSE. Third-party example code keeps its original license, as noted above.

Project tracking: PROGRESS.md (what was done, what is blocked) · DECISIONS.md (choices made autonomously) · docs/install.md