[llvm] [mlir] [mlir][GPU][transform] Add gpu_to_rocdl conversion pattern to transfo… (PR #146962)
Nicolas Vasilache via llvm-commits
llvm-commits at lists.llvm.org
Mon Jul 7 09:15:17 PDT 2025
https://github.com/nicolasvasilache updated https://github.com/llvm/llvm-project/pull/146962
>From 235e608eacd388819ecae01937c601c82708365a Mon Sep 17 00:00:00 2001
From: Son Tuan Vu <vuson at google.com>
Date: Thu, 3 Jul 2025 23:09:00 +0200
Subject: [PATCH 1/2] [mlir][GPU][transform] Add gpu_to_rocdl conversion
pattern to transform dialect
---
.../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 f7fcd99b030dd..f2f797caf942f 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/Arith/IR/Arith.h"
@@ -40,6 +41,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;
@@ -127,6 +129,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 633776398649f..7f59236a3bb27 100644
--- a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
+++ b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
@@ -5505,6 +5505,7 @@ cc_library(
":GPUDialect",
":GPUToGPURuntimeTransforms",
":GPUToNVVMTransforms",
+ ":GPUToROCDLTransforms",
":GPUTransformOpsIncGen",
":GPUTransforms",
":IR",
@@ -5512,6 +5513,7 @@ cc_library(
":MemRefDialect",
":NVGPUDialect",
":NVVMDialect",
+ ":ROCDLDialect",
":SCFDialect",
":Support",
":TransformDialect",
>From ab7fce1f7d935b83404bbaf8b4f9dff8d7feb914 Mon Sep 17 00:00:00 2001
From: Nicolas Vasilache <nico.vasilache at amd.com>
Date: Mon, 7 Jul 2025 18:13:22 +0200
Subject: [PATCH 2/2] Address review
---
mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp | 8 ++++----
1 file changed, 4 insertions(+), 4 deletions(-)
diff --git a/mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp b/mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp
index f2f797caf942f..c9e91535df946 100644
--- a/mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp
+++ b/mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp
@@ -18,6 +18,7 @@
#include "mlir/Dialect/GPU/TransformOps/Utils.h"
#include "mlir/Dialect/GPU/Transforms/Passes.h"
#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
+#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
#include "mlir/Dialect/SCF/IR/DeviceMappingInterface.h"
#include "mlir/Dialect/SCF/IR/SCF.h"
@@ -136,14 +137,13 @@ void transform::ApplyGPUToROCDLConversionPatternsOp::populatePatterns(
llvmTypeConverter, [](AddressSpace space) {
switch (space) {
case AddressSpace::Global:
- return 1;
+ return ROCDL::ROCDLDialect::kGlobalMemoryAddressSpace;
case AddressSpace::Workgroup:
- return 3;
+ return ROCDL::ROCDLDialect::kSharedMemoryAddressSpace;
case AddressSpace::Private:
- return 5;
+ return ROCDL::ROCDLDialect::kPrivateMemoryAddressSpace;
}
llvm_unreachable("unknown address space enum value");
- return 0;
});
FailureOr<amdgpu::Chipset> maybeChipset =
amdgpu::Chipset::parse(getChipset());
More information about the llvm-commits
mailing list