kernels/pr_751/en/llms-full.txt

# Kernels

## Docs

- [Environment variables](https://huggingface.co/docs/kernels/pr_751/env.md)
- [Why kernels?](https://huggingface.co/docs/kernels/pr_751/why_kernels.md)
- [Blog posts](https://huggingface.co/docs/kernels/pr_751/blog.md)
- [Installation](https://huggingface.co/docs/kernels/pr_751/installation.md)
- [FAQ](https://huggingface.co/docs/kernels/pr_751/faq.md)
- [kernels versions](https://huggingface.co/docs/kernels/pr_751/cli-versions.md)
- [kernels info](https://huggingface.co/docs/kernels/pr_751/cli-info.md)
- [Quickstart](https://huggingface.co/docs/kernels/pr_751/basic-usage.md)
- [Projects using kernels](https://huggingface.co/docs/kernels/pr_751/integrating-kernels.md)
- [kernels benchmark](https://huggingface.co/docs/kernels/pr_751/cli-benchmark.md)
- [CLI reference for kernel-builder](https://huggingface.co/docs/kernels/pr_751/builder-cli.md)
- [Install agent skills](https://huggingface.co/docs/kernels/pr_751/cli-skills.md)
- [Kernel requirements](https://huggingface.co/docs/kernels/pr_751/kernel-requirements.md)
- [Migrate from older versions](https://huggingface.co/docs/kernels/pr_751/migration.md)
- [Kernels](https://huggingface.co/docs/kernels/pr_751/index.md)
- [Overview](https://huggingface.co/docs/kernels/pr_751/cli.md)
- [Talks](https://huggingface.co/docs/kernels/pr_751/talks.md)
- [kernels lock](https://huggingface.co/docs/kernels/pr_751/cli-lock.md)
- [Use layers](https://huggingface.co/docs/kernels/pr_751/layers.md)
- [Lock kernel versions](https://huggingface.co/docs/kernels/pr_751/locking.md)
- [kernels download](https://huggingface.co/docs/kernels/pr_751/cli-download.md)
- [kernels verify-signature](https://huggingface.co/docs/kernels/pr_751/cli-verify-signature.md)
- [Why Nix?](https://huggingface.co/docs/kernels/pr_751/builder/why-nix.md)
- [Write kernels](https://huggingface.co/docs/kernels/pr_751/builder/writing-kernels.md)
- [Develop kernels with agents](https://huggingface.co/docs/kernels/pr_751/builder/agents-guide.md)
- [Develop locally](https://huggingface.co/docs/kernels/pr_751/builder/local-dev.md)
- [Building and testing kernels with GitHub Actions](https://huggingface.co/docs/kernels/pr_751/builder/github-actions.md)
- [Build with Nix](https://huggingface.co/docs/kernels/pr_751/builder/build.md)
- [Set up for Metal kernels](https://huggingface.co/docs/kernels/pr_751/builder/metal.md)
- [Nix builder design](https://huggingface.co/docs/kernels/pr_751/builder/design-nix-builder.md)
- [Set up your IDE](https://huggingface.co/docs/kernels/pr_751/builder/ide-setup.md)
- [Build variants](https://huggingface.co/docs/kernels/pr_751/builder/build-variants.md)
- [Architecture overview](https://huggingface.co/docs/kernels/pr_751/builder/design-overview.md)
- [Secure your kernels](https://huggingface.co/docs/kernels/pr_751/builder/security.md)
- [Kernels API Reference](https://huggingface.co/docs/kernels/pr_751/api/kernels.md)
- [Layers API Reference](https://huggingface.co/docs/kernels/pr_751/api/layers.md)

### Environment variables
https://huggingface.co/docs/kernels/pr_751/env.md

# Environment variables

## `KERNELS_CACHE`

The directory to use as the local kernel cache. If not set, the cache
of the `huggingface_hub` package is used.

## `DISABLE_KERNEL_MAPPING`

Disables kernel mappings for [`layers`](layers).

### Why kernels?
https://huggingface.co/docs/kernels/pr_751/why_kernels.md

# Why kernels?

Our goal with the `kernels` package is manifold:

* Establish a standard way of structuring and building kernels. The
`builder` component takes kernel source in a pre-defined layout, with
a declarative build configuration, and produces compiled kernels for
a wide matrix of compute backends (e.g., CUDA, ROCM, and XPU), operating
systems, and architectures. The builder also enforces the reproducibility of
the build environment and build steps.
* Provide a standard way of distributing and loading kernels. Kernels
are distributed through the Hugging Face Hub and can be loaded through
the kernels Python package. `kernels` fetches the right kernel build for
the system that it runs, avoiding long local, sometimes hours-long,
builds. It also supports loading of multiple versions of the same
kernel, effectively eliminating "dependency hell".
* Provide kernel builds across all supported PyTorch versions,
accelerators, and capabilities. This is particularly important
because local kernel builds can become unusable once the base
machine learning framework (e.g., PyTorch) is updated to a recent version.

Additionally, there are several advantages to hosting the pre-built
kernels on the Hugging Face Hub platform:

* Trends -- kinds of models using a kernel more than others,
[applications](https://huggingface.co/spaces) using a particular kernel, etc.
* Users immediately know if their hardware supports a kernel
without having to run any installation locally.
* General features of the Hub platform:
    * Seamless versioning
    * Fast download and upload powered by [XET](https://huggingface.co/docs/hub/xet/index)
    * Visibility into the download stats

### Blog posts
https://huggingface.co/docs/kernels/pr_751/blog.md

# Blog posts

This page lists blog posts on the Kernels project. See the full list
tagged [`kernels`](https://huggingface.co/blog?tag=kernels) on the
Hugging Face blog:

* [🤗 Kernels: Major Updates](https://huggingface.co/blog/revamped-kernels)
* [Custom Kernels for All from Codex and Claude](https://huggingface.co/blog/custom-cuda-kernels-agent-skills)
* [Easily Build and Share ROCm Kernels with Hugging Face](https://huggingface.co/blog/build-rocm-kernels)
* [From Zero to GPU: A Guide to Building and Scaling Production-Ready CUDA Kernels](https://huggingface.co/blog/kernel-builder)
* [Creating custom kernels for the AMD MI300](https://huggingface.co/blog/mi300kernels)
* [Tricks from OpenAI gpt-oss YOU 🫵 can use with transformers](https://huggingface.co/blog/faster-transformers)
* [Mixture of Experts (MoEs) in Transformers](https://huggingface.co/blog/moe-transformers)

This page will be kept updated.

### Installation
https://huggingface.co/docs/kernels/pr_751/installation.md

# Installation

> [!WARNING]
> `kernels` has not reached `1.0` yet. Until then, minor releases may contain
> breaking changes. If you depend on `kernels` in a library or application, we
> **strongly recommend pinning a version range** rather than an unbounded
> dependency. For example, in `pyproject.toml`:
>
> ```toml
> dependencies = [
>     "kernels>=0.15,<0.16",
> ]
> ```
>
> or equivalently `kernels~=0.15` (compatible release). This protects your
> project from unexpected breakage when a new `kernels` version is released.

Install the `kernels` package with `pip` (requires `torch>=2.5` and CUDA):

```bash
pip install kernels
```

or with `uv`

```bash
uv pip install kernels
```

or if you want the latest version from the `main` branch:

```bash
pip install "kernels[benchmark] @ git+https://github.com/huggingface/kernels#subdirectory=kernels"
```

## Curated installations

Some kernels rely on additional packages at runtime (for example,
[CUTLASS DSL](https://docs.nvidia.com/cutlass/), [einops](https://einops.rocks/),
and [Apache TVM FFI](https://github.com/apache/tvm-ffi)). The `curated` extra
installs these commonly-needed dependencies in one go:

```bash
pip install "kernels[curated]"
```

On XPU (Intel GPU) platforms, use the `curated-xpu` extra instead, which omits
the CUDA-only dependencies:

```bash
pip install "kernels[curated-xpu]"
```

> [!IMPORTANT]
> On Windows, we recommend using the Linux version of Torch through
> [WSL 2](https://learn.microsoft.com/en-us/windows/wsl/install), since
> many more kernels support Linux. If you want to use GPU acceleration,
> check out the [CUDA on WSL](https://docs.nvidia.com/cuda/wsl-user-guide/index.html#getting-started-with-cuda-on-wsl-2)
> and [PyTorch with DirectML on WSL 2](https://learn.microsoft.com/en-us/windows/ai/directml/pytorch-wsl)
> guides.

> [!IMPORTANT]
> We strongly recommend not using a free-threaded Python build yet.
> These builds are not only experimental, but do not support the stable ABI
> on Python versions before 3.15. Kernels are compiled with the stable ABI
> to support a wide range of Python versions.

### FAQ
https://huggingface.co/docs/kernels/pr_751/faq.md

# FAQ

## Kernel layers

### Why is the kernelization step needed as a separate step?

In earlier versions of `kernels`, a layer's `forward` method was replaced
by [use_kernel_forward_from_hub()](/docs/kernels/pr_751/en/api/layers#kernels.use_kernel_forward_from_hub) and [replace_kernel_forward_from_hub()](/docs/kernels/pr_751/en/api/layers#kernels.replace_kernel_forward_from_hub).
The new `forward` would dispatch to a kernel based on the device type,
whether a model was training, etc. However, this approach was
fundamentally incompatible with `torch.compile` since it relied
on data-dependent branching.

To avoid branching, we have to make dispatch decisions ahead of time,
which is what the [kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize) function does.

### Why does kernelization only replace `forward` methods?

There are some other possible approaches. The first is to completely
replace existing layers by kernel layers. However, since this would
permit free-form layer classes, it would be much harder to validate
that layers are fully compatible with the layers that they are
replacing. For instance, they could have completely different member
variables. Besides that, we would also need to hold on to the original
layers, in case we need to revert to the base layers when the model
is [kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize)d again with different options.

A second approach would be to make an auxiliary layer that wraps the
original layer and the kernel layer and dispatches to the kernel layer.
This wouldn't have the issues of the first approach, because kernel layers
could be similarly strict as they are now, and we would still have access
to the original layers when [kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize)-ing the model again. However,
this would change the graph structure of the model and would break use
cases where programs access the model internals (e.g.
`model.layers[0].attention.query_weight`) or rely on the graph structure
in other ways.

The approach of `forward`-replacement is the least invasive, because
it preserves the original model graph. It is also reversible, since
even though the `forward` of a layer _instance_ might be replaced,
the corresponding class still has the original `forward`.

## Misc

### How can I disable kernel reporting in the user-agent?

By default, we collect telemetry when a call to [get_kernel()](/docs/kernels/pr_751/en/api/kernels#kernels.get_kernel) is made.
This only includes the `kernels` version, `torch` version, and the build
information for the kernel being requested.

You can disable this by setting `export DISABLE_TELEMETRY=yes`.

### kernels versions
https://huggingface.co/docs/kernels/pr_751/cli-versions.md

# kernels versions

Use `kernels versions` to list all available versions of a kernel on the Hub
and marks compatible versions.

## Usage

```bash
kernels versions <repo_id>
```

## Examples

List versions of a kernel:

```bash
kernels versions kernels-community/activation
```

## Example Output

```text
Version 1: torch210-metal-aarch64-darwin, torch28-cxx11-cu126-aarch64-linux, torch28-cxx11-cu129-aarch64-linux, torch28-cxx11-cu128-aarch64-linux, torch29-cxx11-cu130-x86_64-linux, torch27-cxx11-cu118-x86_64-linux, torch210-cxx11-cu130-x86_64-linux, torch29-cxx11-cu128-aarch64-linux, torch29-cxx11-cu130-aarch64-linux, torch27-cxx11-cu126-x86_64-linux, ✅ torch29-cxx11-cu126-x86_64-linux (compatible), torch27-cxx11-cu128-x86_64-linux, torch210-cxx11-cu126-x86_64-linux, torch29-metal-aarch64-darwin, torch27-cxx11-cu128-aarch64-linux, torch210-cu128-x86_64-windows, torch28-cxx11-cu128-x86_64-linux, torch28-cxx11-cu126-x86_64-linux, torch210-cxx11-cu128-x86_64-linux, torch29-cxx11-cu126-aarch64-linux, ✅ torch29-cxx11-cu128-x86_64-linux (preferred), torch28-cxx11-cu129-x86_64-linux
```

## Use Cases

- Check which versions are available before locking dependencies
- Find the latest version of a kernel
- Identify version SHAs for pinning in `pyproject.toml`

## See Also

- [kernels lock](cli-lock) - Lock kernel versions in your project
- [kernels download](cli-download) - Download locked kernels

### kernels info
https://huggingface.co/docs/kernels/pr_751/cli-info.md

# kernels info

Use `kernels info` to describe a kernel from a Hub repo ID or a local path.
It prints the kernel metadata (name, version, license, upstream/source
repositories, Python dependencies, and supported backends). To list the
build variants of a kernel and their compatibility with the current system,
use [kernels versions](cli-versions).

## Usage

```bash
kernels info <repo_id_or_path>
```

When the argument is an existing local directory, it is treated as a kernel
repository with build variants in its `build/` directory. Otherwise it is
treated as a Hub repo ID.

Options:

- `--revision <revision>`: describe a specific branch, tag, or commit
  (Hub only). Cannot be used together with `--version`.
- `--version <version>`: describe a specific kernel version (Hub only).
  Cannot be used together with `--revision`. When neither is given, the
  latest version is used (or `main` for unversioned repositories).
- `--json`: print the kernel information as JSON for machine-readable output.

## Examples

Describe a kernel on the Hub:

```bash
kernels info kernels-community/activation
```

Describe a local kernel repository (e.g. after `kernel-builder build-and-copy`):

```bash
kernels info ./relu
```

Get machine-readable output:

```bash
kernels info --json kernels-community/activation
```

## Example Output

```text
Repository: kernels-community/activation
Revision: v1
Name: activation
Version: 1
License: Apache-2.0
Upstream: -
Source: -
Python dependencies: -
Backends: cuda, metal
```

## See Also

- [kernels versions](cli-versions) - List available versions of a kernel
- [kernels download](cli-download) - Download locked kernels

### Quickstart
https://huggingface.co/docs/kernels/pr_751/basic-usage.md

# Quickstart

## Loading Kernels

Here is how you would use the [activation](https://huggingface.co/kernels-community/activation) kernels from the Hugging Face Hub:

```python
import torch
from kernels import get_kernel

# Download optimized kernels from the Hugging Face hub
activation = get_kernel("kernels-community/activation", version=1)

# Create a random tensor
x = torch.randn((10, 10), dtype=torch.float16, device="cuda")

# Run the kernel
y = torch.empty_like(x)
activation.gelu_fast(y, x)

print(y)
```

This fetches version `1` of the kernel `kernels-community/activation`.
Kernels are versioned using a major version number. Using `version=1` will
get the latest kernel build from the `v1` branch.

Kernels within a version branch must never break the API or remove builds
for older PyTorch versions. This ensures that your code will continue to work.

Hub kernels must be loaded with either a `version` or an explicit `revision`.

> [!NOTE]
> Version `0` kernels are excluded from the API compatibility requirement,
> since it is used for alpha/beta-quality kernels that may still have
> rapidly changing APIs.

## Checking Kernel Availability

You can check if a particular version of a kernel supports the environment
that the program is running on:

```python
from kernels import has_kernel

# Check if kernel is available for current environment
is_available = has_kernel("kernels-community/activation", version=1)
print(f"Kernel available: {is_available}")
```

When no compatible kernel is found, [has_kernel()](/docs/kernels/pr_751/en/api/kernels#kernels.has_kernel) does not say _why_.
[get_kernel_variants()](/docs/kernels/pr_751/en/api/kernels#kernels.get_kernel_variants) returns the full resolution trace instead: one
decision per build variant in the repository, with compatible variants listed
first. Each decision is a `VariantAccepted` or a `VariantRejected`, and rejected
variants carry a human-readable `reason`:

```python
from kernels import get_kernel_variants, VariantAccepted

for decision in get_kernel_variants("kernels-community/activation", version=1):
    name = decision.variant.variant_str
    if isinstance(decision, VariantAccepted):
        print(f"{name}: compatible")
    else:
        print(f"{name}: rejected ({decision.reason})")
```

## Inspecting Loaded Kernels

[get_loaded_kernels()](/docs/kernels/pr_751/en/api/kernels#kernels.get_loaded_kernels) returns a snapshot of every kernel that has been loaded
into the current process. Each entry is a [LoadedKernel](/docs/kernels/pr_751/en/api/kernels#kernels.LoadedKernel) namedtuple with the
imported `module`, the `package_name`, and `repo_infos` (repo id, resolved
revision, and the backend argument that was passed).

```python
from kernels import get_kernel, get_loaded_kernels

get_kernel("kernels-community/activation", version=1)

for loaded in get_loaded_kernels():
    print(loaded.package_name, loaded.repo_infos)
```

`repo_infos` is populated only for kernels loaded with [get_kernel()](/docs/kernels/pr_751/en/api/kernels#kernels.get_kernel). Kernels
loaded from a local path ([get_local_kernel()](/docs/kernels/pr_751/en/api/kernels#kernels.get_local_kernel)) or via a lockfile
([get_locked_kernel()](/docs/kernels/pr_751/en/api/kernels#kernels.get_locked_kernel), [load_kernel()](/docs/kernels/pr_751/en/api/kernels#kernels.load_kernel)) have `repo_infos=None`.

Browse through different kernels compatible with `kernels` from [here](https://huggingface.co/kernels).

A kernel can provide layers in addition to kernel functions. Refer to [Layers](./layers) to know more.

### Projects using kernels
https://huggingface.co/docs/kernels/pr_751/integrating-kernels.md

# Projects using kernels

This page shows how different projects use `kernels`.

## autoresearch

[karpathy/autoresearch](https://github.com/karpathy/autoresearch) [uses](https://github.com/karpathy/autoresearch/blob/c2450add72cc80317be1fe8111974b892da10944/train.py#L23) `kernels` to
integrate Flash-Attention 3 through the [get_kernel()](/docs/kernels/pr_751/en/api/kernels#kernels.get_kernel) method.

## AReaL

[inclusionAI/AReaL](https://github.com/inclusionAI/AReaL) uses `kernels` in an opt-in manner to integrate
optimized attention mechanisms.

## transformers

[huggingface/transformers](https://github.com/huggingface/transformers/) primarily
depends on `kernels` for all optimizations related to optimized kernels, including
optimized attention implementations, MoE blocks, and quantization. Besides
[get_kernel()](/docs/kernels/pr_751/en/api/kernels#kernels.get_kernel), it also uses [kernel layers](./layers) to optimize the forward passes
of common layers involved in the modeling blocks. Some references are available
[here]()
and [here](https://github.com/search?q=repo%3Ahuggingface%2Ftransformers+use_kernel_forward_from_hub&type=code).

Refer to the following posts to know more:

* [Tricks from OpenAI gpt-oss YOU 🫵 can use with transformers](https://huggingface.co/blog/faster-transformers)
* [Mixture of Experts (MoEs) in Transformers](https://huggingface.co/blog/moe-transformers)

## diffusers

Similar to `transformers`, [huggingface/diffusers](https://github.com/huggingface/diffusers/) uses
`kernels` for integrating optimized kernels to [compute attention](https://github.com/huggingface/diffusers/blob/e5aa719241f9b74d6700be3320a777799bfab70a/src/diffusers/models/attention_dispatch.py).

Besides leveraging pre-built compute kernels, different projects
rely on `kernels` to also package, build, and distribute their
kernels on the Hugging Face Hub platform. This is made possible by the
["builder" component of `kernels`](./builder/writing-kernels).
Visit [huggingface.co/kernels](https://huggingface.co/kernels) to browse
the pre-built compute kernels available on the Hub.

Feel free to open a PR enlisting your project to show how `kernels`
is leveraged there.

### kernels benchmark
https://huggingface.co/docs/kernels/pr_751/cli-benchmark.md

# kernels benchmark

Use `kernels benchmark` to run benchmark scripts shipped with a kernel repository.

The command:

- Downloads the kernel repo at a specific **branch** or **version**
- Runs all `benchmarks/benchmark*.py` scripts
- Times each `benchmark_*` workload and prints a results table
- Optionally saves results as JSON

## Installation

`kernels benchmark` requires extra dependencies:

```bash
uv pip install 'kernels[benchmark]' # or pip install 'kernels[benchmark]'
```

## Example

```bash
kernels benchmark kernels-community/activation --version 1
```

Example output:

```text
Downloading kernels-community/activation@v1...
Running benchmark.py...

  GPU      Apple M3 Max (30 cores)
  CPU      Apple M3 Max
  OS       Darwin 25.2.0
  PyTorch  2.10.0

  Running SiluWorkloads on mps

┌───────────────┬────────────┬─────┬───────────┬────────────┬───────────┬───────────┬───────────┬───────────┬────────────┬───────────┬─────────┐
│ Benchmark     │ Workload   │   N │ Speedup   │   Mean(ms) │   Std(ms) │   Min(ms) │   Max(ms) │   IQR(ms) │   Outliers │   Ref(ms) │ Match   │
├───────────────┼────────────┼─────┼───────────┼────────────┼───────────┼───────────┼───────────┼───────────┼────────────┼───────────┼─────────┤
│ SiluWorkloads │ large      │ 100 │ 1.72x     │     6.5153 │    0.4343 │    6.2883 │    8.4699 │    0.1701 │          8 │   11.2048 │ ✓       │
│ SiluWorkloads │ medium     │ 100 │ 2.48x     │     1.1813 │    0.3976 │    1.04   │    4.2146 │    0.0698 │          5 │    2.9332 │ ✓       │
│ SiluWorkloads │ small      │ 100 │ 1.96x     │     0.4909 │    0.2175 │    0.4407 │    2.6438 │    0.0085 │         16 │    0.9622 │ ✓       │
└───────────────┴────────────┴─────┴───────────┴────────────┴───────────┴───────────┴───────────┴───────────┴────────────┴───────────┴─────────┘

  large: 1.72x faster (95% CI: 6.4302-6.6004ms vs ref 11.2048ms) ✓ significant
  medium: 2.48x faster (95% CI: 1.1034-1.2592ms vs ref 2.9332ms) ✓ significant
  small: 1.96x faster (95% CI: 0.4483-0.5335ms vs ref 0.9622ms) ✓ significant

Kernel: 2385e44  Benchmark: 5b53516
```

## Usage

You must specify which revision to benchmark, either via flags or with `@...` in the repo id:

```bash
kernels benchmark <repo_id> --version <N>
kernels benchmark <repo_id> --branch <name>
kernels benchmark <repo_id>@v<N>
kernels benchmark <repo_id>@<branch>
```

## Examples

Benchmark a tagged kernel version:

```bash
kernels benchmark kernels-community/activation --version 1
```

Equivalent shorthand:

```bash
kernels benchmark kernels-community/activation@v1
```

Benchmark a branch:

```bash
kernels benchmark kernels-community/activation --branch main
```

Tune warmup and iteration count:

```bash
kernels benchmark kernels-community/activation@v1 --warmup 20 --iterations 200
```

Save results to a file (JSON):

```bash
kernels benchmark kernels-community/activation@v1 --output results.json
```

Benchmark a local kernel checkout (must contain `benchmarks/`):

```bash
kernels benchmark ./my_kernel
```

## Output

- By default, a table is printed (timings in ms).
- `--output <file>.json` writes a JSON payload to disk.

## Writing Benchmark Scripts

Benchmark scripts must live under `benchmarks/` in the kernel repository and match `benchmark*.py`.
Each script should define one or more subclasses of `kernels.benchmark.Benchmark`.

Minimal example (`benchmarks/benchmark_activation.py`):

```python
import torch

from kernels.benchmark import Benchmark

class ActivationBenchmark(Benchmark):
    seed = 0

    def setup(self):
        self.x = torch.randn(128, 1024, device=self.device, dtype=torch.float16)
        self.out = torch.empty(128, 512, device=self.device, dtype=torch.float16)

    def benchmark_silu_and_mul(self):
        self.kernel.silu_and_mul(self.out, self.x)

    def verify_silu_and_mul(self):
        # Return reference tensor; runner compares with self.out
        return torch.nn.functional.silu(self.x[..., :512]) * self.x[..., 512:]
```

The runner will:

- Call `setup()` once per workload (or `setup_<workload>()` if present)
- Warm up (`--warmup`)
- Time `benchmark_<workload>()` for `--iterations`
- If `verify_<workload>()` exists, check that outputs match (`torch.allclose(..., atol=1e-2)`) and show a speedup vs the reference computation

## Troubleshooting

- If the repo does not contain a `benchmarks/` directory (or no `benchmark*.py` files), the command exits with an error.
- If a benchmark script defines no `Benchmark` subclasses, the command exits with an error.
- If `verify_<workload>()` exists and the outputs do not match, the command exits with an error.

### CLI reference for kernel-builder
https://huggingface.co/docs/kernels/pr_751/builder-cli.md

# CLI reference for kernel-builder

This document contains the help content for the `kernel-builder` command-line program.

**Command Overview:**

* [`kernel-builder`↴](#kernel-builder)
* [`kernel-builder completions`↴](#kernel-builder-completions)
* [`kernel-builder init`↴](#kernel-builder-init)
* [`kernel-builder build`↴](#kernel-builder-build)
* [`kernel-builder build-and-copy`↴](#kernel-builder-build-and-copy)
* [`kernel-builder build-and-upload`↴](#kernel-builder-build-and-upload)
* [`kernel-builder upload`↴](#kernel-builder-upload)
* [`kernel-builder check-config`↴](#kernel-builder-check-config)
* [`kernel-builder check-abi`↴](#kernel-builder-check-abi)
* [`kernel-builder check-builds`↴](#kernel-builder-check-builds)
* [`kernel-builder create-pyproject`↴](#kernel-builder-create-pyproject)
* [`kernel-builder devshell`↴](#kernel-builder-devshell)
* [`kernel-builder list-variants`↴](#kernel-builder-list-variants)
* [`kernel-builder hash`↴](#kernel-builder-hash)
* [`kernel-builder testshell`↴](#kernel-builder-testshell)
* [`kernel-builder update-build`↴](#kernel-builder-update-build)
* [`kernel-builder skills`↴](#kernel-builder-skills)
* [`kernel-builder skills add`↴](#kernel-builder-skills-add)
* [`kernel-builder clean-pyproject`↴](#kernel-builder-clean-pyproject)

## `kernel-builder`

Build Hugging Face Hub kernels

**Usage:** `kernel-builder <COMMAND>`

###### **Subcommands:**

* `completions` — Generate shell completions
* `init` — Initialize a new kernel project from template
* `build` — Build the kernel locally (alias for build-and-copy)
* `build-and-copy` — Build the kernel and copy artifacts locally
* `build-and-upload` — Build the kernel and upload to Hugging Face Hub
* `upload` — Upload kernel build artifacts to the Hugging Face Hub
* `check-config` — Validate the build.toml file
* `check-abi` — Check the ABI compatibility of a kernel extension
* `check-builds` — Validate kernel builds
* `create-pyproject` — Generate CMake files for a kernel extension build
* `devshell` — Spawn a kernel development shell
* `list-variants` — List build variants
* `hash` — Hash the builds of a kernel
* `testshell` — Spawn a kernel test shell
* `update-build` — Update a `build.toml` to the current format
* `skills` — Install skills for AI coding assistants (Claude, Codex, OpenCode)
* `clean-pyproject` — Clean generated artifacts

## `kernel-builder completions`

Generate shell completions

**Usage:** `kernel-builder completions <SHELL>`

###### **Arguments:**

* `<SHELL>`

  Possible values: `bash`, `elvish`, `fish`, `powershell`, `zsh`

## `kernel-builder init`

Initialize a new kernel project from template

**Usage:** `kernel-builder init [OPTIONS] [PATH]`

###### **Arguments:**

* `<PATH>` — Directory to initialize (defaults to current directory)

###### **Options:**

* `--license <LICENSE>` — The kernel's license

  Default value: `Apache-2.0`
* `--name <OWNER/REPO>` — Name of the kernel repo (e.g. `drbh/my-kernel`)
* `--backends <BACKENDS>` — Backends to enable (`all`, `cpu`, `cuda`, `metal`, `rocm`, `xpu`)

* `--overwrite` — Overwrite existing scaffold files (preserves other files)

## `kernel-builder build`

Build the kernel locally (alias for build-and-copy)

**Usage:** `kernel-builder build [OPTIONS] [KERNEL_DIR]`

###### **Arguments:**

* `<KERNEL_DIR>` — Directory of the kernel project (defaults to current directory)

###### **Options:**

* `--variant <VARIANT>` — Build a specific variant
* `--max-jobs <MAX_JOBS>` — Maximum number of parallel Nix build jobs
* `--cores <CORES>` — Number of CPU cores to use for each build job
* `-L`, `--print-build-logs` — Print full build logs on standard error

## `kernel-builder build-and-copy`

Build the kernel and copy artifacts locally

**Usage:** `kernel-builder build-and-copy [OPTIONS] [KERNEL_DIR]`

###### **Arguments:**

* `<KERNEL_DIR>` — Directory of the kernel project (defaults to current directory)

###### **Options:**

* `--max-jobs <MAX_JOBS>` — Maximum number of parallel Nix build jobs
* `--cores <CORES>` — Number of CPU cores to use for each build job
* `-L`, `--print-build-logs` — Print full build logs on standard error

## `kernel-builder build-and-upload`

Build the kernel and upload to Hugging Face Hub

**Usage:** `kernel-builder build-and-upload [OPTIONS] [KERNEL_DIR]`

###### **Arguments:**

* `<KERNEL_DIR>` — Directory of the kernel project (defaults to current directory)

###### **Options:**

* `--variant <VARIANT>` — Build a specific variant
* `--max-jobs <MAX_JOBS>` — Maximum number of parallel Nix build jobs
* `--cores <CORES>` — Number of CPU cores to use for each build job
* `-L`, `--print-build-logs` — Print full build logs on standard error
* `--repo-id <REPO_ID>` — Repository ID on the Hugging Face Hub (e.g. `user/my-kernel`)
* `--branch <BRANCH>` — Upload to a specific branch (defaults to `v{version}` from metadata)
* `--private` — Create the repository as private
* `--repo-type <REPO_TYPE>` — Repository type on Hugging Face Hub (`kernel` by default, or `model` for legacy repos)

  Default value: `kernel`

  Possible values: `model`, `kernel`

* `--create-pr` — Open a pull request with the changes rather than committing them directly (does not require write access to the repository)
* `--output-json <PATH>` — Write a machine-readable JSON summary of the upload to this path
* `-q`, `--quiet` — Suppress progress output

## `kernel-builder upload`

Upload kernel build artifacts to the Hugging Face Hub

**Usage:** `kernel-builder upload [OPTIONS] [KERNEL_DIR]`

###### **Arguments:**

* `<KERNEL_DIR>` — Directory of the kernel build (defaults to current directory)

###### **Options:**

* `--repo-id <REPO_ID>` — Repository ID on the Hugging Face Hub (e.g. `user/my-kernel`). Defaults to `general.hub.repo-id` from `build.toml`
* `--branch <BRANCH>` — Upload to a specific branch (defaults to `v{version}` from metadata)
* `--private` — Create the repository as private
* `--repo-type <REPO_TYPE>` — Repository type on Hugging Face Hub (`kernel` by default, or `model` for legacy repos)

  Default value: `kernel`

  Possible values: `model`, `kernel`

* `--create-pr` — Open a pull request with the changes rather than committing them directly (does not require write access to the repository)
* `--output-json <PATH>` — Write a machine-readable JSON summary of the upload to this path
* `-q`, `--quiet` — Suppress progress output

## `kernel-builder check-config`

Validate the build.toml file

**Usage:** `kernel-builder check-config [KERNEL_DIR]`

###### **Arguments:**

* `<KERNEL_DIR>`

## `kernel-builder check-abi`

Check the ABI compatibility of a kernel extension

**Usage:** `kernel-builder check-abi [OPTIONS] [KERNEL_DIR]`

###### **Arguments:**

* `<KERNEL_DIR>` — Directory with kernels

###### **Options:**

* `-m`, `--manylinux <VERSION>` — Manylinux version

  Default value: `manylinux_2_28`
* `--macos <VERSION>` — macOS version

  Default value: `15.0`
* `-p`, `--python-abi <VERSION>` — Python ABI version

  Default value: `3.9`
* `--torch-stable-abi <VERSION>` — Torch stable ABI version

## `kernel-builder check-builds`

Validate kernel builds

**Usage:** `kernel-builder check-builds [KERNEL_DIR]`

###### **Arguments:**

* `<KERNEL_DIR>`

## `kernel-builder create-pyproject`

Generate CMake files for a kernel extension build

**Usage:** `kernel-builder create-pyproject [OPTIONS] [KERNEL_DIR] [TARGET_DIR]`

###### **Arguments:**

* `<KERNEL_DIR>`
* `<TARGET_DIR>` — The directory to write the generated files to (directory of `BUILD_TOML` when absent)

###### **Options:**

* `-f`, `--force` — Force-overwrite existing files
* `--unique-id <UNIQUE_ID>` — This is an optional unique identifier that is suffixed to the kernel name to avoid name collisions. (e.g. Git SHA)
* `--kernel-sha <KERNEL_SHA>` — Full commit SHA of the kernel source, recorded in the build metadata. When absent, it is detected from the kernel's git repository (used by Nix builds where the source has no `.git`)
* `--kernel-dirty` — Mark the kernel source as having uncommitted changes in the build metadata. Only meaningful together with `--kernel-sha`

## `kernel-builder devshell`

Spawn a kernel development shell

**Usage:** `kernel-builder devshell [OPTIONS] [KERNEL_DIR]`

###### **Arguments:**

* `<KERNEL_DIR>`

###### **Options:**

* `--variant <VARIANT>` — Use a specific variant
* `--max-jobs <MAX_JOBS>` — Maximum number of parallel Nix build jobs
* `--cores <CORES>` — Number of CPU cores to use for each build job
* `-L`, `--print-build-logs` — Print full build logs on standard error

## `kernel-builder list-variants`

List build variants

**Usage:** `kernel-builder list-variants [OPTIONS] [KERNEL_DIR]`

###### **Arguments:**

* `<KERNEL_DIR>`

###### **Options:**

* `--arch` — Only list variants for the current architecture

## `kernel-builder hash`

Hash the builds of a kernel

**Usage:** `kernel-builder hash [KERNEL_DIR]`

###### **Arguments:**

* `<KERNEL_DIR>`

## `kernel-builder testshell`

Spawn a kernel test shell

**Usage:** `kernel-builder testshell [OPTIONS] [KERNEL_DIR]`

###### **Arguments:**

* `<KERNEL_DIR>`

###### **Options:**

* `--variant <VARIANT>` — Use a specific variant
* `--max-jobs <MAX_JOBS>` — Maximum number of parallel Nix build jobs
* `--cores <CORES>` — Number of CPU cores to use for each build job
* `-L`, `--print-build-logs` — Print full build logs on standard error

## `kernel-builder update-build`

Update a `build.toml` to the current format

**Usage:** `kernel-builder update-build [KERNEL_DIR]`

###### **Arguments:**

* `<KERNEL_DIR>`

## `kernel-builder skills`

Install skills for AI coding assistants (Claude, Codex, OpenCode)

**Usage:** `kernel-builder skills <COMMAND>`

###### **Subcommands:**

* `add` — Install a kernels skill for an AI assistant

## `kernel-builder skills add`

Install a kernels skill for an AI assistant

**Usage:** `kernel-builder skills add [OPTIONS]`

###### **Options:**

* `--skill <SKILL>` — Skill to install

  Default value: `cuda-kernels`

  Possible values: `cuda-kernels`, `cpu-kernels`, `rocm-kernels`, `triton-kernels`, `xpu-kernels`

* `--claude` — Install for Claude
* `--codex` — Install for Codex
* `--opencode` — Install for OpenCode
* `-g`, `--global` — Install globally (user-level) instead of in the current project directory
* `--dest <DEST>` — Install into a custom destination (path to skills directory)
* `--force` — Overwrite existing skills in the destination

## `kernel-builder clean-pyproject`

Clean generated artifacts

**Usage:** `kernel-builder clean-pyproject [OPTIONS] [KERNEL_DIR] [TARGET_DIR]`

###### **Arguments:**

* `<KERNEL_DIR>`
* `<TARGET_DIR>` — The directory to clean from (directory of `BUILD_TOML` when absent)

###### **Options:**

* `-d`, `--dry-run` — Show what would be deleted without actually deleting
* `-f`, `--force` — Force deletion without confirmation
* `--ops-id <OPS_ID>` — This is an optional unique identifier that is suffixed to the kernel name to avoid name collisions. (e.g. Git SHA)

    This document was generated automatically by
    clap-markdown.

### Install agent skills
https://huggingface.co/docs/kernels/pr_751/cli-skills.md

# Install agent skills

Use `kernel-builder skills add` to install the skills for AI coding assistants like Claude, Codex, and OpenCode.
Supported skills include:
- `cuda-kernels` (default)
- `rocm-kernels`
- `xpu-kernels`
- `cpu-kernels`
- `triton-kernels`

Skill files are downloaded from the `huggingface/kernels` directory in this [repository](https://github.com/huggingface/kernels/tree/main/kernel-builder/skills).

Skills instruct agents how to deal with hardware-specific optimizations, integrate with libraries like diffusers and transformers, and benchmark kernel performance in consistent ways.

> [!TIP]
> **When are CPU kernels actually helpful?** Two main cases:
> - **Better performance on Intel Xeon** — custom AVX2/AVX512 kernels (and AMX via brgemm for quantized GEMM) outperform generic PyTorch ops for element-wise and quantized workloads, especially in CPU-only or latency-sensitive serving.
> - **Enabling functionality that otherwise can't run** — some kernels are a hard requirement, e.g. `megablocks` MoE on CPU, where without the kernel you simply cannot run MXFP4.

Example CPU kernels built with this skill (available on the Hub under [`kernels-community`](https://huggingface.co/kernels-community)):

- [`kernels-community/megablocks`](https://huggingface.co/kernels-community/megablocks) — MoE kernels with a CPU backend that enable running MXFP4 MoE models on CPU.
- [`kernels-community/quantization-gptq`](https://huggingface.co/kernels-community/quantization-gptq) — INT4 quantized GEMM using AVX512.
- [`kernels-community/rmsnorm`](https://huggingface.co/kernels-community/rmsnorm) — RMSNorm with AVX2/AVX512 element-wise paths.

> [!TIP]
> **When are Triton kernels useful?** Two main cases:
> - **Portability across NVIDIA and AMD** — a single Triton kernel runs on both CUDA and ROCm without modification. No vendor-specific code needed.
> - **Fusing multiple ops to reduce memory traffic** — operations like softmax (5 PyTorch ops, ~8MN memory ops) become a single kernel (2MN ops) with a ~4x reduction in DRAM round-trips. Any chain of element-wise + reduction ops that PyTorch executes as separate kernels is a fusion opportunity.

Example Triton kernels built with this skill:

- [`jaygala223/triton-layernorm`](https://huggingface.co/kernels/jaygala223/triton-layernorm) — fused LayerNorm with fp32 accumulation, ~1.45x faster than PyTorch on V100.

Examples:

```bash
# install for Claude in the current project
kernel-builder skills add --claude

# install ROCm kernels skill for Codex
kernel-builder skills add --skill rocm-kernels --codex

# install globally for Codex
kernel-builder skills add --codex --global

# install for multiple assistants
kernel-builder skills add --claude --codex --opencode

# install to a custom destination and overwrite if already present
kernel-builder skills add --dest ~/my-skills --force
```

### Kernel requirements
https://huggingface.co/docs/kernels/pr_751/kernel-requirements.md

# Kernel requirements

Kernels on the Hub must fulfill the requirements outlined on this page. By
ensuring kernels are compliant, they can be used on a wide range of Linux
systems and Torch builds.

[Join us on GitHub Discussions](https://github.com/huggingface/kernels/discussions/categories/kernel-builder)
for questions and discussions about building kernels!

## Repository type

Compliant kernels are published as `kernel`-type repositories on the Hub
(the first-class kernel repository type). New uploads via `kernel-builder`
default to this type; see the [migration guide](migration) if you
maintain an older `model`-type kernel repository.

## Trusted publishers

`kernels` only loads kernels from a curated set of trusted publishers by
default. Loading from any other publisher raises an error unless the caller
opts in with `trust_remote_code=True`:

```python
# Trusted publisher: works without opt-in.
get_kernel("kernels-community/activation", version=1)

# Untrusted publisher: must opt in explicitly.
get_kernel("some-other-org/my-kernel", version=1, trust_remote_code=True)
```

The Hub also exposes a `trustedKernelPublisher` flag on the kernel API and
displays a corresponding badge in the UI.

## Directory layout

A kernel repository on the Hub must contain a `build` directory. This
directory contains build variants of a kernel in the form of directories
following the template
`<framework><version>-cxx<abiver>-<cu><cudaver>-<arch>-<os>`.
For example `build/torch26-cxx98-cu118-x86_64-linux`.

The kernel is in the build variant directory and must contain a
`__init__.py` file. For compatibility with older versions of the
`kernels` package, each variant directory must also contain a single
directory with the same name as the repository (replacing `-` by `_`).
For instance, kernels in the `kernels-community/activation` repository
have a directory like `build/<variant>/activation`. This directory
must contain an `__init__.py` file that exports the same symbols as
`__init__.py` in the build variant directory `build/<variant>`.
[This example](https://huggingface.co/kernels-test/flattened-build/blob/main/build/torch-universal/flattened_build/__init__.py)
shows how this can be done. This compatibility directory is
automatically created by `kernel-builder`.

## Build variants

A kernel can be compliant for a specific compute framework (e.g. CUDA) or
architecture (e.g. x86_64). For compliance with a compute framework and
architecture combination, all the variants from the [build variant list](builder/build-variants)
must be available for that combination.

## Kernel metadata

The build variant directory must contain a `metadata.json` file with kernel
metadata. Currently the following top-level keys are supported:

- `id` (`str`, required): a unique identifier for the kernel. This
  identifier must also be a valid Python module name. If the kernel
  registers Torch ops, they must be registered as `torch.ops.<id>`
- `name` (`str`, required): then name of the kernel. Replacing dashes
  by underscores should result in the module name of the kernel.
- `version` (`int`, required): the kernel version number.
- `license` (`str`, required): the kernel license in. Refer to the
  list of [supported license identifiers](https://huggingface.co/docs/hub/repositories-licenses).
- `upstream` (`str`, optional): Git-compatible URL (passable to `git clone`)
  of the original upstream repository where the kernel source code comes from.
- `source` (`str`, optional): Git-compatible URL (passable to `git clone`)
  of the kernel-builder formatted source repository (must contain `build.toml`
  and `flake.nix`).
- `backend` (`dict`, required): information about the compute backend that
  this build variant supports.
- `digest` (`Digest`, required): hash digest of the kernel files.
- `python-depends` (`list[str]`, optional): list of Python dependencies
  from a curated set of Python dependencies.
- `provenance` (`dict`, optional): provenance of the build, used to flag
  non-reproducible (dirty) builds. It contains two optional sub-objects:
  - `kernel-builder`: the `kernel-builder` that produced the build, with its
    `version` (`str`), the `sha` (`str`) of the `kernel-builder` source it was
    built from (when known), and a `dirty` (`bool`) flag that is `true` when
    `kernel-builder` was built from a source tree with uncommitted changes.
  - `kernel`: the kernel source that was built, with its commit `sha` (`str`)
    and a `dirty` (`bool`) flag that is `true` when the kernel source had
    uncommitted changes.

  When either `dirty` flag is set, the kernel was built from uncommitted
  sources and cannot be reliably reproduced.

  > **Note:** For Nix builds, dirtiness follows Nix's flake tree status, which
  > also counts **untracked** files (including an uncommitted `flake.lock`).
  > Commit your `flake.lock` (and avoid stray untracked files) so that clean
  > builds are not flagged as dirty. Local `create-pyproject` runs only
  > consider changes to tracked files.

  > **Note:** The kernel `sha`/`dirty` are captured at the moment
  > `create-pyproject` runs, so they describe the source tree as it was *then*.
  > Running `create-pyproject` and committing afterwards is bad practice: the
  > recorded provenance keeps pointing at the pre-commit state (a stale `sha`,
  > and `dirty: true` if the tree was dirty) even though the committed source
  > differs. Generate the metadata from the final, committed source instead.

Example `metadata.json`:

```json
{
  "name": "mykernel",
  "id": "_mykernel_cuda_7a4e5a7",
  "version": 1,
  "license": "Apache-2.0",
  "python-depends": ["einops"],
  "backend": {
    "type": "cuda",
    "archs": ["7.0", "7.2", "7.5", "8.0", "8.6", "8.7", "8.9", "9.0+PTX"]
  },
  "digest": {
    "algorithm": "sha256",
    "files": {
      "__init__.py": "xLMbARTcTl8L/m1kJLc/h/QL4Kzt772F872a46pfRGI=",
      "_mykernel_cuda_7645816_dirty.abi3.so": "vtdzzToloH38HZkVs7sFEf69QFDxROuPsBAond3Jic0=",
      "_ops.py": "Hrp5aF4o0eHSttw4sQGsbBAXFqvLJ42Y9YJ2KkqvZhg=",
      "mykernel/__init__.py": "DFYPlrhXwYjEqCl/8n0SmWGZV8NFml5DPhMjKfv98GY="
    }
  }
}
```

The `metadata.json` file is generated automatically by `kernel-builder`.

## Backend

The `backend` specifies a dictionary of the following form:

```json
{
  # ...
  "backend": {
    "type": "cuda",
    "archs": ["7.0", "7.2", "7.5", "8.0", "8.6", "8.7", "8.9", "9.0+PTX"]
  }
}
```

The backend `type` must be one of `cann`, `cpu`, `cuda`, `metal`, `neuron`,
`rocm`, `tpu`, or `xpu`. For CUDA and ROCm, the supported architectures must
be specified in the `archs` field.

### Python dependencies

You can specify Python dependencies that your kernel requires. Dependencies can be either general (required for all backends) or backend-specific (required only for certain compute backends like CUDA, ROCm, XPU, Metal, or CPU).

#### General dependencies

For dependencies required regardless of the backend, use the `python-depends` field:

```json
{
  "python-depends": ["einops"]
}
```

#### Backend-specific dependencies

For dependencies that are only needed for specific backends, use the `python-depends-backends` field:

```json
{
  "python-depends-backends": {
    "cuda": ["nvidia-cutlass-dsl"],
    "xpu": ["onednn"]
  }
}
```

#### Combined example

You can specify both general and backend-specific dependencies:

```json
{
  "python-depends": ["einops"],
  "python-depends-backends": {
    "cuda": ["nvidia-cutlass-dsl"]
  },
  "version": 1
}
```

#### Allowed dependencies

The following dependencies are currently allowed:

**General dependencies:**

- `einops`
- `helion`

**Backend-specific dependencies:**

- CUDA: `nvidia-cutlass-dsl`
- XPU: `onednn`

Dependencies are validated based on the backend being used. When a kernel is loaded, only the dependencies relevant to the active backend are checked.

## Versioning

Kernels are versioned using a major version. The kernel revisions of a
version are stored in a branch of the form `v<version>`. Each build
variant will also have the kernel version in `metadata.json`.

The version **must** be bumped in the following cases:

- The kernel API is changed in an incompatible way.
- The API is extended in a compatible way, but not all build variants
  receive the extension (e.g. because they are for older Torch versions
  that are not supported by `kernel-builder` anymore).

In both cases, build variants that are not updated must be removed from
the new version's branch.

> [!IMPORTANT]
> The *kernel API* covered by these versioning guarantees is only the
> public API: the symbols listed in the `__all__` of the kernel's
> top-level `__init__.py`. Anything not in `__all__` (e.g. internal
> helpers or names prefixed with `_`) is considered private and may
> change or be removed at any time without a version bump. Export every
> symbol you intend consumers to rely on in `__all__`.

> [!NOTE]
> By convention, we reserve version `0` for kernels that are still in
> alpha or beta stage and are not recommended for production use (e.g.
> because the API is still changing regularly or there are still too
> many issues).

## Native Python module

Kernels will typically contain a native Python module with precompiled
compute kernels and bindings. This module must fulfill the requirements
outlined in this section. For all operating systems, a kernel must not
have dynamic library dependencies outside:

- Torch;
- CUDA/ROCm libraries installed as dependencies of Torch.

## Compatibility with torch.compile

The Kernel Hub also encourages to write the kernels in a `torch.compile`
compliant way. This helps to ensure that the kernels are compatible with
`torch.compile` without introducing any graph breaks and triggering
recompilation which can limit the benefits of compilation.

[Here](https://github.com/huggingface/kernels/blob/f83b4da6b7f6b171b47bb9bf96271ae2273bc9d3/builder/examples/relu-backprop-compile/tests/test_relu.py#L162)
is a simple test example which checks for graph breaks and
recompilation triggers during `torch.compile`.

### Linux

- Use [ABI3/Limited API](https://docs.python.org/3/c-api/stable.html#stable-application-binary-interface)
  for compatibility with Python 3.9 and later.
- Compatible with [`manylinux_2_28`](https://github.com/pypa/manylinux?tab=readme-ov-file#manylinux_2_28-almalinux-8-based).
  This means that the extension **must not** use symbols versions higher than:
  - GLIBC 2.28
  - GLIBCXX 3.4.24
  - CXXABI 1.3.11
  - GCC 7.0.0

These requirements can be checked with the ABI checker (see below).

### macOS

- Use [ABI3/Limited API](https://docs.python.org/3/c-api/stable.html#stable-application-binary-interface)
  for compatibility with Python 3.9 and later.
- macOS deployment target 26.0.
- Metal 4.0 (`-std=metal4.0`).

The ABI3 requirement can be checked with the ABI checker (see below).

### ABI checker

The manylinux_2_28 and Python ABI 3.9 version requirements can be checked with
`kernel-builder check-abi`:

```bash
$ kernel-builder check-abi examples/kernels/relu
🐍 Checking for compatibility with manylinux_2_28 and Python ABI version 3.9: /home/daniel/git/kernels/examples/kernels/relu/result/torch211-cpu-x86_64-linux/_relu_cpu_30dc0ae_dirty.abi3.so
✅ No compatibility issues found
🐍 Checking for compatibility with manylinux_2_28 and Python ABI version 3.9: /home/daniel/git/kernels/examples/kernels/relu/result/torch211-cu126-x86_64-linux/_relu_cuda_30dc0ae_dirty.abi3.so
✅ No compatibility issues found
🐍 Checking for compatibility with manylinux_2_28 and Python ABI version 3.9: /home/daniel/git/kernels/examples/kernels/relu/result/torch211-cu128-x86_64-linux/_relu_cuda_30dc0ae_dirty.abi3.so
✅ No compatibility issues found
🐍 Checking for compatibility with manylinux_2_28 and Python ABI version 3.9: /home/daniel/git/kernels/examples/kernels/relu/result/torch211-cu130-x86_64-linux/_relu_cuda_30dc0ae_dirty.abi3.so
✅ No compatibility issues found
[...]
```

## Torch extension

Torch native extension functions must be [registered](https://pytorch.org/tutorials/advanced/cpp_custom_ops.html#cpp-custom-ops-tutorial)
in `torch.ops.<namespace>`. Since we allow loading of multiple versions of
a module in the same Python process, `namespace` must be unique for each
version of a kernel. Failing to do so will create clashes when different
versions of the same kernel are loaded. Two suggested ways of doing this
are:

- Appending a truncated SHA-1 hash of the git commit that the kernel was
  built from to the name of the extension.
- Appending random material to the name of the extension.

**Note:** we recommend against appending a version number or git tag.
Version numbers are typically not bumped on each commit, so users
might use two different commits that happen to have the same version
number. Git tags are not stable, so they do not provide a good way
of guaranteeing uniqueness of the namespace.

## Layers

A kernel can provide layers in addition to kernel functions. A layer from
the Hub can replace the `forward` method of an existing layer for a certain
device type. This makes it possible to provide more performant kernels for
existing layers. See the [layers documentation](layers) for more information
on how to use layers.

### Writing layers

To make the extension of layers safe, the layers must fulfill the following
requirements:

- The layers are subclasses of `torch.nn.Module`.
- The layers are pure, meaning that they do not have their own state. This
  means that:
  - The layer must not define its own constructor.
  - The layer must not use class variables.
- No other methods must be defined than `forward`.
- The `forward` method has a signature that is compatible with the
  `forward` method that it is extending.

There are two exceptions to the _no class variables rule_:

1. The `has_backward` variable can be used to indicate whether the layer has
   a backward pass implemented (`True` when absent).
2. The `can_torch_compile` variable can be used to indicate whether the layer
   supports `torch.compile` (`False` when absent).

This is an example of a pure layer:

```python
class SiluAndMul(nn.Module):
    # This layer does not implement backward.
    has_backward: bool = False

    def forward(self, x: torch.Tensor):
        d = x.shape[-1] // 2
        output_shape = x.shape[:-1] + (d,)
        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
        ops.silu_and_mul(out, x)
        return out
```

For some layers, the `forward` method has to use state from the adopting class.
In these cases, we recommend to use type annotations to indicate what member
variables are expected. For instance:

```python
class LlamaRMSNorm(nn.Module):
    weight: torch.Tensor
    variance_epsilon: float

    def forward(self, hidden_states: torch.Tensor) -> torch.Tensor:
        return rms_norm_fn(
            hidden_states,
            self.weight,
            bias=None,
            residual=None,
            eps=self.variance_epsilon,
            dropout_p=0.0,
            prenorm=False,
            residual_in_fp32=False,
        )
```

This layer expects the adopting layer to have `weight` and `variance_epsilon`
member variables and uses them in the `forward` method.

### Exporting layers

To accommodate portable loading, `layers` must be defined in the main
`__init__.py` file. For example:

```python
from . import layers

__all__ = [
  # ...
  "layers"
  # ...
]
```

> [!IMPORTANT]
> Only symbols listed in `__all__` are treated as the kernel's public
> API. This is the surface that consumers can rely on and that the
> [versioning guarantees](#versioning) apply to, so be sure to export
> every function, class, and `layers` module you want to expose.

## Python requirements

- Python code must be compatible with Python 3.9 and later.
- All Python code imports from the kernel itself must be relative. So,
  for instance if in the example kernel `example`,
  `module_b` needs a function from `module_a`, import as:

  ```python
  from .module_a import foo
  ```

  **Never use:**

  ```python
  # DO NOT DO THIS!

  from example.module_a import foo
  ```

  The latter would import from the module `example` that is in Python's
  global module dict. However, since we allow loading multiple versions
  of a module, we uniquely name the module.

- Only modules from the Python standard library, Torch, or the kernel itself
  can be imported.

### Migrate from older versions
https://huggingface.co/docs/kernels/pr_751/migration.md

# Migrate from older versions

## 0.12

### Adopting kernel versions

Before `kernels` 0.12, kernels could be pulled from a repository
without specifying a version. This is deprecated in kernels 0.12
and is an error in kernels 0.15. Instead, use of a kernel should
always specify a version or revision (except for local kernels).

Kernels only use a major version. The kernel maintainer is responsible
for never breaking a kernel within a major version and should bump up
the major version if the kernel API changes and/or when support for
older Torch versions is removed.

> [!NOTE]
> Version `0` kernels are excluded from the API compatibility requirement,
> since it is used for alpha/beta-quality kernels that may still have
> rapidly changing APIs.

You can find the versions that are supported by a kernel using the
`kernels versions command`. For example:

```bash
$ kernels versions kernels-community/activation
Version 1: torch210-cxx11-cu126-x86_64-linux, torch210-cxx11-cu128-x86_64-linux, torch210-cxx11-cu130-x86_64-linux, torch27-cxx11-cu118-x86_64-linux, torch27-cxx11-cu126-x86_64-linux, torch27-cxx11-cu128-aarch64-linux, torch27-cxx11-cu128-x86_64-linux ✅, torch28-cxx11-cu126-aarch64-linux, torch28-cxx11-cu126-x86_64-linux, torch28-cxx11-cu128-aarch64-linux, torch28-cxx11-cu128-x86_64-linux, torch28-cxx11-cu129-aarch64-linux, torch28-cxx11-cu129-x86_64-linux, torch29-cxx11-cu126-aarch64-linux, torch29-cxx11-cu126-x86_64-linux, torch29-cxx11-cu128-aarch64-linux, torch29-cxx11-cu128-x86_64-linux, torch29-cxx11-cu130-aarch64-linux, torch29-cxx11-cu130-x86_64-linux
```

The command lists all available versions (here only version 1) with
all the variants that are supported. A check mark is printed after
the variant that is compatible with your current environment.

Code that uses a kernel can be updated as follows:

```python
# Old:
activation = get_kernel("kernels-community/activation")
activation = get_kernel("kernels-community/activation", version=">=0.0.2 && <0.1.0")

# New:
activation = get_kernel("kernels-community/activation", version=1)

# Old:
kernel_layer_mapping = {
    "SiluAndMul": {
        "cuda": LayerRepository(
            repo_id="kernels-community/activation",
            layer_name="SiluAndMul",
        ),
    }
}
kernel_layer_mapping = {
    "SiluAndMul": {
        "cuda": LayerRepository(
            repo_id="kernels-community/activation",
            layer_name="SiluAndMul",
            version=">=0.0.2 && <0.1.0",
        ),
    }
}

# New:
kernel_layer_mapping = {
    "SiluAndMul": {
        "cuda": LayerRepository(
            repo_id="kernels-community/activation",
            layer_name="SiluAndMul",
            version=1,
        ),
    }
}
```

## 0.14

### `kernel` repo type on the Hub

Kernels are now a first-class repository type on the Hugging Face Hub, and
`kernels` 0.14 loads kernels exclusively from `kernel`-type repositories.
`model`-type kernel repositories are no longer supported by the loader.

New uploads via `kernel-builder build-and-upload` default to
`--repo-type kernel`. To publish, the owning user or org must have
kernel-creation access. Request it from
[huggingface.co/settings/account](https://huggingface.co/settings/account)
("Request Kernels Creation").

To migrate an existing `model`-type kernel repository:

1. Make sure the publishing org has been granted kernel-creation access
   (see above).
2. Re-upload with `kernel-builder build-and-upload` to a `kernel`-type
   repository. Either keep the same `repo-id` in `build.toml` if the
   repository has been migrated to the new type, or point it at a newly
   created `kernel`-type repository.
3. Update consumers' [get_kernel()](/docs/kernels/pr_751/en/api/kernels#kernels.get_kernel) and [LayerRepository](/docs/kernels/pr_751/en/api/layers#kernels.LayerRepository) calls
   to reference the new repository if the `repo-id` changed.

## 0.16

> [!WARNING] **Deprecation of kernel functions.**
> `use_kernel_func_from_hub`, `FuncRepository`, `LocalFuncRepository`, and
`LockedFuncRepository` are now deprecated.

To make a function extensible by a layer, use the same decorator as for
layers ([use_kernel_forward_from_hub()](/docs/kernels/pr_751/en/api/layers#kernels.use_kernel_forward_from_hub)). This makes it clearer
that the function is actually replaced by a layer. You can also use the
[use_kernelized_func()](/docs/kernels/pr_751/en/api/layers#kernels.use_kernelized_func) decorator to attach such a function to
the layer where it is used, making it discoverable by
[kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize).

For example:

```python
import torch
import torch.nn as nn
import torch.nn.functional as F

from kernels import use_kernel_func_from_hub  # old
from kernels import use_kernel_forward_from_hub, use_kernelized_func

# Old:
@use_kernel_func_from_hub("silu_and_mul")
def silu_and_mul(x: torch.Tensor) -> torch.Tensor:
    d = x.shape[-1] // 2
    return F.silu(x[..., :d]) * x[..., d:]

class FeedForward(nn.Module):
    def __init__(self, in_features: int, out_features: int):
        super().__init__()
        self.silu_and_mul = silu_and_mul
        self.linear = nn.Linear(in_features, out_features)

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        return self.silu_and_mul(self.linear(x))

# New:
@use_kernel_forward_from_hub("silu_and_mul")
def silu_and_mul(x: torch.Tensor) -> torch.Tensor:
    d = x.shape[-1] // 2
    return F.silu(x[..., :d]) * x[..., d:]

@use_kernelized_func(silu_and_mul)
class FeedForward(nn.Module):
    def __init__(self, in_features: int, out_features: int):
        super().__init__()
        self.linear = nn.Linear(in_features, out_features)

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        return silu_and_mul(self.linear(x))
```

`FuncRepository`, `LocalFuncRepository`, and `LockedFuncRepository` are
not replaced. They allowed using an arbitrary function from a kernel
as a layer, but this was easily misused and did not have a clean way
of marking such a function as supporting `torch.compile` or backward
passes. Going forward, kernel functions should be exposed as regular
kernel layers and used with [LayerRepository](/docs/kernels/pr_751/en/api/layers#kernels.LayerRepository),
[LocalLayerRepository](/docs/kernels/pr_751/en/api/layers#kernels.LocalLayerRepository), or [LockedLayerRepository](/docs/kernels/pr_751/en/api/layers#kernels.LockedLayerRepository).
For example:

```python
import torch
import torch.nn as nn

# Old:
def fast_silu_and_mul(x: torch.Tensor) -> torch.Tensor:
  ...

# New:
def fast_silu_and_mul(x: torch.Tensor) -> torch.Tensor:
  ...

# Kernel layer that exposes the function.
class FastSiluAndMul(nn.Module):
    def forward(self, x: torch.Tensor) -> torch.Tensor:
        return fast_silu_and_mul(x)
```

For more information, see the [layer documentation](layers).

### Kernels
https://huggingface.co/docs/kernels/pr_751/index.md

# Kernels

The Kernel Hub allows Python libraries and applications to load compute
kernels directly from the [Hub](https://huggingface.co/). Kernels are a first-class
repository type on the Hub, with dedicated pages that surface supported
hardware and versions. To support dynamic loading, Hub kernels differ from
traditional Python kernel packages in that they are made to be:

- **Portable**: a kernel can be loaded from paths outside `PYTHONPATH`.
- **Unique**: multiple versions of the same kernel can be loaded in the
  same Python process.
- **Compatible**: `kernels` must support all recent versions of Python and
  the different PyTorch build configurations (various CUDA versions
  and C++ ABIs). Furthermore, older C library versions must be supported.

Browse available kernels at [huggingface.co/kernels](https://huggingface.co/kernels).

The Kernels project is divided into two parts:

- Builder: [`kernel-builder`](builder-cli) provides utilities to build, package, and distribute compute kernels in a way that is compatible with the Hugging Face Hub and `kernels`.
- `kernels`: The [`kernels`](basic-usage) is a Python package that lets
  users load compatible compute kernels from the Hub. Refer to the [quickstart](basic-usage) to know more.

If you're looking for a more involved "Why kernels?" answer, refer to
[this page](./why_kernels).

The [talks page](./talks) page has links to talks on the
Kernels project. The [blog page](./blog) collects blog posts
on the Kernels project.

### Overview
https://huggingface.co/docs/kernels/pr_751/cli.md

# Overview

The `kernels` CLI provides commands for managing compute kernels.

## Commands

| Command                                     | Description                        |
| ------------------------------------------- | ---------------------------------- |
| [benchmark](cli-benchmark)               | Run benchmark results for a kernel |
| [download](cli-download)                 | Download a kernel                  |
| [info](cli-info)                         | Describe a kernel                  |
| [lock](cli-lock)                         | Lock kernel revisions              |
| [verify-signature](cli-verify-signature) | Verify a kernel signature          |
| [versions](cli-versions)                 | Show kernel versions               |

### Talks
https://huggingface.co/docs/kernels/pr_751/talks.md

# Talks

This page lists talks on the Kernels project, delivered at
different events:

* [Lecture 106: Hugging Face Kernels (GPU Mode)](https://www.youtube.com/watch?v=Ok8vi6JemVQ)
* [PyTorch Day India 2026 Beyond the Brrr: Building a Unified Ecosystem for Optimized Kernels](https://www.youtube.com/watch?v=q0GfzJmuaUM)
* [Build a PyTorch ReLU Kernel with Hugging Face Kernels (CPU + Metal)](https://youtu.be/wQR-QC7pbqQ?is=JY4pvjXpkz8ghgwS)

This page will be kept updated.

### kernels lock
https://huggingface.co/docs/kernels/pr_751/cli-lock.md

# kernels lock

Use `kernels lock` to generate a `kernels.lock` file that pins kernel dependencies to specific revisions.

## Usage

```bash
kernels lock <project_dir>
```

## What It Does

- Reads kernel dependencies from `pyproject.toml` under `[tool.kernels.dependencies]`
- Resolves each kernel to its current revision SHA
- Writes a `kernels.lock` file with pinned versions and variant information

## Examples

Lock kernels in the current project:

```bash
kernels lock .
```

Lock kernels in a specific project:

```bash
kernels lock /path/to/my-project
```

## pyproject.toml Format

Add your kernel dependencies to `pyproject.toml`:

```toml
[tool.kernels.dependencies]
"kernels-community/activation" = 1
```

The version can be:

- A version number (e.g., `1`, `2`)

## kernels.lock Format

The generated lock file contains:

```json
[
  {
    "repo_id": "kernels-community/activation",
    "sha": "ece277f908b9453112722d584fee4b5696f21c49",
    "variants": {
      "torch210-cu128-x86_64-windows": {
        "hash": "sha256-cbf085e1d297d990d9cb074fb5079ff48e9682c729f53a0899a36b5164a6fb45",
        "hash_type": "git_lfs_concat"
      },
      // ...
      "torch29-metal-aarch64-darwin": {
        "hash": "sha256-9f665b54a53246a7d3627422f8a0d41d7956dc5409043dbd14c4ec0327aea310",
        "hash_type": "git_lfs_concat"
      }
    }
  }
]
```

## Workflow

1. Add dependencies to `pyproject.toml`
2. Run `kernels lock .` to generate the lock file
3. Commit both `pyproject.toml` and `kernels.lock`
4. Use `kernels download .` to install locked kernels

## See Also

- [kernels download](cli-download) - Download locked kernels
- [kernels versions](cli-versions) - View available kernel versions

### Use layers
https://huggingface.co/docs/kernels/pr_751/layers.md

# Use layers

A kernel can provide layers in addition to kernel functions. A layer from
the Hub can replace the `forward` method of an existing layer for a certain
device type. This makes it possible to provide more performant kernels for
existing layers.

See [Kernel requirements](kernel-requirements) for more information on the
requirements of Hub layers.

## Making a layer extensible with kernels from the hub

### Using a decorator

A layer can be made extensible with the [use_kernel_forward_from_hub()](/docs/kernels/pr_751/en/api/layers#kernels.use_kernel_forward_from_hub)
decorator. For example:

```python
@use_kernel_forward_from_hub("SiluAndMul")
class SiluAndMul(nn.Module):
    def forward(self, input: torch.Tensor) -> torch.Tensor:
        d = input.shape[-1] // 2
        return F.silu(input[..., :d]) * input[..., d:]
```

The decorator does not change the behavior of the class -- it annotates
the class with the given name (here `SiluAndMul`). The [kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize) function
described below uses this name to look up kernels for the layer.

### External layers

An existing layer that does not (yet) have the [use_kernel_forward_from_hub()](/docs/kernels/pr_751/en/api/layers#kernels.use_kernel_forward_from_hub)
decorator can be made extensible using the [replace_kernel_forward_from_hub()](/docs/kernels/pr_751/en/api/layers#kernels.replace_kernel_forward_from_hub)
function:

```python
from somelibrary import SiluAndMul

replace_kernel_forward_from_hub(SiluAndMul, "SiluAndMul")
```

**Warning:** we strongly recommend using layers with a decorator, since
it signifies that the maintainer intends to keep the `forward` signature
compatible with layers from the hub.

### Using a function as a layer

Sometimes it can be useful to make a function extensible, for example
because the function cannot be replaced by a layer. In such cases, you
can use the [use_kernel_forward_from_hub()](/docs/kernels/pr_751/en/api/layers#kernels.use_kernel_forward_from_hub) decorator on a
function:

```python
@use_kernel_forward_from_hub("silu_and_mul")
def silu_and_mul(x: torch.Tensor) -> torch.Tensor:
    d = x.shape[-1] // 2
    return F.silu(x[..., :d]) * x[..., d:]
```

This will replace the function by an instantiated `torch.nn.Module`
(singleton) that calls the function itself in its forward method. It will
still behave as a function, since `torch.nn.Module` provides an
implementation of `__call__` that delegates to `forward`.

For [kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize) to see the function, you must use
[use_kernelized_func()](/docs/kernels/pr_751/en/api/layers#kernels.use_kernelized_func) on an `torch.nn.Module` that is part
of the to-be kernlized model to make the function visible to
[kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize). The function is typically attached to the module
that uses it. For example:

```python
@use_kernelized_func(silu_and_mul)
class FeedForward(nn.Module):
  def __init__(self, in_features: int, out_features: int):
      self.linear = nn.Linear(in_features, out_features)

  def forward(self, x: torch.Tensor) -> torch.Tensor:
      return silu_and_mul(self.linear(x))
```

Functions used by [use_kernelized_func()](/docs/kernels/pr_751/en/api/layers#kernels.use_kernelized_func) must always have a
[use_kernel_forward_from_hub()](/docs/kernels/pr_751/en/api/layers#kernels.use_kernel_forward_from_hub) decorator.

## Kernelizing a model

A model will not use Hub kernels by default, even if it contains extensible
layers. To enable the use of Hub kernels in the model, it needs to be
'kernelized' using the [kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize) function. This function traverses the
model graph and replaces the `forward` methods of extensible layers for which
Hub kernels are registered. [kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize) can be used as follows:

```python
model = MyModel(...)
model = kernelize(model, mode=Mode.INFERENCE)
```

The [kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize) function modifies the model in-place, the model itself is
returned as a convenience. The `mode` specifies that the model will be used
in inference. Similarly, you can ask [kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize) to prepare the model for
training:

```python
model = MyModel(...)
model = kernelize(model, mode=Mode.TRAINING)
```

A model that is kernelized for training can also be used for inference, but
not the other way around. If you want to change the mode of the kernelized
model, you can just run [kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize) on the model again with the new mode.

If you want to compile a model with `torch.compile`, this should be indicated
in the mode as well. You can do this by combining `Mode.INFERENCE` or
`Mode.TRAINING` with `Mode.TORCH_COMPILE` using the set union (`|`) operator:

```python
model = MyModel(...)

# Inference
model = kernelize(model, mode=Mode.INFERENCE | Mode.TORCH_COMPILE)

# Training
model = kernelize(model, mode=Mode.TRAINING | Mode.TORCH_COMPILE)
```

### Kernel device

Kernels can be registered per device type. For instance, separate `cuda` and
`metal` kernels could be registered for the name `SiluAndMul`. By default,
[kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize) will try to infer the device type from the model's parameters.
You can pass the device type to [kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize) if the device type cannot be
inferred (e.g. because the model has no parameters):

```python
model = MyModel(...)
model = kernelize(model, device="cuda", mode=Mode.INFERENCE)
```

### Fallback `forward`

If the `TRAINING` and/or `TORCH_COMPILE` modes are used, but a registered
kernel does not support backward passes or `torch.compile` respectively,
[kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize) will fall back to the original, non-kernelized, layer. You
can let [kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize) raise an exception instead by using `use_fallback=False`:

```python
model = MyModel(...)
model = kernelize(model, mode=Mode.INFERENCE | Mode.TORCH_COMPILE, use_fallback=False)
```

This can be useful if you want to guarantee that Hub kernels are used.

### Inspecting which kernels are used

The kernels that are used are logged at the `INFO` level by [kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize).
See the [Python logging](https://docs.python.org/3/library/logging.html)
documentation for information on how to configure logging.

## Registering a hub kernel for a layer

[kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize) relies on kernel mappings to find Hub kernels for layers.
Kernel mappings map a kernel name such as `SiluAndMul` to a kernel on
the Hub. For example:

```python
kernel_layer_mapping = {
    "SiluAndMul": {
        "cuda": LayerRepository(
            repo_id="kernels-community/activation",
            layer_name="SiluAndMul",
            version=1,
        ),
        "rocm": LayerRepository(
            repo_id="kernels-community/activation",
            layer_name="SiluAndMul",
            version=1,
        )
    }
}
```

This uses version `1` of the `SiluAndMul` kernel layer from
`kernels-community/activation` for the `cuda` and `rocm` backends. Kernel
layers are versioned using a major version number. Using `version=1`
will get the latest kernel build from the `v1` branch. Kernel layers
within a version branch must never break the API or remove builds for
older PyTorch versions. This ensures that your code will continue to
work.
Hub-backed [LayerRepository](/docs/kernels/pr_751/en/api/layers#kernels.LayerRepository) and [FuncRepository](/docs/kernels/pr_751/en/api/layers#kernels.FuncRepository) entries must specify
either a `version` or an explicit `revision`.

> [!NOTE]
> Version `0` kernels are excluded from the API compatibility requirement,
> since it is used for alpha/beta-quality kernels that may still have
> rapidly changing APIs.

You can register a mapping, like the one above, using [register_kernel_mapping()](/docs/kernels/pr_751/en/api/layers#kernels.register_kernel_mapping):

```python
register_kernel_mapping(kernel_layer_mapping)
```

This will register the kernel mapping in the current context, which is
normally global. It is recommended to scope the mapping to where it is
used with the [use_kernel_mapping()](/docs/kernels/pr_751/en/api/layers#kernels.use_kernel_mapping) context manager:

```python
with use_kernel_mapping(kernel_layer_mapping):
    # Use the layer for which the mapping is applied.
    model = kernelize(model, mode=Mode.TRAINING | Mode.TORCH_COMPILE)
```

This ensures that the mapping is not active anymore outside the
`with`-scope.

If the layer is stateless (it does not use member variables in its forward _or_ it was
originally a function that was converted into a kernel layer with
[use_kernel_func_from_hub()](/docs/kernels/pr_751/en/api/layers#kernels.use_kernel_func_from_hub)), it can also be mapped to a kernel function:

```python
kernel_layer_mapping = {
    "SiluAndMul": {
        "cuda": FuncRepository(
            repo_id="kernels-community/activation",
            func_name="silu_and_mul",
            version=1,
        ),
    }
}
```

### Registering kernels for specific modes

You might want to register two different kernels for a particular layer,
where one kernel is optimized for a specific mode. You can do so by
registering layer repositories for specific modes. For example:

```python
kernel_layer_mapping = {
    "SiluAndMul": {
        "cuda": {
          Mode.INFERENCE: LayerRepository(
              repo_id="kernels-community/activation-inference-optimized",
              layer_name="SiluAndMul",
              version=1,
          ),
          Mode.TRAINING | Mode.TORCH_COMPILE: LayerRepository(
              repo_id="kernels-community/activation-training-optimized",
              layer_name="SiluAndMul",
              version=1,
          ),
      }
    }
}
```

The [kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize) function will attempt to use the following registered
kernels for a given mode:

- `INFERENCE`: `INFERENCE` → `INFERENCE | TORCH_COMPILE` → `TRAINING` →
  `TRAINING | TORCH_COMPILE` → `FALLBACK`
- `INFERENCE | TORCH_COMPILE`: `INFERENCE | TORCH_COMPILE` →
  `TRAINING | TORCH_COMPILE` → `FALLBACK`
- `TRAINING`: `TRAINING` → `TRAINING | TORCH_COMPILE` → `FALLBACK`
- `TRAINING | TORCH_COMPILE`: `TRAINING | TORCH_COMPILE` → `FALLBACK`

`Mode.FALLBACK` is a special mode that is used when no other mode matches. It
is also used when a kernel is registered without a mode, as described in the
previous section.

```python
kernel_layer_mapping = {
    "SiluAndMul": {
        "cuda": {
            Mode.FALLBACK: LayerRepository(
                repo_id="kernels-community/activation",
                layer_name="SiluAndMul",
                version=1,
            ),
            Mode.INFERENCE: LayerRepository(
                repo_id="kernels-community/activation-inference-optimized",
                layer_name="SiluAndMul",
                version=1,
            ),
            Mode.TRAINING: LayerRepository(
                repo_id="kernels-community/activation-training-optimized",
                layer_name="SiluAndMul",
                version=1,
            ),
        }
    }
}
```

In this case, both `Mode.INFERENCE | Mode.TORCH_COMPILE` and
`Mode.TRAINING | Mode.TORCH_COMPILE` will use the `Mode.FALLBACK` kernel,
since the other kernels do not support `torch.compile`.

### Registering kernels for specific CUDA capabilities

Some kernels only work with newer CUDA architectures. For instance, some
kernels require capability 9.0 for the TMA unit on Hopper GPUs. `kernels`
supports registering layers for a range of CUDA capabilities. To do so,
you need to register the layer for a [Device](/docs/kernels/pr_751/en/api/layers#kernels.Device) with type `cuda` and
set the supported range of CUDA capabilities with using `CUDAProperties`:

```python
kernel_layer_mapping = {
    "SiluAndMul": {
        Device(
            type="cuda",
            properties=CUDAProperties(
                min_capability=75, max_capability=89
            ),
        ): LayerRepository(
            repo_id="kernels-community/activation",
            layer_name="SiluAndMul",
            version=1,
        ),
        Device(
            type="cuda",
            properties=CUDAProperties(
                min_capability=90, max_capability=sys.maxsize
            ),
        ): LayerRepository(
            repo_id="kernels-community/activation-hopper",
            layer_name="SiluAndMul",
            version=1,
        ),
    }
}
```

Capabilities behave as follows:

- The minimum and maximum capabilities are inclusive.
- When a new kernel is registered with the same min/max capabilities as
  an existing kernel, the new kernel will replace the old kernel.
- When there are multiple kernels that support a capability, the kernel
  with the smaller capability interval will be used. E.g. given:
  - `KernelA` with `min_capability=80` and `max_capability=89`;
  - `KernelB` with `min_capability=75` and `max_capability=89`;
  - [kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize) runs on a system with capability 8.6.

  Then `KernelA` will be used because the interval 80..89 is smaller
  than 75..89. The motivation is that kernels with smaller ranges
  tend to be more optimized for a specific set of GPUs. **This behavior
  might still change in the future.**

### Registering kernels for specific ROCm capabilities

Registering kernels for the ROCm architecture follows the exact same
pattern as CUDA kernels, using `min_capability` and `max_capability` to restrict
a kernel to a range of ROCm capabilities.

### Loading from a local repository for testing

The [LocalLayerRepository](/docs/kernels/pr_751/en/api/layers#kernels.LocalLayerRepository) class is provided to load a repository from
a local directory. For example:

```python
with use_kernel_mapping(
    {
        "SiluAndMul": {
            "cuda": LocalLayerRepository(
                repo_path="/home/daniel/kernels/activation",
                package_name="activation",
                layer_name="SiluAndMul",
            )
        }
    },
    inherit_mapping=False,
):
    kernelize(linear, mode=Mode.INFERENCE)
```

Similarly, the [LocalFuncRepository](/docs/kernels/pr_751/en/api/layers#kernels.LocalFuncRepository) class can be used to load a kernel
function from a local directory:

```python
with use_kernel_mapping(
    {
        "silu_and_mul": {
            "cuda": LocalFuncRepository(
                repo_path="/home/daniel/kernels/activation",
                package_name="activation",
                func_name="silu_and_mul",
            )
        }
    },
    inherit_mapping=False,
):
    kernelize(model, mode=Mode.INFERENCE)
```

### Lock kernel versions
https://huggingface.co/docs/kernels/pr_751/locking.md

# Lock kernel versions

Projects that use `setuptools` can lock the kernel versions that should be
used. First specify the accepted versions in `pyproject.toml` and make
sure that `kernels` is a build dependency:

```toml
[build-system]
requires = ["kernels", "setuptools"]
build-backend = "setuptools.build_meta"

[tool.kernels.dependencies]
"kernels-community/activation" = 1
```

Then run `kernels lock .` in the project directory. This generates a `kernels.lock` file with
the locked revisions. The locked revision will be used when loading a kernel with
[get_locked_kernel()](/docs/kernels/pr_751/en/api/kernels#kernels.get_locked_kernel):

```python
from kernels import get_locked_kernel

activation = get_locked_kernel("kernels-community/activation")
```

**Note:** the lock file is included in the package metadata, so it will only be visible
to `kernels` after doing an (editable or regular) installation of your project.

## Locked kernel layers

Locking is also supported for kernel layers. To use locked layers, register them
with the [LockedLayerRepository](/docs/kernels/pr_751/en/api/layers#kernels.LockedLayerRepository) class:

```python
kernel_layer_mapping = {
    "SiluAndMul": {
        "cuda": LockedLayerRepository(
            repo_id="kernels-community/activation",
            layer_name="SiluAndMul",
        )
    }
}

register_kernel_mapping(kernel_layer_mapping)
```

Similarly, you can use the [LockedFuncRepository](/docs/kernels/pr_751/en/api/layers#kernels.LockedFuncRepository) class to lock kernel function
versions:

```python
kernel_layer_mapping = {
    "silu_and_mul": {
        "cuda": LockedFuncRepository(
            repo_id="kernels-community/activation",
            func_name="silu_and_mul",
        )
    }
}

register_kernel_mapping(kernel_layer_mapping)
```

## Pre-downloading locked kernels

Locked kernels can be pre-downloaded by running `kernels download .` in your
project directory. This will download the kernels to your local Hugging Face
Hub cache.

The pre-downloaded kernels are used by the [get_locked_kernel()](/docs/kernels/pr_751/en/api/kernels#kernels.get_locked_kernel) function.
[get_locked_kernel()](/docs/kernels/pr_751/en/api/kernels#kernels.get_locked_kernel) will download a kernel when it is not pre-downloaded. If you
want kernel loading to error when a kernel is not pre-downloaded, you can use
the [load_kernel()](/docs/kernels/pr_751/en/api/kernels#kernels.load_kernel) function instead:

```python
from kernels import load_kernel

activation = load_kernel("kernels-community/activation")
```

### kernels download
https://huggingface.co/docs/kernels/pr_751/cli-download.md

# kernels download

Use `kernels download` to download kernels that have been locked in a project's `kernels.lock` file.

## Usage

```bash
kernels download <project_dir> [--all-variants]
```

## What It Does

- Reads the `kernels.lock` file from the specified project directory
- Downloads each locked kernel at its pinned revision (SHA)
- Installs the appropriate variant for your platform (or all variants with `--all-variants`)

## Examples

Download kernels for the current project:

```bash
kernels download .
```

Download all build variants (useful for CI or multi-platform builds):

```bash
kernels download . --all-variants
```

Download kernels for a specific project:

```bash
kernels download /path/to/my-project
```

## Options

| Option           | Description                                                                               |
| ---------------- | ----------------------------------------------------------------------------------------- |
| `--all-variants` | Download all build variants of each kernel instead of just the current platform's variant |

## Prerequisites

Your project directory must contain a `kernels.lock` file. Generate one using [`kernels lock`](cli-lock).

## See Also

- [kernels lock](cli-lock) - Generate the lock file
- [kernels versions](cli-versions) - View available kernel versions

### kernels verify-signature
https://huggingface.co/docs/kernels/pr_751/cli-verify-signature.md

# kernels verify-signature

Use `kernels verify-signature` to verify the metadata signature and check
that kernel files match the digest embedded in the metadata.

## Usage

```bash
kernels verify-signature <repo_id> <version> [--all-variants] \
  [--filter-unsigned] [--filter-no-digest]
```

## What It Does

- Checks that the signing identity in `metadata.json.sigstore` is approved.
- Verifies that `metadata.json` is not tampered with, using the signature
  in `metadata.json.sigstore`.
- Verifies that other kernel files are not tampered with, using the digest
  in `metadata.json`.

## Examples

Verify version `1` of the `kernels-community/relu` kernel. Only checks
the variant that is compatible with the current system:

```bash
kernels verify-signature kernels-community/relu 1
```

Verify all build variants of the same kernel:

```bash
kernels verify-signature kernels-community/relu 1 --all-variants
```

## Example Output

```bash
$ kernels verify-signature kernels-community/relu 1
✅ torch211-cxx11-cu126-x86_64-linux: kernel metadata is correctly signed
$ kernels verify-signature kernels-community/flash-attn2 1
❌ torch211-cxx11-cu126-x86_64-linux: cannot verify kernel integrity, signature not found
```

## Options

| Option               | Description                                                                                                      |
| -------------------- | ---------------------------------------------------------------------------------------------------------------- |
| `--all-variants`     | Verify all build variants of each kernel instead of just the variant that is compatible with the current system. |
| `--filter-no-digest` | Skip variants that do not have a digest in the metadata (typically older builds that precede code signing).      |
| `--filter-unsigned`  | Skip variants that do not have a detached signature (typically older builds that precede code signing).          |

### Why Nix?
https://huggingface.co/docs/kernels/pr_751/builder/why-nix.md

# Why Nix?

The Kernel Builder project uses Nix to build custom kernels designed specifically for PyTorch.

Here’s why we chose Nix and why it's particularly suited to our workflow:

## 1. Consistent and Reproducible Builds

Nix guarantees deterministic evaluation, ensuring that every kernel is built identically, regardless of the host environment. This consistency prevents "it works on my machine" problems, making debugging and deployment straightforward.

## 2. Simplified Dependency Management

Compiling PyTorch kernels often involves complex dependencies such as CUDA versions, PyTorch APIs, and various C++ toolchains. Nix explicitly defines and manages these dependencies, eliminating version conflicts and making maintenance easier.

## 3. Declarative Configuration

Nix’s declarative approach clearly specifies exactly what each kernel build needs. This transparency aids collaboration, speeds up troubleshooting, and makes it easy to document the build process.

## 4. Isolated, Reliable Builds

Each kernel build with Nix runs in a fully isolated sandbox, removing any uncertainty about external state. This isolation ensures clean builds, free of unexpected side effects.

## 5. Efficient Caching and CI Integration

Kernel compilation can be resource-intensive. Nix leverages efficient caching of build artifacts, significantly reducing build times and optimizing continuous integration workflows.

## 6. Easy Experimentation and Rollbacks

Nix allows you to experiment with different kernel configurations, PyTorch versions, or CUDA toolkits easily. If a change introduces an issue, reverting to a previous state is quick and effortless.

Overall, Nix streamlines the Kernel Builder workflow, allowing us to efficiently and reliably manage complex machine learning kernel builds.

## Commonly Asked Questions

**Q. Why not use Docker or other containerization tools instead of Nix?**

While Docker provides isolation and consistent runtime environments, it doesn't guarantee fully reproducible builds. Factors like base image changes or implicit dependencies can still introduce variability.

Nix focuses on reproducibility through deterministic builds, ensuring the same inputs always produce identical outputs. Its declarative configuration, precise dependency management, and efficient caching also make it well-suited for complex environments and CI/CD workflows.

---

If you want to learn more about Nix, check out the following resources:

## References

- **The Official Nix Manual:**
  - The definitive source for all things Nix, providing comprehensive coverage of its features, commands, and ecosystem.
  - Link: [Nix Manual (nixos.org)](https://nixos.org/manual/nix/stable/)
- **Nix Pills:**
  - A series of blog posts breaking down complex Nix concepts into digestible pieces, ideal for a structured, tutorial-style approach.
  - Link: [Nix Pills (nixos.org)](https://nixos.org/guides/nix-pills/)
- **nix.dev**:
  - Home of official documentation for the Nix ecosystem.
  - Link [nix.dev](https://nix.dev/)
- **NixOS Wiki:**
  - A community-driven wiki with a wealth of information, including tips, tricks, and tutorials, covering a wide range of topics, including NixOS-specific information.
  - Link: [NixOS Wiki](https://nixos.wiki/wiki/Main_Page)

### Write kernels
https://huggingface.co/docs/kernels/pr_751/builder/writing-kernels.md

# Write kernels

> [!TIP]
> Refer to the [kernel requirements' page](../kernel-requirements) to
> get an idea of what is expected from the kernel structure and content. 

## Introduction

The Kernel Hub allows Python libraries and applications to load compute
kernels directly from the [Hub](https://hf.co/). To support this kind
of dynamic loading, Hub kernels differ from traditional Python kernel
packages in that they are made to be:

- Portable: a kernel can be loaded from paths outside `PYTHONPATH`.
- Unique: multiple versions of the same kernel can be loaded in the
  same Python process.
- Compatible: kernels must support all recent versions of Python and
  the different PyTorch build configurations (various CUDA versions
  and C++ ABIs). Furthermore, older C library versions must be supported.

`kernel-builder` is a set of tools that can build conforming kernels. It
takes care of:

- Building kernels for all supported PyTorch configurations (C++98/11 and
  different CUDA versions).
- Compatibility with old glibc and libstdc++ versions, so that kernels also
  work on older Linux distributions.
- Registering Torch ops, such that multiple versions the same kernel can be
  loaded without namespace conflicts.

`kernel-builder` builds are configured through a `build.toml` file.
`build.toml` is a simple format that does not require intricate knowledge
of CMake or setuptools.

This page describes the directory layout of a kernel-builder project, the
format of the `build.toml` file, and some additional Python glue that
`kernel-builder` provides. We will use a [simple ReLU kernel](https://github.com/huggingface/kernels/tree/main/examples/kernels/relu)
as the running example. After reading this page, you may also want to have
a look at the more realistic [ReLU kernel with backprop and `torch.compile`](https://github.com/huggingface/kernels/tree/main/examples/kernels/relu-backprop-compile)
support.

> [!TIP]
> We maintain a set of conforming kernels in the
> [kernels-community repository](https://github.com/huggingface/kernels-community).
> We try to keep these kernels synced with upstream as much as possible.

## Setting up environment

### Quick install

The fastest way to get started is to run the install script. This
installs [Determinate Nix](https://docs.determinate.systems/determinate-nix/)
and `kernel-builder` in a single command:

```bash
curl -fsSL https://raw.githubusercontent.com/huggingface/kernels/main/install.sh | bash
```

This will:

1. Install Determinate Nix (if not already installed).
2. Configure the Hugging Face binary cache (to avoid building dependencies from
   source).
3. Install `kernel-builder` via `nix profile install`.

To update `kernel-builder` later:

```bash
nix profile upgrade --all
```

For a step-by-step breakdown of what the script does, see
[Using the kernel builder with Nix](nix).

### Cloud environment

In the [`terraform`](https://github.com/huggingface/kernels/tree/main/terraform) directory, we provide an
example of programatically spinning up an EC2 instance that is ready
with everything needed for you to start developing and building
kernels.

If you use a different provider, the Terraform bridges should be
similar and straightforward to modify.

## Starting a new kernel

The easiest way to start a new kernel is by using the `init` subcommand
of `kernel-builder`. This creates a minimal, compilable kernel:

```bash
$ kernel-builder init --name myorg/mykernel
Initialized `myorg/mykernel` at /home/daniel/git/kernels/examples/kernels/mykernel
```

This creates a kernel named `mykernel` in the directory `mykernel`. The
kernel is configured to upload to the `myorg/mykernel` Hub
repository when an upload command is used.

By default, the `init` subcommand creates a CUDA kernel. You can specify
another backend with the `--backends` option:

```bash
$ kernel-builder init --name myorg/mykernel --backends xpu
```

You can also make a multi-backend kernel by adding all the backends
that you would like to support as arguments to `--backends`:

```bash
$ kernel-builder init --name myorg/mykernel --backends cuda xpu
Initialized `myorg/mykernel` at /home/daniel/git/kernels/examples/kernels/mykernel
```

Finally, if you want to create a kernel for all supported backends, you
can use `--backends all`.

## Kernel project layout

Kernel projects follow this general directory layout:

```text
mykernel
├── benchmarks
│   └── benchmark.py
├── build.toml
├── CARD.md
├── example.py
├── flake.nix
├── mykernel_cuda
│   └── mykernel.cu
├── tests
│   ├── __init__.py
│   └── test_mykernel.py
└── torch-ext
├── mykernel
│   └── __init__.py
├── torch_binding.cpp
└── torch_binding.h
```

In this example we can find:

- The build configuration in `build.toml`.
- One or more top-level directories containing kernels (`mykernel_cuda`).
- The `torch-ext` directory, which contains:
  - `torch_binding.h`: contains declarations for kernel entry points
    (from `kernel_a` and `kernel_b`).
  - `torch_binding.cpp`: registers the entry points as Torch ops.
  - `torch_ext/mykernel`: contains any Python wrapping the kernel needs. At the
    bare minimum, it should contain an `__init__.py` file.
- Kernel tests in the directory `tests`.
- Benchmarks in the directory `benchmarks`.
- A kernel card template in `CARD.md`. This placeholders in the card are filled
  during the kernel build.
- The Nix flake configuration in `flake.nix`.
- An example script that uses the kernel in `example.py`.

## `build.toml`

`build.toml` tells `kernel-builder` what to build and how. It looks as
follows for the `mykernel` kernel:

```toml
[general]
backends = [
  "cuda",
]
name = "mykernel"
version = 1
edition = 5

[general.hub]
repo-id = "myorg/mykernel"

[torch]
src = [
  "torch-ext/torch_binding.cpp",
  "torch-ext/torch_binding.h",
]

[kernel.mykernel]
backend = "cuda"
depends = ["torch"]
src = ["mykernel_cuda/mykernel.cu"]
# If the kernel is only supported on specific capabilities, set the
# cuda-capabilities option:
#
# cuda-capabilities = [ "9.0", "10.0", "12.0" ]
```

The following sections enumerate all supported options for `build.toml`.

### `general`

- `name` (required): the name of the kernel. The Python code for a Torch
  extension must be stored in `torch-ext/<name>`.
- `version` (int): the major version of the kernel.
  The version is written to the kernel's `metadata.json` and is used
  by the `kernel-builder upload` command to upload the kernel to a version
  branch named `v<version>`.
- `edition` (required): the `build.toml` format edition. The current
  edition is `5`. Older `build.toml` files can be migrated with
  `kernel-builder update-build`.
- `backends` (required): a list of supported backends. Must be one or
  more of `cpu`, `cuda`, `metal`, `rocm`, or `xpu`.
- `upstream`: Git-compatible URL (passable to `git clone`) of the original
  upstream repository where the kernel source code comes from.
- `source`: Git-compatible URL (passable to `git clone`) of the kernel-builder
  formatted source repository. This repository must contain a `build.toml` and
  `flake.nix` so that it can be pulled and built with the kernel builder.
- `python-depends` (**experimental**): a list of additional Python dependencies
  that the kernel requires. The only supported dependencies are `einops`,
  `helion`, and `nvidia-cutlass-dsl`.

### `general.hub`

- `repo-id`: the Hub repository to upload the kernel to when the `upload` or
  `build-and-upload` subcommands of `kernel-builder` are used.

### `general.cuda`

- `maxver`: the maximum CUDA toolkit version (inclusive). This option
  _must not_ be set under normal circumstances, since it can exclude Torch
  build variants that are [required for compliant kernels](../kernel-requirements).
  This option is provided for kernels that cause compiler errors on
  newer CUDA toolkit versions.
- `minver`: the minimum required CUDA toolkit version. This option
  _must not_ be set under normal circumstances, since it can exclude Torch
  build variants that are [required for compliant kernels](../kernel-requirements).
  This option is provided for kernels that require functionality only
  provided by newer CUDA toolkits.

### Framework sections

The framework section specifies framework-specific settings. The name of
the section depends on the framework that is used. The currently supported
frameworks are:

- AOT-compiled Torch kernel (`torch`).
- AOT-compiled TVM-FFI kernel (`tvm-ffi`).
- JIT-compiled or not-compiled Torch kernel (`torch-noarch`, experimental).

### `torch`

This framework section is used for AOT-compiled Torch kernels, and has the
following options:

- `src` (required): a list of source files and headers.
- `pyext` (optional): the list of extensions for Python files. Default:
  `["py", "pyi"]`.
- `include` (optional): include directories relative to the project root.
  Default: `[]`.
- `maxver` (optional): only build for this Torch version and earlier. Use cautiously, since this option produces
  non-compliant kernels if the version range does not correspond to the [required variants](build-variants).
- `minver` (optional): only build for this Torch version and later. Use cautiously, since this option produces
  non-compliant kernels if the version range does not correspond to the [required variants](build-variants).
- `stable-abi` (**experimental**): a table mapping backend names to the Torch
  version (e.g. `"2.11"`) that the backend is built against using the Torch
  stable ABI. This requires that the kernel itself only use
  [stable ABI headers](https://docs.pytorch.org/docs/2.12/notes/libtorch_stable_abi.html).
  For an example, see the [`relu-torch-stable-abi`](https://github.com/huggingface/kernels/tree/main/examples/kernels/relu-torch-stable-abi)
  example kernel.

  Backends that are not listed in the table are built normally (without the
  stable ABI), allowing a kernel to mix stable-ABI and full-ABI backends:

  ```toml
  [torch.stable-abi]
  cuda = "2.11"
  rocm = "2.9"
  ```

  Entries for backends that are not in `[general].backends` are ignored, so a
  backend can be commented out of `[general].backends` for testing without
  having to also remove it from the table.

### `tvm-ffi`

This framework section is used for AOT-compiled TVM-FFI kernels.

- `src` (required): a list of source files and headers.
- `pyext` (optional): the list of extensions for Python files. Default:
  `["py", "pyi"]`.
- `include` (optional): include directories relative to the project root.
  Default: `[]`.

### `torch-noarch`

The `torch-noarch` section is used for JIT-compiled kernels or kernels that
do not require any ahead-of-time compilation (e.g. a kernel that packages plain PyTorch
layers).

Normally, it is expected that this type of kernel runs on all CUDA capabilities
or ROCm architectures. However, for kernels that support only a limited range
of archs, the `cuda-capabilites` and `rocm-archs` options can be used to specify
the supported archs. These are then exported to `metadata.json` for consumption
by e.g. the Hugging Face Hub.

- `pyext` (optional): the list of extensions for Python files. Default:
  `["py", "pyi"]`.
- `cuda-capabilities` (optional): a list of CUDA compute capabilities the
  kernel supports (e.g. `["9.0", "10.0"]`).
- `rocm-archs` (optional): a list of ROCm architectures the kernel supports
  (e.g. `["gfx942"]`).

### `kernel.<name>`

Specification of a kernel with the name `<name>`. Multiple `kernel.<name>`
sections can be defined in the same `build.toml`.
See for example [`kernels-community/quantization`](https://huggingface.co/kernels-community/quantization/)
for an example with multiple kernel sections.

The following options can be set for a kernel:

- `backend` (required): the compute backend of the kernel. The currently
  supported backends are `cpu`, `cuda`, `metal`, `rocm`, and `xpu`.
  **The `cpu` backend is currently experimental and might still change.**
- `depends` (required): a list of dependencies. The supported dependencies
  are listed in [`deps.nix`](https://github.com/huggingface/kernels/blob/main/builder/lib/deps.nix).
- `src` (required): a list of source files and headers.
- `include` (optional): include directories relative to the project root.
  Default: `[]`.

Besides these shared options, the following backend-specific options
are available:

#### cuda

- `cuda-capabilities` (optional): a list of CUDA capabilities that the
  kernel should be compiled for. When absent, the kernel will be built
  using all capabilities that the builder supports. The effective
  capabilities are the intersection of this list and the capabilities
  supported by the CUDA compiler. It is recommended to leave this option
  unspecified **unless** a kernel requires specific capabilities.
- `cuda-flags` (optional): additional flags to be passed to `nvcc`.
  **Warning**: this option should only be used in exceptional circumstances.
  Custom compile flags can interfere with the build process or break
  compatibility requirements.

#### rocm

- `rocm-archs`: a list of ROCm architectures that the kernel should be
  compiled for.

#### xpu

- `sycl-flags`: a list of additional flags to be passed to the SYCL
  compiler.

### cpu

- `cxx-flags`: a list of additional flags to be passed to the C++
  compiler.

## Torch bindings

### Defining bindings

Torch bindings are defined in C++, kernels commonly use two files:

- `torch_binding.h` containing function declarations.
- `torch_binding.cpp` registering the functions as Torch ops.

For instance, the `mykernel` kernel discussed above has the following
declaration in `torch_binding.h`:

```cpp
#pragma once

#include <torch/torch.h>

void mykernel(torch::Tensor &out, torch::Tensor const &input);
```

This function is then registered as a Torch op in `torch_binding.cpp`:

```cpp
#include <torch/library.h>

#include "registration.h"
#include "torch_binding.h"

TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
  ops.def("mykernel(Tensor! out, Tensor input) -> ()");
#if defined(CUDA_KERNEL) || defined(ROCM_KERNEL)
  ops.impl("mykernel", torch::kCUDA, &mykernel);
#endif
}

REGISTER_EXTENSION(TORCH_EXTENSION_NAME)
```

This snippet uses macros from `registration.h` to register the function.
`registration.h` is generated by `kernel-builder` itself. A function
is registered through the `def`/`ops` methods. `ops` specifies the
function signature following the [function schema](https://github.com/pytorch/pytorch/blob/main/aten/src/ATen/native/README.md#func).
`impl` associates the function name with the C/C++ function and
the applicable device.

## Using kernel functions from Python

The bindings are typically wrapped in Python code in `torch_ext/<name>`.
The native code is exposed under the `torch.ops` namespace. However,
we add some unique material to the name of the extension to ensure that
different versions of the same extension can be loaded at the same time.
As a result, the extension is registered as
`torch.ops.<name>_<unique_material>`.

To deal with this uniqueness, `kernel_builder` generates a Python module
named `_ops` that contains an alias for the name. This can be used to
refer to the correct `torch.ops` module. For example:

```python
from typing import Optional

import torch

from ._ops import ops

def mykernel(x: torch.Tensor, out: Optional[torch.Tensor] = None) -> torch.Tensor:
    if out is None:
        out = torch.empty_like(x)
    ops.mykernel(out, x)
    return out
```

## Registering Torch operators

You may want to register Torch ops from your kernel's Python code or
[register](https://docs.pytorch.org/tutorials/advanced/custom_ops_landing_page.html) fake ops for `torch.compile` support. It is important to register
such ops in the namespace that kernel-builder makes for your kernel
build. This is required for compliant kernels to ensure that multiple
versions of the same kernel can be loaded at the same time without
namespace conflicts.

You can use the `add_op_namespace_prefix` to prefix an op name with the
correct prefix. So for instance, replace

```python
@torch.library.register_fake("relu::relu_fwd")
def relu_fwd_fake(input: torch.Tensor) -> torch.Tensor:
    return torch.empty_like(input)
```

by

```python
from ._ops import add_op_namespace_prefix

@torch.library.register_fake(add_op_namespace_prefix("relu_fwd"))
def relu_fwd_fake(input: torch.Tensor) -> torch.Tensor:
    return torch.empty_like(input)
```

As mentioned in the above, the `_ops` module is generated by kernel-builder.

kernel-builder uses a hook to reject incorrect usage of Torch op registration
functions. However, it can only catch direct use of certain `torch.library`
decorators. For instance, the hook would not reject the following decorator,
so it should be seen as a last-resort check if human review failed:

```python
@some_indirection_for_register_fake("relu::relu_fwd")
def relu_fwd_fake(input: torch.Tensor) -> torch.Tensor:
    return torch.empty_like(input)
```

> [!WARNING]
> To facilitate static analysis of ops by kernel-builder, `add_op_namespace_prefix` should not be rewrapped. Furthermore, no fallbacks should be added when importing `add_op_namnespace_prefix`, since such fallbacks can mask issues (e.g. incorrect import paths), resulting in non-unique op names. Below is an example of this antipattern:

```py
try:
    from ._ops import add_op_namespace_prefix as _generated_add_op_namespace_prefix
except ImportError:
    def _generated_add_op_namespace_prefix(name: str) -> str:
        return name if "::" in name else f"my_kernel::{name}"

def add_op_namespace_prefix(name: str) -> str:
    return _generated_add_op_namespace_prefix(name)
```

## Kernel tests

### Use `get_kernel` in tests

Kernel tests are stored in the `tests` directory. Tests must not use direct
imports, but instead use `get_kernel` to test the kernel as it will be used.
For example:

```python
import kernels
import torch
import torch.nn.functional as F

relu = kernels.get_kernel("kernels-community/relu", version=1)

def test_relu():
    x = torch.randn(1024, 1024, dtype=torch.float32, device=torch.device("cuda"))
    y = relu.relu(x, torch.empty_like(x))
    y_ref = F.relu(x)
    torch.testing.assert_close(y_ref, y)
```

Development shells (`kernel-builder devshell`/`kernel-builder testshell`)
will set the `LOCAL_KERNELS` variable to ensure that the kernel will be
loaded from the development environment.

### Mark CI tests

Since running all kernel tests in CI may be prohibitively expensive, the
`pyproject.toml` generated by the builder adds support for the special
`kernels_ci` PyTest marker that can be used as follows:

```python
import pytest

@pytest.mark.kernels_ci
def test_mykernel():
  ...
```

We recommend that you to pick tests that together would catch most error
cases while running within 60 seconds.

You can run the tests (e.g. in CI) using:

```bash
$ nix run .#ci-test
```

If the kernel supports multiple backends, it will run the test for the
first supported backend that was found, obeying the following order: CUDA,
ROCm, XPU, Metal, CPU. If you would like to the tests for a specific build
variant, you can use `nix run .#ciTests.<variant>`. For instance:

```bash
$ nix run .#ciTests.torch210-cxx11-cpu-x86_64-linux
```

When running the tests on a non-NixOS systems, make sure that
[the CUDA driver library can be found](https://danieldk.eu/Software/Nix/Nix-CUDA-on-non-NixOS-systems#solutions).

## Kernel docs

We provide a utility to generate a system card for a given kernel, utilizing
information from its `build.toml` and metadata. This system card provides a
reasonable starting point and is meant to be edited afterward by the kernel
developer.

The template card is generated as a part of `kernel-builder init`
command and is serialized in the root directory of the kernel.

The card will be filled automatically by the builder when using the
`build-and-upload` or `build-and-copy` command. It will be serialized
to the `build` sub-directory inside the main kernel directory. It
will be uploaded as `README.md` to the Hub.

### Develop kernels with agents
https://huggingface.co/docs/kernels/pr_751/builder/agents-guide.md

# Develop kernels with agents

Code agents are a good fit to build custom kernels because the hard part is not just writing in Domain Specific Language (DSLs) like CUDA. You also need the right project layout, PyTorch bindings, architecture-specific choices, model-specific integration, and trustworthy benchmarks.

Kernels on Hugging Face are compatible with agents via skills and the `hf` CLI. The `cuda-kernels`, `rocm-kernels`, `xpu-kernels`, `cpu-kernels`, and `triton-kernels` skills contain knowledge so an agent can generate and publish a complete kernel project, instead of isolated snippets.

This guide is for **authoring new kernels**. If you only want to **load an existing precompiled kernel**, use `get_kernel()` instead.

## Before you start

You need:

- a coding agent that supports skills, such as Claude Code, Codex, Cursor, or OpenCode
- a clear target: library, model, operation, GPU, dtype, and representative shapes

The skill currently focuses on NVIDIA GPUs such as **H100**, **A100**, and **T4**, and on integration patterns for **transformers** and **diffusers**.

Install the skill into your agent. If you need the latest version from `main`, use:

```shell
cargo install --git https://github.com/huggingface/kernels hf-kernel-builder

# Install your skills. Use --claude, --codex, or --opencode
kernel-builder skills add --claude
```

> [!NOTE]
> Check [this example](https://github.com/burtenshaw/kernel-skill/tree/main/examples/ltx_video) to see what generated kernels look like.

## 1. Give the agent a precise task prompt

Writing kernels is a hard problem, so be specific to agents. A robust prompt will declare all core attributes, including:

- the library, for example `transformers` or `diffusers`
- the model id, for example `Qwen3-8B` or `LTX-Video`
- the operation, for example `RMSNorm`, attention, RoPE, `GEGLU`, or `AdaLN`
- the target GPU, for example `H100`, `A100`, or `T4`
- the dtype, for example `bfloat16`, `float16`, or `float32`
- the outputs you expect: kernel code, bindings, tests, and benchmarks

In practice, you can often skip some of these and the agent will infer based on common practice, but if you know a detail declare it.

For example:

```
Build a vectorized RMSNorm kernel for H100 targeting Qwen3-8B in transformers.
Create the full kernel-builder project, PyTorch bindings, correctness tests, and benchmark scripts.
```

Or for diffusers:

```
Build an H100 RMSNorm kernel for LTX-Video in diffusers.
Patch the pipeline correctly, benchmark it against the PyTorch baseline, and report end-to-end impact.
```

If you prefer, you can first scaffold a project with `kernel-builder init --name <org>/<kernel>` and then ask the agent to fill in the implementation.

## 2. Verify that the agent produces a complete kernel project

A useful result is a full `kernel-builder` project, not just a `.cu` file. The exact layout can vary, but it should include at least:

```
examples/your_model/
├── kernel_src/
│   └── rmsnorm.cu              # Vectorized CUDA kernel
├── torch-ext/
│   ├── your_kernels/__init__.py
│   └── torch_binding.cpp       # PyTorch C++ bindings
├── benchmark_rmsnorm.py        # Micro-benchmark script
├── build.toml                  # kernel-builder config
├── setup.py                    # python setup.py build_kernel
└── pyproject.toml
```

The agent skills contain example scipts to help you verify the project. So you can briefly test it yourself by running:

```
Verify the kernel project works with a transformers example.
```

## 3. Review the generated files

Let's dive deeper into the generated files, and explore how to validate the project.

### `build.toml`

This is the main configuration file for `kernel-builder`. It tells `kernel-builder` what to build and how so it should contain all the core information about your kernel project.

```
[general]
name = "your_kernels"
backends = ["cuda"]
version = 1

[torch]
src = ["torch-ext/torch_binding.cpp"]

[kernel.rmsnorm]
backend = "cuda"
src = ["kernel_src/rmsnorm.cu"]
depends = ["torch"]
cuda-capabilities = ["9.0"]  # H100
```

First check that:

- `backends = ["cuda"]` is correct for your project
- the kernel source files are listed correctly
- the Torch binding sources are included under `[torch]`
- `cuda-capabilities` is only set when the kernel truly targets specific architectures

For architecture-specific kernels, typical capability values are:

- H100: `9.0`
- A100: `8.0`
- T4: `7.5`

If the kernel does **not** require a specific capability, the kernels docs recommend leaving `cuda-capabilities` unset so the builder can target all supported capabilities. In practice, you can prompt your agent to review the `build.toml` for excessive definitions. Agents have a tendency to over-specify capabilities.

### Torch bindings

The kernel should be registered as Torch ops in `torch-ext/torch_binding.cpp`, with declarations in a header and a small Python wrapper in `torch-ext/<name>/__init__.py`. This is what makes the kernel callable from Python and is the right foundation for `torch.compile` compatibility.

### Model integration code

Make sure the integration matches the library:

- **transformers**: patch the target modules directly, often RMSNorm modules whose class name contains `RMSNorm`
- **diffusers**: inspect the actual pipeline structure before patching, because modules and attention processors can differ across pipelines

> [!NOTE]
> One common issue is that the agent will not integrate the kernel at all. Typically because the project's context is so long.

A few patterns matter in practice for the integration code:

- In **transformers**, RMSNorm modules generally have weights, but epsilon may be exposed as `variance_epsilon` or `eps` depending on the model.
- In **diffusers**, some RMSNorm modules may have `weight=None`, so the integration code needs to handle both weighted and unweighted cases.
- In **diffusers**, checking `type(module).__name__` is often more reliable than `isinstance(...)` for matching RMSNorm modules across implementations.
- If a diffusers pipeline uses CPU offloading, inject custom kernels **before** enabling offload.

For attention, prefer the model library's existing optimized path when one already exists. For example, in `transformers`, Flash Attention 2 is usually the right baseline for attention, while custom kernels are especially useful for operations like RMSNorm and other targeted hotspots.

## 5. Build and test, and benchmark

Kernel Hub kernels must support all recent PyTorch and CUDA configurations. The kernel-builder Nix flake handles this automatically. This step needs both Nix and `kernel-builder`; the easiest way to get both is the [quick install script](writing-kernels#quick-install). Copy the [example `flake.nix`](https://github.com/huggingface/kernels/blob/main/builder/examples/relu/flake.nix) into your project and run:

```shell
nix flake update
kernel-builder build-and-copy -L
```

This builds the kernel for every required PyTorch/CUDA variant and places the results in `build/`. For faster builds, enable the HuggingFace Nix cache:

```shell
nix run nixpkgs#cachix -- use huggingface
```

## 6. Benchmark

There are two main benchmarks to consider:

1. an isolated kernel micro-benchmark
2. an end-to-end benchmark in the real model or pipeline

The agent will generate both benchmarks based on the agent skills examples. Typically as a script called `benchmark_example.py`. If you have access to the target hardware, you can run it to verify the kernel works. For example, the agent will generat a table like this:

```markdown
| Shape         | Custom (ms) | PyTorch (ms) |  Speedup  |
| :------------ | :---------: | :----------: | :-------: |
| [1x128x4096]  |    0.040    |    0.062     | **1.58x** |
| [1x512x4096]  |    0.038    |    0.064     | **1.69x** |
| [1x1024x4096] |    0.037    |    0.071     | **1.90x** |
| [1x2048x4096] |    0.045    |    0.091     | **2.03x** |
| [1x4096x4096] |    0.071    |    0.150     | **2.12x** |
| [4x512x4096]  |    0.056    |    0.093     | **1.67x** |
| [8x256x4096]  |    0.045    |    0.092     | **2.06x** |
| [1x8192x4096] |    0.109    |    0.269     | **2.47x** |
```

Interpret the results carefully. A kernel can show a large isolated speedup but only a modest end-to-end gain if that operation is a small fraction of total runtime. In the LTX-Video example from [the blog we wrote](https://huggingface.co/blog/custom-cuda-kernels-agent-skills), the generated RMSNorm kernel improved the isolated benchmark by about **1.88x** on average, but end-to-end video generation improved by about **6%**, which matched the fact that RMSNorm accounted for only a small share of total compute.

## 7. Publish to the Hub

Once the project is correct and benchmarked, you can build Hub-compatible artifacts and upload them. For this, you should first push to the Hub using the `hf` CLI tool:

```shell
# install the hf CLI tool
hf skills add

# Authenticate
hf auth login

# Push to the Hub
<agent-prompt>
Push the kernel to the Hub.
</agent-prompt>
```

Or, you can manually create the repository and upload the artifacts:

```shell
# Create the repository
hf repo create your-org/your-kernel --type model

# Upload the artifacts
# Run inside the main kernel directory, where build/ is.
kernel-builder upload
```

After pushing to the Hub, users can load the kernel without compiling:

```py
from kernels import get_kernel

kernel = get_kernel("your-org/your-kernel", version=1)
```

Well done! You have now built a custom kernel and published it to the Hub.

### Develop locally
https://huggingface.co/docs/kernels/pr_751/builder/local-dev.md

# Develop locally

## Introduction

`kernel-builder` builds kernels in a sandbox. This has various benefits,
such as building kernels for a wide range of Torch versions, compatibility
with older C library versions and avoiding accidental dependencies.

However, this is not ideal during kernel development, since language
servers and IDEs do not interpret the `build.toml` file. As a result,
code completion will typically not work. `kernel-builder` provides the
`kernel-builder` utility to generate CMake files to build native code and
setuptools files for building the kernel as a regular Python package.
Since CMake and setuptools are widely supported by IDEs, this provides
a much-improved development experience.

## Generating a Python project with `kernel-builder`

`kernel-builder` can create CMake/Python project files for a kernel with
a [`build.toml`](./writing-kernels) file. The `create-pyproject`
command will create the files for the kernel in the current directory:

```bash
$ kernel-builder create-pyproject -f
```

The `-f` flag is optional and instructs `kernel-builder` to overwrite
existing files.

You can build the kernel with

```bash
$ python setup.py build_kernel
```

This builds the kernel and puts the variant that is compatible with the build
environment in `build`. The build can then be loaded directly with `kernels`:

```shell
$ python -c 'import pathlib; import kernels; k = kernels.get_local_kernel(pathlib.Path("build")); print(k)'
```

For AOT kernels, if you want to skip the CMake configuration step in subsequent
builds, you can also run Ninja directly to do incremental builds:

```shell
$ ninja -C _cmake_build local_install
```

You can also create a Python project for a kernel in another directory:

```bash
$ kernel-builder create-pyproject -f path/to/kernel
```

**Warnings:**

- Kernels built in this way should **not** be published on the Kernel
  Hub. They do not fulfill the [kernel requirements](../kernel-requirements).
- Do not add the generated files to Git. `kernel-builder` has regular updates
  and you generally want to use files generated by the latest version.

See [IDE Setup](./ide-setup) for wiring the generated project into
VS Code with direnv.

## Testing kernel builds before publishing

Once you have built a kernel with kernel builder, you may want to test it
locally with software that uses `get_kernel` or `LayerRepository` before
publishing. This can be done using the `LOCAL_KERNELS` variable, which
maps a repository ID to a local kernel directory. For example, you could
use the kernel in `devel/activation` for any use of the
`kernels-community/activation` repository with:

```bash
$ LOCAL_KERNELS="kernels-community/activation=devel/activation" \
  python my_app.py
```

It is also possible to map multiple repositories to local kernel directories
by separating the entries with a colon (`:`):

```bash
$ LOCAL_KERNELS="kernels-community/activation=devel/activation:kernels-community/flash-attn2=devel/flash-attn2" \
  python my_app.py
```

### Building and testing kernels with GitHub Actions
https://huggingface.co/docs/kernels/pr_751/builder/github-actions.md

# Building and testing kernels with GitHub Actions

Compiling a kernel is CPU-intensive and testing it requires an accelerator (such as a GPU), two things that GitHub's standard runners do not provide cheaply. Instead of maintaining self-hosted runners, you can offload both steps to [Hugging Face Jobs](https://huggingface.co/docs/huggingface_hub/guides/jobs) directly from a GitHub Actions workflow.

Two prebuilt actions make this possible:

- [`huggingface/kernel-builder-job`](https://github.com/huggingface/kernel-builder-job) runs the Nix kernel builder on a CPU flavor and publishes the built kernel to the Hub.
- [`huggingface/hf-jobs-action`](https://github.com/huggingface/hf-jobs-action) runs an arbitrary script on any flavor (including GPUs), which is convenient for testing the kernel you just built. This also helps to test the kernel across different hardware.

A typical setup has two workflows: one that **builds** the kernel on push, and one that **tests** it on a GPU. They communicate through the Hub, the build uploads artifacts, the test pulls them back down.

## Prerequisites

- An [HF access token](https://huggingface.co/settings/tokens) with the `job.write` permission. If the test job loads gated models, the token also needs read access to them.
- The token stored as a repository secret named `HF_TOKEN` (**Settings → Secrets and variables → Actions**).
- A kernel repository on the Hub to upload to, with kernel-creation access for the owning user or org (see [Building kernels](build#uploading-your-kernel-to-the-hub)).

> [!NOTE]
> Jobs run under the `namespace` you specify (your username or an org) and
> count against that namespace's compute quota.

## Building on push

The build action checks out nothing by itself, your script clones the exact commit and invokes `kernel-builder`, which comes preinstalled in the default job image. Compilation happens on the HF Jobs CPU flavor, and `build-and-upload` pushes the finished variants to the Hub.

```yaml
# .github/workflows/build-kernel.yml
name: Build Kernel

on:
  push:
    branches: [main]
    paths:
      - "csrc/**"
      - "torch-ext/**"
      - build.toml
      - flake.nix
      - flake.lock
  workflow_dispatch:

jobs:
  build:
    runs-on: ubuntu-latest
    steps:
      - uses: actions/checkout@v4
      - name: Build kernel via HF Jobs
        uses: huggingface/kernel-builder-job@main
        with:
          token: ${{ secrets.HF_TOKEN }}
          namespace: your-username
          flavor: cpu-xl
          timeout: "21600"
          script: |
            # The container starts with `set -x`; disable tracing so the
            # token below is not echoed into the streamed logs.
            set +x
            export HF_TOKEN="${{ secrets.HF_TOKEN }}"
            # Rebuild artifacts, so skip pulling existing LFS blobs.
            export GIT_LFS_SKIP_SMUDGE=1

            git clone "${{ github.server_url }}/${{ github.repository }}" kernel
            cd kernel
            git checkout "${{ github.sha }}"

            kernel-builder build-and-upload \
              --max-jobs 4 \
              --cores 8 \
              --repo-id your-username/your-kernel
```

The path filter keeps the build from running on unrelated commits, and `workflow_dispatch` lets you trigger it by hand from the **Actions** tab. The upload destination is taken from `--repo-id` (or, if omitted, from the `repo-id`/`version` fields in `build.toml`).

> [!TIP]
> Builds can take a long time on the first run because every PyTorch and CUDA
> variant is compiled. Set a generous `timeout` (the example allows six hours)
> and rely on the [Hugging Face binary cache](build#using-the-hugging-face-binary-cache)
> to keep subsequent builds fast.

You can speed up builds by tuning how much work runs in parallel. `--max-jobs`
sets how many kernel variants are built concurrently, while `--cores` sets how
many CPU cores each of those jobs may use. Pick values that fit the chosen CPU
`flavor`: a larger flavor (such as `cpu-xl`) has more cores to spread across
jobs, so raising `--max-jobs` and `--cores` together shortens the total build
time. Setting them too high for the flavor only adds scheduling overhead.

### `kernel-builder-job` inputs

| Input       | Required | Default                                         | Description                                       |
| ----------- | -------- | ----------------------------------------------- | ------------------------------------------------- |
| `token`     | yes      |                                                 | HF token with `job.write` permission.             |
| `namespace` | yes      |                                                 | HF namespace (username or org) that owns the job. |
| `script`    | yes      |                                                 | Shell script to run in the container.             |
| `flavor`    | no       | `cpu-upgrade`                                   | Hardware flavor (e.g. `cpu-xl`).                  |
| `image`     | no       | `ghcr.io/huggingface/kernel-builder-job:latest` | Container image to run the build in.              |
| `timeout`   | no       | `1200`                                          | Maximum seconds to wait for the job.              |

The action exposes `job_id` and `job_url` outputs that link to the run on huggingface.co.

> [!NOTE]
> The default image ships `kernel-builder` on `PATH`, so the script can call it
> directly. Every push to the action's `main` branch republishes `:latest` and
> also pushes an immutable `sha-<commit>` tag, so you can pin a known-good image
> through the `image` input. If you point `image` at your own image that only has
> Nix, either install the CLI first with
> `nix profile install --accept-flake-config github:huggingface/kernels#kernel-builder`
> or invoke it ad-hoc with
> `nix run github:huggingface/kernels#kernel-builder -- build-and-upload ...`.

## Testing on a GPU

Once the kernel is on the Hub, the generic jobs action runs a test script on a GPU flavor. The `files` input copies repository files into the container (under `/tmp/files` by default), and a [`uv` script](https://docs.astral.sh/uv/guides/scripts/) with inline dependencies keeps the environment self-contained.

```yaml
# .github/workflows/run-tests.yml
name: Run tests

on:
  push:
    branches: [main]
    paths:
      - scripts/test.py
  workflow_dispatch:

jobs:
  run:
    runs-on: ubuntu-latest
    steps:
      - uses: actions/checkout@v4
      - name: Run test.py on an HF Jobs GPU
        uses: huggingface/hf-jobs-action@main
        with:
          token: ${{ secrets.HF_TOKEN }}
          namespace: your-username
          flavor: rtx-pro-6000
          image: ghcr.io/astral-sh/uv:python3.10-bookworm
          timeout: "3600"
          files: scripts/test.py
          script: |
            set +x
            export HF_TOKEN="${{ secrets.HF_TOKEN }}"
            uv run /tmp/files/test.py
```

The test script pulls the kernel straight from the Hub with the [`kernels`](../basic-usage) library, so it always runs against the artifacts the build workflow just published:

```python
# scripts/test.py
# /// script
# dependencies = ["kernels", "torch"]
# ///
from kernels import get_kernel

kernel = get_kernel("your-username/your-kernel")
# ... exercise the kernel and assert on the results ...
```

Run the build workflow before the test workflow so the Hub has a fresh kernel to pull. For tightly coupled steps, you can also have one workflow trigger the other, or combine both jobs in a single workflow with a `needs:` dependency.

### `hf-jobs-action` inputs

| Input        | Required | Default      | Description                                             |
| ------------ | -------- | ------------ | ------------------------------------------------------- |
| `token`      | yes      |              | HF token with `job.write` permission.                   |
| `namespace`  | yes      |              | HF namespace (username or org) that owns the job.       |
| `image`      | yes      |              | Container image to run.                                 |
| `script`     | yes      |              | Shell script to execute in the container.               |
| `flavor`     | no       | `cpu-basic`  | Hardware flavor (e.g. `rtx-pro-6000`).                  |
| `files`      | no       |              | Newline-separated repo files to copy into the job.      |
| `files_dest` | no       | `/tmp/files` | Directory the files are copied to inside the container. |
| `env`        | no       | `{}`         | Environment variables as a JSON object.                 |
| `timeout`    | no       | `1200`       | Maximum seconds to wait for the job.                    |

## Choosing a flavor

Flavors map to the machine types available on Hugging Face Jobs, CPU flavors such as `cpu-upgrade` and `cpu-xl` for builds, and GPU flavors such as `l4x1`, `a100-large`, `h200`, or `rtx-pro-6000` for tests. Pick the most reasonable GPU that fits your model to keep jobs low cost. The current list and pricing are in the [Hugging Face Jobs documentation](https://huggingface.co/docs/huggingface_hub/guides/jobs).

> [!NOTE]
> HF Jobs currently only offers a few CPU architectures, so the kernel is built
> for whatever architecture the available CPU flavors provide. This is a current
> limitation to keep in mind if you need to target a specific architecture.

> [!WARNING]
> HF Jobs containers start with shell tracing enabled (`set -x`). Always run
> `set +x` before exporting `HF_TOKEN` so the token does not leak into the
> streamed build logs.

### Build with Nix
https://huggingface.co/docs/kernels/pr_751/builder/build.md

# Build with Nix

## Installation

> [!NOTE]
> The [install script](writing-kernels#quick-install) automates
> the Nix and kernel-builder setup described below. Use these manual
> instructions if you prefer step-by-step control.

### Installing Nix

The kernel builder uses Nix for building kernels. You can build or
run the kernels directly if you have Nix installed on your system.
We recommend installing Nix in the following way:

- Linux: use the [official Nix installer](https://nixos.org/download/).
- macOS: use the [Determinate Nix installer](https://docs.determinate.systems/determinate-nix/).
  In addition, Xcode 16.x is currently required to build kernels.

### Using the Hugging Face binary cache

Since the kernel builder depends on many packages (e.g. every supported
PyTorch version), it is recommended to enable the huggingface cache
to avoid expensive rebuilds.

To use the cache, you can either install cachix and configure it:

```bash
# Install cachix and configure the cache
cachix use huggingface
```

Or run it once without installing cachix permanently:

```bash
# Use cachix without installing it
nix run nixpkgs#cachix -- use huggingface
```

### GPU library configuration

The kernel builder also provides Nix development shells with all Torch
and CUDA/ROCm dependencies needed to develop kernels (see below). If
you want to test your kernels inside a Nix development shell and you
are not using NixOS, [make sure that the CUDA driver is visible](https://danieldk.eu/Nix-CUDA-on-non-NixOS-systems#make-runopengl-driverlib-and-symlink-the-driver-library) to Torch.

## Getting started

The easiest way to start a new kernel is using the `kernel-builder init`
subcommand, which is discussed in [Writing Kernels](writing-kernels).
The commands discussed in the following sections will also work on
existing kernel sources that have `build.toml`/`flake.nix`.

## Building a kernel

A kernel can be built with the `kernel-builder build-and-copy` command.
For example:

```bash
cd examples/relu
kernel-builder build-and-copy -L
```

The `-L` option prints out build logs in the terminal, which can be handy
for monitoring the build. The compiled kernel will then be in the local
`build/` directory.

## Shell for local development

`kernel-builder` provides shells for developing kernels. In such a shell,
all required dependencies are available, as well as `kernel-builder` for generating
project files. For example, you can use the development shell to build a
arch (AOT-compiled) kernel:

```bash
$ kernel-builder devshell
# A devshell is opened in which you can run the following commands:
$ kernel-builder create-pyproject
$ cmake -B build-ext
$ cmake --build build-ext --target local_install
```

This will build the kernel and puts the output in the `build` directory
and can be used with the `kernels` library.

Noarch (JIT-compiled) kernels do not use CMake. For this reason, we also
create a `setup.py` that works both for arch and noarch kernels:

```bash
$ kernel-builder devshell
$ kernel-builder create-pyproject
$ python setup.py build_kernel
```

Development shells are available for every build configuration. For
instance, you can get a Torch 2.11 development shell for ROCm kernels
using:

```bash
$ rm -rf .venv  # Remove existing venv if any.
$ kernel-builder devshell --variant torch212-cxx11-rocm71-x86_64-linux
```

For an editor-driven workflow with `direnv` activating the devshell on
`cd`, see [IDE Setup](./ide-setup).

You can list the variants that the kernel supports with the `list-variants`
subcommand:

```bash
$ kernel-builder list-variants
torch29-cxx11-cu129-x86_64-linux
torch210-cxx11-cu126-x86_64-linux
torch210-cxx11-cu128-x86_64-linux
torch210-cxx11-cu130-x86_64-linux
torch210-cxx11-rocm70-x86_64-linux
torch210-cxx11-rocm71-x86_64-linux
torch210-cxx11-cpu-x86_64-linux
torch210-cxx11-xpu20253-x86_64-linux
torch211-cxx11-cpu-x86_64-linux
torch211-cxx11-cu126-x86_64-linux
torch211-cxx11-cu128-x86_64-linux
torch211-cxx11-cu130-x86_64-linux
torch211-cxx11-rocm71-x86_64-linux
torch211-cxx11-rocm72-x86_64-linux
torch211-cxx11-xpu20253-x86_64-linux
```

## Shell for testing a kernel

You can also start a test shell. This will give you a Python interpreter
with the kernel in Python's search path. This makes it more convenient to run
tests:

```bash
cd examples/relu
kernel-builder testshell
python -m pytest tests
```

`testshell` also supports the `--variant` option, so you can test a particular
kernel variant.

## Adding test dependencies to development shells

You can add test dependencies to a development or testing shell. Adapt
the kernel's `flake.nix` to use the `pythonCheckInputs` option:

```nix
{
  description = "Flake for my kernel";

  inputs = {
    builder.url = "github:huggingface/kernels";
  };

  outputs =
    {
      self,
      builder,
    }:
    builder.lib.genKernelFlakeOutputs {
      inherit self;
      path = ./.;

      # The einops and numpy test dependencies are added here:
      pythonCheckInputs = pkgs: with pkgs; [ numpy ];
    };
}
```

The available packages can be found on [search.nixos.org](https://search.nixos.org/packages?channel=25.05&query=python312Packages).

Keep in mind that these additional dependencies will only be available to
the Nix shells, not the final kernel uploaded to the Hub.

## Uploading your kernel to the Hub

Finally, when you are ready to make a kernel release, you can build and
upload a kernel to the Hub:

```bash
$ cd mykernel
$ kernel-builder build-and-upload
```

> [!NOTE]
> Uploads go to a `kernel`-type Hub repository (the first-class kernel
> repository type). The owning user or org must have kernel-creation
> access. Request it from
> [huggingface.co/settings/account](https://huggingface.co/settings/account)
> ("Request Kernels Creation").

Aside from building and uploading the kernel itself, this will also fill
the card template and upload it as `README.md` to the Hub if the card
template is provided in the source repository as `CARD.md`.

The repository to upload to is determined by the `repo-id` and `version`
fields in `build.toml`. For example, with the following `build.toml`, the
kernel will be uploaded to the repository `kernels-community/flash-attn4`
in the `v1` version branch:

```toml
[general]
# ...
version = 1

[general.hub]
repo-id = "kernels-community/flash-attn4"
```

See [Writing Kernels](writing-kernels) for more details on the `build.toml`
format.

### Uploading through a pull request

If you do not have write access to the target repository — for example,
when proposing builds for a kernel maintained by another user or org —
you can pass `--create-pr` to the `upload` or `build-and-upload`
subcommands:

```bash
$ kernel-builder build-and-upload --create-pr
```

Instead of committing directly, this opens a pull request on the Hub for
each branch that has changes: one against `main` for the kernel card and
one against the `v<version>` branch for the build artifacts and
benchmarks. The pull request URLs are printed after the upload, so that
the repository maintainers can be pointed to them for review.

> [!NOTE]
> A pull request can only target an existing branch. If the `v<version>`
> branch does not exist yet and you do not have write access to create
> it, ask a maintainer of the repository to create the branch first.

> [!TIP]
> You can automate building, uploading, and testing kernels on Hugging Face
> Jobs from CI, see [Building and testing kernels with GitHub Actions](./github-actions).

## Updating the kernel build toolchain

The kernel's dependencies are fully pinned down in the `flake.lock` that
is shipped with the kernel. We periodically release new versions of the
build toolchain that includes bug fixes and supports newer Torch and compute backend
versions. To update the kernel build toolchain, run `nix flake update`
in the kernel directory:

```bash
❯ nix flake update
• Added input 'kernel-builder':
    'github:huggingface/kernels/8ad8a5094f1b3c425f70900699ed690d65d878c3?narHash=sha256-m8tBntCIlH/rY4BcIv5X5%2BdtgSS1yQi883Co%2Bj5cudI%3D' (2026-04-09)
• Added input 'kernel-builder/flake-compat':
    'github:edolstra/flake-compat/5edf11c44bc78a0d334f6334cdaf7d60d732daab?narHash=sha256-vNpUSpF5Nuw8xvDLj2KCwwksIbjua2LZCqhV1LNRDns%3D' (2025-12-29)
• Added input 'kernel-builder/flake-utils':
    'github:numtide/flake-utils/11707dc2f618dd54ca8739b309ec4fc024de578b?narHash=sha256-l0KFg5HjrsfsO/JpG%2Br7fRrqm12kzFHyUHqHCVpMMbI%3D' (2024-11-13)
• Added input 'kernel-builder/flake-utils/systems':
    'github:nix-systems/default/da67096a3b9bf56a91d16901293e51ba5b49a27e?narHash=sha256-Vy1rq5AaRuLzOxct8nz4T6wlgyUR7zLU309k9mBC768%3D' (2023-04-09)
• Added input 'kernel-builder/nixpkgs':
    'github:NixOS/nixpkgs/2f4fd5e1abf9bac8c1d22750c701a7a5e6b524c6?narHash=sha256-Mh6bLcYAcENBAZk3RoMPMFCGGMZmfaGMERE4siZOgP4%3D' (2026-03-31)
• Added input 'kernel-builder/rust-overlay':
    'github:oxalica/rust-overlay/962a0934d0e32f42d1b5e49186f9595f9b178d2d?narHash=sha256-JMdDYn0F%2BswYBILlpCeHDbCSyzqkeSGNxZ/Q5J584jM%3D' (2026-03-31)
• Added input 'kernel-builder/rust-overlay/nixpkgs':
    follows 'kernel-builder/nixpkgs'
```

## Skipping the `get_kernel` check

`kernel-builder` verifies that a kernel can be
imported with the [`kernels`](https://github.com/huggingface/kernels)
package. This check can be disabled by passing `doGetKernelCheck = false`
to `genKernelFlakeOutputs`. **Warning:** it is strongly recommended to keep
this check enabled, as it is one of the checks that validates that a kernel
is compliant. This option is primarily intended for kernels with
`triton.autotune` decorators, which can fail because there is no GPU available
in the build sandbox.

### Set up for Metal kernels
https://huggingface.co/docs/kernels/pr_751/builder/metal.md

# Set up for Metal kernels

Instructions on this page assume that you installed Nix with the
[Determinate Nix installer](https://docs.determinate.systems/determinate-nix/).

## Targeted macOS versions

Since new macOS versions get [adopted quickly](https://telemetrydeck.com/survey/apple/macOS/versions/),
we only support the latest major macOS version except for the first weeks
after a release, when we also support the previous major version.

We currently support macOS 26.0 and later on ARM64 (Apple silicon).

## Requirements

To build a Metal kernel, the following requirements must be met:

- Xcode 26.x must be available on the build machine.
- `xcode-select -p` must point to the Xcode 26 installation, typically
  `/Applications/Xcode.app/Contents/Developer`. If this is not the case,
  you can set the path with:
  `sudo xcode-select -s /path/to/Xcode.app/Contents/Developer`
- The Metal Toolchain must be installed. Starting with macOS 26, this is
  a separate download from Xcode. You can install it with:
  `xcodebuild -downloadComponent MetalToolchain`
- The Nix sandbox should be set to `relaxed`, because the Nix derivation
  that builds the kernel must have access to Xcode and the Metal Toolchain.
  You can verify this by checking that `/etc/nix/nix.custom.conf` contains
  the line:

  ```
  sandbox = relaxed
  ```

  If you had to add the line, make sure to restart the Nix daemon:

  ```
  sudo launchctl kickstart -k system/systems.determinate.nix-daemon
  ```

You can check these requirements as follows. First, you can check the Xcode
version as follows:

```bash
$ xcodebuild -version
Xcode 26.1
Build version 17B55
```

The reported version must be 26.0 or newer. Then you can validate that the
Metal Toolchain is installed with:

```bash
$ xcodebuild -showComponent metalToolchain
Asset Path: /System/Library/AssetsV2/com_apple_MobileAsset_MetalToolchain/68d8db6212b48d387d071ff7b905df796658e713.asset/AssetData
Build Version: 17B54
Status: installed
Toolchain Identifier: com.apple.dt.toolchain.Metal.32023
Toolchain Search Path: /private/var/run/com.apple.security.cryptexd/mnt/com.apple.MobileAsset.MetalToolchain-v17.2.54.0.mDxgz0
```

### Nix builder design
https://huggingface.co/docs/kernels/pr_751/builder/design-nix-builder.md

# Nix builder design

## Introduction

kernel-builder uses a Nix-based builder that orchestrates the build. The Nix
builder provides:

- Reproducible evaluation. The same Nix builder version will always produce
  the same derivations (build recipes).
- Largely reproducible builds by using a build sandbox that only has the
  dependencies specified in a derivation.
- Seamless creation of different build environments (e.g. different Torch
  and CUDA combinations).

## Kernel build steps

A kernel derivation builds a kernel in the following steps:

1. Generate CMake files for the kernel using
   `kernel-builder create-pyproject`.
2. Generate Ninja build files using CMake.
3. Build the kernel using Ninja.
4. Perform various checks on the compiled kernel, such as:
   - Verify that the kernel only uses ABI3/`manylinux_2_28` symbols.
   - Verify that the kernel can be loaded by the `kernels` Python package.
5. Strip runpaths (ELF-embedded library directories) from kernel binaries
   to make the kernel distribution-independent.

## manylinux_2_28 compatibility

To achieve `manylinux_2_28` compatibility, kernels are built using a
toolchain similar to the `manylinux_2_28` Docker images. This toolchain
is based on the gcc toolsets from AlmaLinux 8. `manylinux_2_28` [uses
AlmaLinux 8 as its base](https://github.com/pypa/manylinux#manylinux_2_28-almalinux-8-based),
so we have to compile against the same glibc/libstdc++ versions to
ensure compatibility.

We repackage the AlmaLinux 8 toolsets and libstdc++ as Nix derivations (see
the `nix-builder/packages/manylinux_2_28` source directory). Then we merge
various toolset packages to an unwrapped gcc that resembles unwrapped gcc in
nixpkgs. Finally, we wrap binutils and gcc to combine them into a stdenv.

The stdenv does not reuse glibc from AlmaLinux, since its dynamic loader has
hardcoded FHS paths (`/lib64` etc.) that are not valid in Nix. Using this
dynamic loader results in linking errors, since the paths in the dynamic
loader are used as a last resort (to link glibc libraries). So, instead we
build our own glibc 2.28 package
(see `nix-builder/pkgs/manylinux_2_28/stdenv.nix`) and use that.

## The package set pattern

We repackage various existing package sets as Nix derivations. For instance,
this is done for ROCm, XPU, and manylinux_2_28 packages. We do this because
we want these libraries to be as close as what the user would install. This
avoids compatibility issues between the kernels and the official vendor
packages. For instance, suppose that we built a ROCm library as a shared
library and ROCm provides the same library as a static library, then compiled
kernels could use symbols that cannot be resolved when installing the official
ROCm packages. Similarly, using the official packages allows us to test
against the official upstream packages.

These package sets all follow the same pattern:

```nix
{
  lib,
  callPackage,
  newScope,
  pkgs,
}:

{
  packageMetadata,
}:

let
  inherit (lib.fixedPoints) extends composeManyExtensions;

  fixedPoint = final: {
    inherit lib;
  };
  composed = lib.composeManyExtensions [
    # Base package set.
    (import ./components.nix { inherit packageMetadata; })

    # Package-specific overrides.
    (import ./overrides.nix)

    # Additional overlays that extend the package set.
    (import ./some-overlay.nix)
  ];
in
lib.makeScope newScope (lib.extends composed fixedPoint)
```

We use a fixed point to build up the package set as a list of
[overlays](https://nixos.org/manual/nixpkgs/stable/#sec-overlays-definition).
This has various benefits. For instance, it allows us to refine the
package set incrementally and we can refer to the final versions of
packages in intermediate overlays.

The package sets all use a similar list of overlays:

- An initial overlay (`components.nix`) that applies a generic builder
  to the package set metadata. The metadata typically comes from a Yum/DNF
  repository that contains RPM packages.The generic builder will extract the
  RPMs and move binaries, libraries, and headers to the right location. This
  results in a set of Nix derivations that may or may not build.
- The next overlay (`overrides.nix`) fixes up derivations generated by the
  generic builder in the previous overlay that do not build. Fixing the
  derivations typically consists of adding missing dependencies and changing
  embedded FHS paths to Nix store paths.
- Additional overlays with derivations that combine outputs from previous
  overlays. One typical example are derivations that construct a full compiler
  toolchain (e.g. `nix-builder/pkgs/manylinux_2_28/gcc-unwrapped.nix`).

### Set up your IDE
https://huggingface.co/docs/kernels/pr_751/builder/ide-setup.md

# Set up your IDE

## Introduction

Language servers do not interpret `build.toml`, so IDE completion for
CUDA, ROCm, framework headers, and the kernel's Python wrapper does not
work out of the box. This guide shows how to configure VS Code so that
completion resolves against the same toolchain `kernel-builder`
uses.

The setup has three pieces:

- `kernel-builder create-pyproject` to emit CMake and setuptools files
  the IDE can read (see [Local Development](./local-dev)).
- The kernel-builder devshell, which provides the toolchain (CUDA, ROCm,
  Torch headers, etc.) from the Nix store.
- `direnv` to activate the devshell on `cd`, so VS Code inherits the
  environment through the shell.

Pinning the toolchain through Nix keeps IDE completion aligned with
the build. It also makes switching between CUDA, ROCm, or XPU a
one-line change in `.envrc`.

## Installing direnv and nix-direnv

On non-NixOS systems, install both via `nix profile`:

```bash
$ nix profile install nixpkgs#nix-direnv
```

Add the direnv hook to your shell rc (`~/.bashrc` or
`~/.zshrc`, for example):

```bash
eval "$(direnv hook bash)"    # or: direnv hook zsh
```

Source the rc file (or open a new shell) so the hook is
active in the current session:

```bash
$ source ~/.bashrc            # or: source ~/.zshrc
```

Wire `nix-direnv` into direnv:

```bash
$ mkdir -p ~/.config/direnv
$ echo 'source $HOME/.nix-profile/share/nix-direnv/direnvrc' \
    >> ~/.config/direnv/direnvrc
```

On [NixOS](https://github.com/nix-community/nix-direnv#via-system-configuration-on-nixos)
or with [home-manager](https://github.com/nix-community/nix-direnv#via-home-manager),
enable `programs.direnv` with
`nix-direnv` instead. See
[`terraform/nixos-configuration.nix`](https://github.com/huggingface/kernels/tree/main/terraform/nixos-configuration.nix)
for a working example.

## Activating the devshell with direnv

From the kernel root directory (the one containing `flake.nix` and
`build.toml`), tell direnv to use the flake's default devshell:

```bash
$ echo 'use flake' > .envrc
$ direnv allow
```

`direnv` now activates the default devshell whenever you `cd` into the
project. The devshell's `shellHook` creates and activates a `.venv` on
first entry. Confirm it picked up the toolchain and venv:

```bash
$ which nvcc
/nix/store/.../bin/nvcc
$ ls -ld .venv
drwxr-xr-x ... .venv
$ which python
/path/to/kernel/.venv/bin/python
```

If `.venv` is missing, re-run `direnv reload` and check the output for
the `Creating new venv environment in path: './.venv'` line from the
`shellHook`.

To pin a non-default build variant, name it explicitly:

```bash
$ echo 'use flake .#devShells.torch212-cxx11-rocm71-x86_64-linux' > .envrc
$ direnv allow
```

See [Build Variants](./build-variants) for the variant list.

## Generating IDE-facing project files

direnv puts the toolchain on `PATH`, but the C++ language server still
needs a CMake-derived `compile_commands.json` to resolve per-file
include paths. Generate the CMake/setuptools project and the file:

```bash
$ kernel-builder create-pyproject -f
$ cmake -B build-ext -DCMAKE_EXPORT_COMPILE_COMMANDS=ON
$ ln -sf build-ext/compile_commands.json compile_commands.json
```

`-DCMAKE_EXPORT_COMPILE_COMMANDS=ON` is required: the generated CMake
does not set it. The symlink lets the language server find the file
at the project root.

As noted in [Local Development](./local-dev), do not commit the
generated files.

## Configuring VS Code

Install the [`mkhl.direnv`](https://github.com/direnv/direnv-vscode)
extension. It activates the project's `.envrc` when VS Code opens
the workspace, so language servers and the integrated terminal see
the devshell environment without launching `code` from a shell.

Alternatively, skip the extension and open the project from a
direnv-activated shell — VS Code inherits the environment that way
too:

```bash
$ cd path/to/kernel
$ code .
```

Install one of the following first-party extensions for C++/CUDA
completion:

- `llvm-vs-code-extensions.vscode-clangd` (recommended for CUDA).
- `ms-vscode.cpptools` (Microsoft C/C++).

Add `.vscode/settings.json` (do not commit):

```jsonc
{
  "python.defaultInterpreterPath": "${workspaceFolder}/.venv/bin/python",

  // clangd
  "clangd.arguments": ["--compile-commands-dir=${workspaceFolder}"],

  // Microsoft C/C++ extension
  "C_Cpp.default.compileCommands": "${workspaceFolder}/compile_commands.json",
}
```

Depending on the extension being used, the configuration above behaves
differently:

- With `clangd`, the `clangd.arguments` line is optional. clangd already
  looks in the parent directories of each source file for
  `compile_commands.json` and will find the workspace-root symlink on its
  own ([clangd docs](https://clangd.llvm.org/installation#project-setup)).
  Setting it explicitly does no harm.
- With the Microsoft C/C++ extension, the `C_Cpp.default.compileCommands`
  line is required. The extension does not pick up
  `compile_commands.json` from the workspace root on its own, unless
  another extension (such as CMake Tools) tells it where to look.

To verify, open `torch-ext/torch_binding.cpp` and hover an
`#include <torch/torch.h>` directive. The resolved path should point
into `/nix/store/...`, not a system path.

## Remote development

Use the VS Code Remote-SSH extension and put the direnv hook in the
remote shell's rc. The remote integrated terminal activates the
devshell on `cd`, and VS Code's language servers — which run on the
remote — inherit that environment. The
[`terraform/`](https://github.com/huggingface/kernels/tree/main/terraform)
setup is already configured this way.

## Switching toolchains

Change the `use flake` line in `.envrc` to point at a different
variant. For example:

```bash
# CUDA 13.0
use flake .#devShells.torch212-cxx11-cu130-x86_64-linux

# ROCm 7.1
use flake .#devShells.torch212-cxx11-rocm71-x86_64-linux

# XPU
use flake .#devShells.torch212-cxx11-xpu20253-x86_64-linux
```

Remove `.venv/` first if it was created against a different variant,
then reload direnv to recreate it via the new devshell's `shellHook`:

```bash
$ rm -rf .venv
$ direnv reload
```

## noarch kernels

For Python-only (noarch) kernels, skip the CMake step in "Generating
IDE-facing project files" and the C++ portions of the VS Code
configuration. The `direnv` setup and `python.defaultInterpreterPath`
are all that is needed.

### Build variants
https://huggingface.co/docs/kernels/pr_751/builder/build-variants.md

# Build variants

A kernel can be compliant for a specific compute framework (e.g. CUDA) or
architecture (e.g. x86_64). For compliance with a compute framework and
architecture combination, all the build variants listed below must be
available. This list will be updated as new PyTorch versions are released.

## CPU aarch64-darwin

- `torch212-cpu-aarch64-darwin`
- `torch213-cpu-aarch64-darwin`

## Metal aarch64-darwin

- `torch212-metal-aarch64-darwin`
- `torch213-metal-aarch64-darwin`

## CPU aarch64-linux

- `torch212-cxx11-cpu-aarch64-linux`
- `torch213-cxx11-cpu-aarch64-linux`

## CUDA aarch64-linux

- `torch212-cxx11-cu126-aarch64-linux`
- `torch212-cxx11-cu130-aarch64-linux`
- `torch212-cxx11-cu132-aarch64-linux`
- `torch213-cxx11-cu126-aarch64-linux`
- `torch213-cxx11-cu130-aarch64-linux`
- `torch213-cxx11-cu132-aarch64-linux`

## CPU x86_64-linux

- `torch212-cxx11-cpu-x86_64-linux`
- `torch213-cxx11-cpu-x86_64-linux`

## CUDA x86_64-linux

- `torch212-cxx11-cu126-x86_64-linux`
- `torch212-cxx11-cu130-x86_64-linux`
- `torch212-cxx11-cu132-x86_64-linux`
- `torch213-cxx11-cu126-x86_64-linux`
- `torch213-cxx11-cu130-x86_64-linux`
- `torch213-cxx11-cu132-x86_64-linux`

## ROCm x86_64-linux

- `torch212-cxx11-rocm71-x86_64-linux`
- `torch212-cxx11-rocm72-x86_64-linux`
- `torch213-cxx11-rocm71-x86_64-linux`
- `torch213-cxx11-rocm72-x86_64-linux`

## XPU x86_64-linux

- `torch212-cxx11-xpu20253-x86_64-linux`
- `torch213-cxx11-xpu20260-x86_64-linux`

## Python-only kernels

Kernels that are in pure Python (e.g. Triton kernels) only need to provide
one or more of the following variants:

- `torch-cpu`
- `torch-cuda`
- `torch-metal`
- `torch-rocm`
- `torch-xpu`

### Architecture overview
https://huggingface.co/docs/kernels/pr_751/builder/design-overview.md

# Architecture overview

The Kernel Builder is a Nix flake that combines two components:

![Kernel Builder architecture](https://huggingface.co/datasets/huggingface/documentation-images/resolve/main/kernels/kernel_design_overview.jpg)

- `kernel-builder/`: a Rust-based CLI
  that parses and validates `build.toml`, scaffolds new kernel projects, and drives the
  build. It does not compile kernels itself: every build subcommand executes
  `nix build`/`nix run`/`nix develop` against the kernel project's `flake.nix`.
- `nix-builder`: the Nix builder drives the build itself. It generates a build matrix using the supported build configurations
  (see `nix-builder/versions.nix`) and the kernel's `build.toml`. The build matrix is the cartesian product of
  the backend (one or more of CPU/CUDA/Metal/ROCm/XPU), backend versions, and framework versions.
  Nix builder uses `kernel-builder` to generate the CMake files that drive the build. To do so,
  `kernel-builder` itself is packaged in `nix-builder/pkgs`.

A kernel author uses the Nix builder through the `lib.genKernelFlakeOutputs` that is exposed through the top-level `flake.nix`. This generates the Nix flake outputs for building and developing kernels, such as `bundle` for bundle builds and `devShells` for development shells.

For a deeper look at the design of the Nix builder, see
[Nix Builder design](./design-nix-builder).

### Secure your kernels
https://huggingface.co/docs/kernels/pr_751/builder/security.md

# Secure your kernels

## Introduction

As a kernel builder, you provide code that might be run on thousands or
even millions of machines. This comes with the responsibility of ensuring
no malicious code is distributed.

Below, we provide guidelines to help avoid common attack vectors. These
are _in addition to_ common-sense security practices, such as keeping
your credentials/tokens safe, being vigilant against machine compromise,
and doing proper code reviews.

## Handling pull requests

The Hugging Face Hub allows users to submit pull requests to your
repositories. **Never** merge a pull request that contains a `build/`
directory. The binaries inside the `build/` directory might be compromised
even when the source code looks fine. When a pull request includes
`build/`, ask the submitter to re-submit it without builds. Build the
kernel on your own trusted infrastructure after the PR is merged.

When a PR does not contain build outputs and is ready to review, _carefully_
review every changed line, also taking security into account. Even if the
PR is from a trusted party, review it as if their credentials might have
been compromised.

## Build hygiene

If possible, do builds on a dedicated build machine/VM that is only used
for sandboxed builds (non-macOS kernel-builder builds are sandboxed as
well). Specialized machines are less likely to be compromised, especially
when they are accessed with a hardware-stored SSH key that requires user
interaction.

## Supporting reproducibility

Reproducible builds are very helpful to verify that no malicious code has
slipped into a kernel. If a kernel build is reproducible, then anyone can
rebuild a kernel and verify the binaries match the distributed binaries.
Full reproducibility is a goal we are working toward in `kernel-builder`.

However, this also requires assistance from the kernel builder. This section
describes what you need to do to make reproducible builds possible.

### Only build kernels with Nix sandboxing enabled.

Nix can be used with sandboxing disabled to support systems that do not
support sandboxing (e.g. Linux systems that are configured to disable
mount/network namespaces). **Never** build kernels with sandboxing disabled.
Not only can this cause stray system dependencies to be picked up, but
it can also cause other impurities to slip into the build, making it
impossible to reproduce the build. You can verify that sandboxing is enabled
using `nix-info`:

```bash
$ nix-shell -p nix-info --run "nix-info -m"
 - system: `"x86_64-linux"`
 - host os: `Linux 6.12.39, NixOS, 25.11 (Xantusia), 25.11.20250723.1744f3d`
 - multi-user?: `yes`
 - sandbox: `yes`
 - version: `nix-env (Nix) 2.28.4`
 - nixpkgs: `/nix/store/fqwc3ghi5qfdmzklpwssbamxcqj1vgl3-source`
```

### Do not build from dirty Git trees

Before building a kernel, ensure that all changes are committed. This
makes it possible to reproduce a build from exactly the same source code.
We bake the git shorthash into the ops name, so that it is clear from
which git hash a kernel was built.

The build [`provenance`](kernel-requirements) records a `dirty` flag when
a variant was built from a working tree with uncommitted changes. To keep
these builds from slipping by unnoticed:

- `kernel-builder upload` prints a warning listing every dirty variant it is
  about to publish, and injects a warning banner into the uploaded
  `README.md` (built from `CARD.md`) naming those variants.
- The `kernels` loader emits a warning when a kernel variant it loads was
  built from a dirty tree.

### Kernels API Reference
https://huggingface.co/docs/kernels/pr_751/api/kernels.md

# Kernels API Reference

## Main Functions

### get_kernel[[kernels.get_kernel]]

#### kernels.get_kernel[[kernels.get_kernel]]

```python
kernels.get_kernel(repo_id: str, revision: str | None = None, version: int | None = None, backend: str | None = None, user_agent: str | dict | None = None, trust_remote_code: bool | list[str] = False)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/load.py#L47)

**Parameters:**

repo_id (*str*) : The Hub repository containing the kernel.

revision (*str*, *optional*) : The specific revision (branch, tag, or commit) to download. Cannot be used together with *version*.

version (*int*, *optional*) : The kernel version to download. Cannot be used together with *revision*. Either *version* or *revision* must be specified.

backend (*str*, *optional*) : The backend to load the kernel for. Can only be *cpu* or the backend that Torch is compiled for. The backend will be detected automatically if not provided.

user_agent (*Union[str, dict]*, *optional*) : The *user_agent* info to pass to *snapshot_download()* for internal telemetry.

trust_remote_code (*bool | list[str]*, *optional*, defaults to *False*) : Whether to allow loading kernels from untrusted organisations. When `False`, only kernels from trusted organisations are allowed. When `True`, all repositories are allowed. A list of strings will be used to verify signing identities in a future release; for now it emits a warning and falls back to the default trust check.

**Returns:** `*ModuleType*`

The imported kernel module.

Load a kernel from the kernel hub.

This function downloads a kernel to the local Hugging Face Hub cache directory (if it was not downloaded before)
and then loads the kernel.

Example:
```python
import torch
from kernels import get_kernel

activation = get_kernel("kernels-community/relu", version=1)
x = torch.randn(10, 20, device="cuda")
out = torch.empty_like(x)
result = activation.relu(out, x)
```

### get_local_kernel[[kernels.get_local_kernel]]

#### kernels.get_local_kernel[[kernels.get_local_kernel]]

```python
kernels.get_local_kernel(repo_path: Path, backend: str | None = None)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/load.py#L117)

**Parameters:**

repo_path (`Path`) : The local path to the kernel repository.

backend (`str`, *optional*) : The backend to load the kernel for. Can only be `cpu` or the backend that Torch is compiled for. The backend will be detected automatically if not provided.

**Returns:** `ModuleType`

The imported kernel module.

Import a kernel from a local kernel repository path.

### has_kernel[[kernels.has_kernel]]

#### kernels.has_kernel[[kernels.has_kernel]]

```python
kernels.has_kernel(repo_id: str, revision: str | None = None, version: int | None = None, backend: str | None = None)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/load.py#L153)

**Parameters:**

repo_id (`str`) : The Hub repository containing the kernel.

revision (`str`, *optional*) : The specific revision (branch, tag, or commit) to download. Cannot be used together with `version`.

version (`int`, *optional*) : The kernel version to download. Cannot be used together with `revision`. Either `version` or `revision` must be specified.

backend (`str`, *optional*) : The backend to load the kernel for. Can only be `cpu` or the backend that Torch is compiled for. The backend will be detected automatically if not provided.

**Returns:** `bool`

`True` if a kernel is available for the current environment.

Check whether a kernel build exists for the current environment (Torch version and compute framework).

### get_kernel_variants[[kernels.get_kernel_variants]]

#### kernels.get_kernel_variants[[kernels.get_kernel_variants]]

```python
kernels.get_kernel_variants(repo_id: str, revision: str | None = None, version: int | None = None, backend: str | None = None)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/variants.py#L558)

**Parameters:**

repo_id (`str`) : The Hub repository containing the kernel.

revision (`str`, *optional*) : The specific revision (branch, tag, or commit) to inspect. Cannot be used together with `version`.

version (`int`, *optional*) : The kernel version to inspect. Cannot be used together with `revision`. Either `version` or `revision` must be specified.

backend (`str`, *optional*) : The backend to resolve variants for. Can only be `cpu` or the backend that Torch is compiled for. The backend will be detected automatically if not provided.

**Returns:** `list[Decision]`

One `VariantAccepted` or `VariantRejected` per build variant
in the repository, compatible variants first.

Resolve all build variants of a kernel against the current environment.

The decisions are sorted with compatible variants first, the most preferred
variant leading.

Example:
```python
from kernels import get_kernel_variants, VariantAccepted

for decision in get_kernel_variants("kernels-community/activation", version=1):
    name = decision.variant.variant_str
    if isinstance(decision, VariantAccepted):
        print(f"{name}: compatible")
    else:
        print(f"{name}: rejected ({decision.reason})")

```

### get_loaded_kernels[[kernels.get_loaded_kernels]]

#### kernels.get_loaded_kernels[[kernels.get_loaded_kernels]]

```python
kernels.get_loaded_kernels()
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/importer.py#L45)

**Returns:** `list[LoadedKernel]`

One [LoadedKernel](/docs/kernels/pr_751/en/api/kernels#kernels.LoadedKernel) per distinct kernel variant path
loaded in this process.

Return a snapshot of every kernel that has been loaded into the current process.

The returned list is a new list; mutating it does not affect the registry.

Example:
```python
from kernels import get_kernel, get_loaded_kernels

get_kernel("kernels-community/activation", version=1)
for loaded in get_loaded_kernels():
    print(loaded.metadata.name, loaded.repo_info)
```

## Loading locked kernels

### load_kernel[[kernels.load_kernel]]

#### kernels.load_kernel[[kernels.load_kernel]]

```python
kernels.load_kernel(repo_id: str, lockfile: pathlib.Path | None, backend: str | None = None)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/load.py#L194)

**Parameters:**

repo_id (`str`) : The Hub repository containing the kernel.

lockfile (`Path`, *optional*) : Path to the lockfile. If not provided, the lockfile will be loaded from the caller's package metadata.

backend (`str`, *optional*) : The backend to load the kernel for. Can only be `cpu` or the backend that Torch is compiled for. The backend will be detected automatically if not provided.

**Returns:** `ModuleType`

The imported kernel module.

Get a pre-downloaded, locked kernel.

If `lockfile` is not specified, the lockfile will be loaded from the caller's package metadata.

### get_locked_kernel[[kernels.get_locked_kernel]]

#### kernels.get_locked_kernel[[kernels.get_locked_kernel]]

```python
kernels.get_locked_kernel(repo_id: str)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/load.py#L239)

**Parameters:**

repo_id (`str`) : The Hub repository containing the kernel.

local_files_only (`bool`, *optional*, defaults to `False`) : Whether to only use local files and not download from the Hub.

**Returns:** `ModuleType`

The imported kernel module.

Get a kernel using a lock file.

## Classes

### LoadedKernel[[kernels.LoadedKernel]]

#### kernels.LoadedKernel[[kernels.LoadedKernel]]

```python
kernels.LoadedKernel(metadata: Metadata, module: module, repo_info: kernels.hf_hub.RepoInfo | None)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/importer.py#L14)

This dataclass provides information about a loaded kernel:

- `metadata` (`Metadata`): kernel metadata.
- `module` (`ModuleType`): the imported kernel module.
- `repo_info` (`kernels.hf_hub.RepoInfo | None`): populated only for
  kernels loaded via `get_kernel`. Loaders that work from a local path
  (`get_local_kernel`) or a lockfile (`get_locked_kernel`, `load_kernel`)
  leave this as `None`.

The metadata includes the following properties that describe a kernel:

- `id` (`str`): kernel identifier that is unique to the kernel version + backend.
- `name` (`str`): the name of the kernel.
- `version` (`int`): the version of the kernel.
- `license` (`str`): the license of the kernel.
- `upstream` (`str | None`): the original upstream repository of the kernel.
- `source` (`str | None`): the kernel-builder formatted source repository.
- `python_depends` (`list[str]`): required Python dependencies.
- `backend`: information about the kernel's backend.

### RepoInfo[[kernels.RepoInfo]]

#### kernels.RepoInfo[[kernels.RepoInfo]]

```python
kernels.RepoInfo(repo_id: str, revision: str)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/hf_hub.py#L81)

This dataclass stores the origin of the kernel.

The following fields are available:

- `repo_id` (`str`): the Hub repository containing the kernel.
- `revision` (`str`): the specific revision of the kernel.

### Layers API Reference
https://huggingface.co/docs/kernels/pr_751/api/layers.md

# Layers API Reference

## Making layers kernel-aware

### use_kernel_forward_from_hub[[kernels.use_kernel_forward_from_hub]]

#### kernels.use_kernel_forward_from_hub[[kernels.use_kernel_forward_from_hub]]

```python
kernels.use_kernel_forward_from_hub(layer_name: str)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/layer/layer.py#L272)

**Parameters:**

layer_name (`str`) : The name of the layer to use for kernel lookup in registered mappings.

**Returns:** `Callable`

A decorator function that can be applied to layer classes.

Decorator factory that makes a layer extensible using the specified layer name.

This decorator which prepares a layer class to use kernel layers from the Hugging
Face Hub.

When applied to a function, the function is converted into a layer (`nn.Module`),
made extensible using the given layer name, and then the class is instantiated.
Since `nn.Module` also implements the `__call__` method, the module can still be
used as if it was a function. Note that a decorated function is only visible to
[kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize) if it is attached to a module using [use_kernelized_func()](/docs/kernels/pr_751/en/api/layers#kernels.use_kernelized_func).

Example:
```python
import torch
import torch.nn as nn

from kernels import use_kernel_forward_from_hub
from kernels import use_kernelized_func
from kernels import Mode, kernelize

@use_kernel_forward_from_hub("MyCustomLayer")
class MyCustomLayer(nn.Module):
    def __init__(self, hidden_size):
        super().__init__()
        self.hidden_size = hidden_size

    def forward(self, x: torch.Tensor):
        # original implementation
        return x

model = MyCustomLayer(768)

# The layer can now be kernelized:
# model = kernelize(model, mode=Mode.TRAINING | Mode.TORCH_COMPILE, device="cuda")

# Use on a function (converts the function to `nn.Module`). The function
# can then be replaced with a layer mapping for `MyCustomLayer`.
@use_kernel_forward_from_hub("MyCustomLayer")
def identity(x: torch.Tensor) -> torch.Tensor:
    return x

@use_kernelized_func(identity)
class LayerUsingIdentity(nn.Module):
    def forward(self, x: torch.Tensor) -> torch.Tensor:
        return identity(x)

```

### use_kernel_func_from_hub[[kernels.use_kernel_func_from_hub]]

#### kernels.use_kernel_func_from_hub[[kernels.use_kernel_func_from_hub]]

```python
kernels.use_kernel_func_from_hub(func_name: str)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/layer/func.py#L198)

**Parameters:**

func_name (`str`) : The name of the function name to use for kernel lookup in registered mappings.

**Returns:** `Callable`

A decorator function that can be applied to layer classes.

Decorator that makes a function extensible using the specified function name.

This is a decorator factory that returns a decorator which prepares a function to use kernels from the
Hugging Face Hub.

The function will be exposed as an instance of `torch.nn.Module` in which
the function is called in `forward`. For the function to be properly
kernelized, it **must** be a member of another `torch.nn.Module` that is
part of the model (see the example).

> [!WARNING]
> `use_kernel_func_from_hub` is deprecated and will be removed in kernels 0.17.
> Use [use_kernel_forward_from_hub()](/docs/kernels/pr_751/en/api/layers#kernels.use_kernel_forward_from_hub) instead.

Example:
```python
import torch
import torch.nn as nn

from kernels import use_kernel_func_from_hub
from kernels import Mode, kernelize

@use_kernel_func_from_hub("my_custom_func")
def my_custom_func(x: torch.Tensor):
    # Original implementation
    return x

class MyModel(torch.nn.Module):
    def __init__(self):
        super().__init__()
        self.fn = my_custom_func

    def forward(self, x):
        return self.fn(x)

model = MyModel()

# The layer can now be kernelized:
# model = kernelize(model, mode=Mode.TRAINING | Mode.TORCH_COMPILE, device="cuda")
```

### use_kernelized_func[[kernels.use_kernelized_func]]

#### kernels.use_kernelized_func[[kernels.use_kernelized_func]]

```python
kernels.use_kernelized_func(*args: Callable)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/layer/layer.py#L344)

**Parameters:**

- ***args** (`Callable`) : Kernel functions to attach to the module.

**Returns:** `Callable`

A decorator function that can be applied to modules.

This decorator attaches the target function within the module as a plain
attribute (not as a submodule). This makes the function visible to
[kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize).

Example:
```python
import torch
import torch.nn as nn

from kernels import use_kernel_forward_from_hub
from kernels import use_kernelized_func
from kernels import Mode, kernelize

# Use on a function (converts the function to `nn.Module`). The function
# can then be replaced with a layer mapping for `MyCustomLayer`.
@use_kernel_forward_from_hub("MyCustomLayer")
def identity(x: torch.Tensor) -> torch.Tensor:
    return x

@use_kernelized_func(identity)
class LayerUsingIdentity(nn.Module):
    def forward(self, x: torch.Tensor) -> torch.Tensor:
        return identity(x)

model = LayerUsingIdentity()

# The layer can now be kernelized:
# model = kernelize(model, mode=Mode.TRAINING | Mode.TORCH_COMPILE, device="cuda")
```

### replace_kernel_forward_from_hub[[kernels.replace_kernel_forward_from_hub]]

#### kernels.replace_kernel_forward_from_hub[[kernels.replace_kernel_forward_from_hub]]

```python
kernels.replace_kernel_forward_from_hub(layer_name: str)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/layer/layer.py#L249)

Function that prepares a layer class to use kernels from the Hugging Face Hub.

It is recommended to use [use_kernel_forward_from_hub()](/docs/kernels/pr_751/en/api/layers#kernels.use_kernel_forward_from_hub) decorator instead.
This function should only be used as a last resort to extend third-party layers,
it is inherently fragile since the member variables and `forward` signature
of such a layer can change.

Example:
```python
from kernels import replace_kernel_forward_from_hub
import torch.nn as nn

replace_kernel_forward_from_hub(nn.LayerNorm, "LayerNorm")
```

## Registering kernel mappings

### use_kernel_mapping[[kernels.use_kernel_mapping]]

#### kernels.use_kernel_mapping[[kernels.use_kernel_mapping]]

```python
kernels.use_kernel_mapping(mapping: dict[str, dict[Device | str, RepositoryProtocol | dict[Mode, RepositoryProtocol]]], inherit_mapping: bool = True)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/layer/kernelize.py#L17)

**Parameters:**

mapping (`dict[str, dict[Union[Device, str], Union[LayerRepositoryProtocol, dict[Mode, LayerRepositoryProtocol]]]]`) : The kernel mapping to apply. Maps layer names to device-specific kernel configurations.

inherit_mapping (`bool`, *optional*, defaults to `True`) : When `True`, the current mapping will be extended by `mapping` inside the context. When `False`, only `mapping` is used inside the context.

**Returns:**

Context manager that handles the temporary kernel mapping.

Context manager that sets a kernel mapping for the duration of the context.

This function allows temporary kernel mappings to be applied within a specific context, enabling different
kernel configurations for different parts of your code.

Example:
```python
import torch
import torch.nn as nn
from torch.nn import functional as F

from kernels import use_kernel_forward_from_hub
from kernels import use_kernel_mapping, LayerRepository, Device
from kernels import Mode, kernelize

# Define a mapping
mapping = {
    "SiluAndMul": {
        "cuda": LayerRepository(
            repo_id="kernels-community/activation",
            layer_name="SiluAndMul",
            version=1
        )
    }
}

@use_kernel_forward_from_hub("SiluAndMul")
class SiluAndMul(nn.Module):
    def forward(self, x: torch.Tensor) -> torch.Tensor:
        d = x.shape[-1] // 2
        return F.silu(x[..., :d]) * x[..., d:]

model = SiluAndMul()

# Use the mapping for the duration of the context.
with use_kernel_mapping(mapping):
    # kernelize uses the temporary mapping
    model = kernelize(model, mode=Mode.TRAINING | Mode.TORCH_COMPILE, device="cuda")

# Outside the context, original mappings are restored
```

### register_kernel_mapping[[kernels.register_kernel_mapping]]

#### kernels.register_kernel_mapping[[kernels.register_kernel_mapping]]

```python
kernels.register_kernel_mapping(mapping: dict[str, dict[Device | str, RepositoryProtocol | dict[Mode, RepositoryProtocol]]], inherit_mapping: bool = True)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/layer/kernelize.py#L97)

**Parameters:**

mapping (`dict[str, dict[Union[Device, str], Union[RepositoryProtocol, dict[Mode, RepositoryProtocol]]]]`) : The kernel mapping to register globally. Maps layer names to device-specific kernels. The mapping can specify different kernels for different modes (training, inference, etc.).

inherit_mapping (`bool`, *optional*, defaults to `True`) : When `True`, the current mapping will be extended by `mapping`. When `False`, the existing mappings are erased before adding `mapping`.

Register a global mapping between layer names and their corresponding kernel implementations.

This function allows you to register a mapping between a layer name and the corresponding kernel(s) to use,
depending on the device and mode. This should be used in conjunction with [kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize).

Example:
```python
from kernels import LayerRepository, register_kernel_mapping, Mode

# Simple mapping for a single kernel per device
kernel_layer_mapping = {
    "LlamaRMSNorm": {
        "cuda": LayerRepository(
            repo_id="kernels-community/layer_norm",
            layer_name="LlamaRMSNorm",
            version=1,
        ),
    },
}
register_kernel_mapping(kernel_layer_mapping)

# Advanced mapping with mode-specific kernels
advanced_mapping = {
    "MultiHeadAttention": {
        "cuda": {
            Mode.TRAINING: LayerRepository(
                repo_id="kernels-community/training-kernels",
                layer_name="TrainingAttention",
                version=1,
            ),
            Mode.INFERENCE: LayerRepository(
                repo_id="kernels-community/inference-kernels",
                layer_name="FastAttention",
                version=1,
            ),
        }
    }
}
register_kernel_mapping(advanced_mapping)
```

## Kernelizing a model

### kernelize[[kernels.kernelize]]

#### kernels.kernelize[[kernels.kernelize]]

```python
kernels.kernelize(model: 'nn.Module', mode: Mode, device: str | 'torch.device' | None = None, use_fallback: bool = True)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/layer/kernelize.py#L175)

**Parameters:**

model (`nn.Module`) : The PyTorch model to kernelize.

mode ([Mode](/docs/kernels/pr_751/en/api/layers#kernels.Mode)) : The mode that the kernel is going to be used in. For example, `Mode.TRAINING | Mode.TORCH_COMPILE` kernelizes the model for training with `torch.compile`.

device (`Union[str, torch.device]`, *optional*) : The device type to load kernels for. Supported device types are: "cuda", "mps", "npu", "rocm", "tpu", "xpu". The device type will be inferred from the model parameters when not provided.

use_fallback (`bool`, *optional*, defaults to `True`) : Whether to use the original forward method of modules when no compatible kernel could be found. If set to `False`, an exception will be raised in such cases.

**Returns:** `nn.Module`

The kernelized model with optimized kernel implementations.

Replace layer forward methods with optimized kernel implementations.

This function iterates over all modules in the model and replaces the `forward` method of extensible layers
for which kernels are registered using [register_kernel_mapping()](/docs/kernels/pr_751/en/api/layers#kernels.register_kernel_mapping) or [use_kernel_mapping()](/docs/kernels/pr_751/en/api/layers#kernels.use_kernel_mapping).

Example:
```python
import torch
import torch.nn as nn

from kernels import kernelize, Mode, use_kernel_mapping, LayerRepository
from kernels import use_kernel_forward_from_hub

@use_kernel_forward_from_hub("SiluAndMul")
class SiluAndMul(nn.Module):
    def forward(self, x: torch.Tensor) -> torch.Tensor:
        d = x.shape[-1] // 2
        return F.silu(x[..., :d]) * x[..., d:]

mapping = {
    "SiluAndMul": {
        "cuda": LayerRepository(
            repo_id="kernels-community/activation",
            layer_name="SiluAndMul",
            version=1,
        )
    }
}

# Create and kernelize a model
model = nn.Sequential(
    nn.Linear(1024, 2048, device="cuda"),
    SiluAndMul(),
)

# Kernelize for inference
with use_kernel_mapping(mapping):
    kernelized_model = kernelize(model, mode=Mode.TRAINING | Mode.TORCH_COMPILE)
```

## Classes

### Device[[kernels.Device]]

#### kernels.Device[[kernels.Device]]

```python
kernels.Device(type: str, properties: kernels.layer.device.CUDAProperties | kernels.layer.device.ROCMProperties | None = None)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/layer/device.py#L106)

**Parameters:**

type (`str`) : The device type (e.g., "cuda", "mps", "npu", "rocm", "xpu").

properties ([CUDAProperties](/docs/kernels/pr_751/en/api/layers#kernels.CUDAProperties), *optional*) : Device-specific properties. Currently only [CUDAProperties](/docs/kernels/pr_751/en/api/layers#kernels.CUDAProperties) is supported for CUDA devices.

Represents a compute device with optional properties.

This class encapsulates device information including device type and optional device-specific properties
like CUDA capabilities.

Example:
```python
from kernels import Device, CUDAProperties

# Basic CUDA device
cuda_device = Device(type="cuda")

# CUDA device with specific capability requirements
cuda_device_with_props = Device(
    type="cuda",
    properties=CUDAProperties(min_capability=75, max_capability=90)
)

# MPS device for Apple Silicon
mps_device = Device(type="mps")

# XPU device (e.g., Intel(R) Data Center GPU Max 1550)
xpu_device = Device(type="xpu")

# NPU device (Huawei Ascend)
npu_device = Device(type="npu")
```

#### validate[[kernels.Device.validate]]

```python
validate()
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/huggingface_hub/dataclasses.py#L247)

Run class validators on the instance.

### CUDAProperties[[kernels.CUDAProperties]]

#### kernels.CUDAProperties[[kernels.CUDAProperties]]

```python
kernels.CUDAProperties(min_capability: int, max_capability: int)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/layer/device.py#L8)

**Parameters:**

min_capability (`int`) : Minimum CUDA compute capability required (e.g., 75 for compute capability 7.5).

max_capability (`int`) : Maximum CUDA compute capability supported (e.g., 90 for compute capability 9.0).

CUDA-specific device properties for capability-based kernel selection.

This class defines CUDA compute capability constraints for kernel selection, allowing kernels to specify
minimum and maximum CUDA compute capabilities they support.

Example:
```python
from kernels import CUDAProperties, Device

# Define CUDA properties for modern GPUs (compute capability 7.5 to 9.0)
cuda_props = CUDAProperties(min_capability=75, max_capability=90)

# Create a device with these properties
device = Device(type="cuda", properties=cuda_props)
```

Note:
CUDA compute capabilities are represented as integers where the major and minor versions are concatenated.
For example, compute capability 7.5 is represented as 75, and 8.6 is represented as 86.

#### validate[[kernels.CUDAProperties.validate]]

```python
validate()
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/huggingface_hub/dataclasses.py#L247)

Run class validators on the instance.

### ROCMProperties[[kernels.ROCMProperties]]

#### kernels.ROCMProperties[[kernels.ROCMProperties]]

```python
kernels.ROCMProperties(min_capability: int, max_capability: int)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/layer/device.py#L57)

**Parameters:**

min_capability (`int`) : Minimum ROCM compute capability required (e.g., 75 for compute capability 7.5).

max_capability (`int`) : Maximum ROCM compute capability supported (e.g., 90 for compute capability 9.0).

ROCM-specific device properties for capability-based kernel selection.

This class defines ROCM compute capability constraints for kernel selection, allowing kernels to specify
minimum and maximum ROCM compute capabilities they support.

Example:
```python
from kernels import ROCMProperties, Device

# Define ROCM properties for modern GPUs (compute capability 7.5 to 9.0)
rocm_props = ROCMProperties(min_capability=75, max_capability=90)

# Create a device with these properties
device = Device(type="rocm", properties=rocm_props)
```

Note:
ROCM compute capabilities are represented as integers where the major and minor versions are concatenated.
For example, compute capability 7.5 is represented as 75, and 8.6 is represented as 86.

#### validate[[kernels.ROCMProperties.validate]]

```python
validate()
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/huggingface_hub/dataclasses.py#L247)

Run class validators on the instance.

### Mode[[kernels.Mode]]

#### kernels.Mode[[kernels.Mode]]

```python
kernels.Mode(value, names = None, module = None, qualname = None, type = None, start = 1)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/layer/mode.py#L4)

**Parameters:**

INFERENCE : The kernel is used for inference.

TRAINING : The kernel is used for training.

TORCH_COMPILE : The kernel is used with `torch.compile`.

FALLBACK : In a kernel mapping, this kernel is used when no other mode matches.

Kernelize mode

The `Mode` flag is used by [kernelize()](/docs/kernels/pr_751/en/api/layers#kernels.kernelize) to select kernels for the given mode. Mappings can be registered for
specific modes.

Note:
Different modes can be combined. For instance, `INFERENCE | TORCH_COMPILE` should be used for layers that
are used for inference *with* `torch.compile`.

### FuncRepository[[kernels.FuncRepository]]

#### kernels.FuncRepository[[kernels.FuncRepository]]

```python
kernels.FuncRepository(repo_id: str, func_name: str, revision: str | None = None, version: int | None = None, trust_remote_code: bool | list[str] = False)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/layer/func.py#L28)

**Parameters:**

repo_id (`str`) : The Hub repository containing the layer.

func_name (`str`) : The name of the function within the kernel repository.

revision (`str`, *optional*) : The specific revision (branch, tag, or commit) to download. Cannot be used together with `version`.

version (`int`, *optional*) : The kernel version to download. Cannot be used together with `revision`. Either `version` or `revision` must be specified.

Repository and name of a function for kernel mapping.

> [!WARNING]
> `FuncRepository` is deprecated and will be removed in kernels 0.17.
> Use [LayerRepository](/docs/kernels/pr_751/en/api/layers#kernels.LayerRepository) instead.

Example:
```python
from kernels import FuncRepository

# Reference a specific layer by revision
layer_repo = FuncRepository(
    repo_id="kernels-community/activation",
    func_name="silu_and_mul",
    revision="main",
)

# Reference a layer by version
layer_repo_versioned = FuncRepository(
    repo_id="kernels-community/relu",
    func_name="relu",
    version=1
)
```

### LayerRepository[[kernels.LayerRepository]]

#### kernels.LayerRepository[[kernels.LayerRepository]]

```python
kernels.LayerRepository(repo_id: str, layer_name: str, revision: str | None = None, version: int | None = None, trust_remote_code: bool | list[str] = False)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/layer/layer.py#L35)

**Parameters:**

repo_id (`str`) : The Hub repository containing the layer.

layer_name (`str`) : The name of the layer within the kernel repository.

revision (`str`, *optional*) : The specific revision (branch, tag, or commit) to download. Cannot be used together with `version`.

version (`int`, *optional*) : The kernel version to download. Cannot be used together with `revision`. Either `version` or `revision` must be specified.

trust_remote_code (`bool | list[str]`, *optional*, defaults to `False`) : Whether to allow loading kernels from untrusted organisations. A list of signing identities can be provided for future verification support; until then it warns and falls back to the default trust check.

Repository and name of a layer for kernel mapping.

Example:
```python
from kernels import LayerRepository

# Reference a specific layer by version
layer_repo = LayerRepository(
    repo_id="kernels-community/activation",
    layer_name="SiluAndMul",
    version=1,
)
```

### LocalFuncRepository[[kernels.LocalFuncRepository]]

#### kernels.LocalFuncRepository[[kernels.LocalFuncRepository]]

```python
kernels.LocalFuncRepository(repo_path: Path, func_name: str)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/layer/func.py#L137)

**Parameters:**

repo_path (`Path`) : The local repository containing the layer.

func_name (`str`) : The name of the function within the kernel repository.

Repository and function name from a local directory for kernel mapping.

> [!WARNING]
> `LocalFuncRepository` is deprecated and will be removed in kernels 0.17.
> Use [LocalLayerRepository](/docs/kernels/pr_751/en/api/layers#kernels.LocalLayerRepository) instead.

Example:
```python
from pathlib import Path

from kernels import LocalFuncRepository

# Reference a specific layer by revision
layer_repo = LocalFuncRepository(
    repo_path=Path("/home/daniel/kernels/activation"),
    func_name="silu_and_mul",
)
```

### LocalLayerRepository[[kernels.LocalLayerRepository]]

#### kernels.LocalLayerRepository[[kernels.LocalLayerRepository]]

```python
kernels.LocalLayerRepository(repo_path: Path, layer_name: str)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/layer/layer.py#L131)

**Parameters:**

repo_path (`Path`) : The local repository containing the layer.

layer_name (`str`) : The name of the layer within the kernel repository.

Repository from a local directory for kernel mapping.

Example:
```python
from pathlib import Path

from kernels import LocalLayerRepository

# Reference a specific layer by revision
layer_repo = LocalLayerRepository(
    repo_path=Path("/home/daniel/kernels/activation"),
    layer_name="SiluAndMul",
)
```

### LockedFuncRepository[[kernels.LockedFuncRepository]]

#### kernels.LockedFuncRepository[[kernels.LockedFuncRepository]]

```python
kernels.LockedFuncRepository(repo_id: str, lockfile: pathlib.Path | None = None, func_name: str, trust_remote_code: bool | list[str] = False)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/layer/func.py#L257)

**Parameters:**

repo_id (`str`) : The Hub repository containing the function.

lockfile (`Path`, *optional*) : Path to the lockfile. If not provided, the lockfile will be inferred from the caller's context.

func_name (`str`) : The name of the function within the kernel repository.

trust_remote_code (`bool`, *optional*, defaults to `False`) : Whether to allow loading kernels from untrusted organisations.

Repository and name of a function.

In contrast to `FuncRepository`, this class uses repositories that
are locked inside a project.

> [!WARNING]
> `LockedFuncRepository` is deprecated and will be removed in kernels 0.17.
> Use [LockedLayerRepository](/docs/kernels/pr_751/en/api/layers#kernels.LockedLayerRepository) instead.

### LockedLayerRepository[[kernels.LockedLayerRepository]]

#### kernels.LockedLayerRepository[[kernels.LockedLayerRepository]]

```python
kernels.LockedLayerRepository(repo_id: str, lockfile: Path | None = None, layer_name: str, trust_remote_code: bool | list[str] = False)
```

[Source](https://github.com/huggingface/kernels/blob/vr_751/kernels/src/kernels/layer/layer.py#L182)

Repository and name of a layer.

In contrast to `LayerRepository`, this class uses repositories that
are locked inside a project.