I had an AI agent port Karpathy's GPT-2 trainer from CUDA to SYCL, and made it prove every step

Checo Cadena — Fri, 06 Nov 2026 20:48:55 +0000

A lot of AI software only runs on NVIDIA GPUs, because it's written in CUDA. SYCL is an open standard (from Khronos, the group behind OpenGL and Vulkan) that lets the same C++ code run on NVIDIA, AMD, Intel and plain CPUs.

Moving code from CUDA to SYCL is slow, careful work. AI coding agents are now good enough to do a lot of it. The problem is trust: an agent saying "done" is not the same as the code being right.

So for Meta's Global AI Developer Hackathon I built Muse Port: Meta's Muse model plans and ports real CUDA repositories to SYCL, and nothing counts as done until a check the model never sees confirms the output matches the original.

Code and evidence: https://github.com/checocadena/museport
2-minute demo: https://youtu.be/uBpT5IuxPOY

How it works

Scan, no model involved. museport scan reads the whole CUDA repo and lists every kernel, every NVIDIA library call (cuBLAS, cuDNN, NCCL...) and its oneAPI equivalent, and every CUDA-only assumption.
Plan. Muse writes a migration plan grounded against that scan, so it can't invent files or skip the hard parts.
Port, item by item. An agent loop writes SYCL, compiles it, and runs it. A first pass from Intel's SYCLomatic tool is available, and the agent finishes the job.
Hidden checks. Each item is graded by a check the model never sees, usually built from the project's own CPU code or its own tests. If it fails, the loop keeps going.
Real GPUs. Finished ports run on rented NVIDIA GPUs side by side with the original CUDA, and on AMD Instinct GPUs against the reference.
Report. The write-up is generated only from committed evidence. Every number links to a log.

The test case: llm.c

I used Andrej Karpathy's llm.c, which trains GPT-2 in raw CUDA. Its main trainer is 12 files and 4,884 lines of bf16 mixed precision, with cuBLASLt, warp reductions and stochastic rounding. It's real code, not a toy.

The hidden check was llm.c's own test, at llm.c's own bf16 tolerances, with a guard that fails any port that tries to loosen them.

What happened:

The port, in one run: 125 agent turns and $5.67 of model calls, from the CUDA alone. llm.c's test passes on a CPU, an AMD MI350X and an NVIDIA GPU. That covers logits, all 16 gradient tensors, 10 losses against PyTorch, and a bit-exact checkpoint replay.
It was slow: 36x slower than CUDA.
Then the agent tuned it on real GPUs. I gave it a gpu_test tool that builds its current code on a rented GPU and returns llm.c's test result, milliseconds per step next to llm.c's CUDA on the same GPU, and an nsys kernel profile. It profiled, rewrote the slow kernels, measured each change, and fixed two crashes the GPU exposed.

Where it ended up, as time per training step compared with llm.c's CUDA on the same GPU:

Setup	vs CUDA
Portable SYCL + oneMath, NVIDIA L4 / RTX A4000	1.40-1.55x
Portable SYCL + oneMath, NVIDIA H100	2.90x
With NVIDIA's cuBLAS called through SYCL interop	1.13-1.21x
AMD Instinct MI350X	63 ms/step (llm.c's CUDA doesn't build on AMD)

It also ported 19 of llm.c's standalone kernel experiments, a fp32 version reassembled as a working repo (llm.c-sycl, whose make builds llm.c's own test), and a second project, llama2.cu, in one run for $0.45, with byte-identical stories on CPU, NVIDIA and AMD.

What a port looks like

One real kernel, before and after:

// CUDA (llm.c)
__global__ void residual_forward_kernel(float* out, float* inp1, float* inp2, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        out[idx] = __ldcs(&inp1[idx]) + __ldcs(&inp2[idx]);
    }
}

// SYCL (as Muse wrote it)
static void residual_forward(float *out, const float *a, const float *b, size_t N) {
    launch_1d(N, [=](sycl::id<1> id) { out[id[0]] = a[id[0]] + b[id[0]]; });
}

__ldcs is an NVIDIA-only cache hint, so the port drops it. The result is identical everywhere, but that kind of vendor tuning is exactly what you have to add back per target. Which brings me to the honest part.

What I learned

Portable and correct is not the same as fast. The H100 number is the clearest example. The fastest NVIDIA code uses chip-specific features (fused kernels around tensor cores, warp specialization, Hopper's clusters) that portable code can't fully reach today. My bet is that the realistic version of "portable" is a correct portable baseline plus vendor-specific fast paths for the few kernels that dominate runtime. Agents might be what finally makes that second part affordable.

Hide the checks. An agent that can read its test can end up satisfying the test instead of the program, for example by loosening a tolerance. That's why the grading lives somewhere the model can't see, and why a guard fails any port that touches llm.c's tolerances.

Publish the failures. The one-run bf16 port succeeded in 1 of 2 tries. One tuning run's own loop didn't converge, even though a version it built along the way passes. Speeds are medians of 2-3 runs on rented GPUs whose exact model varies, so compare the ratios. All of that is in the repo, including every mistake and its fix.

Agents find real upstream bugs. Along the way the port hit unimplemented for bf16 gemm in oneMath's cuBLAS backend, oneAPI's math library. The fix was two lines, and a oneMath reviewer has approved it: https://github.com/uxlfoundation/oneMath/pull/793

Try it

git clone https://github.com/checocadena/museport
cd museport/llm.c-sycl
./dev/download_starter_pack.sh
make                      # CPU; or: make TARGET=nvidia / TARGET=amd AMD_ARCH=gfx950
./test_gpt2fp32sycl       # llm.c's own test: ends with "overall okay: 1"

I'll be talking about this at the oneAPI DevSummit, hosted by the UXL Foundation, on December 1 (online): "Evidence, Not Claims: Teaching Coding Agents the Last Mile of CUDA-to-SYCL Migration."
https://uxlfoundation.org/events/oneapi-devsummit-hosted-by-uxl-foundation-2026/