kernel-builder/skills/cuda-kernels/SKILL.md
Provides guidance for writing and benchmarking optimized CUDA kernels for NVIDIA GPUs (H100, A100, T4) targeting HuggingFace diffusers and transformers libraries. Kernels must be kernel-builder/ABI3-compliant: no pybind11, no setup.py, TORCH_LIBRARY_EXPAND bindings only. Supports models like LTX-Video, Stable Diffusion, LLaMA, Mistral, and Qwen. Includes integration with HuggingFace Kernels Hub (get_kernel) for loading pre-compiled kernels. Includes benchmarking scripts to compare kernel performance against baseline implementations.
npx skillsauth add huggingface/kernels cuda-kernelsInstall this skill globally with one command. Works with Claude Code, Cursor, and Windsurf.
3 of 9 scanners reported clean
Some scanners were skipped, did not run, or reported a non-clean status. Review each row below.
This skill provides patterns and guidance for developing optimized CUDA kernels targeting NVIDIA GPUs (H100, A100, T4) for use with HuggingFace diffusers and transformers libraries.
Kernels MUST build with kernel-builder and meet the Kernel Hub requirements. kernel-builder compiles against the Python limited API (ABI3) so a single binary works for Python 3.9+ across versions. Several patterns that are standard in generic PyTorch-extension tutorials are therefore hard build failures here. Do not use them, even if PyTorch documentation or your training data suggests them.
| ❌ Never use | Why it fails | ✅ Use instead |
|---|---|---|
| pybind11 in any form: #include <torch/extension.h>, #include <pybind11/...>, PYBIND11_MODULE(...), py::arg, any py:: symbol | pybind11 is incompatible with the limited API (ABI3); the build does not compile | TORCH_LIBRARY_EXPAND in torch-ext/torch_binding.cpp (see below). Note: torch/extension.h transitively includes pybind11 — include torch/torch.h + torch/library.h instead |
| Hand-written setup.py / pyproject.toml using torch.utils.cpp_extension (CUDAExtension, BuildExtension, cpp_extension.load, load_inline) | setuptools extensions are not ABI3 and bypass build.toml; kernel-builder owns the build | build.toml + nix run .#build-and-copy -L. For an editable dev install, generate the project files with kernel-builder create-pyproject -f — never write them by hand |
| TORCH_LIBRARY(my_kernel, m), TORCH_LIBRARY_FRAGMENT(...), or TORCH_LIBRARY_IMPL(...) with a hardcoded namespace | kernel-builder suffixes the op namespace with a per-build hash (e.g. _my_kernel_a1b2c3d); a hardcoded name never resolves | TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) from the generated registration.h |
| Hardcoded torch.ops.my_kernel.fn(...) calls in Python | Same namespace mangling — the op namespace name is only known at build time | from ._ops import ops then ops.fn(...) |
| Hand-written PyMODINIT_FUNC PyInit__... or any manual CPython module init | Generated by REGISTER_EXTENSION; duplicating it breaks module loading | REGISTER_EXTENSION(TORCH_EXTENSION_NAME) exactly once, in torch_binding.cpp |
| Non-limited CPython API calls (PyArg_ParseTuple, direct PyObject* manipulation) | Violates ABI3 | Stay within the torch C++ API: torch::Tensor, TORCH_CHECK, at::cuda::* |
| Absolute imports of your own package inside torch-ext/ (from my_kernel.utils import x) | The package directory is renamed when loaded from the Hub; absolute imports break | Relative imports only: from .utils import x, from ._ops import ops |
| Runtime Python deps beyond torch (and einops if truly needed) | Hub compliance restricts kernel dependencies; imports of numpy, triton, packaging, etc. are rejected | Standard library + torch only |
| Python-side @torch.library.custom_op as the primary binding | The op must be registered in C++ so it ships in the compiled extension | C++ registration via TORCH_LIBRARY_EXPAND; Python-side torch.library.register_fake is only for adding a fake/meta impl (see torch.compile section) |
registration.h and _ops.py are generated by kernel-builder — reference them, never write them yourself.
torch-ext/torch_binding.h:
#pragma once
#include <torch/torch.h>
void my_kernel_forward(torch::Tensor &out, torch::Tensor const &input);
torch-ext/torch_binding.cpp:
#include <torch/library.h>
#include "registration.h"
#include "torch_binding.h"
TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
ops.def("my_kernel_forward(Tensor! out, Tensor input) -> ()");
ops.impl("my_kernel_forward", torch::kCUDA, &my_kernel_forward);
}
REGISTER_EXTENSION(TORCH_EXTENSION_NAME)
torch-ext/my_kernel/__init__.py:
import torch
from ._ops import ops
def my_kernel(x: torch.Tensor) -> torch.Tensor:
out = torch.empty_like(x)
ops.my_kernel_forward(out, x)
return out
grep -rn "pybind11\|PYBIND11\|torch/extension.h\|py::" torch-ext/ returns nothing.grep -rn "TORCH_LIBRARY(\|TORCH_LIBRARY_FRAGMENT\|PyInit" torch-ext/ returns nothing (only TORCH_LIBRARY_EXPAND is allowed).setup.py exists unless generated by kernel-builder create-pyproject.kernel-builder check-config passes — [general] needs a dash-separated name (never underscores) and a license, plus [torch] (binding sources) and [kernel.<name>] sections.nix run .#build-and-copy -L.kernel-builder check-abi (after building).For benchmarking kernel performance:
# Benchmark with optimized kernels (6% end-to-end speedup)
python generate_video.py --use-optimized-kernels
# Benchmark baseline with torch.compile (34% speedup)
python generate_video.py --no-optimized-kernels --compile
# Compare configurations (note: --compile and --use-optimized-kernels are mutually exclusive)
python generate_video.py --use-optimized-kernels && \
python generate_video.py --no-optimized-kernels --compile
For a minimal diffusers integration example (~150 lines):
python scripts/ltx_kernel_injection_example.py
For a minimal transformers integration example (~120 lines):
python scripts/transformers_injection_example.py
Load pre-compiled kernels from HuggingFace Hub (no local compilation):
from kernels import get_kernel
# Load optimized activation kernels
activation = get_kernel("kernels-community/activation", version=1)
# Use the kernel
y = torch.empty_like(x)
activation.gelu_fast(y, x)
For a complete HuggingFace Kernels example:
python scripts/huggingface_kernels_example.py
python benchmark_rmsnorm.py
| Library | Supported Models | Key Kernels | |---------|------------------|-------------| | diffusers | LTX-Video, Stable Diffusion, FLUX, DiT | RMSNorm, GEGLU, RoPE, AdaLN | | transformers | LLaMA, Mistral, Qwen, Falcon | RMSNorm, Attention |
| GPU | Compute Capability | Guide | |-----|-------------------|-------| | H100 | sm_90 | h100-optimization-guide.md | | A100 | sm_80 | a100-optimization-guide.md | | T4 | sm_75 | t4-optimization-guide.md |
Use this skill when:
Complete working examples ship with the kernels repo under examples/kernels/ (also at github.com/huggingface/kernels):
relu/ — the canonical minimal kernel: build.toml, flake.nix, TORCH_LIBRARY_EXPAND bindings, Python API, layers/, testsrelu-backprop-compile/ — backward pass + torch.compile support (fake-op registration)silu-and-mul/ — activation kernel following the same layoutUse the benchmark script to measure kernel performance:
# Full benchmark with all options
python scripts/benchmark_example.py \
--use-optimized-kernels \
--compile \
--batch-size 1 \
--num-frames 161 \
--height 512 \
--width 768 \
--steps 50 \
--warmup-iterations 2
| Option | Default | Description |
|--------|---------|-------------|
| --use-optimized-kernels | auto | Use custom H100 CUDA kernels |
| --no-optimized-kernels | - | Use baseline implementation |
| --compile | false | Enable torch.compile on transformer |
| --batch-size | 1 | Number of videos per prompt |
| --num-frames | 161 | Number of frames to generate |
| --height | 512 | Video height in pixels |
| --width | 768 | Video width in pixels |
| --steps | 50 | Denoising steps |
| --warmup-iterations | 2 | Warmup runs before benchmark |
End-to-End Video Generation (49 frames, 30 steps, H100 80GB):
| Configuration | Time (s) | it/s | Speedup | Notes | |:---|:---:|:---:|:---:|:---| | Baseline (no compile) | 2.87 | 12.58 | 1.00x | Reference | | Optimized Kernels | 2.70 | 13.52 | 1.06x | 6% faster | | Baseline + torch.compile | 2.14 | 19.05 | 1.34x | 34% faster |
Important: --use-optimized-kernels and --compile are currently mutually exclusive. Custom kernels require PyTorch custom op registration to work with torch.compile.
Key metrics to capture:
The vectorized RMSNorm kernel achieves 2.67x average speedup over PyTorch baseline:
| Shape | Custom (ms) | PyTorch (ms) | Speedup | |:---|:---:|:---:|:---:| | [1×1024×2048] | 0.019 | 0.065 | 3.37x | | [2×1024×2048] | 0.024 | 0.073 | 3.04x | | [4×1024×2048] | 0.036 | 0.093 | 2.58x | | [2×4096×3072] | 0.087 | 0.208 | 2.41x | | [4×4096×3072] | 0.157 | 0.392 | 2.49x |
Bandwidth efficiency: 38% of H100's theoretical 3.35 TB/s
Why end-to-end speedup is smaller: RMSNorm accounts for ~5% of total compute in LTX-Video. The remaining time is spent in attention (Flash Attention/SDPA), linear projections, and VAE decode.
.claude/skills/cuda-kernels/
├── scripts/
│ ├── benchmark_example.py # End-to-end video generation benchmark
│ ├── benchmark_rmsnorm.py # Isolated RMSNorm micro-benchmark
│ ├── ltx_kernel_injection_example.py # Minimal diffusers integration (~150 lines)
│ ├── transformers_injection_example.py # Minimal transformers integration (~120 lines)
│ └── huggingface_kernels_example.py # HuggingFace Kernels Hub integration
├── references/
│ ├── diffusers-integration.md # Complete diffusers integration guide
│ ├── transformers-integration.md # Complete transformers integration guide
│ ├── huggingface-kernels-integration.md # HuggingFace Kernels Hub (get_kernel) guide
│ ├── troubleshooting.md # Common issues and solutions
│ ├── kernel-templates.md # CUDA kernel templates (includes vectorized)
│ ├── h100-optimization-guide.md # H100 (Hopper) optimization deep dive
│ ├── a100-optimization-guide.md # A100 (Ampere) optimization deep dive
│ └── t4-optimization-guide.md # T4 (Turing) optimization deep dive
└── SKILL.md # This file
examples/kernels/relu/ # Canonical working example (kernels repo)
├── build.toml # kernel-builder build configuration
├── flake.nix # Nix build entry point
├── CARD.md # Kernel card template (becomes README.md)
├── relu_cuda/relu.cu # CUDA kernel source
├── torch-ext/
│ ├── torch_binding.h / .cpp # TORCH_LIBRARY_EXPAND bindings
│ └── relu/__init__.py # Python API (+ optional layers/)
└── tests/test_relu.py # Kernel tests (nix run .#ci-test)
| Spec | Value | Optimization Impact | |------|-------|---------------------| | SMs | 132 | Grid sizing: aim for multiples of 132 | | Threads/SM | 2048 | Max 16 blocks of 128 threads per SM | | Shared Memory | 192 KB/SM | Large tiles possible | | L2 Cache | 50 MB | Reuse across blocks | | Memory BW | 3.35 TB/s | Coalesced access critical | | Warp Size | 32 | All reductions use warp shuffles |
| Spec | H100 | A100 | T4 | |------|------|------|-----| | SMs | 132 | 108 | 40 | | Memory BW | 3.35 TB/s | 2.0 TB/s | 320 GB/s | | Shared Mem/SM | 192 KB | 164 KB | 64 KB | | BF16 Support | Yes | Yes | No (FP16 only) | | Compute Cap | sm_90 | sm_80 | sm_75 |
See detailed guides: H100 | A100 | T4
BFloat16 vectorization using __nv_bfloat162:
// Load 2 bfloat16 elements at once (32-bit load)
const __nv_bfloat162* vec_input = reinterpret_cast<const __nv_bfloat162*>(row_input);
#pragma unroll 4
for (int i = tid; i < vec_hidden; i += stride) {
__nv_bfloat162 v = vec_input[i];
float v0 = __bfloat162float(v.x);
float v1 = __bfloat162float(v.y);
sum_sq += v0 * v0 + v1 * v1;
}
FP16 vectorization using __half2:
const __half2* vec_input = reinterpret_cast<const __half2*>(row_input);
__half2 v = vec_input[i];
float v0 = __half2float(v.x);
float v1 = __half2float(v.y);
FP32 vectorization using float4:
const float4* vec_input = reinterpret_cast<const float4*>(row_input);
float4 v = vec_input[i];
sum_sq += v.x * v.x + v.y * v.y + v.z * v.z + v.w * v.w;
template <typename T>
__device__ __forceinline__ T warp_reduce_sum(T val) {
#pragma unroll
for (int offset = 16; offset > 0; offset >>= 1) {
val += __shfl_xor_sync(0xffffffff, val, offset);
}
return val;
}
BLOCK_SIZE_M = 128, BLOCK_SIZE_N = 64, BLOCK_SIZE_K = 64NUM_WARPS = 8For element-wise ops (RoPE, GEGLU):
constexpr int BLOCK_SIZE = 256;
int num_blocks = (total_elements + BLOCK_SIZE - 1) / BLOCK_SIZE;
For reduction ops (LayerNorm, RMSNorm) with vectorization:
// Divide by 2 for bf16/fp16 vectorized access
int threads = min(hidden_size / 2, MAX_THREADS);
threads = max(threads, WARP_SIZE);
threads = (threads + 32 - 1) / 32 * 32; // Round to warp boundary
All kernels support three precision modes:
__half (FP16) - Default for inference__nv_bfloat16 (BF16) - Preferred for trainingfloat (FP32) - Reference/debuggingStart new kernels with kernel-builder init instead of creating files by hand — it generates the compliant layout in one shot:
kernel-builder init --name my-username/my-kernel
This creates build.toml (valid dash-separated name, license, [general.hub] repo-id already wired), flake.nix, torch-ext/ with compilable torch_binding.{h,cpp} and the Python package, a <name>_cuda/ kernel source dir, tests/, benchmarks/, example.py, and CARD.md — and it initializes a git repository (required for builds). Then replace the stub kernel with your own sources and update the src lists in build.toml.
nix run .#build-and-copy --max-jobs 2 --cores 8 -L
kernel-builder build-and-upload
The target repo is set by repo-id under [general.hub] and version under [general] in build.toml. Uploads go to a kernel-type Hub repository (not a model repo); the owning user/org needs kernel-creation access ("Request Kernels Creation" at huggingface.co/settings/account).
Never hand-write a setup.py (it leads to torch.utils.cpp_extension/pybind11, which cannot build under ABI3). Let kernel-builder generate the project files, then build with setup.py build_kernel (no pip install/editable install needed):
kernel-builder create-pyproject -f
python setup.py build_kernel
This builds the kernel and puts the output in build, which can be loaded directly with kernels.get_local_kernel(Path("build")). Inside kernel-builder devshell/testshell, LOCAL_KERNELS is set automatically so get_kernel("<repo-id>") resolves to this local build.
[general]
# Name MUST be dash-separated lowercase (my-kernel), never underscores —
# `kernel-builder check-config` rejects underscores. The Python package
# lives at torch-ext/<name with dashes replaced by underscores>.
name = "ltx-kernels"
backends = ["cuda"]
version = 1
license = "Apache-2.0" # required field
[general.hub]
# Hub repo for `kernel-builder build-and-upload`; with `version` this
# selects the version branch (e.g. v1).
repo-id = "my-username/ltx-kernels"
[torch]
src = [
"torch-ext/torch_binding.cpp",
"torch-ext/torch_binding.h"
]
[kernel.your_kernel]
backend = "cuda"
src = ["kernel_src/your_kernel.cu"]
depends = ["torch"]
# Only constrain cuda-capabilities when the kernel truly requires it —
# do not over-specify.
The kernel directory must be a git repository with files committed (git init && git add -A && git commit) — Nix refuses to build non-git kernels ("Kernel is not in a git repository").
See huggingface-kernels-integration.md for the complete guide.
Load pre-compiled, optimized kernels directly from HuggingFace Hub without local compilation:
from kernels import get_kernel, has_kernel
# Check availability and load — Hub loads REQUIRE version= (or revision=);
# a bare get_kernel(repo_id) raises ValueError.
if has_kernel("kernels-community/activation", version=1):
activation = get_kernel("kernels-community/activation", version=1)
# Use the kernel
x = torch.randn((4, 4), dtype=torch.float16, device="cuda")
y = torch.empty_like(x)
activation.gelu_fast(y, x)
Key functions:
get_kernel(repo_id, version=N) - Download and load kernel from Hub; version= (major version) or revision= (branch/tag/commit) is requiredhas_kernel(repo_id, version=N) - Check if compatible build existsget_local_kernel(Path("path/to/kernel-project")) - Load a local build (looks in <path> and <path>/build) — use during developmentTesting local builds through the get_kernel() code path: set LOCAL_KERNELS="org/name=/path/to/kernel-project" and call get_kernel("org/name") unchanged — the override short-circuits the Hub entirely (no download, no version needed), so integration code can be tested verbatim against a local build.
Popular community kernels:
kernels-community/activation - GELU, SiLU, etc.kernels-community/flash-attn - Flash Attention 2kernels-community/triton-layer-norm - LayerNorm, RMSNormSee diffusers-integration.md for the complete guide.
See transformers-integration.md for the complete guide.
Key differences from diffusers:
elementwise_affine=False)'RMSNorm' in class_name to match LlamaRMSNorm, MistralRMSNorm, etc.variance_epsilon (LLaMA) or eps (others) for epsilonset_processor() pattern - use Flash Attention 2 insteadMinimal transformers pattern:
from transformers import AutoModelForCausalLM
from ltx_kernels import rmsnorm
def patch_rmsnorm(model):
for name, module in model.named_modules():
if 'RMSNorm' in type(module).__name__:
eps = getattr(module, 'variance_epsilon', None) or getattr(module, 'eps', 1e-6)
def make_forward(mod, epsilon):
def forward(x):
return rmsnorm(x, mod.weight, eps=epsilon)
return forward
module.forward = make_forward(module, eps)
model = AutoModelForCausalLM.from_pretrained("meta-llama/Llama-2-7b-hf", torch_dtype=torch.bfloat16)
patch_rmsnorm(model)
LTX-Video uses elementwise_affine=False for some RMSNorm modules:
# Transformer blocks: NO WEIGHT
self.norm1 = RMSNorm(dim, elementwise_affine=False)
# Attention modules: HAS WEIGHT
self.norm_q = torch.nn.RMSNorm(..., elementwise_affine=True)
Solution: Handle both cases:
has_weight = hasattr(module, 'weight') and module.weight is not None
if has_weight:
output = rmsnorm(x, module.weight, eps=eps)
else:
weight = torch.ones(x.shape[-1], device=x.device, dtype=x.dtype)
output = rmsnorm(x, weight, eps=eps)
# WRONG - misses diffusers RMSNorm
if isinstance(module, torch.nn.RMSNorm):
# CORRECT - catches all RMSNorm variants
if type(module).__name__ == 'RMSNorm':
LTX-Video uses activation_fn="gelu-approximate". Don't patch GEGLU for LTX-Video.
pipe = LTXPipeline.from_pretrained(...)
pipe.to("cuda")
inject_optimized_kernels(pipe) # BEFORE offloading
pipe.enable_model_cpu_offload() # Now safe
from diffusers import LTXPipeline
from ltx_kernels import rmsnorm
def patch_rmsnorm_modules(model):
"""Patch all RMSNorm modules to use custom kernel."""
for name, module in model.named_modules():
if type(module).__name__ == 'RMSNorm':
eps = getattr(module, 'eps', 1e-6)
has_weight = hasattr(module, 'weight') and module.weight is not None
if has_weight:
def make_forward(mod, epsilon):
def forward(x):
return rmsnorm(x, mod.weight, eps=epsilon)
return forward
module.forward = make_forward(module, eps)
else:
def make_forward(epsilon):
def forward(x):
w = torch.ones(x.shape[-1], device=x.device, dtype=x.dtype)
return rmsnorm(x, w, eps=epsilon)
return forward
module.forward = make_forward(eps)
# Usage
pipe = LTXPipeline.from_pretrained("Lightricks/LTX-Video", torch_dtype=torch.bfloat16)
pipe.to("cuda")
patch_rmsnorm_modules(pipe.transformer)
pipe.enable_model_cpu_offload()
[..., hidden_size]elementwise_affine=False__nv_bfloat162 for BF16, __half2 for FP16, float4 for FP32[batch, seq, heads, head_dim] - for text[batch, t*h*w, heads, head_dim] - for videoLTXVideoRotaryPosEmbed[batch, seq, 2*hidden] -> Output [batch, seq, hidden]norm(x) * weight * (1 + scale) + shift# NVIDIA Nsight Systems
nsys profile -o profile python your_script.py
# NVIDIA Nsight Compute
ncu --set full -o metrics python your_script.py
See troubleshooting.md for all common issues and solutions.
Quick fixes:
type(module).__name__ insteadenable_model_cpu_offload()Custom CUDA kernels and torch.compile are mutually exclusive unless you register the kernel as a PyTorch custom op.
Error message:
torch._dynamo.exc.Unsupported: Attempted to call function marked as skipped
Workaround options:
--use-optimized-kernels without --compile (6% speedup)--compile without custom kernels (34% speedup)To make the op torch.compile-compatible: ops registered via TORCH_LIBRARY_EXPAND in C++ are already proper custom ops — do NOT re-wrap them with @torch.library.custom_op in Python. Just register a fake (meta) implementation using the generated _ops.py helpers:
import torch
from ._ops import ops, add_op_namespace_prefix
@torch.library.register_fake(add_op_namespace_prefix("rmsnorm_forward"))
def _(out, input, weight, eps):
return None # out-variant op: no shape changes
development
Provides guidance for writing, optimizing, and benchmarking Triton kernels for Intel XPU GPUs (Battlemage/Arc Pro B50) using the Xe-Forge optimization framework. Includes an LLM-driven trial-loop workflow (analyze, validate, benchmark, profile, finalize), XPU-specific patterns (tensor descriptors, GRF mode, tile swizzling), KernelBench fused kernels, and Flash Attention.
development
| name | triton-kernels | | --- | --- | | description | Provides guidance for writing and benchmarking portable Triton kernels targeting NVIDIA and AMD GPUs. Covers core DSL patterns, @triton.autotune, numerics (fp16/bf16/fp8), masked loads, reductions, tiling, benchmarking harness, correctness testing, and integration with HuggingFace Kernels Hub (get_kernel). Vendor-neutral: points to rocm-kernels and xpu-kernels for backend-specific tuning. | | disable-model-invocation | false | | user-invoca
tools
Provides guidance for writing, optimizing, and benchmarking C++ CPU kernels with SIMD intrinsics (AVX2/AVX512) for the Hugging Face kernels ecosystem. Includes a two-phase workflow: Phase 1 correctness (generic → AVX2) and Phase 2 performance exploration (AVX512 with branching trial loop), runtime CPU dispatch, OpenMP threading, and brgemm integration for GEMM-heavy kernels.
devops
Provides guidance for writing and benchmarking optimized Triton kernels for AMD GPUs (MI355X, R9700) on ROCm, targeting HuggingFace diffusers (LTX-Video, SD3, FLUX) and transformers. Core kernels: RMSNorm, RoPE 3D, GEGLU, AdaLN. Includes XCD swizzle, autotune, diffusers integration patterns, and LTX-Video pipeline injection.