KernelBench: Can LLMs Write Efficient GPU Kernels?

A benchmark for evaluating the ability of LLMs to write efficient GPU kernels.

Contributed By

@jiito

Code

src/inspect_evals/kernelbench

Paper

https://arxiv.org/html/2502.10517v1

Overview

KernelBench is a benchmark for evaluating whether LLMs can generate efficient GPU kernels.

This implementation follows KernelBench 0.1, an improved version of KernelBench by the original authors. See this blog post for more information: https://scalingintelligence.stanford.edu/blogs/kernelbenchv01/

Usage

Note: This eval has dependencies that conflict with other evals in this repo. It ships its own isolated environment under packages/kernelbench/ rather than sharing the repo’s main one.

Installation

From the repository (recommended)

Sync the eval’s environment from its package directory:

cd packages/kernelbench
uv sync

Then run evaluations from that directory:

uv run inspect eval inspect_evals/kernelbench --model openai/gpt-5-nano

Into an existing environment

If you already have a Python environment, install the eval’s package directly — it brings in inspect-evals alongside the eval-specific dependencies:

pip install /path/to/inspect_evals/packages/kernelbench/

Then run without uv run:

inspect eval inspect_evals/kernelbench --model openai/gpt-5-nano

You can also import the task directly from Python:

from inspect_ai import eval
from inspect_evals.kernelbench import kernelbench
eval(kernelbench)

From a Jupyter notebook (using Inspect Flow)

Inspect Flow manages the isolated environment for you. Install it once:

pip install inspect-flow

Then from a notebook cell, point Flow at this eval’s pyproject.toml so it builds the isolated venv from the same dependency set as uv sync:

from inspect_flow import FlowSpec, FlowTask, FlowDependencies
from inspect_flow.api import run

spec = FlowSpec(
    execution_type="venv",
    log_dir="logs/",
    dependencies=FlowDependencies(
        dependency_file="path/to/inspect_evals/packages/kernelbench/pyproject.toml",
    ),
    tasks=[FlowTask(name="inspect_evals/kernelbench", model="openai/gpt-5-nano")],
)

run(spec)  # blocking; builds isolated venv on first run, cached thereafter

After the run completes, read and analyse logs in the same notebook — the eval’s heavy dependencies are not required for log reading.

Running evaluations

After installing via any of the methods above, you can evaluate models. The examples below assume the repository approach; drop uv run if you are managing dependencies yourself.

uv run inspect eval inspect_evals/kernelbench --model openai/gpt-5-nano

Log viewing (inspect view) and default-model setup are documented in the Inspect Evals README.

Maintenance tasks

The tasks below score the evaluation harness rather than the model, for example a judge meta-evaluation or a sandbox image check. Their results are not benchmark results.

kernelbench_sandbox_check: Runs the real scorer over fixture kernels with known verdicts to verify a sandbox image on GPU hardware; accuracy 1.0 certifies the image.

uv run inspect eval inspect_evals/kernelbench_sandbox_check

Options

You can control a variety of options from the command line. For example:

uv run inspect eval inspect_evals/kernelbench --limit 10 --sample-shuffle
uv run inspect eval inspect_evals/kernelbench --max-connections 10
uv run inspect eval inspect_evals/kernelbench --temperature 0.5

See uv run inspect eval --help for all available options.

Parameters

`kernelbench`

splits (list[Literal['level_1', 'level_2', 'level_3', 'level_4']]): Dataset splits to use for evaluation (default: ['level_1', 'level_2', 'level_3'])
limit (int | None): Maximum number of samples to evaluate (None for all) (default: None)
option (Literal['zero_shot', 'one_shot', 'few_shot']): Prompt option for kernel generation (default: 'one_shot')
precision (Literal['fp32', 'fp16', 'bf16']): Computation precision (torch.dtype) (default: 'fp32')
include_hardware (bool): Whether to include hardware-specific information in prompts (default: False)
gpu_name (Literal['L40S', 'H100', 'A100', 'A100-80GB', 'L4', 'T4', 'A10G']): Target GPU name for hardware-specific prompts (default: 'L40S')
backend (Literal['cuda', 'triton', 'tilelang', 'cute', 'thunderkittens']): Backend to use (‘cuda’, ‘triton’, ‘tilelang’, or ‘cute’) (default: 'cuda')
check_kernel (bool): Whether to perform static kernel validation before scoring (default: True)
max_regenerations (int): Maximum number of regeneration attempts for invalid kernels (default: 10)
num_correct_trials (int): Number of trials with different random inputs for correctness testing; correctness passes only if all trials pass (default: 5)
num_perf_trials (int): Number of performance trials to run for averaging timing results (default: 100)
measure_performance (bool): Whether to measure kernel performance (default: True)
timing_method (Literal['cuda_event', 'do_bench', 'do_bench_impl', 'host_time']): Method to time kernel execution (see timing.py for details) (default: 'cuda_event')
verbose (bool): Whether to enable verbose output during evaluation (default: False)
build_dir (pathlib.Path | None): Directory for building kernels (default: None)
device (int | None): GPU (cuda) device to run the evaluation on (default: None)
p (float): Threshold for fast_p metric (fraction of samples that are correct and have speedup > p) (default: 1.0)
temperature (float): Sampling temperature for generation, defaults to 0.0 for this task (default: 0.0)
sandbox_type (str): Sandbox provider type. “docker” (default, uses the pinned published image) and “k8s” have built-in configs; any other registered provider can be used by also passing sandbox_config. (default: 'docker')
sandbox_config (str | None): Custom sandbox config. For k8s, this can be a runner-local Helm values file path or inline YAML mapping; for other providers, it must be a runner-local config path and cannot be combined with sandbox_image or sandbox_node_selector. (default: None)
sandbox_image (str | None): Pullable image reference overriding the sandbox image. Docker and Kubernetes default to the pinned public KernelBench image; see “Running on Kubernetes” in the README. (default: None)
sandbox_node_selector (dict[str, str] | None): Kubernetes node labels used to select a GPU node. These are merged over any nodeSelector in sandbox_config. (default: None)

`kernelbench_sandbox_check` (maintenance task)

backend (Literal['cuda', 'triton', 'tilelang', 'cute', 'thunderkittens']): Backend whose passing fixture kernel to use (‘cuda’ or ‘triton’) (default: 'cuda')
precision (Literal['fp32', 'fp16', 'bf16']): Precision passed to the scorer (fp32, fp16, bf16) (default: 'fp32')
timing_method (Literal['cuda_event', 'do_bench', 'do_bench_impl', 'host_time']): Timing method passed to the scorer (default: 'cuda_event')
num_correct_trials (int): Correctness trials per kernel (default: 5)
num_perf_trials (int): Performance trials per kernel (default: 20)
device (int | None): GPU (cuda) device index passed to the scorer (default: None)
sandbox_type (str): Sandbox provider type, as for kernelbench (default: 'docker')
sandbox_config (str | None): Custom sandbox config, as for kernelbench (default: None)
sandbox_image (str | None): Image reference to check instead of the pinned default (default: None)
sandbox_node_selector (dict[str, str] | None): Kubernetes node labels, as for kernelbench (default: None)

Running with Docker

The default Docker configuration pulls the published, digest-pinned linux/amd64 image. The task records the resolved image reference in its metadata so timing results retain their measurement environment.

To exercise Dockerfile or sandbox dependency changes before publishing a new image, use the packaged local-build variant from a repository checkout:

uv run inspect eval inspect_evals/kernelbench -T sandbox_config=src/inspect_evals/kernelbench/compose.local.yaml

The published image is amd64-only. On another architecture, the local-build variant is the available route when the NVIDIA CUDA base image and locked dependencies support that host.

Updating sandbox dependencies

The Dockerfile installs the standalone sandbox uv project with uv sync --locked --no-dev --no-install-project. Its committed uv.lock pins direct and transitive Python dependencies, artifact hashes, and the full KernelBench commit. The sandbox project has no path dependency on inspect_evals. Host dependencies continue to use packages/kernelbench/uv.lock; changes to that host lock, including Dependabot updates and make lock-isolated, do not require a sandbox image rebuild.

The sandbox project declares KernelBench and Torch as its direct dependencies. Its lock contains 148 runtime packages for Linux amd64/Python 3.12, removing 34 packages inherited from the former host dependency graph, including inspect-ai. All retained versions, sources, and artifact hashes match the original 182-package snapshot from the 2026-09-16 image (packages/kernelbench/uv.lock at commit 23e2462449a550f11e8e259f28aa8cac7e903a7a, SHA-256 cfc90d957b738184d6b142b0beda1cc082107e3bcae33fc34a25964d8a006c36). KernelBench’s upstream dependency declarations are unchanged. They still include inference and development tooling; its eager imports require OpenAI and LiteLLM even when model calls run on the host.

The pinned default is the 2026-09-17 image built from this lock by the repository’s image workflow. Its digest and the dependency checksum below were updated together, so the image dependency checksum guard passes. Date tags are re-pushed by later rebuilds on the same day, so always pin the digest, not the tag alone.

The Dockerfile pins Python 3.12.12 and the CUDA base digest used by that published image. Lock generation and future builds use uv 0.12.5, independently of the host project’s uv version. The sandbox project also pins setuptools for KernelBench’s isolated source build. OS packages installed with apt are not locked, so rebuilding the Dockerfile is not guaranteed to produce an identical image. Evaluation runs use the published image digest.

When deliberately updating the sandbox environment, edit sandbox/pyproject.toml as needed and update selected packages from the repository root:

uvx --from uv==0.12.5 uv lock --directory src/inspect_evals/kernelbench/sandbox --upgrade-package PACKAGE
uvx --from uv==0.12.5 uv lock --check --directory src/inspect_evals/kernelbench/sandbox

Replace PACKAGE with the package to update. Use --upgrade only when intending to update the whole sandbox. CI and pre-commit check freshness without changing the lockfile. Keep the sandbox’s uv version pin aligned in its manifest, Dockerfile, CI check, pre-commit hook, and these commands.

Review the lock diff, build the image, run the sandbox check below against it on a GPU, and publish the image. Then update the digest in DEFAULT_SANDBOX_IMAGE and compose.yaml, and set DEFAULT_SANDBOX_DEPENDENCIES_SHA256 to the SHA-256 of the manifest bytes followed by the lockfile bytes:

cat src/inspect_evals/kernelbench/sandbox/pyproject.toml src/inspect_evals/kernelbench/sandbox/uv.lock | shasum -a 256

The checksum test requires acknowledgement of sandbox dependency changes independently of the host lock; it does not verify an image’s contents. Assess task versioning when the measurement environment changes. Host dependency changes still need compatibility testing with the pinned sandbox.

Checking a sandbox image on a GPU

kernelbench_sandbox_check is a maintenance task, not a benchmark. It runs the eval’s own scorer over three fixture kernels with known verdicts inside the sandbox: a passing kernel for the chosen backend (load_inline CUDA or Triton), a kernel that compiles but returns wrong output, and a source that is not Python. A sample passes only if the scorer’s verdict matches, a passing kernel gets finite positive timings, and the sandbox’s installed packages, Python version and manifest-plus-lock checksum match the committed sandbox project. Accuracy 1.0 means the image is fit for kernelbench. No model is called. The task is built on the shared inspect_evals.utils.sandbox_check helper; the KernelBench-specific parts are the fixture kernels, the verdict predicates and the GPU inventory checks.

Run it against the pinned image, or a candidate via sandbox_image, on a host with an NVIDIA container runtime:

uv run inspect eval inspect_evals/kernelbench_sandbox_check --model mockllm/model
uv run inspect eval inspect_evals/kernelbench_sandbox_check --model mockllm/model -T sandbox_image=ghcr.io/generality-labs/inspect-eval-kernelbench:2026-09-17@sha256:...
uv run inspect eval inspect_evals/kernelbench_sandbox_check --model mockllm/model -T backend=triton -T precision=fp16 -T timing_method=do_bench

tests/kernelbench/test_sandbox_check.py covers the same task. Its unit tests run in CI without a GPU. The GPU run is marked gpu and slow, so it is skipped unless both gates are opened; it sweeps both backends, fp32 and fp16, and all four timing methods:

uv run --no-sync tox -e kernelbench -- --runslow --rungpu -k test_pinned_sandbox_image

Set KERNELBENCH_SANDBOX_TYPE=k8s on a runner with cluster access (the default is docker), and KERNELBENCH_SANDBOX_IMAGE to check an image other than the pinned default. On Hawk, express the same sweep as an eval set; each item runs the task on a GPU node with zero LLM spend:

name: kernelbench-sandbox-check
tasks:
  - package: inspect_evals @ git+https://github.com/UKGovernmentBEIS/inspect_evals@main
    name: inspect_evals
    items:
      - name: kernelbench_sandbox_check
        args: { sandbox_type: k8s, backend: cuda, precision: fp32, timing_method: cuda_event }
      - name: kernelbench_sandbox_check
        args: { sandbox_type: k8s, backend: triton, precision: fp16, timing_method: do_bench }
packages:
  - inspect-evals-kernelbench @ git+https://github.com/UKGovernmentBEIS/inspect_evals@main#subdirectory=packages/kernelbench
models:
  - package: inspect-test-utils
    name: hardcoded
    items:
      - name: hardcoded
        args: { answer: unused }
retry_attempts: 0

Python compatibility

The pinned upstream KernelBench commit declares requires-python = "==3.10.*". uv ignores dependency Python upper bounds during resolution, so this metadata does not prevent locking or installing it on Python 3.12. No metadata override or source patch is applied. The published image already uses Python 3.12.12, and our uploaded runner imports typing.NotRequired, which requires Python 3.11 or newer.

The earlier 182-package migration passed fresh dependency resolution, runner import, 38 upstream CPU-only tests covering static checks, code extraction and score helpers, and 90 Inspect KernelBench tests on Python 3.12.10, with three dataset-dependent tests skipped. A complete Linux amd64 Docker build passed on Python 3.12.12; runner startup, KernelBench imports, and CPU tensor execution passed inside that image with networking disabled. These checks establish compatibility for the exercised Python paths; CUDA compilation and GPU timing still require validation on a GPU when updating the image. See uv’s Python resolution behaviour.

A fresh build of the reduced image passed on Linux amd64/Python 3.12.12. Its installed inventory matches all 148 locked versions, and host-only packages including inspect-ai are absent. With networking disabled, the uploaded runner started, CPU tensor operations passed for fp32/fp16/bf16, all four supported timing functions resolved, and Ninja and the CUDA compiler were available. The same build was then exercised on a Tesla T4 (CUDA driver 580, Torch 2.9.0+cu128) through the unchanged runner with networking disabled: a load_inline CUDA kernel and a Triton kernel compiled and passed five correctness trials in fp32 and fp16, all four timing methods returned finite runtimes for both the kernel and the reference, an incorrect kernel was reported as correctness=False, and a non-Python source as compiled=False, with no error payloads. The same probe was then repeated against the published 2026-09-17 digest pinned below, on a T4 node in Kubernetes, with identical verdicts and an identical installed inventory.

Running on Kubernetes

Install the Kubernetes sandbox provider, then select it when running the eval:

uv pip install inspect_k8s_sandbox

uv run inspect eval inspect_evals/kernelbench -T sandbox_type=k8s

This uses the pinned public linux/amd64 image published by the repository’s image workflow to the Generality Labs GHCR namespace and renders the packaged values.yaml, which requests one nvidia.com/gpu and tolerates the conventional nvidia.com/gpu GPU-node taint. Use -T sandbox_image=<registry-ref> to pull from a mirror or a custom image. An earlier revision completed two samples end to end on EKS with genuine T4 timing scores. The pinned 2026-09-17 digest has since run the current runner on an EKS T4 node through the k8s sandbox provider, covering CUDA and Triton compilation, correctness trials and all timing methods, but the full Inspect task has not been rerun on it end to end, so supports_k8s remains unset pending that check.

Each concurrent Kubernetes sample reserves an entire GPU for its lifetime. On node pools with fixed GPU counts, increasing concurrency can sharply increase the selected instance size and cost; start with --max-samples 1 and scale deliberately.

sandbox_node_selector is only needed when a cluster mixes GPU models and the prompt must name the selected model. On an autoscaled cluster, the selector value must match both the label declared by the autoscaler node-pool template and the label ultimately applied to the running node. Those values can differ: selecting only the observed node label can prevent scale-up, while selecting only the template label can provision nodes that are immediately relabelled out from under the pod. Check both the node-pool configuration and running-node labels (for example, kubectl get nodes -L nvidia.com/gpu.product); if they disagree, fix the node pool rather than adding a selector here.

After confirming a consistent label, pass it alongside the prompt’s hardware parameters:

uv run inspect eval inspect_evals/kernelbench -T sandbox_type=k8s -T sandbox_node_selector='{"nvidia.com/gpu.product":"Tesla-T4"}' -T gpu_name=T4 -T include_hardware=true

For other Kubernetes customisation, copy values.yaml, edit it, and pass either its runner-local path or its YAML content through sandbox_config. Both block-style and single-line flow-style YAML mappings are accepted. sandbox_image and sandbox_node_selector override the corresponding fields in a custom Kubernetes config. Passing a file path requires that file to exist on the process running Inspect; inline YAML is portable to remote runners. For non-k8s providers, sandbox_config remains a runner-local file path and cannot be combined with those k8s overrides.

Docker and Kubernetes use the same digest-pinned image by default, keeping their dependency environment aligned. sandbox_image can override that default for either provider. Because KernelBench scores runtime speedup, treat results from different image digests or CPU/GPU hardware as different measurement environments. The resolved image is recorded in task metadata.

The scorer uploads the current checkout’s eval_runner.py before each sample and executes it with the sandbox environment’s python. This keeps the scorer and runner CLI in sync with the host eval code while the image pins KernelBench, PyTorch, and the remaining runtime dependencies.

Installing on remote runners

Remote runners can install the eval’s isolated package directly from git:

pip install "git+https://github.com/UKGovernmentBEIS/inspect_evals@<ref>#subdirectory=packages/kernelbench"

Installed this way, the package’s inspect_evals dependency resolves from PyPI at the latest release, which can skew from the eval code at <ref>. To keep them in lockstep, install inspect_evals from the same ref in the same command:

pip install "inspect_evals @ git+https://github.com/UKGovernmentBEIS/inspect_evals@<ref>" "git+https://github.com/UKGovernmentBEIS/inspect_evals@<ref>#subdirectory=packages/kernelbench"

Dataset

copied from: https://github.com/ScalingIntelligence/KernelBench

KernelBench is a comprehensive benchmark designed to evaluate LLMs’ ability to generate efficient GPU kernels. Each task provides a PyTorch operator implementation as a reference and asks the model to transpile or generate a corresponding GPU kernel (in CUDA or related DSLs). Tasks are organized into several levels of increasing complexity:

Level 1: Single-kernel operators (e.g., matrix multiplications, convolutions, layer normalization).
Level 2: Simple fusion patterns (e.g., combining operations like conv+bias+ReLU).
Level 3: Full model architectures (e.g., MobileNet, VGG, MiniGPT, Mamba).
Level 4: Hugging Face model-level optimizations.

The problems can be sourced either from a local dataset or the Hugging Face KernelBench dataset.

Scoring

copied from: https://github.com/ScalingIntelligence/KernelBench

KernelBench evaluates two core aspects of model-generated kernels:

Correctness: Each generated kernel is run against randomized test cases and checked for consistency with the corresponding PyTorch reference implementation.
Performance (Speedup): The runtime of the generated kernel is compared to the PyTorch baseline over multiple trials to compute the speedup.

The main benchmark metric is fast_p: the fraction of tasks for which the model output is both correct and achieves a speedup greater than a threshold p. For example:

fast_1 measures the fraction of tasks where the kernel is correct and faster than PyTorch.
fast_2 measures the fraction of tasks where the kernel is correct and at least 2x faster than PyTorch.
fast_0 is simply the overall correctness rate.

Notes

KernelBench uses temperature 0.0 as a default value, however this can be overridden using the available -T temperature parameter

There is a fixed set of GPUs that this can be run on while using the original prompts. You can track the development of fix in this issue.

Due to dependencies (uvloop) in the original repository that we depend on, this eval is not compatible on Windows.

Evaluation Report

Following the pattern from the original paper and METRs write-up, the evaluation tested on Levels 1 - 3. The explanation from METR:

As with the original KernelBench work, we did not evaluate our models on level 4. Level 4 of KernelBench consists of 20 tasks that require models to directly optimize the performance of 20 medium-sized 2019-2022 language models. This is mainly because the KernelBench PyTorch implementation for Level 4 does not provide information about the internals of the models but instead loads the model using the HuggingFace transformer package – fairly evaluating frontier model performance in writing these agents would require providing agent scaffolding that efficiently parses the HuggingFace codebase, just to provide the agents with the information required to perform the task.

Inspect evals version 1.0.0 of Kernelbench used (this implementation is Kernelbench 0.1)

Model	Implementation	Level 1 Fast_1 (n=100)	Level 1 Correct	Level 2 Fast_1 (n=100)	Level 2 Correct	Level 3 Fast_1 (n=50)	Level 3 Correct
gpt-4o	Original (temp=0)	0.04	30.0%	0.07	8.0%	0.04	8.0%
gpt-4o	inspect (temp=0)	0.07	30.0%	0.06	7.0%	0.02	6.0%

Changelog

[6-C] - 2026-08-27

Added k8s sandbox support (#2116) with a built-in values template, pinned public image, custom or inline sandbox_config, and sandbox_image and sandbox_node_selector overrides. An earlier revision completed an end-to-end EKS run; verification of the final image and runner path remains pending.
Fixed Docker and k8s scoring to upload the current eval runner and execute it with the sandbox environment’s Python. Previously, every sample could fail with AttributeError: 'NoneType' object has no attribute 'get_torch_dtype_from_string' because uv run selected an interpreter without the KernelBench dependencies.
Changed the Docker default from an eight-CPU local image build to the same four-CPU, digest-pinned linux/amd64 image used by k8s, making the dependency environment reproducible across both providers. Added compose.local.yaml for Dockerfile development and compatible non-amd64 hosts, and record the resolved image in task metadata.

[5-B] - 2026-05-17

Reworked the scorer to distinguish infrastructure failures from verdicts on the generated kernel. See #1679.
- Outcomes attributable to the model’s kernel are scored: static-check failure, no code block, segfault, timeout, and any verdict encoded in output.json (including compiled=False / correctness=False). A timeout is a documented judgment call: it is ambiguous (hanging kernel vs. slow compile under contention), but upstream KernelBench treats it as a kernel failure.
- Failures of the evaluation machinery raise instead, so the sample is marked ERRORED and stays eligible for --retry-on-error: eval_runner exiting non-zero without a segfault, unparseable output.json (the runner was killed mid-write), and "error" payloads, which mean no verdict was reached (an exception escaped eval_kernel_against_ref, or compilation lock retries were exhausted). Previously these paths were scored 0.0, conflating them with kernel failures.
- eval_runner now tags error payloads with error_kind ("exception" with a traceback, or "retry_exhausted") so the failure cause is visible in the error message and log.

[4-B] - 2026-04-13

Pin HuggingFace dataset to revision ca1464e5.

[3-B] - 2026-03-02

Scorer refactored to execute kernel evaluation inside a Docker sandbox instead of via subprocess on the host. Results are not directly comparable with v2-A due to the changed execution environment.
Added sandbox_type task parameter.

[2-A] - 2026-02-16

Migrate version to new scheme. See #907.

Static checks

Results of inspect-evals-lint at the commit this documentation was built from. Rule names link to their documentation; findings link to the file at that commit. These checks describe structure and conventions, not whether the evaluation measures what it claims.

Loading lint results…