oneapi-agent-toolkit

AGENTS.md: porting CUDA to SYCL with oneapi-agent-toolkit

Instructions for any coding agent working in this repository or using its tools. This is a plain-text version of the oneapi-port skill; if they ever differ, the skill is the reference.

What has been verified, and what has not

• Ports measured end to end (every result in docs/) were made by the bundled loop, oneapi-agent loop, driving Claude models (claude-sonnet-5-5, claude-haiku-5-5) through the Anthropic SDK.
• MCP registration with oneapi-agent setup was checked against the real clients: Claude Code, Gemini CLI (Connected), Cursor (ready, all tools listed) and Codex (entry parsed). See docs/agent-registration.md. That shows the server starts and its tools are listed; no port made by those clients over MCP has been measured.
• Any other agent (or the CLI route below) is untested. The tools are ordinary commands that print JSON, so they should work anywhere a shell does, but nothing here claims that.

Two ways to call the tools

MCP (Claude Code, Codex, Cursor, Gemini CLI, …): oneapi-agent setup --agent NAME registers the server. Tools: migrate, compile, scan_helpers, verify, benchmark, gpu_test, doctor.

CLI (any agent with a shell): the same functions as commands. Each prints one ToolResult as JSON on stdout (schema) and exits with 0 = ok, 1 = fail, 2 = anything else (error, toolchain_missing, timeout). Put --summary before the command for a one-line result instead of JSON.

MCP tool	CLI equivalent (flags as in oneapi-agent <cmd> --help)
doctor	oneapi-agent doctor
migrate(cuda_path, out_dir)	oneapi-agent migrate CUDA_PATH --out OUT_DIR [--in-root DIR] [--extra-arg=ARG ...]
compile(sources, flags, output)	oneapi-agent compile SRC [SRC ...] [--flags "-O2 -Iinc"] [--output BIN] [--no-dpct-include]
scan_helpers(path)	oneapi-agent scan-helpers PATH
verify(reference_cmd, candidate_cmd, ...)	oneapi-agent verify --ref "REF_CMD {output}" --cand "CAND_CMD {output}" [--input FILE ...] [--atol X] [--rtol X] [--max-ulp N] [--mode auto\|array\|audio] [--threshold-dbfs DB] [--work-dir DIR]
benchmark(cmd, runs, devices)	oneapi-agent benchmark "CMD ARGS" [--runs N] [--warmup N] [--device SELECTOR ...] [--report-regex RE]
gpu_test(task_file, out_dir)	oneapi-agent gpu-test TASK_JSON [--out DIR] [--vendor nvidia\|amd\|intel] [--device SELECTOR] [--runs N] [--no-native]

migrate needs SYCLomatic (c2s), compile needs icpx, and gpu-test needs a SYCL GPU visible to sycl-ls; without them the result is toolchain_missing and its evidence names what is missing. scan-helpers is a text scan and needs nothing; verify and benchmark run whatever commands you give them (benchmark uses sycl-ls, when present, to time each SYCL device). See docs/install.md.

Example: oneapi-agent --summary verify --ref "./ref {output}" --cand "./port {output}" --atol 1e-6 prints e.g. [verify] ok: PASS: 1/1 case(s) match (worst max abs err 4.77e-07, atol 1e-06).

Ground rules

1. Target SYCL 2020 for Intel oneAPI DPC++ (icpx -fsycl). Do not keep CUDA idioms: no hard-coded warp size 32, no error-code plumbing, no 3-D launches for 1-D work. Use sycl::queue, USM, parallel_for, nd_range, group algorithms, sycl::atomic_ref.
2. Every claim cites a tool result: quote status, summary and the key evidence numbers. Never write “should work”, “is equivalent” or “is faster” without one.
3. Never report success without verify returning ok on the current code. Editing a file after the last passing verify makes it unverified again.
4. Never fabricate. On toolchain_missing, say so, quote evidence.failed_command, and continue only with steps that do not need that tool.
5. Read limitations on every result and carry the relevant ones into your report.
6. Do not widen tolerances to pass. If a numerical reason forces it, report old and new.

Procedure

migrate -> compile -> [fix errors] -> scan_helpers -> refactor to idiomatic SYCL
        -> compile (no dpct include) -> [fix] -> verify -> [fix] -> scan_helpers (= 0)
        -> benchmark -> report

1. migrate. Read evidence.warnings: each DPCT code marks something c2s could not translate faithfully (DPCT1009/1010 error codes, DPCT1049 work-group size, DPCT1065 barrier hint, DPCT1083 kernel parameter layout). On fail, c2s could not parse something; pass include paths as --extra-arg=-I<dir> (repeatable; the = is required because the value starts with -).
2. compile. Fix the first error first; later ones are often cascades. Repeat until ok.
3. scan_helpers on the migrated tree gives the work list: by_category, occurrences, and suggestions per category.
4. Refactor, one category at a time:
   • device/queue management: one sycl::queue (in-order if the code relies on ordering, as CUDA’s default stream does), passed explicitly;
   • memory: sycl::malloc_device / malloc_shared, q.memcpy, sycl::free(p, q);
   • error handling: remove dpct::err0 and DPCT_CHECK_ERROR; use try { } catch (const sycl::exception&);
   • timing: the sycl::event from submit with property::queue::enable_profiling;
   • sub-groups: sycl::shift_group_left, select_from_group, reduce_over_group; never assume a sub-group size of 32;
   • atomics: sycl::atomic_ref; math: sycl:: builtins;
   • act on advisories too (e.g. nd_item3_for_1d: use range<1> / nd_range<1>);
   • remove #include <dpct/...> last.
5. compile again without the dpct include (--no-dpct-include) to prove the helpers are gone. Fix until ok.
6. verify against the reference. Arrays: start near atol 1e-5 for float32. Audio: --mode audio, pass when the residual peak is at most -96 dBFS. On fail, use first_mismatches and max_abs_err to localise the bug (indexing and range rounding, missing bounds checks, reduction order, float vs double literals, denormals/FTZ).
7. scan_helpers must be ok (0 helpers) on the final code.
8. benchmark the final candidate; report only the numbers returned.

Loop discipline

• One iteration per fix attempt after a failing compile or verify; stop at the iteration limit (default 8) and report the best version so far.
• Keep the best passing version: snapshot files whenever verify passes; never overwrite a passing snapshot with an unverified edit.
• Stuck rule: the same error three times in a row means stop trial-and-error; re-read the evidence and change approach, or report that you are stuck, quoting the error.

Final report

## Port report: <name>
Result: SUCCESS | PARTIAL (best passing version kept) | BLOCKED (<reason>)
Iterations: N (compile fixes: a, verify fixes: b)
migrate:      <status> - <k> DPCT warnings (<codes>)
compile:      <status> - <errors/warnings>, dpct include needed: yes/no
scan_helpers: <status> - <before> -> <after> helpers
verify:       <status> - <cases>, max_abs_err=<x> / residual peak=<y> dBFS (threshold <t>)
benchmark:    <device>: median <m> ms, p95 <p> ms (runs=<n>)
What failed and how it was fixed:
- <error> -> <fix> (iteration i)
Limitations: <carried over from tool results>

SUCCESS only when verify is ok and scan_helpers is ok on the same final code.

This site is open source. Improve this page.