diff --git a/modules/sdnq/triton_mm.py b/modules/sdnq/triton_mm.py index beda09f03..955761e88 100644 --- a/modules/sdnq/triton_mm.py +++ b/modules/sdnq/triton_mm.py @@ -12,16 +12,31 @@ import torch import triton import triton.language as tl +try: + from .common import is_rdna2_and_older +except Exception: + is_rdna2_and_older = False -matmul_configs = [ - triton.Config({'BLOCK_SIZE_M': BM, 'BLOCK_SIZE_N': BN, "BLOCK_SIZE_K": BK, "GROUP_SIZE_M": GM}, num_warps=w, num_stages=s) - for BM in [32, 64, 128, 256] - for BN in [32, 64, 128, 256] - for BK in [32, 64, 128] - for GM in [4, 8] - for w in [4, 8] - for s in [2] -] +if is_rdna2_and_older: + matmul_configs = [ + triton.Config({'BLOCK_SIZE_M': BM, 'BLOCK_SIZE_N': BN, "BLOCK_SIZE_K": BK, "GROUP_SIZE_M": GM}, num_warps=w, num_stages=s) + for BM in [64, 128] + for BN in [64, 128] + for BK in [64] + for GM in [2, 4] + for w in [2, 4] + for s in [2] + ] +else: + matmul_configs = [ + triton.Config({'BLOCK_SIZE_M': BM, 'BLOCK_SIZE_N': BN, "BLOCK_SIZE_K": BK, "GROUP_SIZE_M": GM}, num_warps=w, num_stages=s) + for BM in [32, 64, 128, 256] + for BN in [32, 64, 128, 256] + for BK in [32, 64, 128] + for GM in [4, 8] + for w in [4, 8] + for s in [2] + ] @triton.autotune(configs=matmul_configs, key=["M", "N", "K", "stride_bk", "ACCUMULATOR_DTYPE"], cache_results=True)