Skip to content

Commit 25f11bb

Browse files
committed
refactor(triton): share INT8 autotune kwargs
Keep both INT8 matmul kernels on one configs, cache-key, and pruning contract so future device-scoping changes cannot diverge between kernel paths. Add a focused assertion covering the shared mapping. Signed-off-by: liminfei-amd <91481003+liminfei-amd@users.noreply.github.com>
1 parent 0036a3e commit 25f11bb

2 files changed

Lines changed: 14 additions & 10 deletions

File tree

comfy_kitchen/backends/triton/quantization.py

Lines changed: 8 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -906,11 +906,13 @@ def _prune_int8_autotune_configs(configs, named_args, **kwargs):
906906

907907
_INT8_PRUNE_CONFIGS_BY = {'early_config_prune': _prune_int8_autotune_configs}
908908

909-
@triton.autotune(
910-
configs=_INT8_MATMUL_CONFIGS,
911-
key=['m', 'n', 'k', 'device_index'],
912-
prune_configs_by=_INT8_PRUNE_CONFIGS_BY,
913-
)
909+
_INT8_AUTOTUNE_KWARGS = {
910+
'configs': _INT8_MATMUL_CONFIGS,
911+
'key': ['m', 'n', 'k', 'device_index'],
912+
'prune_configs_by': _INT8_PRUNE_CONFIGS_BY,
913+
}
914+
915+
@triton.autotune(**_INT8_AUTOTUNE_KWARGS)
914916
@triton.jit
915917
def _int8_matmul_dequant_kernel(
916918
# Pointers
@@ -982,11 +984,7 @@ def _int8_matmul_dequant_kernel(
982984
c_mask = (offs_am[:, None] < m) & (offs_bn[None, :] < n)
983985
tl.store(c_ptrs, c, mask=c_mask)
984986

985-
@triton.autotune(
986-
configs=_INT8_MATMUL_CONFIGS,
987-
key=['m', 'n', 'k', 'device_index'],
988-
prune_configs_by=_INT8_PRUNE_CONFIGS_BY,
989-
)
987+
@triton.autotune(**_INT8_AUTOTUNE_KWARGS)
990988
@triton.jit
991989
def _int8_matmul_dequant_per_row_kernel(
992990
# Pointers

tests/test_int8.py

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -204,6 +204,12 @@ def test_int8_autotune_pool_creation_does_not_probe_devices(monkeypatch):
204204
def test_int8_autotune_cache_key_includes_device():
205205
from comfy_kitchen.backends.triton import quantization
206206

207+
assert quantization._INT8_AUTOTUNE_KWARGS == {
208+
"configs": quantization._INT8_MATMUL_CONFIGS,
209+
"key": ["m", "n", "k", "device_index"],
210+
"prune_configs_by": quantization._INT8_PRUNE_CONFIGS_BY,
211+
}
212+
207213
assert "device_index" in quantization._int8_matmul_dequant_kernel.keys
208214
assert "device_index" in quantization._int8_matmul_dequant_per_row_kernel.keys
209215

0 commit comments

Comments
 (0)