Give it a PyTorch model on any one machine. It profiles, generates, validates, and benchmarks optimized kernels on the hardware you have — and qualitatively ranks the same kernels for hardware you don't, by researching each target's ISA and climbing a confidence ladder of static analysis, analytical models, and simulators.
VladBud = AutoKernel (autonomous GPU
kernel optimization — the agent edits one file, runs a fixed evaluation, keeps/reverts,
forever) + crosskernel/, an additive layer that makes it
vendor-portable: AMD ROCm, Intel GPUs, CPU SIMD (AVX-512/AMX/NEON/SVE), and named
accelerators (TPU, Hexagon, Ascend, Cerebras, Trainium, Tenstorrent).
Built and verified live on an NVIDIA GB10 (Grace-Blackwell, DGX Spark, sm_121).
Two hard problems sit under "make a kernel optimizer cross-platform":
- Portability (P1) — one agent loop, many targets, each with a different language (Triton / HIP / SYCL / C++ SIMD / Pallas / CSL), compiler, and execution model.
- The hardware-you-don't-have problem (P2) — you may hold a single NVIDIA card, yet must still create, validate, and rank kernels for an AMD MI300X or a TPU you cannot run. This is the crux — it's what separates real cross-platform optimization from "just add more backends."
AutoKernel only judges a kernel by running it on the local NVIDIA GPU. VladBud answers P2 with a confidence ladder: cheap static signals reject bad candidates, analytical models give a relative ranking, simulators give occasional high-trust validation, and real hardware (when present) calibrates everything above it — so one card earns trustworthy estimates for targets it can't execute.
- 8/9 kernel types generated + correctness/stability PASS on the physical GB10; the
9th (
fused_mlp) is correctly flagged — its fp32 path uses TF32, which the strict gate catches (a true positive, not a miss). - The same kernels ranked for AMD MI300X / Intel Max / TPU v6e with no vendor hardware
(confidence-ladder
RANKED), and for unowned NVIDIA archs via compile-onlyptxas. - The autotuned Triton matmul beats cuBLAS on the brand-new Blackwell by up to 1.43× (vendor libraries lag new silicon).
- A 9-kernel reference corpus (web-researched, cited) is the comparison ground truth.
- The evaluator survived a brutal multi-agent code review (9 confirmed bugs incl. two false-pass holes, all fixed) and a 12-guard adversarial regression suite.
- AutoKernel researches hardware/ISA/docs itself (e.g. it discovered GB10 lacks
TMEM/
tcgen05and folded that into its generation knowledge).
| Target | Author surface | Validate (no HW) | Rank (no HW) | Status here |
|---|---|---|---|---|
| NVIDIA (incl. GB10) | Triton / CUDA-C++ | Triton interpreter | ptxas regs/smem, roofline | measured on GB10 |
| AMD ROCm | Triton / HIP (auto-hipify) | Triton interpreter / CPU ref | roofline (gfx), re-autotune | generated + ladder-ranked |
| Intel XPU | Triton / SYCL | Triton interpreter | roofline (Xe) | generated + ladder-ranked |
| CPU (x86 + ARM) | C++ SIMD / Triton-CPU | runs natively (universal) | llvm-mca / OSACA + roofline | measured on Grace CPU |
| Google TPU | Pallas | Pallas interpret=True |
Timeloop/MAESTRO | modelled + ladder-ranked |
| AWS Trainium | NKI | nki.simulate_kernel |
analytical | modelled |
| Qualcomm Hexagon | HVX / hexagon-mlir | hexagon-sim ISS |
cycle-approx PMU | modelled |
| Huawei Ascend | Ascend C | CPU-twin (ICPU_RUN_KF) |
CAModel | modelled |
| Cerebras | CSL | cycle-accurate fabric sim | same (perf too!) | modelled |
| Tenstorrent | TT-Metalium | ttsim (executes kernels) |
ttsim + analytical | modelled |
"measured" = generated, built, validated, and timed on local hardware. "ladder-ranked" = correctness path + analytical roofline with no vendor hardware present. "modelled" = the target carries its specs, validation method, and simulator in the spec DB.
Requirements: a GPU with a matching PyTorch build, Python 3.10+, a C/C++ compiler
(g++/clang). Optional: llvm-mca, ptxas/nvcc, ninja (for the CPU load_inline
build).
GB10 / Blackwell note (important): the GB10's
sm_121needs CUDA 13 wheels (cu130) —torch 2.12+cu130+triton 3.7. The base repo'spyproject.tomlpins thecu128index, which is wrong for Blackwell (it stalls on multi-GB downloads and, even if it completed, is the wrong CUDA). On a DGX Spark, use the system Python that already has the cu130 stack and run scripts withpython3 …(notuv run). Hardware-aware dependency selection is part of the cross-platform story — seeCROSSKERNEL_DESIGN.md§2.
git clone https://github.com/BudEcosystem/VladBud.git
cd VladBud
# Base AutoKernel deps (on non-Blackwell NVIDIA, the documented path works):
uv sync # or use a torch build matching your GPU
pip install ninja # for the crosskernel CPU backend's load_inline build
# One-time AutoKernel setup (test data + baselines):
python3 prepare.pyexport PATH=$HOME/.local/bin:$PATH # so load_inline finds ninja
python3 -m crosskernel.run --all # REAL verification: all 9 kernels on every local target
python3 -m crosskernel.run matmul # one kernel across local (measured) + remote (ladder-estimated)
python3 -m crosskernel.tests.adversarial # evaluator trust-anchor regression suite (12 guards)
python3 -m crosskernel.research # hardware/ISA knowledge cards (incl. unowned targets)
python3 -m crosskernel.demo # foundational thesis demo (create+validate+rank, incl. unowned HW)
python3 crosskernel/targets.py # hardware model + local detection + roofline# Profile any PyTorch model, rank kernels by GPU time:
python3 profile.py --model models/llama_7b.py --class-name LlamaModel --input-shape 1,512 --dtype float16
python3 extract.py --top 5 # extract bottleneck kernels into workspace/
python3 bench.py # 5-stage correctness + roofline on one kernel
# Then run an agent (Claude Code / Codex) in this dir: "Read program.md and start a new experiment."Additive (crosskernel/, ~3k lines) — it imports nothing from AutoKernel's frozen harness
(bench.py, reference.py, …) and never modifies it, so the base NVIDIA+Triton flow keeps
working unchanged.
| Module | Role |
|---|---|
targets.py |
vendor-neutral HardwareTarget + data-driven spec DB (15 targets, 11 vendors) + local detection + roofline |
research.py |
autokernel researches hardware/ISA/architecture → KnowledgeCard (curated playbook + corpus + live-web hook) that steers generation |
corpus/ |
the reference-kernel corpus — 9 web-researched, cited cross-platform references + per-hardware accuracy/stability/perf baselines |
generate.py |
research-informed kernel generation per (kernel, target): Triton (NVIDIA/AMD/Intel) + C++/SIMD (CPU); materialize() → callable |
evaluate.py |
target-parametric evaluator: fp64-oracle accuracy + adversarial stability + measured/ladder-estimated performance vs corpus |
ladder.py |
the confidence ladder: Rung 0 reject → 1 static → 2 correctness → 3 analytical → 4 sim → 5 real-HW |
analysis/static.py |
hardware-free signals: compiler vectorization report, llvm-mca, ptxas -v |
backends/ |
Backend registry (build·validate·time·static_analyze); cpu-cpp universal local backend |
run.py |
end-to-end orchestration: research → generate → evaluate → compare across local + remote targets |
tests/adversarial.py |
the evaluator's adversarial regression suite (12 guards against false-pass / reward-hacks) |
1. The confidence ladder (answer to P2). No single technique both validates and ranks a kernel for absent hardware, so candidates climb a ladder of ascending cost and trust:
| Rung | Technique | Needs target HW? | Signal |
|---|---|---|---|
| 0 Static reject | anti-pattern + reward-hack screen | no | reject only |
| 1 Static signals | CPU: vec-report + llvm-mca; GPU: ptxas -v regs/spills, Triton IR | no (compile only) | perf proxy |
| 2 Correctness | CPU ref / Triton interpreter / functional sim (NKI, Pallas, ttsim, …) | no | correctness |
| 3 Analytical | roofline; Timeloop/MAESTRO/SCALE-Sim for systolic | no (spec sheet) | perf rank |
| 4 Cycle sim | Accel-Sim (NVIDIA), Cerebras fabric sim, … | no, slow | perf |
| 5 Real hardware | do_bench / measure on the device |
yes | ground truth |
Two principles: optimize for relative rank, not absolute latency; and calibrate down the ladder whenever real hardware is touched.
2. The fp64 oracle (correctness done right). AutoKernel's elementwise allclose(atol=1e-4)
mis-FAILs correct matmuls at near-zero output elements (two valid reduction orders disagree
by a few ULP). VladBud compares against the true fp64 result rounded to the kernel's output
dtype with a relative-Frobenius metric — robust to near-zero cancellation and to overflow
(where inf is the correct answer and must match), while still catching silent precision
downgrades (the fused_mlp TF32 case).
3. The evaluator is an adversarial boundary. Every documented kernel-gen system that
trusted a naive harness produced inflated or fraudulent speedups (Sakana's recycled-output
"150×"; METR's timer monkey-patching). The evaluator never benchmarks a kernel that failed
correctness, flags physically-impossible >100%-of-peak readings, carries ±2× uncertainty on
theoretical CPU peaks, tags every cross-target result with its confidence rung — and is
locked down by tests/adversarial.py (broken / cheating / unverifiable kernels all FAIL).
python3 -m crosskernel.run --all — generated, built, validated, and benchmarked on the
physical device (% of the relevant roofline ceiling: compute for matmul-class, bandwidth
for normalization/reduction-class):
| kernel | accuracy | stability | performance (GB10) | bound | corpus reference |
|---|---|---|---|---|---|
| matmul | PASS | PASS | ~69 TFLOPS, 55% | compute | A100 cuBLAS ~250 TFLOPS |
| softmax | PASS | PASS | ~217 GB/s, 79% | memory | H100 Triton ~820 GB/s |
| layernorm | PASS | PASS | ~203 GB/s, 74% | memory | H100 Triton 765 GB/s |
| rmsnorm | PASS | PASS | ~211 GB/s, 77% | memory | A100 1365 GB/s |
| reduce | PASS | PASS | ~208 GB/s, 76% | memory | (bandwidth-bound) |
| rotary_embedding | PASS | PASS | ~111 GB/s, 41% | memory | A100 ~8× fused |
| cross_entropy | PASS | PASS | ~226 GB/s, 83% | memory | A100 >2× |
| fused_mlp | FAIL | PASS | flagged (TF32 fp32) | compute | GB10 ~91 TFLOPS |
| flash_attention | PASS | PASS | ~75 TFLOPS, 60% | compute | H100 FlashAttention-3 |
GB10 (DGX Spark) is power-capped and thermally throttles under sustained load (matmul measured 69 TFLOPS cool → 50 TFLOPS after a long run). Correctness is stable; absolute TFLOPS vary — which is exactly why the harness reports confidence and the design calls for clock-locking. Full evidence:
crosskernel/VERIFICATION.md.
Research → design → build → review → fix → verify, with dynamic multi-agent workflows for the fan-out phases:
- Research (5 parallel agents, ~40 cited sources) — AMD/Intel GPU stacks, non-GPU accelerators + simulators, CPU SIMD + static analysis, hardware-free perf estimation, portable IRs + agentic-kernel-gen prior art.
- Reference corpus (10-agent workflow) — canonical implementations + baselines per kernel.
- Build — the 16-module system, GB10-verified at each step.
- Brutal review (multi-agent workflow, dimensions → find → adversarially verify each finding) → 9 confirmed bugs incl. two evaluator false-passes → all fixed.
- Fix-validation (second review workflow) → confirmed fixes correct, 2 LOW refinements → fixed. Locked down by the adversarial regression suite.
VladBud/
crosskernel/ ← the cross-platform system (additive)
targets.py hardware model + data spec DB + detection + roofline
research.py hardware/ISA research → KnowledgeCard
generate.py per-target kernel generation
evaluate.py target-parametric evaluator (fp64 oracle + ladder)
ladder.py the confidence ladder
run.py end-to-end orchestration (run --all)
analysis/static.py llvm-mca / ptxas / vectorization-report
backends/ Backend registry + cpu-cpp universal backend
corpus/data/*.json 9 researched reference kernels + baselines
tests/adversarial.py evaluator trust-anchor regression suite
demo.py foundational thesis demo
README.md VERIFICATION.md
CROSSKERNEL_DESIGN.md ← architecture + methodology (design doc)
kernel.py program.md bench.py reference.py prepare.py ← AutoKernel base (frozen harness)
profile.py extract.py orchestrate.py verify.py
kernels/ kernels/cuda/ kernelbench/ models/
See CROSSKERNEL_DESIGN.md for the full architecture and the
portable-representation decision (tile-DSL source + a PJRT/IREE-HAL-style HAL, not a single
IR), and program.md for the AutoKernel agent's operating manual.
Built on AutoKernel by RightNow AI —
"autoresearch for GPU kernels," itself inspired by Andrej Karpathy's
autoresearch. KernelBench integration is based on
Stanford Scaling Intelligence Lab's
KernelBench (Ouyang, Guo, et al.). The
cross-platform layer is grounded in ~40 cited primary sources across the AMD ROCm/HIP, Intel
oneAPI/SYCL, Triton, TPU/Pallas, Hexagon, Ascend, Cerebras, and CPU-SIMD ecosystems
(summarized in CROSSKERNEL_DESIGN.md).
MIT (inherited from AutoKernel).