# Kernels

## Docs

- [Environment variables](https://huggingface.co/docs/kernels/main/env.md)
- [Quickstart](https://huggingface.co/docs/kernels/main/basic-usage.md)
- [Installation](https://huggingface.co/docs/kernels/main/installation.md)
- [Why kernels?](https://huggingface.co/docs/kernels/main/why_kernels.md)
- [Kernels CLI Reference](https://huggingface.co/docs/kernels/main/cli.md)
- [kernels benchmark](https://huggingface.co/docs/kernels/main/cli-benchmark.md)
- [Kernels](https://huggingface.co/docs/kernels/main/index.md)
- [FAQ](https://huggingface.co/docs/kernels/main/faq.md)
- [kernels lock](https://huggingface.co/docs/kernels/main/cli-lock.md)
- [CLI reference for kernel-builder](https://huggingface.co/docs/kernels/main/builder-cli.md)
- [Layers](https://huggingface.co/docs/kernels/main/layers.md)
- [Kernel requirements](https://huggingface.co/docs/kernels/main/kernel-requirements.md)
- [Locking kernel/layer versions](https://huggingface.co/docs/kernels/main/locking.md)
- [Integrating kernels](https://huggingface.co/docs/kernels/main/integrating-kernels.md)
- [kernels versions](https://huggingface.co/docs/kernels/main/cli-versions.md)
- [kernel-builder skills add](https://huggingface.co/docs/kernels/main/cli-skills.md)
- [Talks](https://huggingface.co/docs/kernels/main/talks.md)
- [Migrating from older versions](https://huggingface.co/docs/kernels/main/migration.md)
- [kernels download](https://huggingface.co/docs/kernels/main/cli-download.md)
- [Kernels API Reference](https://huggingface.co/docs/kernels/main/api/kernels.md)
- [Layers API Reference](https://huggingface.co/docs/kernels/main/api/layers.md)
- [IDE setup with direnv and the kernel devshell](https://huggingface.co/docs/kernels/main/builder/ide-setup.md)
- [Writing custom kernels with code agents](https://huggingface.co/docs/kernels/main/builder/agents-guide.md)
- [Writing Hub kernels with kernel-builder](https://huggingface.co/docs/kernels/main/builder/writing-kernels.md)
- [Using the kernel builder with Nix](https://huggingface.co/docs/kernels/main/builder/build.md)
- [Build variants](https://huggingface.co/docs/kernels/main/builder/build-variants.md)
- [Why Nix?](https://huggingface.co/docs/kernels/main/builder/why-nix.md)
- [Nix Builder design](https://huggingface.co/docs/kernels/main/builder/design-nix-builder.md)
- [Metal kernels 🤘](https://huggingface.co/docs/kernels/main/builder/metal.md)
- [Security](https://huggingface.co/docs/kernels/main/builder/security.md)
- [Local development of kernels](https://huggingface.co/docs/kernels/main/builder/local-dev.md)

### Environment variables
https://huggingface.co/docs/kernels/main/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).

### Quickstart
https://huggingface.co/docs/kernels/main/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`.

## 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/main/en/api/kernels#kernels.has_kernel) does not say *why*.
[get_kernel_variants()](/docs/kernels/main/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/main/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/main/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/main/en/api/kernels#kernels.get_kernel). Kernels
loaded from a local path ([get_local_kernel()](/docs/kernels/main/en/api/kernels#kernels.get_local_kernel)) or via a lockfile
([get_locked_kernel()](/docs/kernels/main/en/api/kernels#kernels.get_locked_kernel), [load_kernel()](/docs/kernels/main/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.

### Installation
https://huggingface.co/docs/kernels/main/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"
```

> [!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.

### Why kernels?
https://huggingface.co/docs/kernels/main/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

### Kernels CLI Reference
https://huggingface.co/docs/kernels/main/cli.md

# Kernels CLI Reference

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

## Commands

| Command                                                 | Description                                              |
| ------------------------------------------------------- | -------------------------------------------------------- |
| [benchmark](cli-benchmark)                           | Run benchmark results for a kernel                       |
| [check](cli-check)                                   | Check a kernel for compliance                            |
| [versions](cli-versions)                             | Show kernel versions                                     |
| [lock](cli-lock)                                     | Lock kernel revisions                                    |

## Quick Start

For building and writing kernels, please refer [building kernels](./builder/build) and 
[writing kernels](./builder/writing-kernels).

### Use kernels in your project

#### Directly from the Hub

```python
import torch

from kernels import get_kernel

# Download optimized kernels from the Hugging Face hub
my_kernel = get_kernel("my-username/my-kernel", version=1)

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

# Run the kernel
y = torch.empty_like(x)
my_kernel.my_kernel_function(y, x)

print(y)
```

or

#### Locked and downloaded

Add to `pyproject.toml`:

```toml
[tool.kernels.dependencies]
"my-username/my-kernel" = "1"
```

Then lock and download:

```bash
kernels lock .
kernels download .
```

### See help

```bash
kernels --help
```

### kernels benchmark
https://huggingface.co/docs/kernels/main/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.

### Kernels
https://huggingface.co/docs/kernels/main/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.

### FAQ
https://huggingface.co/docs/kernels/main/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/main/en/api/layers#kernels.use_kernel_forward_from_hub) and [replace_kernel_forward_from_hub()](/docs/kernels/main/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/main/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/main/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/main/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/main/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 lock
https://huggingface.co/docs/kernels/main/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

### CLI reference for kernel-builder
https://huggingface.co/docs/kernels/main/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 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
* `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`, `neuron`, `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`

* `-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`

* `-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-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 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`, `rocm-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.

### Layers
https://huggingface.co/docs/kernels/main/layers.md

# 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/main/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/main/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/main/en/api/layers#kernels.use_kernel_forward_from_hub)
decorator can be made extensible using the [replace_kernel_forward_from_hub()](/docs/kernels/main/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 annotate the function with the [use_kernel_func_from_hub()](/docs/kernels/main/en/api/layers#kernels.use_kernel_func_from_hub) decorator:

```python
@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:]
```

This will replace the function by an instantiated `torch.nn.Module`
(singleton) that calls the function itself in its forward method.

**Note:** for kernelization to see the function, it must be a member of
another `torch.nn.Module` that is part of the model. For example:

```python
class FeedForward(nn.Module):
  def __init__(self, in_features: int, out_features: int):
      self.linear = nn.Linear(in_features, out_features)
      # Note: silu_and_mul is a Torch module.
      self.silu_and_mul = silu_and_mul

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

## 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/main/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/main/en/api/layers#kernels.kernelize) can be used as follows:

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

The [kernelize()](/docs/kernels/main/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/main/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/main/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/main/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/main/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/main/en/api/layers#kernels.kernelize) will fall back to the original, non-kernelized, layer. You
can let [kernelize()](/docs/kernels/main/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/main/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/main/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/main/en/api/layers#kernels.LayerRepository) and [FuncRepository](/docs/kernels/main/en/api/layers#kernels.FuncRepository) entries must specify
either a `version` or an explicit `revision`.

You can register a mapping, like the one above, using [register_kernel_mapping()](/docs/kernels/main/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/main/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/main/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/main/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/main/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/main/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/main/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/main/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)
```

### Kernel requirements
https://huggingface.co/docs/kernels/main/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 Discord](https://discord.gg/H6Tkmd88N3) 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).
- `backend` (`dict`, required): information about the compute backend that
  this build variant supports.
- `python-depends` (`list[str]`, optional): list of Python dependencies
  from a curated set of Python dependencies.

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"]
  }
}
```

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`, 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`

**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.

## 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 15.0.
- Metal 3.0 (`-std=metal3.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"
  # ...
]
```

## 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.

### Locking kernel/layer versions
https://huggingface.co/docs/kernels/main/locking.md

# Locking kernel/layer 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/main/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/main/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/main/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/main/en/api/kernels#kernels.get_locked_kernel) function.
[get_locked_kernel()](/docs/kernels/main/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/main/en/api/kernels#kernels.load_kernel) function instead:

```python
from kernels import load_kernel

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

### Integrating kernels
https://huggingface.co/docs/kernels/main/integrating-kernels.md

# Integrating 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/main/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/main/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 versions
https://huggingface.co/docs/kernels/main/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

### kernel-builder skills add
https://huggingface.co/docs/kernels/main/cli-skills.md

### kernel-builder skills add

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`

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.

Examples:

```bash
<CopyLLMTxtMenu containerStyle="float: right; margin-left: 10px; display: inline-flex; position: relative; z-index: 10;"></CopyLLMTxtMenu>

# 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
```

### Talks
https://huggingface.co/docs/kernels/main/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.

### Migrating from older versions
https://huggingface.co/docs/kernels/main/migration.md

# Migrating 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.

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/main/en/api/kernels#kernels.get_kernel) and [LayerRepository](/docs/kernels/main/en/api/layers#kernels.LayerRepository) calls
   to reference the new repository if the `repo-id` changed.

### kernels download
https://huggingface.co/docs/kernels/main/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 API Reference
https://huggingface.co/docs/kernels/main/api/kernels.md

# Kernels API Reference

## Main Functions

### get_kernel[[kernels.get_kernel]]

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

[Source](https://github.com/huggingface/kernels/blob/main/kernels/src/kernels/utils.py#L415)

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)
```

**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.

### get_local_kernel[[kernels.get_local_kernel]]

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

[Source](https://github.com/huggingface/kernels/blob/main/kernels/src/kernels/utils.py#L483)

Import a kernel from a local kernel repository path.

**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.

### has_kernel[[kernels.has_kernel]]

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

[Source](https://github.com/huggingface/kernels/blob/main/kernels/src/kernels/utils.py#L516)

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

**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.

### get_kernel_variants[[kernels.get_kernel_variants]]

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

[Source](https://github.com/huggingface/kernels/blob/main/kernels/src/kernels/utils.py#L557)

Resolve all build variants of a kernel against the current environment.

The decisions are sorted with compatible variants first, the most preferred
variant leading.

**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.

### get_loaded_kernels[[kernels.get_loaded_kernels]]

#### kernels.get_loaded_kernels[[kernels.get_loaded_kernels]]

[Source](https://github.com/huggingface/kernels/blob/main/kernels/src/kernels/utils.py#L142)

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)
```

**Returns:**

``list[LoadedKernel]``

One [LoadedKernel](/docs/kernels/main/en/api/kernels#kernels.LoadedKernel) per distinct kernel variant path
loaded in this process.

## Loading locked kernels

### load_kernel[[kernels.load_kernel]]

#### kernels.load_kernel[[kernels.load_kernel]]

[Source](https://github.com/huggingface/kernels/blob/main/kernels/src/kernels/utils.py#L593)

Get a pre-downloaded, locked kernel.

If `lockfile` is not specified, the lockfile will be loaded from the caller's package metadata.

**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.

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

**Returns:**

``ModuleType``

The imported kernel module.

### get_locked_kernel[[kernels.get_locked_kernel]]

#### kernels.get_locked_kernel[[kernels.get_locked_kernel]]

[Source](https://github.com/huggingface/kernels/blob/main/kernels/src/kernels/utils.py#L649)

Get a kernel using a lock file.

**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.

## Classes

### LoadedKernel[[kernels.LoadedKernel]]

#### kernels.LoadedKernel[[kernels.LoadedKernel]]

[Source](https://github.com/huggingface/kernels/blob/main/kernels/src/kernels/utils.py#L112)

This dataclass provides information about a loaded kernel:

- `metadata` (`Metadata`): kernel metadata.
- `module` (`ModuleType`): the imported kernel module.
- `repo_info` (`kernels.utils.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 upstream repository of the kernel.
- `python_depends` (`list[str]`): required Python dependencies.
- `backend`: information about the kernel's backend.

### RepoInfo[[kernels.RepoInfo]]

#### kernels.RepoInfo[[kernels.RepoInfo]]

[Source](https://github.com/huggingface/kernels/blob/main/kernels/src/kernels/utils.py#L97)

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/main/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]]

[Source](https://github.com/huggingface/kernels/blob/main/kernels/src/kernels/layer/layer.py#L269)

Decorator factory that makes a layer extensible using the specified layer name.

This is a decorator factory that returns a decorator which prepares a layer class to use kernels from the
Hugging Face Hub.

Example:
```python
import torch
import torch.nn as nn

from kernels import use_kernel_forward_from_hub
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")
```

**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.

### use_kernel_func_from_hub[[kernels.use_kernel_func_from_hub]]

#### kernels.use_kernel_func_from_hub[[kernels.use_kernel_func_from_hub]]

[Source](https://github.com/huggingface/kernels/blob/main/kernels/src/kernels/layer/func.py#L167)

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).

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")
```

**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.

### replace_kernel_forward_from_hub[[kernels.replace_kernel_forward_from_hub]]

#### kernels.replace_kernel_forward_from_hub[[kernels.replace_kernel_forward_from_hub]]

[Source](https://github.com/huggingface/kernels/blob/main/kernels/src/kernels/layer/layer.py#L246)

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/main/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]]

[Source](https://github.com/huggingface/kernels/blob/main/kernels/src/kernels/layer/kernelize.py#L17)

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
```

**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.

### register_kernel_mapping[[kernels.register_kernel_mapping]]

#### kernels.register_kernel_mapping[[kernels.register_kernel_mapping]]

[Source](https://github.com/huggingface/kernels/blob/main/kernels/src/kernels/layer/kernelize.py#L97)

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/main/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)
```

**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`.

## Kernelizing a model

### kernelize[[kernels.kernelize]]

#### kernels.kernelize[[kernels.kernelize]]

[Source](https://github.com/huggingface/kernels/blob/main/kernels/src/kernels/layer/kernelize.py#L175)

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/main/en/api/layers#kernels.register_kernel_mapping) or [use_kernel_mapping()](/docs/kernels/main/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)
```

**Parameters:**

model (`nn.Module`) : The PyTorch model to kernelize.

mode ([Mode](/docs/kernels/main/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", "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.

## Classes

### Device[[kernels.Device]]

#### kernels.Device[[kernels.Device]]

[Source](https://github.com/huggingface/kernels/blob/main/kernels/src/kernels/layer/device.py#L106)

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")
```

validatekernels.Device.validatehttps://github.com/huggingface/kernels/blob/main/kernels/src/huggingface_hub/dataclasses.py#L247[]
Run class validators on the instance.

**Parameters:**

type (`str`) : The device type (e.g., "cuda", "mps", "npu", "rocm", "xpu").

properties (`CUDAProperties`, *optional*) : Device-specific properties. Currently only `CUDAProperties` is supported for CUDA devices.

### Mode[[kernels.Mode]]

#### kernels.Mode[[kernels.Mode]]

[Source](https://github.com/huggingface/kernels/blob/main/kernels/src/kernels/layer/mode.py#L4)

Kernelize mode

The `Mode` flag is used by [kernelize()](/docs/kernels/main/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`.

**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.

### FuncRepository[[kernels.FuncRepository]]

#### kernels.FuncRepository[[kernels.FuncRepository]]

[Source](https://github.com/huggingface/kernels/blob/main/kernels/src/kernels/layer/func.py#L27)

Repository and name of a function for kernel mapping.

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
)
```

**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.

### LayerRepository[[kernels.LayerRepository]]

#### kernels.LayerRepository[[kernels.LayerRepository]]

[Source](https://github.com/huggingface/kernels/blob/main/kernels/src/kernels/layer/layer.py#L32)

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,
)
```

**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.

### LocalFuncRepository[[kernels.LocalFuncRepository]]

#### kernels.LocalFuncRepository[[kernels.LocalFuncRepository]]

[Source](https://github.com/huggingface/kernels/blob/main/kernels/src/kernels/layer/func.py#L116)

Repository and function name from a local directory for kernel mapping.

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",
)
```

**Parameters:**

repo_path (`Path`) : The local repository containing the layer.

func_name (`str`) : The name of the function within the kernel repository.

### LocalLayerRepository[[kernels.LocalLayerRepository]]

#### kernels.LocalLayerRepository[[kernels.LocalLayerRepository]]

[Source](https://github.com/huggingface/kernels/blob/main/kernels/src/kernels/layer/layer.py#L128)

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",
)
```

**Parameters:**

repo_path (`Path`) : The local repository containing the layer.

layer_name (`str`) : The name of the layer within the kernel repository.

### LockedFuncRepository[[kernels.LockedFuncRepository]]

#### kernels.LockedFuncRepository[[kernels.LockedFuncRepository]]

[Source](https://github.com/huggingface/kernels/blob/main/kernels/src/kernels/layer/func.py#L222)

Repository and name of a function.

In contrast to `FuncRepository`, this class uses repositories that
are locked inside a project.

### LockedLayerRepository[[kernels.LockedLayerRepository]]

#### kernels.LockedLayerRepository[[kernels.LockedLayerRepository]]

[Source](https://github.com/huggingface/kernels/blob/main/kernels/src/kernels/layer/layer.py#L179)

Repository and name of a layer.

In contrast to `LayerRepository`, this class uses repositories that
are locked inside a project.

### IDE setup with direnv and the kernel devshell
https://huggingface.co/docs/kernels/main/builder/ide-setup.md

# IDE setup with direnv and the kernel devshell

## 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.torch211-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.torch211-cxx11-cu130-x86_64-linux

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

# XPU
use flake .#devShells.torch211-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.

### Writing custom kernels with code agents
https://huggingface.co/docs/kernels/main/builder/agents-guide.md

# Writing custom kernels with code 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`, and `xpu-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                    # pip install -e .
└── 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. 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
nix run .#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.

### Writing Hub kernels with kernel-builder
https://huggingface.co/docs/kernels/main/builder/writing-kernels.md

# Writing Hub kernels with kernel-builder

## 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

[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 `kernels upload` command to upload the kernel to a version
  branch named `v<version>`.
- `backends` (required): a list of supported backends. Must be one or
  more of `cpu`, `cuda`, `metal`, `rocm`, or `xpu`.
- `python-depends` (**experimental**): a list of additional Python dependencies
  that the kernel requires. The only supported dependencies are `einops`
  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.

### `torch`

This section describes the Torch extension. In the future, there may be
similar sections for other frameworks. This section 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**): when set to a Torch version (e.g.
  `"2.11"`), the kernel is built 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.

### `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 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)
```

## Kernel tests

Kernel tests are stored in the `tests` directory. 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.

### Using the kernel builder with Nix
https://huggingface.co/docs/kernels/main/builder/build.md

# Using the kernel builder 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:

```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
```

If you want to test the kernel as a Python package, you can do so.
`kernel-builder devshell` will automatically create a virtual environment in
the `.venv` and activate it. You can install the kernel as a regular
Python package in this virtual environment:

```bash
$ kernel-builder devshell
$ kernel-builder create-pyproject
$ pip install --no-build-isolation -e .
```

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 torch211-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.

## 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.

### Build variants
https://huggingface.co/docs/kernels/main/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

- `torch211-cpu-aarch64-darwin`
- `torch212-cpu-aarch64-darwin`

## Metal aarch64-darwin

- `torch211-metal-aarch64-darwin`
- `torch212-metal-aarch64-darwin`

## CPU aarch64-linux

- `torch211-cxx11-cpu-aarch64-linux`
- `torch212-cxx11-cpu-aarch64-linux`

## CUDA aarch64-linux

- `torch211-cxx11-cu126-aarch64-linux`
- `torch211-cxx11-cu128-aarch64-linux`
- `torch211-cxx11-cu130-aarch64-linux`
- `torch212-cxx11-cu126-aarch64-linux`
- `torch212-cxx11-cu130-aarch64-linux`
- `torch212-cxx11-cu132-aarch64-linux`

## CPU x86_64-linux

- `torch211-cxx11-cpu-x86_64-linux`
- `torch212-cxx11-cpu-x86_64-linux`

## CUDA x86_64-linux

- `torch211-cxx11-cu126-x86_64-linux`
- `torch211-cxx11-cu128-x86_64-linux`
- `torch211-cxx11-cu130-x86_64-linux`
- `torch212-cxx11-cu126-x86_64-linux`
- `torch212-cxx11-cu130-x86_64-linux`
- `torch212-cxx11-cu132-x86_64-linux`

## ROCm x86_64-linux

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

## XPU x86_64-linux

- `torch211-cxx11-xpu20253-x86_64-linux`
- `torch212-cxx11-xpu20253-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`

### Why Nix?
https://huggingface.co/docs/kernels/main/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)

### Nix Builder design
https://huggingface.co/docs/kernels/main/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 upstram 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`).

### Metal kernels 🤘
https://huggingface.co/docs/kernels/main/builder/metal.md

# 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
```

### Security
https://huggingface.co/docs/kernels/main/builder/security.md

# Security

## 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.

### Local development of kernels
https://huggingface.co/docs/kernels/main/builder/local-dev.md

# Local development of kernels

## 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.

It is recommended to do an editable install of the generated project into
your Python virtual environment for development:

```bash
$ pip install wheel # Needed once to enable bdist_wheel.
$ pip install --no-build-isolation -e .
```

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
```
