How to Profile DeepSeek-V3 Inference Kernels Using Triton Autotuning

You can profile DeepSeek-V3's FP8 inference kernels by setting TRITON_DEBUG=1 to capture autotune selections at runtime, or by using triton.testing.benchmark to isolate and validate specific kernel configurations from inference/kernel.py.

DeepSeek-V3 achieves high-throughput FP8 inference through custom Triton kernels that are automatically tuned for specific GPU architectures. Understanding how to profile these kernels is essential for optimizing latency and verifying that the autotuned configurations selected in the DeepSeek-V3 codebase are optimal for your hardware.

Understanding the Autotuned Kernels in DeepSeek-V3

DeepSeek-V3 relies on two performance-critical Triton kernels wrapped with the autotune decorator:

• act_quant_kernel – Performs block-wise activation quantization, producing FP8 tensors and per-block scaling factors.
• fp8_gemm_kernel – Executes matrix multiplication on FP8 matrices with per-block scaling.

Both kernels are defined in [inference/kernel.py](https://github.com/deepseek-ai/DeepSeek-V3/blob/main/inference/kernel.py) and use the @triton.autotune decorator to generate multiple launch configurations.

Autotune Configuration Structure

In inference/kernel.py, the fp8_gemm_kernel is decorated with a predefined configuration list that Triton searches at runtime:

fp8_gemm_configs = [
    Config({'BLOCK_SIZE_M': block_m,
            'BLOCK_SIZE_N': block_n,
            'BLOCK_SIZE_K': 128},
           num_stages=num_stages,
           num_warps=8)
    for block_m in [16, 32, 64]
    for block_n in [32, 64, 128]
    for num_stages in [3, 4, 5, 6]
]

@triton.autotune(configs=fp8_gemm_configs, key=['N', 'K'])
@triton.jit
def fp8_gemm_kernel(...):
    ...

The key=['N', 'K'] parameter instructs Triton to cache and select configurations based on matrix dimensions, ensuring the fastest combination of block size and pipeline stage count is chosen for each problem size.

Method 1: Runtime Profiling with TRITON_DEBUG

The simplest way to profile DeepSeek-V3 inference kernels is to enable Triton's built-in debug output. Set the environment variable TRITON_DEBUG=1 before running the model to emit the selected configuration for every autotuned kernel launch.

import os
import torch
from inference.model import Transformer, ModelArgs

os.environ["TRITON_DEBUG"] = "1"

args = ModelArgs()
model = Transformer(args).cuda()
tokens = torch.randint(0, args.vocab_size, (1, 32), device='cuda')
logits = model(tokens)

When executed, Triton prints the selected configuration for each kernel:

[triton] selected config for fp8_gemm_kernel:
{'BLOCK_SIZE_M': 32, 'BLOCK_SIZE_N': 64, 'BLOCK_SIZE_K': 128, 'num_stages': 4, 'num_warps': 8}

This output reveals which block dimensions and stage counts the autotuner selected for your specific GPU and input shapes, allowing you to verify the decisions made by the decorator in inference/kernel.py.

Method 2: Isolated Benchmarking with triton.testing

For granular performance analysis, use triton.testing.benchmark to measure kernel latency in isolation. This approach allows you to verify autotune selections or test custom configurations without running the full model through inference/generate.py.

Benchmarking the FP8 GEMM Kernel

First, prepare FP8 tensors and scaling factors matching the expected layout in inference/kernel.py:

import torch
import triton
import triton.testing as testing
from inference.kernel import fp8_gemm_kernel

M, N, K = 2048, 2048, 2048
a = torch.randn(M, K, dtype=torch.float8_e4m3fn, device='cuda')
b = torch.randn(K, N, dtype=torch.float8_e4m3fn, device='cuda')
a_s = torch.full((M // 128, K // 128), 1.0, dtype=torch.float32, device='cuda')
b_s = torch.full((K // 128, N // 128), 1.0, dtype=torch.float32, device='cuda')
c = torch.empty(M, N, dtype=torch.float32, device='cuda')

Then benchmark a specific configuration by explicitly defining the grid and block parameters:

grid = (triton.cdiv(M, 32), triton.cdiv(N, 64))

latency_us = testing.benchmark(
    lambda: fp8_gemm_kernel[grid](
        a, b, c, a_s, b_s,
        M, N, K,
        BLOCK_SIZE_M=32, BLOCK_SIZE_N=64, BLOCK_SIZE_K=128
    ),
    rep=30,
)

print(f"FP8 GEMM latency: {latency_us.mean():.2f} µs ± {latency_us.stdev():.2f} µs")

By varying BLOCK_SIZE_M, BLOCK_SIZE_N, num_stages, and other parameters, you can reproduce the full autotune search space and identify hardware-specific sweet spots that may differ from Triton's cached selection.

Benchmarking the Activation Quantization Kernel

Similarly, profile the act_quant_kernel to measure quantization overhead separately from matrix multiplication:

from inference.kernel import act_quant_kernel, act_quant

x = torch.randn(1024, 128, dtype=torch.float32, device='cuda')
y, s = act_quant(x, block_size=128)

lat_us = testing.benchmark(
    lambda: act_quant_kernel[(triton.cdiv(x.numel(), 128),)](
        x, y, s,
        BLOCK_SIZE=128, scale_fmt="ue8m0"
    ),
    rep=50,
)
print(f"act_quant latency: {lat_us.mean():.2f} µs")

Analyzing FP8 Quantization Sensitivity

Profiling both kernels helps identify bottlenecks in the FP8 inference pipeline. The act_quant_kernel generates scaling factors (a_s, b_s) that feed into fp8_gemm_kernel. If quantization latency dominates—particularly at small batch sizes—optimizing the BLOCK_SIZE in the activation kernel may yield better end-to-end performance than tuning the GEMM alone.

Hardware-specific considerations matter when profiling DeepSeek-V3:

• On NVIDIA Hopper GPUs, larger block sizes (e.g., BLOCK_SIZE_M=64) often increase occupancy due to enhanced shared memory bandwidth.
• On Ampere GPUs, reducing BLOCK_SIZE_K may alleviate register pressure and improve utilization.

Complete Profiling Workflow

A systematic approach to optimizing DeepSeek-V3 inference involves:

1. Capture autotune selections – Run export TRITON_DEBUG=1 and execute python -m inference.generate to log which configurations Triton selects for your specific batch sizes and sequence lengths.
2. Isolate kernels – Create standalone benchmark scripts using triton.testing.benchmark for fp8_gemm_kernel and act_quant_kernel imported from inference/kernel.py.
3. Sweep parameters – Test all combinations of block sizes, pipeline stages, and warp counts to verify the autotuned choice is optimal for your GPU architecture.
4. Validate changes – After modifying kernel code (e.g., adding vectorized loads or changing data layouts), repeat the benchmark to quantify performance gains before committing changes.

Summary

• DeepSeek-V3's inference performance depends on autotuned Triton kernels in inference/kernel.py, specifically fp8_gemm_kernel and act_quant_kernel.
• Set TRITON_DEBUG=1 to view which kernel configurations Triton selects at runtime without modifying code.
• Use triton.testing.benchmark to isolate kernels and measure latency for specific block sizes and pipeline stages.
• Profile both quantization and GEMM kernels to identify whether data type conversion or matrix multiplication dominates latency.
• Benchmark across different configurations to find hardware-specific optimizations for NVIDIA Hopper or Ampere GPUs.

Frequently Asked Questions

How do I know which autotune configuration is being selected during inference?

Set the environment variable TRITON_DEBUG=1 before launching your script. Triton will print the selected configuration—including block sizes, number of stages, and warps—to stdout every time an autotuned kernel is launched. This works immediately with the existing code in inference/kernel.py and requires no code changes.

Can I override the autotune configuration chosen by Triton?

Yes. While Triton caches the best configuration based on the key parameters (like ['N', 'K']), you can bypass autotuning by removing the @triton.autotune decorator and manually specifying block dimensions when launching the kernel grid. For testing purposes, use triton.testing.benchmark with hardcoded BLOCK_SIZE_M, BLOCK_SIZE_N, and num_stages values to test specific hardware characteristics.

Why is profiling the act_quant_kernel separately important?

The act_quant_kernel produces scaling factors required by the fp8_gemm_kernel. If activation quantization becomes a bottleneck—particularly at small batch sizes or on GPUs with limited shared memory—optimizing the quantization block size can improve overall inference latency more than tuning the GEMM alone. Isolated profiling reveals this split in execution time and helps balance the pipeline.

What hardware considerations affect the optimal kernel configuration?

NVIDIA Hopper GPUs (H100) generally benefit from larger block sizes (64×64 or 128×128) and more pipeline stages due to increased shared memory bandwidth. Ampere GPUs (A100) may perform better with smaller BLOCK_SIZE_K values to reduce register pressure. Always profile on your target hardware using the full autotune configuration grid defined in inference/kernel.py to identify the actual fastest parameters for your specific device.

{
  "mcpServers": {
    "instagit": {
      "command": "npx",
      "args": ["-y", "instagit@latest"]
    }
  }
}