cuda-kernels

Diffusers と Transformers 向け CUDA カーネル

このスキルでは、HuggingFace の diffusers および transformers ライブラリで使用する、NVIDIA GPU（H100、A100、T4）向けの最適化された CUDA カーネルを開発するためのパターンとガイダンスを提供します。

厳守事項 — コードを書く前にお読みください

カーネルは 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）を op 名前空間の末尾に付加します。名前をハードコードすると解決できません	生成された registration.h の TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops)
Python での torch.ops.my_kernel.fn(...) のハードコード	同じ名前空間のマングリングが行われるため、op 名前空間名はビルド時までわかりません	from ._ops import ops の後に ops.fn(...) を呼び出します
手書きの PyMODINIT_FUNC PyInit__...、または手動による CPython モジュール初期化	REGISTER_EXTENSION が生成します。重複させるとモジュールの読み込みに失敗します	torch_binding.cpp 内で REGISTER_EXTENSION(TORCH_EXTENSION_NAME) を一度だけ使用します
Limited API 非対応の 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	コンパイル済み拡張機能に含めるため、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] には、license と、ハイフン区切りの name（アンダースコアは使用不可）が必要です。また、[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（LLM）

最小限の 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

このスキルを使用する場面

次のような場合にこのスキルを使用します。

• ベースライン実装と比較して、カーネルのパフォーマンスをベンチマークする
• 拡散モデルまたは LLM 用の新しい CUDA カーネルを作成する
• H100、A100、T4 アーキテクチャ向けに既存のカーネルを最適化する
• カスタムのアテンション、正規化、活性化レイヤーを実装する
• 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	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 は現在、同時に使用できません。カスタムカーネルを torch.compile で動作させるには、PyTorch のカスタム op 登録が必要です。

記録する主な指標：

• デバイス： GPU モデル（例：NVIDIA H100 80GB HBM3）
• 精度： 使用するデータ型（例：bfloat16）
• 解像度： 幅 x 高さ（例：768x512）
• フレーム数： 生成するフレーム数（例：49、161）

RMSNorm マイクロベンチマーク

ベクトル化された RMSNorm カーネルは、PyTorch のベースラインと比較して平均 2.67 倍高速です。

形状	カスタム（ms）	PyTorch（ms）	高速化率
[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%

エンドツーエンドの高速化率が小さい理由： LTX-Video では、RMSNorm が全体の計算量の約 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	結合アクセスが重要
ワープサイズ	32	すべてのリダクションでワープシャッフルを使用

クイック比較（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;

ワープシャッフルによるリダクション

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

サポートされるデータ型

すべてのカーネルは、次の 3 つの精度モードをサポートします。

• __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 リポジトリ（モデルリポジトリではありません）です。所有するユーザーまたは組織には、カーネル作成権限（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

これによりカーネルがビルドされ、出力が 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.

カーネルのディレクトリは、ファイルがコミット済みの git リポジトリである必要があります（git init && git add -A && git commit）。Nix は git リポジトリではないカーネルのビルドを拒否します（「Kernel is not in a git repository」）。

ライブラリとの統合

HuggingFace Kernels Hub（get_kernel）

完全なガイドについては、huggingface-kernels-integration.md を参照してください。

ローカルでコンパイルせずに、最適化済みのコンパイル済みカーネルを HuggingFace Hub から直接読み込みます。

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 からカーネルをダウンロードして読み込みます。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 へのアクセスが完全にスキップされるため（ダウンロード不要、バージョン指定不要）、統合コードをローカルビルドに対してそのままテストできます。

人気のコミュニティカーネル：

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

Diffusers との統合（動画／画像生成）

完全なガイドについては、diffusers-integration.md を参照してください。

Transformers との統合（LLM）

完全なガイドについては、transformers-integration.md を参照してください。

Diffusers との主な違い：

• Transformers の RMSNorm には常に重みがあります（elementwise_affine=False はありません）
• LlamaRMSNorm、MistralRMSNorm などに一致させるには、'RMSNorm' in class_name を使用します
• イプシロンの取得には、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 は GEGLU ではなく GELU を使用する

LTX-Video は activation_fn="gelu-approximate" を使用します。LTX-Video に対して GEGLU をパッチしないでください。

4. CPU オフロードの前にカーネルを注入する

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

カーネル固有のガイドライン

RMSNorm

• 入力レイアウト：[..., hidden_size]
• イプシロンのデフォルト：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 の weight が None です。ones を作成してください
• isinstance() が一致しない：代わりに type(module).__name__ を使用してください
• GEGLU が呼び出されない：モデルが使用するのは GELU であり、GEGLU ではありません
• パッチが保持されない：enable_model_cpu_offload() の前に注入してください
• カスタムカーネルで torch.compile が失敗する：以下を参照してください

torch.compile との互換性

カスタム CUDA カーネルと torch.compile は、カーネルを PyTorch カスタム op として登録しない限り、併用できません。

エラーメッセージ：

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

回避策：

1. --compile を指定せずに --use-optimized-kernels を使用する（6% の高速化）
2. カスタムカーネルを使わずに --compile を使用する（34% の高速化）
3. C++ で登録された op に fake/meta 実装を追加する（以下を参照）

op を torch.compile と互換性のあるものにするには： C++ の TORCH_LIBRARY_EXPAND 経由で登録された op は、すでに適切なカスタム op です。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 上のコミュニティカーネル