# 用于扩散器和变压器的 CUDA 内核此技能提供了开发针对 NVIDIA GPU（H100、A100、T4）的优化 CUDA 内核的模式和指导，以便与 HuggingFace **diffusers** 和 **transformers** 库一起使用。 ## 快速入门 ### 扩散器（视频/图像生成） **用于基准测试内核性能：** # 优化内核的基准（6% 端到端加速） pythongenerate_video.py --use-optimized-kernels # 使用 torch.compile 进行基准测试（加速 34%） pythongenerate_video.py --no-optimized-kernels --compile # 比较配置（注意：--compile 和 --use-optimized-kernels 是互斥的） pythongenerate_video.py --use-optimized-kernels && \ pythongenerate_video.py --no-optimized-kernels --compile **对于最小扩散器集成示例（约 150 行）：** python 脚本/ltx_kernel_injection_example.py ### 变形金刚（法学硕士） **对于最小变压器集成示例（约 120 行）：** python 脚本/transformers_injection_example.py ### HuggingFace 内核中心 **从 HuggingFace Hub 加载预编译内核（无本地编译）：** 从内核导入 get_kernel # 加载优化的激活内核激活= get_kernel（“内核社区/激活”，版本= 1） # 使用内核 y = torch.empty_like(x) 激活.gelu_fast(y, x) **完整的 HuggingFace 内核示例：** python 脚本/huggingface_kernels_example.py ### 独立内核微基准测试 python benchmark_rmsnorm.py ## 支持的库和模型图书馆 |支持的型号 |关键内核 ---|---|--- **扩散器** | LTX-视频、稳定扩散、FLUX、DiT | RMSNorm、GEGLU、RoPE、AdaLN **变压器** | LLaMA、米斯特拉尔、Qwen、Falcon | RMSNorm，注意图形处理器 |计算能力|指南 ---|---|--- 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/ltx_video/”中找到。这表明： * 自定义 CUDA 内核（RMSNorm、RoPE 3D、GEGLU、AdaLN） * 使用 setup.py、build.toml 和 flake.nix 构建系统设置 * PyTorch C++ 绑定和 Python API * 用于比较优化性能与基准性能的基准测试脚本 ## 基准测试内核使用基准测试脚本来测量内核性能： # 包含所有选项的完整基准测试 python 脚本/benchmark_example.py \ --使用优化内核 \ --编译\ --批量大小 1 \ --num-frames 161 \ --高度512 \ --宽度768 \ --步骤 50 \ --预热迭代 2 ### 基准脚本选项选项 |默认 |描述 ---|---|--- `--使用优化内核` |汽车 |使用自定义 H100 CUDA 内核 `--无优化内核` | - |使用基线实施 `--编译` |假 |在变压器上启用 torch.compile `--批量大小` | 1 |每个提示的视频数量 `--num-frames` | 161 | 161要生成的帧数 `--高度` | 512 | 512视频高度（以像素为单位） `--宽度` | 768 | 768视频宽度（以像素为单位） `--步骤` | 50 | 50去噪步骤 `--预热迭代` | 2 |在基准测试之前运行预热 ### 基准结果示例 **端到端视频生成（49 帧，30 步，H100 80GB）：** 配置|时间|它/它 |加速|注释 ---|---|---|---|--- 基线（未编译）| 2.87 | 2.87 12.58 | 12.58 1.00 倍 |参考 **优化内核** | 2.70 | 2.70 13.52 | 13.52 **1.06x** |速度提高 6% 基线 + torch.compile | 2.14 | 2.14 19.05 | 19.05 1.34 倍 |速度提高 34% **重要提示：** `--use-optimized-kernels` 和 `--compile` 目前是互斥的。自定义内核需要 PyTorch 自定义操作注册才能与 torch.compile 一起使用。 **要捕获的关键指标：** * **设备：** GPU 型号（例如 NVIDIA H100 80GB HBM3） * **精度：** 使用的数据类型（例如，bfloat16） * **分辨率：** 宽度 x 高度（例如，768x512） * **帧：** 生成的帧数（例如 49、161） ### RMSNorm 微基准测试与 PyTorch 基线相比，矢量化 RMSNorm 内核实现了 **2.67 倍平均加速**：形状|自定义（毫秒）| PyTorch（毫秒）|加速比 ---|---|---|--- [1×1024×2048] | 0.019 | 0.019 0.065 | 0.065 **3.37x** [2×1024×2048] | 0.024 | 0.024 0.073 | 0.073 **3.04x** [4×1024×2048] | 0.036 | 0.036 0.093 | 0.093 **2.58x** [2×4096×3072] | 0.087 | 0.087 0.208 | 0.208 **2.41x** [4×4096×3072] | 0.157 | 0.157 0.392 | 0.392 **2.49x** **带宽效率：** H100 理论 3.35 TB/s 的 38% **为什么端到端加速较小：** RMSNorm 约占 LTX-Video 总计算的 5%。剩下的时间花在注意力（Flash Attention/SDPA）、线性投影和 VAE 解码上。 ## 项目结构 .claude/技能/cuda-kernels/ ├── 脚本/ │ ├── benchmark_example.py # 端到端视频生成基准测试 │ ├── benchmark_rmsnorm.py # 孤立的 RMSNorm 微基准测试 │ ├── ltx_kernel_injection_example.py # 最小扩散器集成（~150 行） │ ├── Transformers_injection_example.py # 最小化 Transformer 集成（~120 行） │ └── Huggingface_kernels_example.py # HuggingFace Kernels Hub 集成 ├── 参考资料/ │ ├── differs-integration.md # 完整的扩散器集成指南 │ ├── Transformers-integration.md # 完整的 Transformer 集成指南 │ ├── Huggingface-kernels-integration.md # HuggingFace Kernels Hub (get_kernel) 指南 │ ├── Troubleshooting.md # 常见问题及解决方案 │ ├── kernel-templates.md # CUDA 内核模板（包括矢量化） │ ├── h100-optimization-guide.md # H100（Hopper）优化深入探讨 │ ├── a100-optimization-guide.md # A100（安培）优化深入探讨 │ └── t4-optimization-guide.md # T4（图灵）优化深入探讨 └── SKILL.md # 这个文件 example/ltx_video/ # 完整的工作示例 ├── kernel_src/ │ └── rmsnorm.cu # 矢量化 RMSNorm 内核（快 2.67 倍） ├── torch-ext/ # PyTorch 绑定 ├──generate_video.py # 完整的基准测试脚本 ├── benchmark_rmsnorm.py # 隔离内核基准测试 └── setup.py # pip install -e . ## GPU 架构参考 ### H100（漏斗）- 主要目标规格|价值|优化影响 ---|---|--- 短信 | 132 | 132网格大小：目标是 132 的倍数线程/SM | 2048 | 2048每个 SM 最多 16 个块，每块 128 个线程共享内存| 192 KB/SM |可以使用大瓷砖二级缓存 | 50 MB |跨块重用内存带宽 | 3.35 TB/秒 |合并访问至关重要经纱尺寸| 32 | 32所有缩减均使用扭曲洗牌 ### 快速比较（H100 vs A100 vs T4）规格| H100 | A100 | T4 ---|---|---|--- 短信 | 132 | 132 108 | 108 40 内存带宽 | 3.35 TB/秒 | 2.0 TB/秒 | 320GB/秒共享内存/SM | 192 KB | 164 KB | 64KB BF16 支持 |是的 |是的 | **否（仅限 FP16）** 计算上限 | sm_90 | sm_80 | sm_75 > 查看详细指南：[H100]() | [A100]() | [T4]() ## 核心内核模式 ### 向量化内存访问（对于性能至关重要） **使用`__nv_bfloat162`进行BFloat16矢量化：** // 一次加载 2 个 bfloat16 元素（32 位加载） const __nv_bfloat162* vec_input = reinterpret_cast(row_input); #pragma 展开 4 for (int i = tid; i < vec_hidden; i += stride) { __nv_bfloat162 v = vec_input[i]; 浮动 v0 = __bfloat162float(v.x); 浮动 v1 = __bfloat162float(v.y); sum_sq += v0 * v0 + v1 * v1; } **使用`__half2`进行FP16矢量化：** const __half2* vec_input = reinterpret_cast(row_input); __half2 v = vec_input[i]; 浮点 v0 = __half2float(v.x); 浮点 v1 = __half2float(v.y); **使用 `float4` 进行 FP32 矢量化：** const float4* vec_input = reinterpret_cast(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; ### 扭曲洗牌减少模板 <类型名称 T> __device__ __forceinline__ T warp_reduce_sum(T val) { #pragma展开 for (int offset = 16; offset > 0; offset >>= 1) { val += __shfl_xor_sync(0xffffffff, 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）： // 除以 2 进行 bf16/fp16 向量化访问 int 线程 = min(hidden_size / 2, MAX_THREADS); 线程数 = max(线程数, WARP_SIZE); 线程 = (线程 + 32 - 1) / 32 * 32; // 舍入到扭曲边界 ## 支持的数据类型所有内核都支持三种精度模式： * `__half` (FP16) - 默认推理 * `__nv_bfloat16` (BF16) - 首选训练 * `float` (FP32) - 参考/调试 ## 构建内核 ### 与 Nix 一起（推荐） nix run .#build-and-copy --max-jobs 2 --cores 8 -L ### 使用 pip/uv uv pip install -e 。 ### build.toml 配置 [一般] 名称 =“ltx_kernels” 后端= [“cuda”] [内核.your_kernel] 后端=“cuda” src = ["kernel_src/your_kernel.cu"] cuda-功能= [“9.0”] ## 库集成 ### HuggingFace 内核中心 (get_kernel) > **请参阅 [huggingface-kernels-integration.md]() 了解完整指南。** 直接从 HuggingFace Hub 加载预编译、优化的内核，无需本地编译：从内核导入 get_kernel, has_kernel # 检查可用性和负载 if has_kernel("kernels-community/activation"): 激活= get_kernel（“内核社区/激活”，版本= 1） # 使用内核 x = torch.randn((4, 4), dtype=torch.float16, device="cuda") y = torch.empty_like(x) 激活.gelu_fast(y, x) **主要功能：** * `get_kernel(repo_id, version=None)` \- 从 Hub 下载并加载内核 * `has_kernel(repo_id)` \- 检查是否存在兼容的版本 * `get_local_kernel(path)` \- 从本地目录加载（开发） **流行的社区内核：** * `kernels-community/activation` \- GELU、SiLU 等。 * `kernels-community/flash-attn` \- Flash Attention 2 * `kernels-community/triton-layer-norm` \- LayerNorm, RMSNorm ### 扩散器集成（视频/图像生成） > **请参阅 [diffusers-integration.md]() 了解完整指南。** ### 变压器集成（法学硕士） > **请参阅 [transformers-integration.md]() 了解完整指南。** **与扩散器的主要区别：** * Transformers RMSNorm **始终**具有权重（无 `elementwise_affine=False`） * 在 class_name 中使用 `'RMSNorm' 来匹配 LlamaRMSNorm、MistralRMSNorm 等。 * 检查 epsilon 的“variance_epsilon”（LLaMA）或“eps”（其他） * 没有 `set_processor()` 模式 - 使用 Flash Attention 2 代替 **最小变形金刚图案：** 从变压器导入 AutoModelForCausalLM 从 ltx_kernels 导入 rmsnorm def patch_rmsnorm(模型): 对于名称，model.named_modules()中的模块：如果类型（模块）.__name__ 中为“RMSNorm”： eps = getattr(模块, 'variance_epsilon', None) 或 getattr(模块, 'eps', 1e-6) def make_forward(mod, epsilon): def 前向(x): 返回 rmsnorm(x, mod.weight, eps=epsilon) 向前返回 module.forward = make_forward(模块, eps) 模型 = AutoModelForCausalLM.from_pretrained("meta-llama/Llama-2-7b-hf", torch_dtype=torch.bfloat16) patch_rmsnorm（模型） ### 扩散器的关键陷阱 #### 1\. RMSNorm 重量可能为无 LTX-Video 对某些 RMSNorm 模块使用 `elementwise_affine=False`： # 变压器块：无重量 self.norm1 = RMSNorm(dim, elementwise_affine=False) # 注意力模块：有重量 self.norm_q = torch.nn.RMSNorm(..., elementwise_affine=True) **解决方案：** 处理这两种情况： has_weight = hasattr(module, 'weight') 并且 module.weight 不是 None 如果有_权重：输出 = rmsnorm(x, module.weight, eps=eps) 其他：重量 = torch.ones(x.shape[-1], device=x.device, dtype=x.dtype) 输出 = rmsnorm(x, 重量, eps=eps) #### 2\.扩散器 RMSNorm != torch.nn.RMSNorm # 错误 - 错过扩散器 RMSNorm if isinstance(模块, torch.nn.RMSNorm): # 正确 - 捕获所有 RMSNorm 变体如果类型（模块）.__name__ == 'RMSNorm'： #### 3\. LTX-Video 使用 GELU，而不是 GEGLU LTX-Video 使用 `activation_fn="gelu-approximate"`。不要为 LTX-Video 修补 GEGLU。 #### 4\.在 CPU 卸载之前注入内核管道 = LTXPipeline.from_pretrained(...) 管道.to(“cuda”) ject_optimized_kernels(pipe) # 卸载之前 pipeline.enable_model_cpu_offload() # 现在安全了 ### 最小集成模式从扩散器导入 LTXPipeline 从 ltx_kernels 导入 rmsnorm def patch_rmsnorm_modules（模型）： """修补所有 RMSNorm 模块以使用自定义内核。""" 对于名称，model.named_modules()中的模块：如果类型（模块）.__name__ == 'RMSNorm'： eps = getattr(模块, 'eps', 1e-6) has_weight = hasattr(module, 'weight') 并且 module.weight 不是 None 如果有_权重： def make_forward(mod, epsilon): def 前向(x): 返回 rmsnorm(x, mod.weight, eps=epsilon) 向前返回 module.forward = make_forward(模块, eps) 其他： def make_forward(epsilon): def 前向(x): w = torch.ones(x.shape[-1], device=x.device, dtype=x.dtype) 返回 rmsnorm(x, w, eps=epsilon) 向前返回 module.forward = make_forward(eps) # 用法管道 = LTXPipeline.from_pretrained("Lightricks/LTX-Video", torch_dtype=torch.bfloat16) 管道.to(“cuda”) patch_rmsnorm_modules(管道.变压器) pipeline.enable_model_cpu_offload() ## 内核特定指南 ### 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% ### 绳索 * 1D：`[batch，seq，heads，head_dim]` \- 用于文本 * 3D：`[batch，t*h*w，heads，head_dim]` \- 用于视频 * LTX-Video 通过“LTXVideoRotaryPosEmbed”计算自己的 RoPE ### 格鲁 vs 格鲁 * **GEGLU** ：输入 `[batch, seq, 2*hidden]` -> 输出 `[batch, seq, hide]` * **GELU** ：标准激活 * **LTX-Video 使用 GELU，而不是 GEGLU** ### 阿达恩 * 公式：`norm(x) * 权重 * (1 + scale) + shift` * 用于 DiT 块中进行调节 ## 性能分析 # NVIDIA Nsight 系统 nsys 配置文件 -o 配置文件 python your_script.py # NVIDIA Nsight 计算 ncu --set full -o 指标 python your_script.py ## 常见问题 > **请参阅 [troubleshooting.md]() 了解所有常见问题和解决方案。** 快速修复： * **“NoneType没有连续的属性”**：RMSNorm权重为None，创建一个 * **isinstance() 不匹配** : 使用 `type(module).__name__` 代替 * **GEGLU not called** ：模型使用 GELU，而不是 GEGLU * **修补不会持续** ：在 `enable_model_cpu_offload()` 之前注入 * **torch.compile 使用自定义内核失败**：见下文 ### torch.compile 兼容性自定义 CUDA 内核和 `torch.compile` 是**互斥的**，除非您将内核注册为 PyTorch 自定义操作。 **错误信息：** torch._dynamo.exc.Unsupported：尝试调用标记为已跳过的函数 **解决方法选项：** 1. 使用“--use-optimized-kernels”而不使用“--compile”（加速 6%） 2. 使用 `--compile` 而不使用自定义内核（加速 34%） 3. 将内核注册为自定义操作（高级，需要 `torch.library`） **注册为自定义操作（为了 torch.compile 兼容性）：** 进口火炬 @torch.library.custom_op("ltx_kernels::rmsnorm", mutates_args={"out"}) def rmsnorm(输出：torch.Tensor，输入：torch.Tensor，权重：torch.Tensor，eps：float) -> 无： ops.rmsnorm_forward（输出，输入.连续（），重量.连续（），eps） @rmsnorm.register_fake def _(输出、输入、重量、eps)： pass # 没有形状变化 ## 另请参阅 ### 脚本 * [benchmark_example.py]() \- **用于比较优化与基线的基准测试脚本 - 从这里开始** * [ltx_kernel_injection_example.py]() \- 最小扩散器集成（~150 行） * [transformers_injection_example.py]() \- 最小转换器/LLM 集成（~120 行） * [huggingface_kernels_example.py]() \- HuggingFace Kernels Hub 集成 ### 集成指南 * [huggingface-kernels-integration.md]() \- **HuggingFace Kernels Hub (get_kernel) - 加载预编译内核** * [diffusers-integration.md]() \- 完整的扩散器管道集成 * [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/ltx_video/]() \- 完整的 LTX-Video 示例目录 ### 外部资源 * [HuggingFace 内核文档]() * [HuggingFace 内核 GitHub]() * [集线器上的社区内核]()

CUDA内核

包含技能