skills.rest
tensormux
tensormuxOFFICIAL
Explore all skills by this author →
WebsiteGitHub

Write an FP8 Kernel

Design FP8 compute kernels with E4M3/E5M2 scaling for NVIDIA Hopper and Ada GPUs.

54|7|Updated Apr 10, 2026
Search Tags:#gpu#kernel#quantization#fp8#fp8-kernel#wgmma
One-click install
npx skills add https://github.com/tensormux/kernel-skills --skill write-an-fp8-kernel
Or copy as Structured Prompt for Agent▼
Please help me install this Agent Skill.
Skill: Write an FP8 Kernel
Source: https://github.com/tensormux/kernel-skills/tree/main/skills/quantization/write-fp8-kernel
Command: npx skills add https://github.com/tensormux/kernel-skills --skill write-an-fp8-kernel

SYSTEM DOCUMENTATION & REQUIREMENTS

What problem does it solve?

This Skill guides architects and engineers to design and implement FP8 compute kernels for inference and training on NVIDIA Hopper (sm_90) and Ada Lovelace (sm_89) hardware, addressing FP8 format selection, scaling strategy, tensor core usage via WGMMA or cuBLAS, and dequantization epilogue design.

Core Features & Use Cases

  • FP8 format mapping and per-tensor / per-channel scaling strategies for weights and activations, including E4M3 for forward paths and E5M2 for gradients.
  • Choice of implementation path (cuBLAS FP8, CUTLASS FP8, or custom WGMMA) with guidance on when to prefer each.
  • Epilogue design that dequantizes FP8 results back to fp16/bf16/fp32, with proper saturation (satfinite) and fp32 accumulation in the GEMM.
  • Validation and benchmarking workflow to compare FP8 results against fp32 references and monitor overflow conditions.

Quick Start

Implement a small FP8 GEMM by selecting E4M3 for inputs, setting per-tensor scaling, using cuBLAS FP8 paths, and verifying results against a fp32 baseline.

Frequently Asked Questions about Write an FP8 Kernel

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

FAQPage Schema
How do I implement an FP8 GEMM kernel for Hopper GPUs?▼

To implement an FP8 GEMM kernel for Hopper GPUs, select E4M3 format for inputs, apply per-tensor scaling, utilize cuBLAS FP8 paths or WGMMA instructions, and validate output against an fp32 baseline.

When should I use E4M3 versus E5M2 formats in FP8 compute kernels?▼

Use the E4M3 format for forward pass weights and activations, and use the E5M2 format for gradients in training to maintain accuracy while maximizing throughput across GEMM workloads.

Does cuBLAS support FP8 tensor core operations with per-tensor scaling?▼

Yes, cuBLAS supports FP8 tensor core operations with per-tensor scaling on Hopper and Ada GPUs, enabling high-throughput GEMM execution with fp32 accumulation and dequantization epilogues.

What is the best way to dequantize FP8 results back to fp32 during the GEMM epilogue?▼

The best way to dequantize FP8 results is to craft a precise scaled epilogue that applies satfinite saturation and leverages fp32 accumulation to convert outputs back to fp16, bf16, or fp32 formats.

Why does my FP8 kernel overflow during training on sm_90 hardware?▼

FP8 kernel overflow during training typically occurs when scaling factors are improperly calibrated or satfinite saturation is not applied, requiring rigorous validation against fp32 references to monitor conditions.

Can I use WGMMA for custom FP8 kernels on Ada Lovelace sm_89 GPUs?▼

Yes, you can use WGMMA for custom FP8 kernels on Ada Lovelace sm_89 GPUs, choosing between CUTLASS FP8 templates or custom WGMMA paths depending on your specific throughput and accuracy requirements.

Related Skills

Explore More Software Engineering Skills→
tile-aitile-ai
168

tune-multiplication

Optimize GEMV, GEMM, and grouped GEMM kernels on NVIDIA Ampere and Hopper GPUs.

Official
Advanced
drunkcodingdrunkcoding
4

benchmark-kernel

Benchmark FlashInfer kernel performance across backends using CUPTI timing.

Community
Intermediate
KrxGuKrxGu
54

Write CUDA GEMM Kernel

Design and implement a CUDA GEMM kernel with FP32 accumulation and tiling.

Community
Advanced
jayll1303jayll1303
17

tensorrt-llm

Convert HuggingFace checkpoints into optimized TensorRT-LLM engines with quantization.

Community
Advanced