Comfy-Org / Comfy-Org/comfy-kitchen

INT8 matmul autotune uses NVIDIA-tuned configs on AMD RDNA

Open
#68 1 comment 0 reactions 0 assignees View on GitHub
Dominant language
Python
Stars
220
Forks
91
Avg merge
1d 7h
Merged PRs (30d)
12

Description

The two INT8 matmul Triton kernels ship 6 autotune configs tuned for NVIDIA
(`num_stages=3–4`, `block_n=256`). On AMD RDNA these are suboptimal — RDNA prefers shallow
software pipelining (`num_stages=2`) and deeper `block_k`.

A `triton.testing.do_bench` sweep on gfx1100 (RDNA3) and gfx1201 (RDNA4) shows an RDNA-tuned
`num_stages=2` pool is ~1.03–1.11× faster per shape across typical DiT/SD INT8 GEMM shapes.

**Proposal:** add `_int8_autotune_configs()` that returns an RDNA pool when `torch.version.hip`
is set, keeping the NVIDIA list unchanged for CUDA. INT8 GEMM is int32-exact, so this changes
speed only (outputs byte-identical). PR ready.

Contributor guide

Open the contributing guide

Research direction

Locate the two INT8 matmul Triton kernels and inspect their existing six NVIDIA autotune configs, then trace the torch.version.hip entry point. Done means an RDNA pool is selected on HIP while the NVIDIA list remains unchanged on CUDA, with byte-identical outputs and the reported gfx1100/gfx1201 benchmark improvement.

Written by the indexing model from the issue text.

Assessment

Tech stack
python, pytorch
Domain
machine-learning, performance
Issue type
Feature
Difficulty
3/5
Estimated time
1-2 days
Activity status
Quiet
Clarity
Mostly clear
Newbie friendliness
55/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.