skills.rest
huggingface
huggingfaceOFFICIAL
Explore all skills by this author →
Website𝕏@huggingfaceGitHub

triton-kernels

Guides writing, testing, and benchmarking portable Triton GPU kernels for NVIDIA and AMD hardware.

736|125|Updated Nov 29, 2024
Search Tags:#benchmarking#huggingface#CUDA#autotune#gpu kernels#pyTorch#Triton
One-click install
npx skills add https://github.com/huggingface/kernels --skill triton-kernels
Or copy as Structured Prompt for Agent▼
Please help me install this Agent Skill.
Skill: triton-kernels
Source: https://github.com/huggingface/kernels/tree/main/kernel-builder/skills/triton-kernels
Command: npx skills add https://github.com/huggingface/kernels --skill triton-kernels

SYSTEM DOCUMENTATION & REQUIREMENTS

💡 This Skill requires torch, triton, and includes scripts (resource) and references (resource) components.

What problem does it solve?

Writing correct and fast Triton GPU kernels requires navigating subtle pitfalls like masked load fill values, fp32 accumulation, autotune configuration, and reduction block sizing, where mistakes silently produce wrong results or poor performance.

Core Features & Use Cases

  • Core DSL Patterns: Provides reference implementations for softmax, tiled matmul, and RMSNorm with correct masking, pointer arithmetic, and numerics handling.
  • Autotune & Benchmarking Guidance: Covers @triton.autotune configuration, num_warps/num_stages heuristics, GB/s and TFLOPS measurement, and correctness testing against PyTorch references.
  • Hub Integration: Explains publishing Triton kernels to the HuggingFace Kernels Hub and patching Transformers models with custom kernels.
  • Use Case: An engineer optimizing an LLM's RMSNorm layer uses the skill's patterns to write a fused Triton kernel, validates it against PyTorch across irregular shapes, benchmarks throughput in GB/s, and publishes it via kernel-builder.

Quick Start

Ask the AI to write a Triton softmax kernel with correctness tests and a benchmark comparing it against PyTorch.

Frequently Asked Questions about triton-kernels

High-intent search queries and answers about installing and using this skill.

FAQPage Schema
How do I write a Triton kernel for softmax or RMSNorm?▼

Use one program per row with BLOCK_SIZE set to triton.next_power_of_2 of the row dimension, load with masked offsets, accumulate in fp32, and store back in the input dtype. The skill provides complete reference implementations for softmax, RMSNorm, and matmul.

How do I benchmark a Triton kernel against PyTorch?▼

Use triton.testing.do_bench for timing and triton.testing.perf_report with Benchmark to generate comparison plots. Report GB/s for memory-bound kernels like softmax and TFLOPS for compute-bound kernels like matmul.

Can Triton kernels run on both NVIDIA and AMD GPUs?▼

Yes, Triton kernels written with portable patterns run on both NVIDIA and AMD GPUs without modification. Backend-specific tuning for AMD ROCm or Intel XPU is covered by separate companion skills.

Why does my Triton kernel give wrong results after autotuning?▼

Autotuning a BLOCK_SIZE that controls the reduction dimension causes silent wrong results when the chosen value is smaller than the actual dimension. Compute reduction block sizes with triton.next_power_of_2 in the wrapper instead of autotuning them.

When should I not use a custom Triton kernel?▼

Avoid custom Triton kernels when PyTorch already has highly optimized implementations like cuBLAS matmul or Flash Attention, when inputs are small enough that launch overhead dominates, or for simple element-wise chains that torch.compile already fuses.

How do I publish a Triton kernel to the HuggingFace Hub?▼

Create a build.toml with backends set to triton, a repo-id, and version, then run kernel-builder build-and-upload. Others can load it with get_kernel from the kernels package without local compilation.

Related Skills

Explore More Software Engineering Skills→
sgl-projectsgl-project
31.6k

diffusion-kernel

Integrate custom JIT CUDA and Triton kernels to accelerate diffusion model inference.

Official
Advanced
m0atm0at
1

cuda-kernels

Develop and benchmark optimized CUDA kernels for HuggingFace diffusers and transformers.

Community
Advanced
hsliuustc0106hsliuustc0106
84

vllm-omni-hardware

Configure vLLM-Omni hardware backends across CUDA, ROCm, NPU, and XPU.

Community
Advanced
TorchedHatTorchedHat
6

compile-trace-inductor

Diagnose PyTorch Inductor compilation and performance issues by tracing IR transformations.

Official
Advanced