[llvm-branch-commits] [llvm] [mlir] [mlir][GPU][transform] Add gpu_to_rocdl conversion pattern to transfo… (PR #146962)
Nicolas Vasilache via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Thu Jul 3 14:15:31 PDT 2025
https://github.com/nicolasvasilache created https://github.com/llvm/llvm-project/pull/146962
…rm dialect
Authored-by: Son Tuan Vu <vuson at google.com>
>From d8730eb667660782ec1dce6e9cdea020c5821300 Mon Sep 17 00:00:00 2001
From: Nicolas Vasilache <nico.vasilache at amd.com>
Date: Thu, 3 Jul 2025 23:09:00 +0200
Subject: [PATCH] [mlir][GPU][transform] Add gpu_to_rocdl conversion pattern to
transform dialect
Authored-by: Son Tuan Vu <vuson at google.com>
---
.../GPU/TransformOps/GPUTransformOps.td | 14 +++++++
.../Dialect/GPU/TransformOps/CMakeLists.txt | 1 +
.../GPU/TransformOps/GPUTransformOps.cpp | 38 +++++++++++++++++++
.../llvm-project-overlay/mlir/BUILD.bazel | 2 +
4 files changed, 55 insertions(+)
diff --git a/mlir/include/mlir/Dialect/GPU/TransformOps/GPUTransformOps.td b/mlir/include/mlir/Dialect/GPU/TransformOps/GPUTransformOps.td
index 36b579485fc04..87423c639945f 100644
--- a/mlir/include/mlir/Dialect/GPU/TransformOps/GPUTransformOps.td
+++ b/mlir/include/mlir/Dialect/GPU/TransformOps/GPUTransformOps.td
@@ -54,6 +54,20 @@ def ApplyGPUSubgroupReduceToNVVMConversionPatternsOp : Op<Transform_Dialect,
let assemblyFormat = "attr-dict";
}
+def ApplyGPUToROCDLConversionPatternsOp : Op<Transform_Dialect,
+ "apply_conversion_patterns.gpu.gpu_to_rocdl",
+ [DeclareOpInterfaceMethods<ConversionPatternDescriptorOpInterface,
+ ["verifyTypeConverter"]>]> {
+ let description = [{
+ Collects patterns that convert GPU dialect ops to ROCDL dialect ops. These
+ patterns require an "LLVMTypeConverter".
+ }];
+ let arguments = (ins StrAttr:$chipset);
+ let assemblyFormat = [{
+ `chipset` `=` $chipset attr-dict
+ }];
+}
+
//===----------------------------------------------------------------------===//
// Apply...PatternsOp
//===----------------------------------------------------------------------===//
diff --git a/mlir/lib/Dialect/GPU/TransformOps/CMakeLists.txt b/mlir/lib/Dialect/GPU/TransformOps/CMakeLists.txt
index b26788f675ce5..e5cc0254f1ffe 100644
--- a/mlir/lib/Dialect/GPU/TransformOps/CMakeLists.txt
+++ b/mlir/lib/Dialect/GPU/TransformOps/CMakeLists.txt
@@ -24,4 +24,5 @@ add_mlir_dialect_library(MLIRGPUTransformOps
# ConversionPatterns
MLIRNVGPUToNVVM
MLIRGPUToNVVMTransforms
+ MLIRGPUToROCDLTransforms
)
diff --git a/mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp b/mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp
index a86fc47947130..b764a72529f8f 100644
--- a/mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp
+++ b/mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp
@@ -10,6 +10,7 @@
#include "mlir/Conversion/GPUCommon/GPUCommonPass.h"
#include "mlir/Conversion/GPUToNVVM/GPUToNVVMPass.h"
+#include "mlir/Conversion/GPUToROCDL/GPUToROCDLPass.h"
#include "mlir/Conversion/LLVMCommon/TypeConverter.h"
#include "mlir/Dialect/AMDGPU/IR/AMDGPUDialect.h"
#include "mlir/Dialect/Affine/IR/AffineOps.h"
@@ -42,6 +43,7 @@
#include "llvm/Support/Debug.h"
#include "llvm/Support/ErrorHandling.h"
#include "llvm/Support/InterleavedRange.h"
+#include "llvm/Support/LogicalResult.h"
#include <type_traits>
using namespace mlir;
@@ -129,6 +131,42 @@ LogicalResult transform::ApplyGPUSubgroupReduceToNVVMConversionPatternsOp::
return success();
}
+void transform::ApplyGPUToROCDLConversionPatternsOp::populatePatterns(
+ TypeConverter &typeConverter, RewritePatternSet &patterns) {
+ auto &llvmTypeConverter = static_cast<LLVMTypeConverter &>(typeConverter);
+ populateGpuMemorySpaceAttributeConversions(
+ llvmTypeConverter, [](AddressSpace space) {
+ switch (space) {
+ case AddressSpace::Global:
+ return 1;
+ case AddressSpace::Workgroup:
+ return 3;
+ case AddressSpace::Private:
+ return 5;
+ }
+ llvm_unreachable("unknown address space enum value");
+ return 0;
+ });
+ FailureOr<amdgpu::Chipset> maybeChipset =
+ amdgpu::Chipset::parse(getChipset());
+ assert(llvm::succeeded(maybeChipset) && "expected valid chipset");
+ populateGpuToROCDLConversionPatterns(
+ llvmTypeConverter, patterns, mlir::gpu::amd::Runtime::HIP, *maybeChipset);
+}
+
+LogicalResult
+transform::ApplyGPUToROCDLConversionPatternsOp::verifyTypeConverter(
+ transform::TypeConverterBuilderOpInterface builder) {
+ FailureOr<amdgpu::Chipset> maybeChipset =
+ amdgpu::Chipset::parse(getChipset());
+ if (failed(maybeChipset)) {
+ return emitOpError("Invalid chipset name: " + getChipset());
+ }
+ if (builder.getTypeConverterType() != "LLVMTypeConverter")
+ return emitOpError("expected LLVMTypeConverter");
+ return success();
+}
+
//===----------------------------------------------------------------------===//
// Apply...PatternsOp
//===----------------------------------------------------------------------===//s
diff --git a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
index cc266c2fe3a77..79f2cd5ea71db 100644
--- a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
+++ b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
@@ -5502,6 +5502,7 @@ cc_library(
":GPUDialect",
":GPUToGPURuntimeTransforms",
":GPUToNVVMTransforms",
+ ":GPUToROCDLTransforms",
":GPUTransformOpsIncGen",
":GPUTransforms",
":IR",
@@ -5509,6 +5510,7 @@ cc_library(
":MemRefDialect",
":NVGPUDialect",
":NVVMDialect",
+ ":ROCDLDialect",
":SCFDialect",
":Support",
":TransformDialect",
More information about the llvm-branch-commits
mailing list