llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-mlir-gpu Author: Nicolas Vasilache (nicolasvasilache) <details> <summary>Changes</summary> …rm dialect Authored-by: Son Tuan Vu <vuson@<!-- -->google.com> --- Full diff: https://github.com/llvm/llvm-project/pull/146962.diff 4 Files Affected: - (modified) mlir/include/mlir/Dialect/GPU/TransformOps/GPUTransformOps.td (+14) - (modified) mlir/lib/Dialect/GPU/TransformOps/CMakeLists.txt (+1) - (modified) mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp (+38) - (modified) utils/bazel/llvm-project-overlay/mlir/BUILD.bazel (+2) ``````````diff 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", `````````` </details> https://github.com/llvm/llvm-project/pull/146962 _______________________________________________ llvm-branch-commits mailing list llvm-branch-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
