oneapi-agent

SYCLomatic gets CUDA code most of the way to SYCL, but leaves output full of dpct:: helpers and no proof the port still behaves like the original, and coding agents asked to finish the job tend to stop as soon as it compiles. oneapi-agent-toolkit gives an agent five tools that return evidence instead of claims (migrate, compile, scan_helpers, verify, benchmark), through an MCP server or a bundled agent loop, and counts a port as done only when verification passes and no helpers remain on the same source. In measured runs with Claude models, agents using it finished every port correctly, while the same agents without its checks often stopped early; the agent-written ports then passed unedited on rented NVIDIA and AMD GPUs.

Results at a glance

Claude models only, graded on a CPU device by a hidden check the agent never saw: indicative, not a benchmark. Every number links to the committed traces, final sources and tool results.

Demo: 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).

Try it in a minute, no toolchain or API key: clone, pip install -e ".[all]", then scripts/demo.sh. It runs the agent loop end to end through both model providers (Claude, and Meta's client against a scripted server) with canned tool results, clearly marked MOCK. The real tools need the oneAPI toolchain: see the Quickstart.

More: the full story and its limits in docs/overview.md · GPU testing on four clouds in docs/gpu-testing.md · a real model porting the original examples in docs/real-runs/ · cuBLAS → oneMath in examples/gemm_cublas/ · CI re-checks the claims on every push with scripts/check_claims.py (runs) · agents without MCP: AGENTS.md.

"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.

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: flameox

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]
    end
    subgraph Models["model providers (models.py)"]
        ANT[anthropic<br/>Claude, Anthropic SDK]
        META[meta<br/>Muse, OpenAI-compatible API]
    end
    CC <-- MCP stdio --> SRV[oneapi-agent serve<br/>MCP server]
    LOOP --> ANT & META
    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
oneapi-agent loop --dry-run --mock-tools # works now: no toolchain, no API key (results marked MOCK)
scripts/demo.sh                          # the same through both providers (anthropic, meta)

# Everything below needs the oneAPI toolchain (icpx, SYCLomatic c2s). Without it each tool
# returns status "toolchain_missing" (exit code 2) and names what is missing.
# The route verified here (see docs/install.md):
scripts/install-toolchain.sh /opt && source scripts/oneapi-env.sh

# Use the tools directly (each prints a JSON ToolResult; add --summary for one line)
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     # expect "fail" (exit 1): raw c2s output keeps dpct helpers (11 here with c2s 21.0.0); closing that gap is the agent's job
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 (needs an installed agent)
oneapi-agent setup --dry-run --agent claude   # preview for a named agent, installed or not
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 (needs the toolchain)
ONEAPI_AGENT_API_KEY=... oneapi-agent loop --config examples/fir/task.json   # real model (paid API calls)
oneapi-agent loop --dry-run --mock-tools --provider meta   # Meta client, scripted server: no key

The loop's model comes from --model or ONEAPI_AGENT_MODEL (default claude-sonnet-5-5), and its provider from --provider (anthropic or meta; by default inferred from the model name). Meta's Muse models: see docs/meta-model-api.md. All settings are listed in .env.example.

• 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.

• Cost cap: each run stops at max_cost_usd ($5 by default), priced from a built-in table of Claude models. For any other model set ONEAPI_AGENT_PRICE_IN and ONEAPI_AGENT_PRICE_OUT (USD per million tokens; optionally ONEAPI_AGENT_PRICE_CACHED_IN). Without a price the loop refuses to start, because the cap could not apply; ONEAPI_AGENT_ALLOW_UNPRICED=1 runs anyway with only the token cap, and says so.

• A missing API key is reported before any work starts.

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

• PLANNED, no results yet: evaluate non-Claude models, including Meta's Muse models through the Meta Model API, with the same hidden checks, workloads, arms and run counts as the Claude comparisons. The provider layer is in place (--provider meta, tested offline); prices for the cost cap come from ONEAPI_AGENT_PRICE_IN / ONEAPI_AGENT_PRICE_OUT.

• 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 is the source of the six examples/*_hecbench programs: five under BSD-3-Clause (Zheming Jin et al.; LICENSE.HeCBench), and atomic_reduction under MIT (AMD; LICENSE.AMD). See examples/README.md.

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

Third-party components

Everything else in this repository is original work by the author.

License

Copyright (C) 2026 Checo Cadena.

Open source: AGPL-3.0. Use it, study it, change it and share it freely. If you distribute a modified version, or let others use one over a network (for example as a hosted service), you must publish your source under the AGPL too.

Commercial licenses are available for using it in a closed-source product or service without the AGPL's terms: contact checo@innercircuit.com.

• Earlier versions: everything up to and including commit b08b687 was released under the MIT license, and copies of those versions remain under MIT.

• Third-party example code keeps its own license (BSD-3-Clause or MIT, see Credits); both are compatible with the AGPL.

• Contributing: see CONTRIBUTING.md; outside code contributions need a contributor license agreement, which is not in place yet.

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