# How to Profile DeepSeek-V3 Inference Kernels Using Triton Autotuning

> Profile DeepSeek-V3 inference kernels efficiently with Triton autotuning. Learn to capture runtime selections or benchmark specific kernel configurations for optimized performance.

- Repository: [DeepSeek/DeepSeek-V3](https://github.com/deepseek-ai/DeepSeek-V3)
- Tags: performance
- Published: 2026-02-26

---

**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`](https://github.com/deepseek-ai/DeepSeek-V3/blob/main/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)](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`](https://github.com/deepseek-ai/DeepSeek-V3/blob/main/inference/kernel.py), the `fp8_gemm_kernel` is decorated with a predefined configuration list that Triton searches at runtime:

```python
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.

```python
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`](https://github.com/deepseek-ai/DeepSeek-V3/blob/main/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`](https://github.com/deepseek-ai/DeepSeek-V3/blob/main/inference/generate.py).

### Benchmarking the FP8 GEMM Kernel

First, prepare FP8 tensors and scaling factors matching the expected layout in [`inference/kernel.py`](https://github.com/deepseek-ai/DeepSeek-V3/blob/main/inference/kernel.py):

```python
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:

```python
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:

```python
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`](https://github.com/deepseek-ai/DeepSeek-V3/blob/main/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`](https://github.com/deepseek-ai/DeepSeek-V3/blob/main/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`](https://github.com/deepseek-ai/DeepSeek-V3/blob/main/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`](https://github.com/deepseek-ai/DeepSeek-V3/blob/main/inference/kernel.py) to identify the actual fastest parameters for your specific device.