cuda-kernels

Diffusers 与 Transformers 的 CUDA 内核

此技能提供开发优化 CUDA 内核的模式和指导，适用于 NVIDIA GPU（H100、A100、T4），供 HuggingFace diffusers 和 transformers 库使用。

硬性限制 — 编写任何代码前请阅读

内核必须使用 kernel-builder 构建，并符合 Kernel Hub 要求。kernel-builder 使用 Python limited API（ABI3） 进行编译，因此同一个二进制文件可跨版本用于 Python 3.9 及更高版本。通用 PyTorch 扩展教程中一些常见的模式会导致构建彻底失败。请勿使用这些模式，即使 PyTorch 文档或训练数据中推荐了它们。

禁止的模式 — 绝不可生成以下内容

❌ 绝不可使用	失败原因	✅ 请改用
任何形式的 pybind11：#include <torch/extension.h>、#include <pybind11/...>、PYBIND11_MODULE(...)、py::arg、任何 py:: 符号	pybind11 与 limited API（ABI3）不兼容；构建无法通过	在 torch-ext/torch_binding.cpp 中使用 TORCH_LIBRARY_EXPAND（见下文）。注意：torch/extension.h 会间接包含 pybind11 — 改为包含 torch/torch.h + torch/library.h
使用 torch.utils.cpp_extension（CUDAExtension、BuildExtension、cpp_extension.load、load_inline）手写 setup.py / pyproject.toml	setuptools 扩展不支持 ABI3，且会绕过 build.toml；构建由 kernel-builder 管理	使用 build.toml + nix run .#build-and-copy -L。如需可编辑的开发安装，请使用 kernel-builder create-pyproject -f 生成项目文件 — 绝不可手动编写
使用硬编码命名空间的 TORCH_LIBRARY(my_kernel, m)、TORCH_LIBRARY_FRAGMENT(...) 或 TORCH_LIBRARY_IMPL(...)	kernel-builder 会为算子命名空间追加每次构建专属的哈希值（例如 _my_kernel_a1b2c3d）；硬编码名称永远无法解析	使用生成的 registration.h 中的 TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops)
在 Python 中硬编码调用 torch.ops.my_kernel.fn(...)	同样存在命名空间混淆问题 — 只有在构建时才能确定算子命名空间名称	使用 from ._ops import ops，然后调用 ops.fn(...)
手写 PyMODINIT_FUNC PyInit__... 或任何手动 CPython 模块初始化代码	REGISTER_EXTENSION 会生成此代码；重复定义会导致模块加载失败	在 torch_binding.cpp 中恰好使用一次 REGISTER_EXTENSION(TORCH_EXTENSION_NAME)
使用非 limited CPython API 调用（PyArg_ParseTuple、直接操作 PyObject*）	违反 ABI3 要求	使用 torch C++ API：torch::Tensor、TORCH_CHECK、at::cuda::*
在 torch-ext/ 中绝对导入自有包（from my_kernel.utils import x）	从 Hub 加载时，包目录会被重命名；绝对导入会失效	仅使用相对导入：from .utils import x、from ._ops import ops
除 torch（以及确有需要时的 einops）之外的运行时 Python 依赖	Hub 合规要求限制内核依赖；导入 numpy、triton、packaging 等会被拒绝	仅使用标准库 + torch
将 Python 侧 @torch.library.custom_op 用作主要绑定方式	算子必须在 C++ 中注册，才能打包到编译后的扩展中	通过 TORCH_LIBRARY_EXPAND 在 C++ 中注册；Python 侧的 torch.library.register_fake 仅用于添加 fake/meta 实现（见 torch.compile 部分）

唯一受支持的绑定模式

registration.h 和 _ops.py 由 kernel-builder 生成 — 请引用它们，切勿自行编写。

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

声明内核已完成前的检查清单

1. grep -rn "pybind11\|PYBIND11\|torch/extension.h\|py::" torch-ext/ 无任何返回结果。
2. grep -rn "TORCH_LIBRARY(\|TORCH_LIBRARY_FRAGMENT\|PyInit" torch-ext/ 无任何返回结果（仅允许使用 TORCH_LIBRARY_EXPAND）。
3. 除非由 kernel-builder create-pyproject 生成，否则不得存在 setup.py。
4. kernel-builder check-config 检查通过 — [general] 需要使用以连字符分隔的 name（不能使用下划线）和 license，以及 [torch]（绑定源文件）和 [kernel.<name>] 部分。
5. 内核目录是一个 Git 仓库，且所有文件均已提交（Nix 拒绝非 Git 构建）。
6. 构建成功：nix run .#build-and-copy -L。
7. ABI 兼容性检查通过：构建完成后运行 kernel-builder check-abi。

快速入门

Diffusers（视频／图像生成）

用于对内核性能进行基准测试：

# 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

用于最精简的 diffusers 集成示例（约 150 行）：

python scripts/ltx_kernel_injection_example.py

Transformers（大模型）

用于最精简的 transformers 集成示例（约 120 行）：

python scripts/transformers_injection_example.py

HuggingFace Kernels Hub

从 HuggingFace Hub 加载预编译的内核（无需本地编译）：

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)

用于完整的 HuggingFace Kernels 示例：

python scripts/huggingface_kernels_example.py

隔离的内核微基准测试

python benchmark_rmsnorm.py

支持的库和模型

库	支持的模型	关键内核
diffusers	LTX-Video、Stable Diffusion、FLUX、DiT	RMSNorm、GEGLU、RoPE、AdaLN
transformers	LLaMA、Mistral、Qwen、Falcon	RMSNorm、Attention

GPU	计算能力	指南
H100	sm_90	h100-optimization-guide.md
A100	sm_80	a100-optimization-guide.md
T4	sm_75	t4-optimization-guide.md

适用场景

在以下情况下使用此技能：

• 对照基线实现对内核性能进行基准测试
• 为扩散模型或大模型编写新的 CUDA 内核
• 针对 H100、A100 或 T4 架构优化现有内核
• 实现自定义 Attention、归一化或激活层
• 将内核集成到 diffusers 流水线中（LTX-Video、Stable Diffusion、FLUX、DiT）
• 将内核集成到 transformers 模型中（LLaMA、Mistral、Qwen）
• 调试 NVIDIA GPU 上的内核性能问题

可运行示例

内核仓库的 examples/kernels/ 目录中提供了完整的可运行示例（也可在 github.com/huggingface/kernels 查看）：

• relu/ — 典型的最简内核：build.toml、flake.nix、TORCH_LIBRARY_EXPAND 绑定、Python API、layers/、测试
• relu-backprop-compile/ — 反向传播 + torch.compile 支持（fake-op 注册）
• silu-and-mul/ — 遵循相同布局的激活内核

内核基准测试

使用基准测试脚本测量内核性能：

# 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

基准测试脚本选项

选项	默认值	说明
--use-optimized-kernels	auto	使用自定义 H100 CUDA 内核
--no-optimized-kernels	-	使用基线实现
--compile	false	在 transformer 上启用 torch.compile
--batch-size	1	每个提示词生成的视频数量
--num-frames	161	要生成的帧数
--height	512	视频高度（像素）
--width	768	视频宽度（像素）
--steps	50	去噪步数
--warmup-iterations	2	基准测试前的预热运行次数

示例基准测试结果

端到端视频生成（49 帧、30 步、H100 80GB）：

配置	时间（秒）	it/s	加速比	备注
基线（不编译）	2.87	12.58	1.00x	参考值
优化内核	2.70	13.52	1.06x	快 6%
基线 + torch.compile	2.14	19.05	1.34x	快 34%

重要提示： --use-optimized-kernels 和 --compile 当前不能同时使用。自定义内核需要注册 PyTorch 自定义算子，才能与 torch.compile 配合使用。

需要记录的关键指标：

• 设备： GPU 型号（例如，NVIDIA H100 80GB HBM3）
• 精度： 使用的数据类型（例如，bfloat16）
• 分辨率： 宽度 x 高度（例如，768x512）
• 帧数： 生成的帧数（例如，49、161）

RMSNorm 微基准测试

向量化 RMSNorm 内核相比 PyTorch 基线，平均加速 2.67x：

形状	自定义（毫秒）	PyTorch（毫秒）	加速比
[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

带宽效率： 达到 H100 理论带宽 3.35 TB/s 的 38%

端到端加速比较小的原因： RMSNorm 约占 LTX-Video 总计算量的 5%。其余时间用于注意力（Flash Attention/SDPA）、线性投影和 VAE 解码。

项目结构

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

GPU 架构参考

H100（Hopper）——主要目标

规格	值	优化影响
SM 数量	132	网格大小：尽量设为 132 的倍数
每个 SM 的线程数	2048	每个 SM 最多可容纳 16 个含 128 个线程的块
共享内存	192 KB/SM	可使用较大的块
L2 缓存	50 MB	可在多个块间复用
内存带宽	3.35 TB/s	关键在于合并访问
Warp 大小	32	所有归约都使用 warp shuffle

快速对比（H100、A100 与 T4）

规格	H100	A100	T4
SM 数量	132	108	40
内存带宽	3.35 TB/s	2.0 TB/s	320 GB/s
每个 SM 的共享内存	192 KB	164 KB	64 KB
BF16 支持	是	是	否（仅 FP16）
计算能力	sm_90	sm_80	sm_75

参阅详细指南：H100 | A100 | T4

核心内核模式

向量化内存访问（对性能至关重要）

使用 __nv_bfloat162 进行 BFloat16 向量化：

// 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;
}

使用 __half2 进行 FP16 向量化：

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

使用 float4 进行 FP32 向量化：

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;

Warp Shuffle 归约

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 = 64
• NUM_WARPS = 8

线程配置

对于逐元素操作（RoPE、GEGLU）：

constexpr int BLOCK_SIZE = 256;
int num_blocks = (total_elements + BLOCK_SIZE - 1) / BLOCK_SIZE;

对于使用向量化的归约操作（LayerNorm、RMSNorm）：

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

支持的数据类型

所有内核支持三种精度模式：

• __half（FP16）——推理默认值
• __nv_bfloat16（BF16）——训练首选
• float（FP32）——参考/调试

构建内核

搭建新的内核项目

新建内核时，请使用 kernel-builder init，而不是手动创建文件——它会一次性生成符合要求的目录结构：

kernel-builder init --name my-username/my-kernel

此命令会创建 build.toml（包含有效的连字符分隔名称、许可证，以及已配置好 [general.hub] repo-id）、flake.nix、包含可编译的 torch_binding.{h,cpp} 和 Python 包的 torch-ext/、内核源代码目录 <name>_cuda/、tests/、benchmarks/、example.py 和 CARD.md，并初始化 git 仓库（构建所必需）。然后，用你自己的源代码替换内核桩代码，并更新 build.toml 中的 src 列表。

使用 Nix（推荐）

nix run .#build-and-copy --max-jobs 2 --cores 8 -L

一次性构建并发布到 Hub

kernel-builder build-and-upload

目标仓库由 build.toml 中 [general.hub] 下的 repo-id 和 [general] 下的 version 指定。上传内容会进入 kernel 类型的 Hub 仓库（而非模型仓库）；所属用户或组织需要具备创建 kernel 的权限（在 huggingface.co/settings/account 中点击“Request Kernels Creation”）。

用于开发的本地构建

切勿手动编写 setup.py（这会导致使用 torch.utils.cpp_extension/pybind11，而它们无法在 ABI3 下构建）。让 kernel-builder 生成项目文件，然后使用 setup.py build_kernel 构建（无需 pip install/可编辑安装）：

kernel-builder create-pyproject -f
python setup.py build_kernel

此操作会构建 kernel 并将输出放入 build，可通过 kernels.get_local_kernel(Path("build")) 直接加载。在 kernel-builder devshell/testshell 中，LOCAL_KERNELS 会自动设置，因此 get_kernel("<repo-id>") 会解析到此本地构建。

build.toml 配置

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

kernel 目录必须是包含已提交文件的 git 仓库（git init && git add -A && git commit）——Nix 拒绝构建非 git kernel（“Kernel is not in a git repository”）。

库集成

HuggingFace Kernels Hub（get_kernel）

完整指南请参阅 huggingface-kernels-integration.md。

无需本地编译，即可直接从 HuggingFace Hub 加载预编译的优化 kernel：

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)

主要函数：

• get_kernel(repo_id, version=N) - 从 Hub 下载并加载 kernel；必须指定 version=（主版本号）或 revision=（分支/标签/提交）
• has_kernel(repo_id, version=N) - 检查是否存在兼容的构建
• get_local_kernel(Path("path/to/kernel-project")) - 加载本地构建（会在 <path> 和 <path>/build 中查找）——开发期间使用

**通过 get_kernel() 代码路径测试本地构建：**设置 LOCAL_KERNELS="org/name=/path/to/kernel-project" 并原样调用 get_kernel("org/name")——此覆盖设置会完全绕过 Hub（无需下载，也无需指定版本），因此可以直接使用集成代码测试本地构建。

热门社区 kernel：

• kernels-community/activation - GELU、SiLU 等
• kernels-community/flash-attn - Flash Attention 2
• kernels-community/triton-layer-norm - LayerNorm、RMSNorm

Diffusers 集成（视频/图像生成）

完整指南请参阅 diffusers-integration.md。

Transformers 集成（大模型）

完整指南请参阅 transformers-integration.md。

与 diffusers 的主要区别：

• Transformers 的 RMSNorm 始终具有权重（没有 elementwise_affine=False）
• 使用 'RMSNorm' in class_name 来匹配 LlamaRMSNorm、MistralRMSNorm 等
• epsilon 可通过检查 variance_epsilon（LLaMA）或 eps（其他模型）获取
• 没有 set_processor() 模式——改用 Flash Attention 2

Transformers 最简模式：

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)

Diffusers 的关键注意事项

1. RMSNorm 权重可能为 None

LTX-Video 对部分 RMSNorm 模块使用 elementwise_affine=False：

# Transformer blocks: NO WEIGHT
self.norm1 = RMSNorm(dim, elementwise_affine=False)

# Attention modules: HAS WEIGHT
self.norm_q = torch.nn.RMSNorm(..., elementwise_affine=True)

**解决方案：**兼容这两种情况：

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)

2. Diffusers RMSNorm ≠ torch.nn.RMSNorm

# WRONG - misses diffusers RMSNorm
if isinstance(module, torch.nn.RMSNorm):

# CORRECT - catches all RMSNorm variants
if type(module).__name__ == 'RMSNorm':

3. LTX-Video 使用 GELU，而非 GEGLU

LTX-Video 使用 activation_fn="gelu-approximate"。不要为 LTX-Video 替换 GEGLU。

4. 在 CPU 卸载之前注入 kernel

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

针对 kernel 的指南

RMSNorm

• 输入布局：[..., hidden_size]
• epsilon 默认值：1e-6
• 如果 elementwise_affine=False，权重可能为 None
• **向量化：**BF16 使用 __nv_bfloat162，FP16 使用 __half2，FP32 使用 float4
• **性能：**向量化实现比 PyTorch 快 2.67 倍
• **带宽：**达到 H100 理论带宽 3.35 TB/s 的约 38%

RoPE

• 1D：[batch, seq, heads, head_dim] - 用于文本
• 3D：[batch, t*h*w, heads, head_dim] - 用于视频
• LTX-Video 通过 LTXVideoRotaryPosEmbed 计算自身的 RoPE

GEGLU 与 GELU

• GEGLU：输入 [batch, seq, 2*hidden] -> 输出 [batch, seq, hidden]
• GELU：标准激活函数
• LTX-Video 使用 GELU，而非 GEGLU

AdaLN

• 公式：norm(x) * weight * (1 + scale) + shift
• 在 DiT 块中用于条件控制

性能分析

# NVIDIA Nsight Systems
nsys profile -o profile python your_script.py

# NVIDIA Nsight Compute
ncu --set full -o metrics python your_script.py

常见问题

所有常见问题和解决方案请参见 troubleshooting.md。

快速修复：

• “NoneType has no attribute contiguous”：RMSNorm 权重为 None，请创建 ones
• isinstance() 不匹配：改用 type(module).__name__
• 未调用 GEGLU：模型使用 GELU，而非 GEGLU
• 补丁未持久生效：在 enable_model_cpu_offload() 之前注入
• torch.compile 与自定义内核不兼容：见下文

torch.compile 兼容性

除非将内核注册为 PyTorch 自定义算子，否则自定义 CUDA 内核与 torch.compile 互斥。

错误消息：

torch._dynamo.exc.Unsupported: Attempted to call function marked as skipped

解决方法：

1. 使用 --use-optimized-kernels，不使用 --compile（提速 6%）
2. 使用 --compile，不使用自定义内核（提速 34%）
3. 为 C++ 注册的算子添加 fake/meta 实现（见下文）

**要让算子兼容 torch.compile：**通过 C++ 中的 TORCH_LIBRARY_EXPAND 注册的算子已经是正规的自定义算子——不要在 Python 中用 @torch.library.custom_op 再次封装它们。只需使用生成的 _ops.py 辅助函数注册 fake（meta）实现：

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

另请参阅

脚本

• benchmark_example.py - 用于比较优化版本与基准版本的基准测试脚本 - 从这里开始
• ltx_kernel_injection_example.py - 精简的 diffusers 集成（约 150 行）
• transformers_injection_example.py - 精简的 transformers/LLM 集成（约 120 行）
• huggingface_kernels_example.py - HuggingFace Kernels Hub 集成

集成指南

• huggingface-kernels-integration.md - HuggingFace Kernels Hub（get_kernel）- 加载预编译内核
• diffusers-integration.md - 完整的 diffusers 流水线集成
• transformers-integration.md - 完整的 transformers/LLM 集成

GPU 优化指南

• h100-optimization-guide.md - H100（Hopper，sm_90）深度解析
• a100-optimization-guide.md - A100（Ampere，sm_80）深度解析
• t4-optimization-guide.md - T4（Turing，sm_75）深度解析

参考资料

• troubleshooting.md - 常见问题和解决方案
• kernel-templates.md - 完整的内核模板
• examples/kernels/relu/ - 标准可运行的内核示例（绑定、层、测试）

外部资源

• HuggingFace Kernels 文档
• HuggingFace Kernels GitHub
• Hub 上的社区内核