diff --git a/mlir/lib/Dialect/Rock/Tuning/RockTuningImpl.cpp b/mlir/lib/Dialect/Rock/Tuning/RockTuningImpl.cpp
index 355c87585c20..922b237d1109 100644
--- a/mlir/lib/Dialect/Rock/Tuning/RockTuningImpl.cpp
+++ b/mlir/lib/Dialect/Rock/Tuning/RockTuningImpl.cpp
@@ -1738,6 +1738,24 @@ static LogicalResult getTuningProblemStr(rock::RockGemmWrapperInterface gemmIF,
 // since it can store each field separately.
 // Currently serialize the problem in MIOpenDriver command friendly format
 LogicalResult getTuningProblemStr(ModuleOp mod, SmallVectorImpl<char> &out) {
+  auto serializeWithSplitKSupport = [&](auto tuningOp) -> LogicalResult {
+    if (failed(getTuningProblemStr(tuningOp, out)))
+      return failure();
+
+    // Legality must be judged on the function holding the op we just
+    // serialized. The ModuleOp overload always inspects the module's first
+    // function, which is not necessarily that one.
+    auto func = tuningOp->template getParentOfType<func::FuncOp>();
+    if (!func)
+      return failure();
+
+    llvm::raw_svector_ostream problemOS(out);
+    problemOS << " -supportsSplitK "
+              << (succeeded(rock::testFusionLegalitySplitK(func)) ? "true"
+                                                                  : "false");
+    return success();
+  };
+
   {
     rock::RockGemmWrapperInterface gemmIF;
     WalkResult findPrimary =
@@ -1746,7 +1764,7 @@ LogicalResult getTuningProblemStr(ModuleOp mod, SmallVectorImpl<char> &out) {
           return WalkResult::interrupt();
         });
     if (findPrimary.wasInterrupted())
-      return getTuningProblemStr(gemmIF, out);
+      return serializeWithSplitKSupport(gemmIF);
   }
   {
     rock::RockGemmGemmWrapperInterface gemmGemmOp;
@@ -1756,7 +1774,7 @@ LogicalResult getTuningProblemStr(ModuleOp mod, SmallVectorImpl<char> &out) {
           return WalkResult::interrupt();
         });
     if (findGemmGemm.wasInterrupted())
-      return getTuningProblemStr(gemmGemmOp, out);
+      return serializeWithSplitKSupport(gemmGemmOp);
   }
   return failure();
 }
diff --git a/mlir/test/fusion/fusability-conv-add-max.mlir b/mlir/test/fusion/fusability-conv-add-max.mlir
index 67c758cf787d..7f01309e1b33 100644
--- a/mlir/test/fusion/fusability-conv-add-max.mlir
+++ b/mlir/test/fusion/fusability-conv-add-max.mlir
@@ -2,6 +2,8 @@
 // CHECK-SPLITK: fusible:0
 // RUN: rocmlir-gen -emit-module-fusibility-for=v3:16,16,4,16,16,1,1,1,2,1,1 - < %s | FileCheck %s --check-prefixes=CHECK-NONSPLITK
 // CHECK-NONSPLITK: fusible:1
+// RUN: rocmlir-gen --emit-tuning-key - < %s | FileCheck %s --check-prefix=CHECK-TUNING-KEY
+// CHECK-TUNING-KEY: -supportsSplitK false
 module {
   func.func @mlir_convolution_add_relu(%arg0: memref<64x1x1x1xf32>, %arg1: memref<1x256x56x56xf32>, %arg2: memref<64x256x1x1xf32>, %arg3: memref<1x64x56x56xf32>) attributes {rock.enable_splitk_for_tuning, rock.kernel, mhal.arch = "amdgcn-amd-amdhsa:gfx90a:sramecc+:xnack-"} {
     %cst = arith.constant 0.000000e+00 : f32
diff --git a/mlir/test/fusion/fusability-dot-add.mlir b/mlir/test/fusion/fusability-dot-add.mlir
index b0d2380125ea..41195525c295 100644
--- a/mlir/test/fusion/fusability-dot-add.mlir
+++ b/mlir/test/fusion/fusability-dot-add.mlir
@@ -2,6 +2,8 @@
 // CHECK-SPLITK: fusible:1
 // RUN: rocmlir-gen -emit-module-fusibility-for=v3:16,16,4,16,16,1,1,1,2,1,1 - < %s | FileCheck %s --check-prefixes=CHECK-NONSPLITK
 // CHECK-NONSPLITK: fusible:1
+// RUN: rocmlir-gen --emit-tuning-key - < %s | FileCheck %s --check-prefix=CHECK-TUNING-KEY
+// CHECK-TUNING-KEY: -supportsSplitK true
 module {
   func.func @mlir_dot_add(%arg0: memref<1x2x320xf32>, %arg1: memref<1x2x1280xf32>, %arg2: memref<1x1280x320xf32>, %arg3: memref<1x2x320xf32>) attributes {rock.enable_splitk_for_tuning, rock.kernel, mhal.arch = "amdgcn-amd-amdhsa:gfx90a:sramecc+:xnack-"} {
     %alloc = memref.alloc() {alignment = 64 : i64} : memref<1x2x320xf32>
diff --git a/mlir/test/fusion/fusability-dot-mul.mlir b/mlir/test/fusion/fusability-dot-mul.mlir
index 2c87ad7a2bd5..8236f6301017 100644
--- a/mlir/test/fusion/fusability-dot-mul.mlir
+++ b/mlir/test/fusion/fusability-dot-mul.mlir
@@ -2,6 +2,8 @@
 // CHECK-SPLITK: fusible:1
 // RUN: rocmlir-gen -emit-module-fusibility-for=v3:16,16,4,16,16,1,1,1,2,1,1 - < %s | FileCheck %s --check-prefixes=CHECK-NONSPLITK
 // CHECK-NONSPLITK: fusible:1
+// RUN: rocmlir-gen --emit-tuning-key - < %s | FileCheck %s --check-prefix=CHECK-TUNING-KEY
+// CHECK-TUNING-KEY: -supportsSplitK true
 module {
   func.func @mlir_dot_mul(%arg0: memref<1x2x320xf32>, %arg1: memref<1x2x1280xf32>, %arg2: memref<1x1280x320xf32>, %arg3: memref<1x2x320xf32>) attributes {rock.enable_splitk_for_tuning, rock.kernel, mhal.arch = "amdgcn-amd-amdhsa:gfx90a:sramecc+:xnack-"} {
     %alloc = memref.alloc() {alignment = 64 : i64} : memref<1x2x320xf32>
diff --git a/mlir/test/rocmlir-gen/gemm-misc-options.mlir b/mlir/test/rocmlir-gen/gemm-misc-options.mlir
index 67e5d65406cf..f2e1b23e7ce5 100644
--- a/mlir/test/rocmlir-gen/gemm-misc-options.mlir
+++ b/mlir/test/rocmlir-gen/gemm-misc-options.mlir
@@ -9,7 +9,9 @@
 // RUN: rocmlir-gen --emit-tuning-key -p -t fp8_fp8 --arch gfx1201 | FileCheck %s --check-prefix=CONVOCPFP8
 // CONVOCPFP8: amdgcn-amd-amdhsa:gfx1201   {{.*}}     convfp8_fp8 -F 1 -f GNC01 -I NGC01 -O NGC01 -n 128 -c 8 -H 32 -W 32 -k 128 -y 3 -x 3 -p 0 -q 0 -u 1 -v 1 -l 1 -j 1 -g 1
 // RUN: rocmlir-gen --arch gfx908 --operation gemm -p --emit-tuning-key | FileCheck %s --check-prefix=GEMM
-// GEMM: amdgcn-amd-amdhsa:gfx908   {{.*}}     -t f32 -out_datatype f32 -transA false -transB false -g 1 -m 1024 -n 512 -k 769
+// GEMM: amdgcn-amd-amdhsa:gfx908   {{.*}}     -t f32 -out_datatype f32 -transA false -transB false -g 1 -m 1024 -n 512 -k 769 -supportsSplitK true
+// RUN: rocmlir-gen --arch gfx908 --operation gemm -t i8 -p --emit-tuning-key | FileCheck %s --check-prefix=GEMM_I8
+// GEMM_I8: amdgcn-amd-amdhsa:gfx908   {{.*}}     -t i8 -out_datatype i32 -transA false -transB false -g 1 -m 1024 -n 512 -k 769 -supportsSplitK false{{$}}
 // RUN: rocmlir-gen --emit-tuning-key -p -t fp8_fp8 --arch gfx950 | FileCheck %s --check-prefix=CONVOCPFP8_GFX950
 // CONVOCPFP8_GFX950: amdgcn-amd-amdhsa:gfx950   {{.*}}     convfp8_fp8 -F 1 -f GNC01 -I NGC01 -O NGC01 -n 128 -c 8 -H 32 -W 32 -k 128 -y 3 -x 3 -p 0 -q 0 -u 1 -v 1 -l 1 -j 1 -g 1
 // RUN: rocmlir-gen --arch gfx942 --operation gemm -p --num_cu 40 --num_chiplets 20 | FileCheck %s --check-prefix=NUM_CHIPLETS
diff --git a/mlir/test/rocmlir-gen/problem-key.mlir b/mlir/test/rocmlir-gen/problem-key.mlir
index d0068d4d742e..497986e6a6b2 100644
--- a/mlir/test/rocmlir-gen/problem-key.mlir
+++ b/mlir/test/rocmlir-gen/problem-key.mlir
@@ -1,5 +1,5 @@
 // RUN: rocmlir-gen --arch gfx942 --operation attention -seq_len_q 256 -seq_len_k 512 -head_dim_qk 64 -head_dim_v 32 -t f32 -g 1 | rocmlir-gen --emit-tuning-key - | FileCheck %s  --check-prefixes=CHECK_1
-// CHECK_1: -t f32 -transQ false -transK false -transV false -transO false -causal false -return_lse false -split_kv 1 -num_heads_q 1 -num_heads_kv 1 -g 1 -seq_len_q 256 -seq_len_k 512 -head_dim_qk 64 -head_dim_v 32 -with-attn-scale false -with-attn-bias false -transBias false
+// CHECK_1: -t f32 -transQ false -transK false -transV false -transO false -causal false -return_lse false -split_kv 1 -num_heads_q 1 -num_heads_kv 1 -g 1 -seq_len_q 256 -seq_len_k 512 -head_dim_qk 64 -head_dim_v 32 -with-attn-scale false -with-attn-bias false -transBias false -supportsSplitK true
 // RUN: rocmlir-gen --arch gfx942 --operation attention -seq_len_q 256 -seq_len_k 512 -head_dim_qk 64 -head_dim_v 32 -t f16 -g 4 | rocmlir-gen --emit-tuning-key - | FileCheck %s  --check-prefixes=CHECK_2
 // CHECK_2: -t f16 -transQ false -transK false -transV false -transO false -causal false -return_lse false -split_kv 1 -num_heads_q 1 -num_heads_kv 1 -g 4 -seq_len_q 256 -seq_len_k 512 -head_dim_qk 64 -head_dim_v 32 -with-attn-scale false -with-attn-bias false -transBias false
 // RUN: rocmlir-gen --arch gfx942 --operation attention -seq_len_q 256 -seq_len_k 512 -head_dim_qk 64 -head_dim_v 32 -t i8 -g 8 | rocmlir-gen --emit-tuning-key - | FileCheck %s  --check-prefixes=CHECK_3
diff --git a/mlir/utils/performance/perfRunner.py b/mlir/utils/performance/perfRunner.py
index 060f87b33480..ee633c563e8e 100644
--- a/mlir/utils/performance/perfRunner.py
+++ b/mlir/utils/performance/perfRunner.py
@@ -355,7 +355,8 @@ def get_bank_conflict(filename):
 
 
 # Tuning databases
-MaybeTuningDb = Optional[Dict[Tuple[str, int, int, str], str]]
+TuningDb = Dict[Tuple[str, int, int, str], str]
+MaybeTuningDb = Optional[TuningDb]
 
 
 def parse_tuning_db_line(
@@ -399,6 +400,49 @@ def parse_tuning_db_line(
 PARSER_EXCEPTIONS = (ValueError, IndexError, KeyError, NameError)
 
 
+def extract_tuning_key_metadata(argv: list) -> Tuple[list, Optional[bool]]:
+    """Extract metadata that identifies a tuning problem but is not a rocmlir-gen option."""
+    filtered = []
+    supports_split_k = None
+    i = 0
+    while i < len(argv):
+        if argv[i] == '-supportsSplitK':
+            if i + 1 >= len(argv):
+                raise ValueError("Missing value for tuning-key metadata -supportsSplitK")
+            value = argv[i + 1].lower()
+            if value not in ("true", "false"):
+                raise ValueError(f"Invalid value for tuning-key metadata -supportsSplitK: {value}")
+            supports_split_k = value == "true"
+            i += 2
+            continue
+        filtered.append(argv[i])
+        i += 1
+    return filtered, supports_split_k
+
+
+def infer_split_k_support(arch: str, output_dtype: str) -> bool:
+    """Mirror validOutputAtomicAdd for keys without explicit split-K metadata."""
+    required_feature = {
+        'f32': GemmFeatures.ATOMIC_ADD,
+        'f16': GemmFeatures.ATOMIC_ADD_F16,
+        'bf16': GemmFeatures.ATOMIC_ADD_BF16,
+    }.get(output_dtype.lower())
+    if required_feature is None:
+        return False
+    if not arch:
+        # Some callers construct configurations before assigning a target.
+        # Preserve the historical permissive behavior until an arch is known.
+        return True
+    features = lookup_arch_info(arch).default_features
+    return has_feature(features, required_feature)
+
+
+def resolve_split_k_support(explicit_support: Optional[bool], arch: str, output_dtype: str) -> bool:
+    if explicit_support is not None:
+        return explicit_support
+    return infer_split_k_support(arch, output_dtype)
+
+
 def canonicalize_config(config_str: str, conf_class: type, arch: str, num_cu: int,
                         num_chiplets: int) -> str:
     """Canonicalize a config by round-tripping through conf_class.from_command_line/to_command_line.
@@ -569,6 +613,10 @@ def run_pipeline(proc_specs):
 
 class PerfConfiguration:
     TABLE_COLUMNS = []
+    supports_split_k = True
+
+    def tuning_key_metadata(self) -> str:
+        return f"-supportsSplitK {str(self.supports_split_k).lower()}"
 
     def compute_tflops(self, ns: int) -> float:
         raise NotImplementedError()
@@ -730,6 +778,7 @@ def generate_mlir_driver_commandline(self, rocmlir_gen_flags, kernel_repeats=MLI
 
     @classmethod
     def from_command_line(cls, argv, arch, num_cu, num_chiplets):
+        argv, supports_split_k = extract_tuning_key_metadata(argv)
         # determine datatype from argv[1]
         # Please keep this in sync with mlir::rock::getTuningProblemStr()
         if argv[0] == 'conv':
@@ -815,9 +864,12 @@ def from_command_line(cls, argv, arch, num_cu, num_chiplets):
             else:
                 continue
 
-        return cls(datatype, direction, filter_layout, input_layout, output_layout, n, c, hi, wi, k,
-                   y, x, conv_stride_h, conv_stride_w, padding_h, padding_w, dilation_h, dilation_w,
-                   group, arch, num_cu, num_chiplets)
+        config = cls(datatype, direction, filter_layout, input_layout, output_layout, n, c, hi, wi,
+                     k, y, x, conv_stride_h, conv_stride_w, padding_h, padding_w, dilation_h,
+                     dilation_w, group, arch, num_cu, num_chiplets)
+        output_dtype = OUTPUT_DATA_TYPES_MAP.get(datatype, datatype)
+        config.supports_split_k = resolve_split_k_support(supports_split_k, arch, output_dtype)
+        return config
 
     def to_command_line(self):
         return (
@@ -828,7 +880,8 @@ def to_command_line(self):
             f"-n {self.n} -c {self.c} -H {self.hi} -W {self.wi} -k {self.k} " +
             f"-y {self.y} -x {self.x} -p {self.padding_h} -q {self.padding_w} " +
             f"-u {self.conv_stride_h} -v {self.conv_stride_w} -l {self.dilation_h} " +
-            f"-j {self.dilation_w} -m conv -g {self.group} -t 1")
+            f"-j {self.dilation_w} -m conv -g {self.group} -t 1 "
+            f"{self.tuning_key_metadata()}")
 
     def __init__(self, dtype: str, direction: str, filter_layout: str, input_layout: str,
                  output_layout: str, n: int, c: int, hi: int, wi: int, k: int, y: int, x: int,
@@ -886,10 +939,11 @@ def benchmark_external(cls, commandline, paths: Paths, arch, num_cu, num_chiplet
         # rocMLIR configs use layout names (e.g. GNC01) that MIOpenDriver rejects.
         # Translate them to NCHW/NHWC; skip configs that have no faithful MIOpen
         # equivalent instead of forcing an unfair comparison.
-        miopen_commandline = conv_commandline_to_miopen_layouts(commandline)
+        config_args, _ = extract_tuning_key_metadata(commandline)
+        miopen_commandline = conv_commandline_to_miopen_layouts(config_args)
         if miopen_commandline is None:
             print("Skipping MIOpen benchmark: conv layout has no equivalent MIOpen "
-                  f"NCHW/NHWC representation: {' '.join(commandline)}")
+                  f"NCHW/NHWC representation: {' '.join(config_args)}")
             return config.table_entry(np.nan)
         miopen_driver_cmd = [MIOPENDRIVER, *miopen_commandline, '-V', '0', '-t', '1']
         print("Running MIOpen Benchmark: ", ' '.join(miopen_driver_cmd))
@@ -1216,6 +1270,7 @@ def generate_mlir_driver_commandline(self, rocmlir_gen_flags, kernel_repeats=MLI
 
     @classmethod
     def from_command_line(cls, argv, arch, num_cu, num_chiplets):
+        argv, supports_split_k = extract_tuning_key_metadata(argv)
         # Please keep this in sync with mlir::rock::getTuningProblemStr()
         dtype = None
         g = None
@@ -1276,9 +1331,11 @@ def from_command_line(cls, argv, arch, num_cu, num_chiplets):
             if v is None:
                 raise ValueError("Incomplete GEMM configuration")
 
-        return cls(dtype, out_dtype, g, m, k, n, trans_a, trans_b, scaled_gemm, scale_a_dtype,
-                   scale_b_dtype, trans_scale_a, trans_scale_b, arch, num_cu, num_chiplets,
-                   perf_config)
+        config = cls(dtype, out_dtype, g, m, k, n, trans_a, trans_b, scaled_gemm, scale_a_dtype,
+                     scale_b_dtype, trans_scale_a, trans_scale_b, arch, num_cu, num_chiplets,
+                     perf_config)
+        config.supports_split_k = resolve_split_k_support(supports_split_k, arch, out_dtype)
+        return config
 
     def to_command_line(self):
         result = (f"-t {self.datatype} -out_datatype {self.out_dtype} " +
@@ -1294,7 +1351,7 @@ def to_command_line(self):
             result += f" -transScaleA {str(self.trans_scale_a).lower()}"
         if self.trans_scale_b:
             result += f" -transScaleB {str(self.trans_scale_b).lower()}"
-        return result
+        return f"{result} {self.tuning_key_metadata()}"
 
     def __init__(self,
                  dtype: str,
@@ -1466,6 +1523,7 @@ def generate_mlir_driver_commandline(self, rocmlir_gen_flags, kernel_repeats=MLI
 
     @classmethod
     def from_command_line(cls, argv, arch, num_cu, num_chiplets):
+        argv, supports_split_k = extract_tuning_key_metadata(argv)
         # optional defaults
         perf_config = ''
         dtype = None
@@ -1543,9 +1601,11 @@ def from_command_line(cls, argv, arch, num_cu, num_chiplets):
             if v is None:
                 raise ValueError("Incomplete conv+gemm configuration")
 
-        return cls(dtype, filter_layout, input_layout, trans_c, trans_o, n, c, hi, wi, k, y, x, o,
-                   conv_stride_h, conv_stride_w, padding_h, padding_w, dilation_h, dilation_w,
-                   group, arch, num_cu, num_chiplets, perf_config)
+        config = cls(dtype, filter_layout, input_layout, trans_c, trans_o, n, c, hi, wi, k, y, x, o,
+                     conv_stride_h, conv_stride_w, padding_h, padding_w, dilation_h, dilation_w,
+                     group, arch, num_cu, num_chiplets, perf_config)
+        config.supports_split_k = resolve_split_k_support(supports_split_k, arch, dtype)
+        return config
 
     def to_command_line(self):
         return (f"-t {self.datatype} " +
@@ -1554,7 +1614,8 @@ def to_command_line(self):
                 f"-n {self.n} -c {self.c} -H {self.hi} -W {self.wi} -k {self.k} " +
                 f"-y {self.y} -x {self.x} -p {self.padding_h} -q {self.padding_w} " +
                 f"-u {self.conv_stride_h} -v {self.conv_stride_w} -l {self.dilation_h} " +
-                f"-j {self.dilation_w} -g {self.group} -gemmO {str(self.o)}")
+                f"-j {self.dilation_w} -g {self.group} -gemmO {str(self.o)} "
+                f"{self.tuning_key_metadata()}")
 
 
 class GemmGemmConfiguration(PerfConfiguration):
@@ -1640,6 +1701,7 @@ def generate_mlir_driver_commandline(self, rocmlir_gen_flags, kernel_repeats=MLI
 
     @classmethod
     def from_command_line(cls, argv, arch, num_cu, num_chiplets):
+        argv, supports_split_k = extract_tuning_key_metadata(argv)
         # optional defaults
         perf_config = ''
         dtype = None
@@ -1684,15 +1746,18 @@ def from_command_line(cls, argv, arch, num_cu, num_chiplets):
             if v is None:
                 raise ValueError("Incomplete gemm+gemm configuration")
 
-        return cls(dtype, g, m, k, n, o, trans_a, trans_b, trans_c, trans_o, arch, num_cu,
-                   num_chiplets, perf_config)
+        config = cls(dtype, g, m, k, n, o, trans_a, trans_b, trans_c, trans_o, arch, num_cu,
+                     num_chiplets, perf_config)
+        config.supports_split_k = resolve_split_k_support(supports_split_k, arch, dtype)
+        return config
 
     def to_command_line(self):
         return (f"-t {self.datatype} " +
                 f"-transA {str(self.trans_a).lower()} -transB {str(self.trans_b).lower()} " +
                 f"-transC {str(self.trans_c).lower()} -transO {str(self.trans_o).lower()} " +
                 f"-g {self.g} " +
-                f"-m {str(self.m)} -k {str(self.k)} -n {str(self.n)} -gemmO {str(self.o)}")
+                f"-m {str(self.m)} -k {str(self.k)} -n {str(self.n)} -gemmO {str(self.o)} "
+                f"{self.tuning_key_metadata()}")
 
 
 class AttentionConfiguration(PerfConfiguration):
@@ -1830,6 +1895,7 @@ def generate_mlir_driver_commandline(self, rocmlir_gen_flags, kernel_repeats=MLI
 
     @classmethod
     def from_command_line(cls, argv, arch, num_cu, num_chiplets):
+        argv, supports_split_k = extract_tuning_key_metadata(argv)
         # optional defaults
         perf_config = ''
         dtype = None
@@ -1908,30 +1974,33 @@ def from_command_line(cls, argv, arch, num_cu, num_chiplets):
             if v is None:
                 raise ValueError("Incomplete Attention configuration")
 
-        return cls(dtype,
-                   g,
-                   seq_len_q,
-                   seq_len_k,
-                   num_heads_q,
-                   num_heads_kv,
-                   head_dim_qk,
-                   head_dim_v,
-                   with_attn_scale,
-                   with_attn_bias,
-                   trans_q,
-                   trans_k,
-                   trans_v,
-                   trans_o,
-                   causal,
-                   return_lse,
-                   split_kv,
-                   arch,
-                   num_cu,
-                   num_chiplets,
-                   perf_config,
-                   trans_bias=trans_bias,
-                   current_seqlen=current_seqlen,
-                   sliding_window_size=sliding_window_size)
+        config = cls(dtype,
+                     g,
+                     seq_len_q,
+                     seq_len_k,
+                     num_heads_q,
+                     num_heads_kv,
+                     head_dim_qk,
+                     head_dim_v,
+                     with_attn_scale,
+                     with_attn_bias,
+                     trans_q,
+                     trans_k,
+                     trans_v,
+                     trans_o,
+                     causal,
+                     return_lse,
+                     split_kv,
+                     arch,
+                     num_cu,
+                     num_chiplets,
+                     perf_config,
+                     trans_bias=trans_bias,
+                     current_seqlen=current_seqlen,
+                     sliding_window_size=sliding_window_size)
+        output_dtype = 'f32' if dtype == 'i8' else dtype
+        config.supports_split_k = resolve_split_k_support(supports_split_k, arch, output_dtype)
+        return config
 
     def to_command_line(self):
         return (
@@ -1945,7 +2014,7 @@ def to_command_line(self):
             f"-seq_len_q {str(self.seq_len_q)} -seq_len_k {str(self.seq_len_k)} -num_heads_q {str(self.num_heads_q)} -num_heads_kv {str(self.num_heads_kv)} -head_dim_qk {str(self.head_dim_qk)} -head_dim_v {str(self.head_dim_v)} "
             + f"-with-attn-scale {str(self.with_attn_scale).lower()} " +
             f"-with-attn-bias {str(self.with_attn_bias).lower()} " +
-            f"-transBias {str(self.trans_bias).lower()}")
+            f"-transBias {str(self.trans_bias).lower()} {self.tuning_key_metadata()}")
 
 
 class HipBLASLtGemmConfig(GemmConfiguration):
@@ -2344,6 +2413,27 @@ def run_fusion_kernel(filename, rocmlir_gen_args, paths: Paths):
     return nanoseconds
 
 
+def lookup_fusion_tuning_config(tuning_db: TuningDb, arch: str, num_cu: int, num_chiplets: int,
+                                config: PerfConfiguration) -> Optional[str]:
+    """Find a fusion perf config, allowing a split-K-disabled key to use a base-op row.
+
+    benchmark_fusion_kernels forces every candidate's split-K factor to one before calling this
+    helper, so retrying a split-K-capable key cannot select an illegal split-K configuration.
+    """
+    config_str = config.to_command_line()
+    key = (arch, num_cu, num_chiplets, config_str)
+    if key in tuning_db:
+        return tuning_db[key]
+
+    if config.supports_split_k:
+        return None
+
+    false_suffix = "-supportsSplitK false"
+    assert config_str.endswith(false_suffix)
+    fallback_config_str = config_str[:-len(false_suffix)] + "-supportsSplitK true"
+    return tuning_db.get((arch, num_cu, num_chiplets, fallback_config_str))
+
+
 # Generate fusion vs. gemm/conv performance results
 def benchmark_fusion_kernels(test_dir,
                              paths: Paths,
@@ -2401,9 +2491,10 @@ def benchmark_fusion_kernels(test_dir,
         # Find the best perf_config
         best_perf = ""
         if tuning_db:
-            config_str = config.to_command_line()
-            if (arch, num_cu, num_chiplets, config_str) in tuning_db:
-                best_perf = tuning_db[arch, num_cu, num_chiplets, config_str]
+            tuned_config = lookup_fusion_tuning_config(tuning_db, arch, num_cu, num_chiplets,
+                                                       config)
+            if tuned_config is not None:
+                best_perf = tuned_config
                 config.set_perfconfig(best_perf)
             else:  # Tuning DB present but doesn't contain config, add a NaN entry
                 if test_vector not in perf_results:
@@ -2455,8 +2546,10 @@ def tune_mlir_kernels(configs, arch, num_cu, num_chiplets):
         if config.datatype not in ConvConfiguration.MIOPEN_SUPPORTED_DTYPES:
             print(f"Skipping MIOpen tuning for unsupported datatype: {config.datatype}")
             continue
+        # Tuning-key metadata is not a MIOpenDriver option and must not reach its argv.
+        config_args, _ = extract_tuning_key_metadata(commandline)
         if config.input_layout == 'nchw':
-            miopen_driver_cmd = [MIOPENDRIVER, *commandline, '-V', '0']
+            miopen_driver_cmd = [MIOPENDRIVER, *config_args, '-V', '0']
             print(' '.join(miopen_driver_cmd))
             p1 = subprocess.Popen(miopen_driver_cmd,
                                   stdout=subprocess.PIPE,
diff --git a/mlir/utils/performance/tests/mock_hip.py b/mlir/utils/performance/tests/mock_hip.py
index cd637869464a..283d4a9b356d 100644
--- a/mlir/utils/performance/tests/mock_hip.py
+++ b/mlir/utils/performance/tests/mock_hip.py
@@ -90,11 +90,20 @@ def __init__(self, **kwargs):
         self.has_lds_transpose_load = kwargs.get("has_lds_transpose_load", False)
 
 
-_DEFAULT_MOCK_INFO = _MockAmdArchInfo()
+_ATOMIC_ADD_F32_F16 = _MockGemmFeatures.ATOMIC_ADD | _MockGemmFeatures.ATOMIC_ADD_F16
+_ARCH_DEFAULT_FEATURES = {
+    "gfx900": _MockGemmFeatures.NONE,
+    "gfx908": _ATOMIC_ADD_F32_F16,
+    "gfx90a": _ATOMIC_ADD_F32_F16,
+    "gfx942": _ATOMIC_ADD_F32_F16,
+    "gfx950": _ATOMIC_ADD_F32_F16 | _MockGemmFeatures.ATOMIC_ADD_BF16,
+}
 
 
 def _mock_lookup_arch_info(arch):
-    return _DEFAULT_MOCK_INFO
+    chip = next((part for part in arch.split(":") if part.startswith("gfx")), arch)
+    return _MockAmdArchInfo(
+        default_features=_ARCH_DEFAULT_FEATURES.get(chip, _MockGemmFeatures.NONE))
 
 
 def _mock_has_feature(features, flag) -> bool:
diff --git a/mlir/utils/performance/tests/test_perfRunner.py b/mlir/utils/performance/tests/test_perfRunner.py
index 66898c6dbe23..53c83c3e5597 100644
--- a/mlir/utils/performance/tests/test_perfRunner.py
+++ b/mlir/utils/performance/tests/test_perfRunner.py
@@ -73,9 +73,9 @@ def test_read_with_header_and_comments(self):
                   "-g 1 -m 256 -n 128 -k 64")
         with tempfile.NamedTemporaryFile(mode="w", suffix=".tsv", delete=False) as f:
             f.write("# arch\tconfig\tperfconfig\n")
-            f.write(f"gfx900\t{gemm_a}\tperf_1\n")
+            f.write(f"gfx908\t{gemm_a}\tperf_1\n")
             f.write("\n")
-            f.write(f"gfx900\t{gemm_b}\tperf_2\n")
+            f.write(f"gfx908\t{gemm_b}\tperf_2\n")
             path = f.name
         try:
             db = perfRunner.read_tuning_db(path,
@@ -83,8 +83,38 @@ def test_read_with_header_and_comments(self):
                                            fallback_num_cu=120,
                                            fallback_num_chiplets=1)
             assert len(db) == 2
-            assert db[("gfx900", 120, 1, gemm_a)] == "perf_1"
-            assert db[("gfx900", 120, 1, gemm_b)] == "perf_2"
+            # Legacy rows describe unfused generated kernels, which support split-K.
+            assert db[("gfx908", 120, 1, f"{gemm_a} -supportsSplitK true")] == "perf_1"
+            assert db[("gfx908", 120, 1, f"{gemm_b} -supportsSplitK true")] == "perf_2"
+        finally:
+            os.unlink(path)
+
+    def test_read_distinguishes_split_k_support(self):
+        gemm = ("-t f32 -out_datatype f32 -transA false -transB false "
+                "-g 1 -m 1024 -n 512 -k 769")
+        with tempfile.NamedTemporaryFile(mode="w", suffix=".tsv", delete=False) as f:
+            f.write(f"gfx900\t120\t1\t{gemm} -supportsSplitK true\tperf_split_k\n")
+            f.write(f"gfx900\t120\t1\t{gemm} -supportsSplitK false\tperf_no_split_k\n")
+            path = f.name
+        try:
+            db = perfRunner.read_tuning_db(path, perfRunner.GemmConfiguration)
+
+            assert len(db) == 2
+            assert db[("gfx900", 120, 1, f"{gemm} -supportsSplitK true")] == "perf_split_k"
+            assert db[("gfx900", 120, 1, f"{gemm} -supportsSplitK false")] == "perf_no_split_k"
+        finally:
+            os.unlink(path)
+
+    def test_legacy_i8_gemm_does_not_gain_split_k_support(self):
+        gemm = ("-t i8 -out_datatype i32 -transA false -transB false "
+                "-g 1 -m 1024 -n 512 -k 769")
+        with tempfile.NamedTemporaryFile(mode="w", suffix=".tsv", delete=False) as f:
+            f.write(f"gfx908\t120\t1\t{gemm}\tperf_i8\n")
+            path = f.name
+        try:
+            db = perfRunner.read_tuning_db(path, perfRunner.GemmConfiguration)
+
+            assert db[("gfx908", 120, 1, f"{gemm} -supportsSplitK false")] == "perf_i8"
         finally:
             os.unlink(path)
 
@@ -102,10 +132,10 @@ def test_read_skips_unparseable_entries(self):
         # An .mlir path written by `tuningRunner --config foo.mlir`.
         mlir_path = "/path/to/fusion_kernel.mlir"
         with tempfile.NamedTemporaryFile(mode="w", suffix=".tsv", delete=False) as f:
-            f.write(f"gfx900\t{valid_gemm}\tperf_ok\n")
-            f.write(f"gfx900\t{conv_entry}\tperf_conv\n")
-            f.write(f"gfx900\t{malformed_gemm}\tperf_bad\n")
-            f.write(f"gfx900\t{mlir_path}\tperf_mlir\n")
+            f.write(f"gfx908\t{valid_gemm}\tperf_ok\n")
+            f.write(f"gfx908\t{conv_entry}\tperf_conv\n")
+            f.write(f"gfx908\t{malformed_gemm}\tperf_bad\n")
+            f.write(f"gfx908\t{mlir_path}\tperf_mlir\n")
             path = f.name
         try:
             db = perfRunner.read_tuning_db(path,
@@ -113,7 +143,7 @@ def test_read_skips_unparseable_entries(self):
                                            fallback_num_cu=120,
                                            fallback_num_chiplets=1)
             assert len(db) == 1
-            assert db[("gfx900", 120, 1, valid_gemm)] == "perf_ok"
+            assert db[("gfx908", 120, 1, f"{valid_gemm} -supportsSplitK true")] == "perf_ok"
         finally:
             os.unlink(path)
 
@@ -196,6 +226,55 @@ def test_pair_notation(self):
         assert out_map["fp8"] == "fp8"
 
 
+class TestSplitKSupport:
+    GEMM = ("-t {dtype} -out_datatype {out_dtype} -transA false -transB false "
+            "-g 1 -m 64 -n 64 -k 64")
+
+    def test_rejects_non_atomic_output_type(self):
+        key = self.GEMM.format(dtype="i8", out_dtype="i32")
+        config = perfRunner.GemmConfiguration.from_command_line(key.split(), "gfx908", 120, 1)
+
+        assert config.supports_split_k is False
+        assert config.to_command_line().endswith("-supportsSplitK false")
+
+    def test_mock_arch_database_matches_atomic_capabilities(self):
+        assert perfRunner.infer_split_k_support("gfx900", "f32") is False
+        assert perfRunner.infer_split_k_support("gfx908", "f16") is True
+        assert perfRunner.infer_split_k_support("gfx942", "bf16") is False
+        assert perfRunner.infer_split_k_support("gfx950:sramecc+:xnack-", "bf16") is True
+
+    def test_checks_arch_atomic_add_feature(self, monkeypatch):
+        monkeypatch.setattr(
+            perfRunner, "lookup_arch_info",
+            lambda arch: types.SimpleNamespace(default_features=perfRunner.GemmFeatures.ATOMIC_ADD))
+        key = self.GEMM.format(dtype="f16", out_dtype="f16")
+        config = perfRunner.GemmConfiguration.from_command_line(key.split(), "gfx1030", 2, 1)
+
+        assert config.supports_split_k is False
+
+    def test_explicit_metadata_is_authoritative(self, monkeypatch):
+        monkeypatch.setattr(
+            perfRunner, "lookup_arch_info",
+            lambda arch: types.SimpleNamespace(default_features=perfRunner.GemmFeatures.NONE))
+        key = self.GEMM.format(dtype="f16", out_dtype="f16")
+        config = perfRunner.GemmConfiguration.from_command_line(
+            f"{key} -supportsSplitK true".split(), "gfx1030", 2, 1)
+
+        assert config.supports_split_k is True
+
+    def test_fusion_lookup_falls_back_to_split_k_capable_base_key(self):
+        key = self.GEMM.format(dtype="f32", out_dtype="f32")
+        config = perfRunner.GemmConfiguration.from_command_line(
+            f"{key} -supportsSplitK false".split(), "gfx908", 120, 1)
+        capable_key = key + " -supportsSplitK true"
+        tuning_db = {("gfx908", 120, 1, capable_key): "v2:64,64,32,32,1,1,1"}
+
+        perf_config = perfRunner.lookup_fusion_tuning_config(tuning_db, "gfx908", 120, 1, config)
+
+        assert perf_config == "v2:64,64,32,32,1,1,1"
+        assert config.supports_split_k is False
+
+
 class TestLayoutHelpers:
     """Tests for input/output/filter layout conversion."""
 
diff --git a/mlir/utils/performance/tests/test_tuningRunner.py b/mlir/utils/performance/tests/test_tuningRunner.py
index 5ed9f03029da..ae2e2f3cf8e0 100644
--- a/mlir/utils/performance/tests/test_tuningRunner.py
+++ b/mlir/utils/performance/tests/test_tuningRunner.py
@@ -165,11 +165,15 @@ class TestTuningStateFile:
     """Tests for TuningStateFile (persisted state, no GPU)."""
 
     _CONF_CLASS = GemmConfiguration
-    _ARCH = "gfx900"
+    _ARCH = "gfx908"
     _NUM_CU = 64
     _NUM_CHIPLETS = 1
-    _TV_A = "-t f32 -out_datatype f32 -transA false -transB false -g 1 -m 1024 -n 512 -k 769"
-    _TV_B = "-t f16 -out_datatype f16 -transA false -transB true -g 1 -m 256 -n 128 -k 64"
+    # Canonical form, i.e. what canonicalize_test_vector() produces on load. It includes the
+    # tuning-key metadata that a config file would not spell out.
+    _TV_A = ("-t f32 -out_datatype f32 -transA false -transB false -g 1 -m 1024 -n 512 -k 769"
+             " -supportsSplitK true")
+    _TV_B = ("-t f16 -out_datatype f16 -transA false -transB true -g 1 -m 256 -n 128 -k 64"
+             " -supportsSplitK true")
 
     def _make_state_file(self, filepath, **kwargs):
         return TuningStateFile(filepath,
@@ -236,7 +240,7 @@ def test_old_state_file_configs_are_canonicalized(self):
 class TestTunedConfigsCache:
     """Tests for TunedConfigsCache.from_output_file (parsing only, no GPU)."""
 
-    def _options(self, output_path, arch="gfx900", num_cu=64, num_chiplets=1, tuning_space="full"):
+    def _options(self, output_path, arch="gfx908", num_cu=64, num_chiplets=1, tuning_space="full"):
         return Options(
             chip=arch,
             arch=arch,
@@ -274,18 +278,20 @@ def test_stdout_output_returns_empty(self):
 
     def test_parse_new_format_tsv(self):
         tv = "-t f32 -out_datatype f32 -transA false -transB false -g 1 -m 1024 -n 512 -k 769"
+        # Rows are canonicalized on load, which fills in the absent tuning-key metadata.
+        tv_canonical = f"{tv} -supportsSplitK true"
         with tempfile.NamedTemporaryFile(mode="w", suffix=".tsv", delete=False) as f:
             f.write(
                 "# arch\tnumCUs\tnumChiplets\ttestVector\tperfConfig\tTFlops\ttuningSpace\tcommitId\ttimestamp\tdurationSec\n"
             )
             f.write(
-                f"gfx900\t64\t1\t{tv}\tperf_best\t1.5\tfull\tabc123\t2025-01-01T00:00:00Z\t10.0\n")
+                f"gfx908\t64\t1\t{tv}\tperf_best\t1.5\tfull\tabc123\t2025-01-01T00:00:00Z\t10.0\n")
             path = f.name
         try:
             opts = self._options(path)
             cache = TunedConfigsCache.from_output_file(opts, GemmConfiguration)
             assert cache.count() == 1
-            r = cache.get(tv)
+            r = cache.get(tv_canonical)
             assert r is not None
             assert r.success
             assert r.winning_config == "perf_best"
@@ -379,80 +385,96 @@ def test_arch_mismatch_not_loaded(self):
 
 class TestCanonicalizeTestVector:
     """Tests for canonicalize_config and canonicalize_test_vector across all ops."""
+    _ARCH = "gfx908"
 
     @pytest.mark.parametrize("op", _ALL_OPS)
     def test_reorders_flags(self, op):
         tv = _SAMPLE_TEST_VECTORS[op]
         conf_class = tv["conf_class"]
-        canonical = canonicalize_config(tv["raw"], conf_class, "gfx900", 64, 1)
-        assert canonical == tv["canonical"]
+        canonical = canonicalize_config(tv["raw"], conf_class, self._ARCH, 64, 1)
+        assert canonical == f'{tv["canonical"]} -supportsSplitK true'
 
     @pytest.mark.parametrize("op", _ALL_OPS)
     def test_idempotent(self, op):
         tv = _SAMPLE_TEST_VECTORS[op]
         conf_class = tv["conf_class"]
         idempotent_form = tv["idempotent"]
-        result = canonicalize_config(idempotent_form, conf_class, "gfx900", 64, 1)
-        assert result == idempotent_form
+        result = canonicalize_config(idempotent_form, conf_class, self._ARCH, 64, 1)
+        assert result == f"{idempotent_form} -supportsSplitK true"
 
     @pytest.mark.parametrize("op", _ALL_OPS)
     def test_round_trip_preserves_data(self, op):
         """Canonicalize twice and verify the result is stable."""
         tv = _SAMPLE_TEST_VECTORS[op]
         conf_class = tv["conf_class"]
-        first = canonicalize_config(tv["raw"], conf_class, "gfx900", 64, 1)
-        second = canonicalize_config(first, conf_class, "gfx900", 64, 1)
+        first = canonicalize_config(tv["raw"], conf_class, self._ARCH, 64, 1)
+        second = canonicalize_config(first, conf_class, self._ARCH, 64, 1)
         assert first == second
 
+    @pytest.mark.parametrize("op", _ALL_OPS)
+    def test_split_k_support_metadata_round_trips(self, op):
+        tv = _SAMPLE_TEST_VECTORS[op]
+        conf_class = tv["conf_class"]
+        # Use the non-default value so this actually exercises parsing rather than the fallback.
+        raw = f'{tv["raw"]} -supportsSplitK false'
+
+        canonical = canonicalize_config(raw, conf_class, self._ARCH, 64, 1)
+        config = conf_class.from_command_line(canonical.split(), self._ARCH, 64, 1)
+
+        assert canonical.endswith("-supportsSplitK false")
+        assert config.supports_split_k is False
+        assert "-supportsSplitK" not in config.generate_mlir_driver_commandline("")
+
     def test_mlir_path_passthrough(self):
         path = "/some/test.mlir"
-        assert canonicalize_test_vector(path, GemmConfiguration, "gfx900", 64, 1) == path
+        assert canonicalize_test_vector(path, GemmConfiguration, self._ARCH, 64, 1) == path
 
     def test_invalid_config_raises_valueerror(self):
         with pytest.raises(ValueError, match="Failed to parse"):
-            canonicalize_config("not a valid config", GemmConfiguration, "gfx900", 64, 1)
+            canonicalize_config("not a valid config", GemmConfiguration, self._ARCH, 64, 1)
 
     def test_wrong_op_raises_valueerror(self):
         gemm_tv = "-t f32 -out_datatype f32 -transA false -transB false -g 1 -m 64 -n 128 -k 256"
         with pytest.raises(ValueError, match="Failed to parse"):
-            canonicalize_config(gemm_tv, ConvConfiguration, "gfx900", 64, 1)
+            canonicalize_config(gemm_tv, ConvConfiguration, self._ARCH, 64, 1)
 
     def test_fusion_dispatches_to_conv(self):
         """Fusion path (PerfConfiguration base class) routes 'conv*' prefix to ConvConfiguration."""
         raw = _SAMPLE_TEST_VECTORS["conv"]["raw"]
-        expected = canonicalize_config(raw, ConvConfiguration, "gfx900", 64, 1)
-        result = canonicalize_config(raw, PerfConfiguration, "gfx900", 64, 1)
+        expected = canonicalize_config(raw, ConvConfiguration, self._ARCH, 64, 1)
+        result = canonicalize_config(raw, PerfConfiguration, self._ARCH, 64, 1)
         assert result == expected
 
     def test_fusion_dispatches_to_gemm(self):
         """Fusion path (PerfConfiguration base class) routes non-'conv' prefix to GemmConfiguration."""
         raw = _SAMPLE_TEST_VECTORS["gemm"]["raw"]
-        expected = canonicalize_config(raw, GemmConfiguration, "gfx900", 64, 1)
-        result = canonicalize_config(raw, PerfConfiguration, "gfx900", 64, 1)
+        expected = canonicalize_config(raw, GemmConfiguration, self._ARCH, 64, 1)
+        result = canonicalize_config(raw, PerfConfiguration, self._ARCH, 64, 1)
         assert result == expected
 
     def test_fusion_invalid_raises_valueerror_with_resolved_class(self):
         """Errors from fusion dispatch should name the resolved concrete class, not the base."""
         with pytest.raises(ValueError, match="ConvConfiguration"):
-            canonicalize_config("convfp16 not a real config", PerfConfiguration, "gfx900", 64, 1)
+            canonicalize_config("convfp16 not a real config", PerfConfiguration, self._ARCH, 64, 1)
         with pytest.raises(ValueError, match="GemmConfiguration"):
-            canonicalize_config("not a real config", PerfConfiguration, "gfx900", 64, 1)
+            canonicalize_config("not a real config", PerfConfiguration, self._ARCH, 64, 1)
 
     def test_cache_loaded_with_canonical_key(self):
         """Verify that from_output_file canonicalizes test vectors so cache lookups match."""
         raw = "-g 1 -m 1024 -k 769 -n 512 -t f32 -out_datatype f32 -transA false -transB false"
-        canonical = canonicalize_config(raw, GemmConfiguration, "gfx900", 64, 1)
+        canonical = canonicalize_config(raw, GemmConfiguration, self._ARCH, 64, 1)
         with tempfile.NamedTemporaryFile(mode="w", suffix=".tsv", delete=False) as f:
             f.write(
                 "# arch\tnumCUs\tnumChiplets\ttestVector\tperfConfig\tTFlops\ttuningSpace\tcommitId\ttimestamp\tdurationSec\n"
             )
             f.write(
-                f"gfx900\t64\t1\t{raw}\tperf_best\t1.5\tfull\tabc123\t2025-01-01T00:00:00Z\t10.0\n")
+                f"{self._ARCH}\t64\t1\t{raw}\tperf_best\t1.5\tfull\tabc123\t2025-01-01T00:00:00Z\t10.0\n"
+            )
             path = f.name
         try:
             opts = Options(
-                chip="gfx900",
-                arch="gfx900",
+                chip=self._ARCH,
+                arch=self._ARCH,
                 num_cu=64,
                 num_chiplets=1,
                 debug=False,