[GPUHeuristics] Refactor MMA heuristic seeds to be architecture-specific (#23717)
Each GPU architecture now defines a constexpr ArchSeedSet containing
gemm, scaled gemm, and convolution seed arrays indexed by GemmSize. The
architecture is determined from the target and the corresponding seed
set is returned (e.g. RDNA4 uses tuned seeds from RX 9070 XT
benchmarking, others use the default). GemmSize enum values are changed
to start at 0 for direct array indexing.
Add RDNA4 (gfx1201) config tests covering small, medium, and large
matmul and convolution shapes.
RDNA4 benchmark results (RX 9070 XT, vs default seeds):
- GEMM shapes: ~40% of shapes improved, 15.9% geometric mean speedup.
- Prod Conv shapes: ~25% of shapes improved, 6% geometric mean speedup.
- Proxy Conv shapes: ~30% of shapes improved, 14.5% geometric mean
speedup.
---------
Signed-off-by: yzhang93 <zhyuhang88@gmail.com>
Co-authored-by: Claude <noreply@anthropic.com>
diff --git a/compiler/src/iree/compiler/Codegen/Common/GPU/GPUHeuristics.cpp b/compiler/src/iree/compiler/Codegen/Common/GPU/GPUHeuristics.cpp
index 01a5a0c..f5422af 100644
--- a/compiler/src/iree/compiler/Codegen/Common/GPU/GPUHeuristics.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/GPU/GPUHeuristics.cpp
@@ -50,15 +50,16 @@
return os;
}
-llvm::raw_ostream &operator<<(llvm::raw_ostream &os, const GemmSize &gemmSize) {
+llvm::raw_ostream &operator<<(llvm::raw_ostream &os,
+ const GemmSizeKind &gemmSize) {
switch (gemmSize) {
- case GemmSize::SmallGemm:
+ case GemmSizeKind::SmallGemm:
return os << "SmallGemm";
- case GemmSize::MediumGemm:
+ case GemmSizeKind::MediumGemm:
return os << "MediumGemm";
- case GemmSize::LargeGemm:
+ case GemmSizeKind::LargeGemm:
return os << "LargeGemm";
- case GemmSize::VeryLargeGemm:
+ case GemmSizeKind::VeryLargeGemm:
return os << "VeryLargeGemm";
default:
assert(false && "Unhandled gemm size");
@@ -648,7 +649,7 @@
// (compute=8192, area=512) because throughput matters more. Among
// 16x16x32 and 32x32x16 (both area=1024), prefer smaller K (16 vs 32)
// for less operand staging pressure.
- if (problem.gemmSize == GemmSize::VeryLargeGemm) {
+ if (problem.gemmSize == GemmSizeKind::VeryLargeGemm) {
int64_t lhsCompute = intrinsicCompute(lhs);
int64_t rhsCompute = intrinsicCompute(rhs);
if (lhsCompute != rhsCompute) {
@@ -700,8 +701,7 @@
return bestMNTileCountPerSubgroup;
}
- if (problem.gemmSize == GemmSize::NotSet ||
- problem.gemmSize == GemmSize::SmallGemm) {
+ if (!problem.gemmSize || problem.gemmSize == GemmSizeKind::SmallGemm) {
LDBG() << "Arithmetic intensity is too low, "
<< "skipping adjustment of seeds for workgroup count.";
return bestMNTileCountPerSubgroup;
diff --git a/compiler/src/iree/compiler/Codegen/Common/GPU/GPUHeuristics.h b/compiler/src/iree/compiler/Codegen/Common/GPU/GPUHeuristics.h
index 78d7003..c4fe42a 100644
--- a/compiler/src/iree/compiler/Codegen/Common/GPU/GPUHeuristics.h
+++ b/compiler/src/iree/compiler/Codegen/Common/GPU/GPUHeuristics.h
@@ -4,15 +4,26 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+#ifndef IREE_COMPILER_CODEGEN_COMMON_GPU_GPUHEURISTICS_H_
+#define IREE_COMPILER_CODEGEN_COMMON_GPU_GPUHEURISTICS_H_
+
#include <cstdint>
+#include <optional>
#include "iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUInterfaces.h"
#include "mlir/IR/Types.h"
namespace mlir::iree_compiler {
-enum class GemmSize { NotSet, SmallGemm, MediumGemm, LargeGemm, VeryLargeGemm };
+enum class GemmSizeKind : int {
+ SmallGemm,
+ MediumGemm,
+ LargeGemm,
+ VeryLargeGemm,
+ Count, // Must be last — used for static array sizes.
+};
-llvm::raw_ostream &operator<<(llvm::raw_ostream &os, const GemmSize &gemmSize);
+llvm::raw_ostream &operator<<(llvm::raw_ostream &os,
+ const GemmSizeKind &gemmSize);
/// Struct containing information about a matmul's shape and type.
struct GPUMatmulShapeType {
@@ -27,7 +38,7 @@
Type aScaleType;
Type bScaleType;
- GemmSize gemmSize = GemmSize::NotSet;
+ std::optional<GemmSizeKind> gemmSize;
// Number of horizontally fused operations.
// Horizontal fusion: C1,C2 = fused_matmul(A, B1, B2) where A is shared.
@@ -175,3 +186,5 @@
const GPUMMASchedule &schedule);
} // namespace mlir::iree_compiler
+
+#endif // IREE_COMPILER_CODEGEN_COMMON_GPU_GPUHEURISTICS_H_
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/BUILD.bazel b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/BUILD.bazel
index d7adda1..b0207a8 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/BUILD.bazel
+++ b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/BUILD.bazel
@@ -27,6 +27,7 @@
"//compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR:IREECodegenDialect",
"//compiler/src/iree/compiler/Codegen/Dialect/GPU/IR:IREEGPUDialect",
"//compiler/src/iree/compiler/Codegen/Dialect/GPU/IR:IREEGPUInterfaces",
+ "//compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils:KnownTargets",
"//compiler/src/iree/compiler/Codegen/Interfaces:PartitionableLoopsInterface",
"//compiler/src/iree/compiler/Codegen/Utils",
"//compiler/src/iree/compiler/Dialect/LinalgExt/IR",
@@ -52,8 +53,10 @@
"KnownTargets.h",
],
deps = [
+ "//compiler/src/iree/compiler/Codegen/Common/GPU:GPUHeuristics",
"//compiler/src/iree/compiler/Codegen/Dialect/GPU/IR:IREEGPUDialect",
"@llvm-project//llvm:Support",
+ "@llvm-project//mlir:AMDGPUUtils",
"@llvm-project//mlir:IR",
"@llvm-project//mlir:Support",
],
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/CMakeLists.txt b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/CMakeLists.txt
index 2555441..76eb0e5 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/CMakeLists.txt
+++ b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/CMakeLists.txt
@@ -31,6 +31,7 @@
iree::compiler::Codegen::Dialect::Codegen::IR::IREECodegenDialect
iree::compiler::Codegen::Dialect::GPU::IR::IREEGPUDialect
iree::compiler::Codegen::Dialect::GPU::IR::IREEGPUInterfaces
+ iree::compiler::Codegen::Dialect::GPU::TargetUtils::KnownTargets
iree::compiler::Codegen::Interfaces::PartitionableLoopsInterface
iree::compiler::Codegen::Utils
iree::compiler::Dialect::LinalgExt::IR
@@ -48,8 +49,10 @@
"KnownTargets.cpp"
DEPS
LLVMSupport
+ MLIRAMDGPUUtils
MLIRIR
MLIRSupport
+ iree::compiler::Codegen::Common::GPU::GPUHeuristics
iree::compiler::Codegen::Dialect::GPU::IR::IREEGPUDialect
PUBLIC
)
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/ConfigUtils.cpp b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/ConfigUtils.cpp
index 25646f4..990233d 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/ConfigUtils.cpp
+++ b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/ConfigUtils.cpp
@@ -14,6 +14,7 @@
#include "iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUAttrs.h"
#include "iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUEnums.h"
#include "iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUInterfaces.h"
+#include "iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.h"
#include "iree/compiler/Codegen/Interfaces/PartitionableLoopsInterface.h"
#include "iree/compiler/Codegen/Utils/GPUUtils.h"
#include "iree/compiler/Codegen/Utils/Utils.h"
@@ -258,89 +259,24 @@
return {smallGemmCutoff, largeGemmCutoff, veryLargeGemmCutoff};
}
+/// Returns architecture-specific GPUMMAHeuristicSeeds for the given target,
+/// problem size category (GemmSizeKind), and operation type (gemm, scaled
+/// gemm, or convolution).
static std::optional<GPUMMAHeuristicSeeds>
-getGemmHeuristicSeeds(GemmSize gemmSize, int64_t inBitWidth, bool scaled) {
- switch (gemmSize) {
- case GemmSize::SmallGemm:
- return GPUMMAHeuristicSeeds(
- {/*bestSubgroupCountPerWorkgroup=*/2,
- /*bestMNTileCountPerSubgroup=*/2,
- /*bestKTileCountPerSubgroup=*/4,
- /*bestKElementCountPerSubgroup=*/2 * kCacheLineSizeBits / inBitWidth});
- case GemmSize::MediumGemm:
- if (scaled) {
- return GPUMMAHeuristicSeeds(
- {/*bestSubgroupCountPerWorkgroup=*/8,
- /*bestMNTileCountPerSubgroup=*/32,
- /*bestKTileCountPerSubgroup=*/4,
- /*bestKElementCountPerSubgroup=*/kCacheLineSizeBits / 2 /
- inBitWidth});
- }
- return GPUMMAHeuristicSeeds(
- {/*bestSubgroupCountPerWorkgroup=*/4,
- /*bestMNTileCountPerSubgroup=*/8,
- /*bestKTileCountPerSubgroup=*/4,
- /*bestKElementCountPerSubgroup=*/2 * kCacheLineSizeBits / inBitWidth});
- case GemmSize::LargeGemm:
- case GemmSize::VeryLargeGemm:
- if (scaled) {
- return GPUMMAHeuristicSeeds(
- {/*bestSubgroupCountPerWorkgroup=*/8,
- /*bestMNTileCountPerSubgroup=*/32,
- /*bestKTileCountPerSubgroup=*/2,
- /*bestKElementCountPerSubgroup=*/kCacheLineSizeBits / 2 /
- inBitWidth});
- }
- return GPUMMAHeuristicSeeds(
- {/*bestSubgroupCountPerWorkgroup=*/4,
- /*bestMNTileCountPerSubgroup=*/16,
- /*bestKTileCountPerSubgroup=*/2,
- /*bestKElementCountPerSubgroup=*/kCacheLineSizeBits / 2 / inBitWidth});
- default:
- assert(false && "Unhandled gemm size");
- return std::nullopt;
- }
-}
-
-static std::optional<GPUMMAHeuristicSeeds>
-getConvolutionHeuristicSeeds(GemmSize gemmSize, int64_t inBitWidth) {
- switch (gemmSize) {
- case GemmSize::SmallGemm:
- return GPUMMAHeuristicSeeds(
- {/*bestSubgroupCountPerWorkgroup=*/2,
- /*bestMNTileCountPerSubgroup=*/2,
- /*bestKTileCountPerSubgroup=*/4,
- /*bestKElementCountPerSubgroup=*/kCacheLineSizeBits / inBitWidth});
- case GemmSize::MediumGemm:
- return GPUMMAHeuristicSeeds(
- {/*bestSubgroupCountPerWorkgroup=*/8,
- /*bestMNTileCountPerSubgroup=*/4,
- /*bestKTileCountPerSubgroup=*/4,
- /*bestKElementCountPerSubgroup=*/2 * kCacheLineSizeBits / inBitWidth});
- case GemmSize::LargeGemm:
- case GemmSize::VeryLargeGemm:
- // Favor more subgroups for convolution to help latency hiding from global
- // loads.
- return GPUMMAHeuristicSeeds(
- {/*bestSubgroupCountPerWorkgroup=*/8,
- /*bestMNTileCountPerSubgroup=*/8,
- /*bestKTileCountPerSubgroup=*/2,
- /*bestKElementCountPerSubgroup=*/kCacheLineSizeBits / 2 / inBitWidth});
- default:
- assert(false && "Unhandled convolution gemm size");
- return std::nullopt;
- }
-}
-
-static std::optional<GPUMMAHeuristicSeeds>
-getContractionHeuristicSeeds(GPUMatmulShapeType problem, bool isGemm,
+getContractionHeuristicSeeds(IREE::GPU::TargetAttr target,
+ GPUMatmulShapeType problem, bool isGemm,
bool scaled) {
- GemmSize gemmSize = problem.gemmSize;
- int64_t inBitWidth = problem.aType.getIntOrFloatBitWidth();
- if (isGemm) {
- return getGemmHeuristicSeeds(gemmSize, inBitWidth, scaled);
- }
- return getConvolutionHeuristicSeeds(gemmSize, inBitWidth);
+ const ArchSeedSet &archSeeds = getArchSeedSet(target);
+ assert(problem.gemmSize.has_value() && "GemmSizeKind must be set");
+
+ // Pick the right category, index by GemmSizeKind, then convert the stored
+ // K-element bits to an actual element count.
+ const GPUMMAHeuristicSeeds *table =
+ !isGemm ? archSeeds.conv
+ : (scaled ? archSeeds.scaledGemm : archSeeds.gemm);
+ GPUMMAHeuristicSeeds result = table[static_cast<int>(*problem.gemmSize)];
+ result.bestKElementCountPerSubgroup /= problem.aType.getIntOrFloatBitWidth();
+ return result;
}
/// Given a target and a matmul problem, try to find an MMA schedule for the
@@ -421,25 +357,25 @@
if (computeIntensity <= gemmCutoffs.smallGemmCutoff) {
// For matmuls with small arithmetic intensity, use small
// bestMNTileCountPerSubgroup and large bestKTileCountPerSubgroup.
- problem.gemmSize = GemmSize::SmallGemm;
+ problem.gemmSize = GemmSizeKind::SmallGemm;
} else if (computeIntensity >= gemmCutoffs.veryLargeGemmCutoff) {
// For very large matmuls, prefer low-VGPR-pressure intrinsics (e.g.,
// 32x32x16 over 16x16x32) which provide higher compute throughput per
// register.
- problem.gemmSize = GemmSize::VeryLargeGemm;
+ problem.gemmSize = GemmSizeKind::VeryLargeGemm;
} else if (computeIntensity >= gemmCutoffs.largeGemmCutoff) {
// For matmuls with large arithmetic intensity, use large
// bestMNTileCountPerSubgroup and small bestKTileCountPerSubgroup to
// amortize launch/memory costs and maximize throughput.
- problem.gemmSize = GemmSize::LargeGemm;
+ problem.gemmSize = GemmSizeKind::LargeGemm;
} else {
// Choose balanced tile shapes. Empirically, medium-AI workloads can favor
// either small or large tiles depending on kernel details.
- problem.gemmSize = GemmSize::MediumGemm;
+ problem.gemmSize = GemmSizeKind::MediumGemm;
}
- LDBG() << "This config is " << problem.gemmSize;
+ LDBG() << "This config is " << *problem.gemmSize;
std::optional<GPUMMAHeuristicSeeds> maybeSeeds =
- getContractionHeuristicSeeds(problem, isGemm, scaled);
+ getContractionHeuristicSeeds(target, problem, isGemm, scaled);
assert(maybeSeeds.has_value() && "expected seeds to be found");
GPUMMAHeuristicSeeds seeds = maybeSeeds.value();
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.cpp b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.cpp
index bd8db2b..025e952 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.cpp
+++ b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.cpp
@@ -12,6 +12,7 @@
#include "llvm/ADT/DenseMap.h"
#include "llvm/ADT/STLExtras.h"
#include "llvm/ADT/StringSwitch.h"
+#include "mlir/Dialect/AMDGPU/Utils/Chipset.h"
#include "mlir/IR/Attributes.h"
#include "mlir/IR/BuiltinAttributes.h"
#include "mlir/IR/BuiltinTypes.h"
@@ -1253,4 +1254,76 @@
.Default(nullptr);
}
+//===----------------------------------------------------------------------===//
+// Architecture-specific heuristic seed tables
+//===----------------------------------------------------------------------===//
+
+constexpr int64_t kCacheLineSizeBits = 128 * 8;
+
+// clang-format off
+
+/// Default seeds (CDNA and other architectures).
+static constexpr ArchSeedSet kDefaultSeeds = {
+ /*gemm=*/{
+ /*SmallGemm=*/ {2, 2, 4, 2 * kCacheLineSizeBits},
+ /*MediumGemm=*/ {4, 8, 4, 2 * kCacheLineSizeBits},
+ /*LargeGemm=*/ {4, 16, 2, kCacheLineSizeBits / 2},
+ /*VeryLargeGemm=*/ {4, 16, 2, kCacheLineSizeBits / 2},
+ },
+ /*scaledGemm=*/{
+ /*SmallGemm=*/ {2, 2, 4, 2 * kCacheLineSizeBits},
+ /*MediumGemm=*/ {8, 32, 4, kCacheLineSizeBits / 2},
+ /*LargeGemm=*/ {8, 32, 2, kCacheLineSizeBits / 2},
+ /*VeryLargeGemm=*/ {8, 32, 2, kCacheLineSizeBits / 2},
+ },
+ /*conv=*/{
+ /*SmallGemm=*/ {2, 2, 4, kCacheLineSizeBits},
+ /*MediumGemm=*/ {8, 4, 4, 2 * kCacheLineSizeBits},
+ /*LargeGemm=*/ {8, 8, 2, kCacheLineSizeBits / 2},
+ /*VeryLargeGemm=*/ {8, 8, 2, kCacheLineSizeBits / 2},
+ },
+};
+
+/// RDNA4 seeds (tuned based on RX 9070 XT benchmarking data).
+static constexpr ArchSeedSet kRDNA4Seeds = {
+ /*gemm=*/{
+ /*SmallGemm=*/ {2, 2, 4, 2 * kCacheLineSizeBits},
+ /*MediumGemm=*/ {4, 4, 4, kCacheLineSizeBits},
+ /*LargeGemm=*/ {8, 16, 4, kCacheLineSizeBits},
+ /*VeryLargeGemm=*/ {8, 16, 4, kCacheLineSizeBits},
+ },
+ /*scaledGemm=*/{
+ /*SmallGemm=*/ {2, 2, 4, 2 * kCacheLineSizeBits},
+ /*MediumGemm=*/ {8, 32, 4, kCacheLineSizeBits / 2},
+ /*LargeGemm=*/ {8, 32, 2, kCacheLineSizeBits / 2},
+ /*VeryLargeGemm=*/ {8, 32, 2, kCacheLineSizeBits / 2},
+ },
+ /*conv=*/{
+ /*SmallGemm=*/ {2, 2, 4, kCacheLineSizeBits},
+ /*MediumGemm=*/ {4, 4, 4, kCacheLineSizeBits},
+ /*LargeGemm=*/ {4, 8, 4, kCacheLineSizeBits},
+ /*VeryLargeGemm=*/ {4, 8, 4, kCacheLineSizeBits},
+ },
+};
+
+// clang-format on
+
+/// Look up the seed set for the given target architecture.
+const ArchSeedSet &getArchSeedSet(TargetAttr target) {
+ if (!target) {
+ return kDefaultSeeds;
+ }
+
+ StringRef arch = target.getArch();
+ // RDNA4 is gfx1200/gfx1201 (major=12, minor=0). Note: gfx1250 (minor=5)
+ // is a separate experimental target and should not use RDNA4 seeds.
+ FailureOr<amdgpu::Chipset> chipset = amdgpu::Chipset::parse(arch);
+ bool isRDNA4 = succeeded(chipset) && chipset->majorVersion == 12 &&
+ chipset->minorVersion == 0;
+ if (isRDNA4 || arch == "rdna4") {
+ return kRDNA4Seeds;
+ }
+ return kDefaultSeeds;
+}
+
} // namespace mlir::iree_compiler::IREE::GPU
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.h b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.h
index ccb5531..4168b00 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.h
+++ b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.h
@@ -7,11 +7,28 @@
#ifndef IREE_COMPILER_CODEGEN_DIALECT_GPU_TARGETUTILS_KNOWNTARGETS_H_
#define IREE_COMPILER_CODEGEN_DIALECT_GPU_TARGETUTILS_KNOWNTARGETS_H_
+#include "iree/compiler/Codegen/Common/GPU/GPUHeuristics.h"
#include "iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUAttrs.h"
#include "llvm/ADT/StringRef.h"
namespace mlir::iree_compiler::IREE::GPU {
+//===----------------------------------------------------------------------===//
+// Architecture-specific heuristic seed tables
+//===----------------------------------------------------------------------===//
+
+/// Complete seed set for a GPU architecture, with separate tables for
+/// gemm, scaled gemm, and convolution — each indexed by GemmSizeKind.
+struct ArchSeedSet {
+ static constexpr int numKinds = static_cast<int>(GemmSizeKind::Count);
+ GPUMMAHeuristicSeeds gemm[numKinds] = {};
+ GPUMMAHeuristicSeeds scaledGemm[numKinds] = {};
+ GPUMMAHeuristicSeeds conv[numKinds] = {};
+};
+
+/// Look up the heuristic seed set for the given target architecture.
+const ArchSeedSet &getArchSeedSet(TargetAttr target);
+
constexpr char kNoEncodingLayoutResolverName[] = "none";
constexpr char kPadEncodingLayoutResolverName[] = "pad";
constexpr char kDataTilingEncodingLayoutResolverName[] = "data-tiling";
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/BUILD.bazel b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/BUILD.bazel
index 389e518..9bacaa2 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/BUILD.bazel
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/BUILD.bazel
@@ -24,6 +24,7 @@
"config_direct_conv_tile_and_fuse.mlir",
"config_igemm_tile_and_fuse.mlir",
"config_tile_and_fuse.mlir",
+ "config_tile_and_fuse_gfx1201.mlir",
"config_tile_and_fuse_gfx950.mlir",
"config_user_vector_distribute.mlir",
"config_vector_distribute_gfx1100.mlir",
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/CMakeLists.txt b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/CMakeLists.txt
index c78c28f..6a6ca51 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/CMakeLists.txt
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/CMakeLists.txt
@@ -19,6 +19,7 @@
"config_direct_conv_tile_and_fuse.mlir"
"config_igemm_tile_and_fuse.mlir"
"config_tile_and_fuse.mlir"
+ "config_tile_and_fuse_gfx1201.mlir"
"config_tile_and_fuse_gfx950.mlir"
"config_user_vector_distribute.mlir"
"config_vector_distribute_gfx1100.mlir"
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/config_tile_and_fuse_gfx1201.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/config_tile_and_fuse_gfx1201.mlir
new file mode 100644
index 0000000..020c81d
--- /dev/null
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/config_tile_and_fuse_gfx1201.mlir
@@ -0,0 +1,158 @@
+// RUN: iree-opt --mlir-print-local-scope --split-input-file --iree-gpu-test-target=gfx1201 \
+// RUN: --pass-pipeline="builtin.module(iree-llvmgpu-select-lowering-strategy)" %s \
+// RUN: | FileCheck %s --check-prefix=GEMM
+// RUN: iree-opt --mlir-print-local-scope --split-input-file --iree-gpu-test-target=gfx1201 \
+// RUN: --iree-codegen-llvmgpu-use-igemm=false --iree-codegen-llvmgpu-use-direct-convolution=true \
+// RUN: --pass-pipeline="builtin.module(iree-llvmgpu-select-lowering-strategy)" %s \
+// RUN: | FileCheck %s --check-prefix=CONV
+
+// Verify RDNA4-specific heuristic seed selection produces expected configs for
+// matmul and convolution operations at different arithmetic intensity levels.
+
+// ============================================================================
+// Matmul — small (low arithmetic intensity, memory-bound)
+// ============================================================================
+
+func.func @matmul_small_f16(%arg0: tensor<64x1280xf16>, %arg1: tensor<1280x1280xf16>) -> tensor<64x1280xf32> {
+ %cst = arith.constant 0.000000e+00 : f32
+ %empty = tensor.empty() : tensor<64x1280xf32>
+ %fill = linalg.fill ins(%cst : f32) outs(%empty : tensor<64x1280xf32>) -> tensor<64x1280xf32>
+ %result = linalg.matmul ins(%arg0, %arg1 : tensor<64x1280xf16>, tensor<1280x1280xf16>) outs(%fill : tensor<64x1280xf32>) -> tensor<64x1280xf32>
+ return %result : tensor<64x1280xf32>
+}
+
+// GEMM-LABEL: func.func @matmul_small_f16
+// GEMM-SAME: #iree_codegen.translation_info<pipeline = LLVMGPUTileAndFuse
+// GEMM-SAME: workgroup_size = [128, 1, 1] subgroup_size = 32
+// GEMM: linalg.matmul {{.*}}lowering_config = #iree_gpu.lowering_config
+// GEMM-SAME: mma_kind = #iree_gpu.mma_layout<WMMAR4_F32_16x16x16_F16>
+// GEMM-SAME: promote_operands = [0, 1]
+// GEMM-SAME: reduction = [0, 0, 4]
+// GEMM-SAME: subgroup = [2, 2, 0]
+// GEMM-SAME: workgroup = [64, 64, 0]
+
+// -----
+
+// ============================================================================
+// Matmul — medium (moderate arithmetic intensity)
+// ============================================================================
+
+func.func @matmul_medium_f16(%arg0: tensor<2048x1280xf16>, %arg1: tensor<1280x1280xf16>) -> tensor<2048x1280xf32> {
+ %cst = arith.constant 0.000000e+00 : f32
+ %empty = tensor.empty() : tensor<2048x1280xf32>
+ %fill = linalg.fill ins(%cst : f32) outs(%empty : tensor<2048x1280xf32>) -> tensor<2048x1280xf32>
+ %result = linalg.matmul ins(%arg0, %arg1 : tensor<2048x1280xf16>, tensor<1280x1280xf16>) outs(%fill : tensor<2048x1280xf32>) -> tensor<2048x1280xf32>
+ return %result : tensor<2048x1280xf32>
+}
+
+// GEMM-LABEL: func.func @matmul_medium_f16
+// GEMM-SAME: #iree_codegen.translation_info<pipeline = LLVMGPUTileAndFuse
+// GEMM-SAME: workgroup_size = [128, 1, 1] subgroup_size = 32
+// GEMM: linalg.matmul {{.*}}lowering_config = #iree_gpu.lowering_config
+// GEMM-SAME: mma_kind = #iree_gpu.mma_layout<WMMAR4_F32_16x16x16_F16>
+// GEMM-SAME: promote_operands = [0, 1]
+// GEMM-SAME: reduction = [0, 0, 4]
+// GEMM-SAME: subgroup = [2, 2, 0]
+// GEMM-SAME: workgroup = [64, 64, 0]
+
+// -----
+
+// ============================================================================
+// Matmul — large (high arithmetic intensity, compute-bound)
+// ============================================================================
+
+func.func @matmul_large_f16(%arg0: tensor<4096x4096xf16>, %arg1: tensor<4096x4096xf16>) -> tensor<4096x4096xf32> {
+ %cst = arith.constant 0.000000e+00 : f32
+ %empty = tensor.empty() : tensor<4096x4096xf32>
+ %fill = linalg.fill ins(%cst : f32) outs(%empty : tensor<4096x4096xf32>) -> tensor<4096x4096xf32>
+ %result = linalg.matmul ins(%arg0, %arg1 : tensor<4096x4096xf16>, tensor<4096x4096xf16>) outs(%fill : tensor<4096x4096xf32>) -> tensor<4096x4096xf32>
+ return %result : tensor<4096x4096xf32>
+}
+
+// GEMM-LABEL: func.func @matmul_large_f16
+// GEMM-SAME: #iree_codegen.translation_info<pipeline = LLVMGPUTileAndFuse
+// GEMM-SAME: workgroup_size = [256, 1, 1] subgroup_size = 32
+// GEMM: linalg.matmul {{.*}}lowering_config = #iree_gpu.lowering_config
+// GEMM-SAME: mma_kind = #iree_gpu.mma_layout<WMMAR4_F32_16x16x16_F16>
+// GEMM-SAME: promote_operands = [0, 1]
+// GEMM-SAME: reduction = [0, 0, 4]
+// GEMM-SAME: subgroup = [4, 4, 0]
+// GEMM-SAME: workgroup = [256, 128, 0]
+
+// -----
+
+// ============================================================================
+// Convolution — small (low arithmetic intensity)
+// ============================================================================
+
+func.func @conv_small_f16(%arg0: tensor<1x18x18x16xf16>, %arg1: tensor<32x3x3x16xf16>) -> tensor<1x16x16x32xf32> {
+ %cst = arith.constant 0.000000e+00 : f32
+ %empty = tensor.empty() : tensor<1x16x16x32xf32>
+ %fill = linalg.fill ins(%cst : f32) outs(%empty : tensor<1x16x16x32xf32>) -> tensor<1x16x16x32xf32>
+ %0 = linalg.conv_2d_nhwc_fhwc {dilations = dense<1> : tensor<2xi64>, strides = dense<1> : tensor<2xi64>}
+ ins(%arg0, %arg1 : tensor<1x18x18x16xf16>, tensor<32x3x3x16xf16>)
+ outs(%fill : tensor<1x16x16x32xf32>) -> tensor<1x16x16x32xf32>
+ return %0 : tensor<1x16x16x32xf32>
+}
+
+// CONV-LABEL: func.func @conv_small_f16
+// CONV-SAME: #iree_codegen.translation_info<pipeline = LLVMGPUTileAndFuse
+// CONV-SAME: workgroup_size = [128, 1, 1] subgroup_size = 32
+// CONV: linalg.conv_2d_nhwc_fhwc {{.*}}lowering_config = #iree_gpu.lowering_config
+// CONV-SAME: mma_kind = #iree_gpu.mma_layout<WMMAR4_F32_16x16x16_F16>
+// CONV-SAME: promote_operands = [0, 1]
+// CONV-SAME: reduction = [0, 0, 0, 0, 1, 1, 1]
+// CONV-SAME: subgroup = [1, 2, 1, 1, 0, 0, 0]
+// CONV-SAME: workgroup = [1, 4, 16, 32, 0, 0, 0]
+
+// -----
+
+// ============================================================================
+// Convolution — medium (moderate arithmetic intensity)
+// ============================================================================
+
+func.func @conv_medium_f16(%arg0: tensor<2x66x66x128xf16>, %arg1: tensor<128x3x3x128xf16>) -> tensor<2x64x64x128xf32> {
+ %cst = arith.constant 0.000000e+00 : f32
+ %empty = tensor.empty() : tensor<2x64x64x128xf32>
+ %fill = linalg.fill ins(%cst : f32) outs(%empty : tensor<2x64x64x128xf32>) -> tensor<2x64x64x128xf32>
+ %0 = linalg.conv_2d_nhwc_fhwc {dilations = dense<1> : tensor<2xi64>, strides = dense<1> : tensor<2xi64>}
+ ins(%arg0, %arg1 : tensor<2x66x66x128xf16>, tensor<128x3x3x128xf16>)
+ outs(%fill : tensor<2x64x64x128xf32>) -> tensor<2x64x64x128xf32>
+ return %0 : tensor<2x64x64x128xf32>
+}
+
+// CONV-LABEL: func.func @conv_medium_f16
+// CONV-SAME: #iree_codegen.translation_info<pipeline = LLVMGPUTileAndFuse
+// CONV-SAME: workgroup_size = [128, 1, 1] subgroup_size = 32
+// CONV: linalg.conv_2d_nhwc_fhwc {{.*}}lowering_config = #iree_gpu.lowering_config
+// CONV-SAME: mma_kind = #iree_gpu.mma_layout<WMMAR4_F32_16x16x16_F16>
+// CONV-SAME: promote_operands = [0, 1]
+// CONV-SAME: reduction = [0, 0, 0, 0, 1, 1, 4]
+// CONV-SAME: subgroup = [1, 1, 2, 2, 0, 0, 0]
+// CONV-SAME: workgroup = [1, 1, 64, 64, 0, 0, 0]
+
+// -----
+
+// ============================================================================
+// Convolution — large (high arithmetic intensity, compute-bound)
+// ============================================================================
+
+func.func @conv_large_f16(%arg0: tensor<16x50x34x576xf16>, %arg1: tensor<576x3x3x576xf16>) -> tensor<16x48x32x576xf32> {
+ %cst = arith.constant 0.000000e+00 : f32
+ %empty = tensor.empty() : tensor<16x48x32x576xf32>
+ %fill = linalg.fill ins(%cst : f32) outs(%empty : tensor<16x48x32x576xf32>) -> tensor<16x48x32x576xf32>
+ %0 = linalg.conv_2d_nhwc_fhwc {dilations = dense<1> : tensor<2xi64>, strides = dense<1> : tensor<2xi64>}
+ ins(%arg0, %arg1 : tensor<16x50x34x576xf16>, tensor<576x3x3x576xf16>)
+ outs(%fill : tensor<16x48x32x576xf32>) -> tensor<16x48x32x576xf32>
+ return %0 : tensor<16x48x32x576xf32>
+}
+
+// CONV-LABEL: func.func @conv_large_f16
+// CONV-SAME: #iree_codegen.translation_info<pipeline = LLVMGPUTileAndFuse
+// CONV-SAME: workgroup_size = [128, 1, 1] subgroup_size = 32
+// CONV: linalg.conv_2d_nhwc_fhwc {{.*}}lowering_config = #iree_gpu.lowering_config
+// CONV-SAME: mma_kind = #iree_gpu.mma_layout<WMMAR4_F32_16x16x16_F16>
+// CONV-SAME: promote_operands = [0, 1]
+// CONV-SAME: reduction = [0, 0, 0, 0, 1, 1, 4]
+// CONV-SAME: subgroup = [1, 2, 1, 2, 0, 0, 0]
+// CONV-SAME: workgroup = [1, 2, 32, 64, 0, 0, 0]
