[GH-ISSUE #68] INT8 matmul autotune uses NVIDIA-tuned configs on AMD RDNA #23

New issue

Open

opened 2026-10-26 22:19:14 +00:00 by bsakaguchi · 1 comment

bsakaguchi commented

2026-10-26 22:19:14 +00:00

Owner

Originally created by @liminfei-amd on GitHub (Jul 10, 2026).
Original GitHub issue: https://github.com/Comfy-Org/comfy-kitchen/issues/68

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.

Originally created by @liminfei-amd on GitHub (Jul 10, 2026). Original GitHub issue: https://github.com/Comfy-Org/comfy-kitchen/issues/68 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.

bsakaguchi commented

2026-10-26 22:19:15 +00:00

Author

Owner

@Ph0rk0z commented on GitHub (Jul 13, 2026):

The configs for turning weren't optimal either. I end up adding them myself.

@Ph0rk0z commented on GitHub (Jul 13, 2026): The configs for turning weren't optimal either. I end up adding them myself.

bsakaguchi referenced this issue

2026-10-26 22:20:13 +00:00

[PR #23] [CLOSED] Fix FP8 tensor support on MPS backend for Apple Silicon Macs #85

bsakaguchi referenced this issue

2026-10-26 22:20:17 +00:00

[PR #34] [CLOSED] Fix FP8 dequantization on MPS by falling back to CPU #91