simplifying pattern application

Author: Muzammiluddin-Syed-ECE

mlir/include/mlir/Conversion/GPUToAMDGPU/GPUToAMDGPU.h

@@ -23,8 +23,7 @@ class Pass;
 #define GEN_PASS_DECL_CONVERTGPUTOAMDGPUPASS
 #include "mlir/Conversion/Passes.h.inc"
 
-void populateSubgroupReduceLoweringPatterns(LLVMTypeConverter &converter,
-                                            RewritePatternSet &patterns,
+void populateSubgroupReduceLoweringPatterns(RewritePatternSet &patterns,
                                             unsigned subgroupSize,
                                             PatternBenefit benefit);
 // void populateGPUToAMDGPUConversionPatterns(LLVMTypeConverter &converter,

mlir/lib/Conversion/GPUToAMDGPU/GPUToAMDGPU.cpp

@@ -8,11 +8,7 @@
 
 #include "mlir/Conversion/GPUToAMDGPU/GPUToAMDGPU.h"
 
-#include "mlir/Conversion/LLVMCommon/ConversionTarget.h"
-#include "mlir/Conversion/LLVMCommon/Pattern.h"
-#include "mlir/Conversion/LLVMCommon/TypeConverter.h"
 #include "mlir/Dialect/AMDGPU/IR/AMDGPUDialect.h"
-#include "mlir/Dialect/AMDGPU/Utils/Chipset.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
 #include "mlir/IR/BuiltinTypes.h"
@@ -23,15 +19,8 @@
 #include "mlir/Dialect/GPU/IR/GPUDialect.h"
 #include "mlir/Dialect/Vector/IR/VectorOps.h"
 
+#include "mlir/Transforms/WalkPatternRewriteDriver.h"
 #include "llvm/Support/FormatVariadic.h"
-#include "llvm/Support/MathExtras.h"
-#include <cassert>
-#include <cstdint>
-
-#include "../LLVMCommon/MemRefDescriptor.h"
-
-#include "llvm/ADT/STLExtras.h"
-#include <optional>
 
 namespace mlir {
 #define GEN_PASS_DEF_CONVERTGPUTOAMDGPUPASS
@@ -180,24 +169,17 @@ struct ConvertGPUToAMDGPUPass
 
   void runOnOperation() override {
     RewritePatternSet patterns(&getContext());
-    LLVMTypeConverter converter(&getContext());
-    LLVMConversionTarget target(getContext());
-    target.addLegalDialect<::mlir::LLVM::LLVMDialect>();
-    target.addLegalDialect<::mlir::amdgpu::AMDGPUDialect>();
-    target.addLegalDialect<::mlir::ROCDL::ROCDLDialect>();
-
     int subgroupSizeInt = static_cast<int>(subgroupSize);
-    populateSubgroupReduceLoweringPatterns(converter, patterns, subgroupSizeInt,
+    populateSubgroupReduceLoweringPatterns(patterns, subgroupSizeInt,
                                            PatternBenefit(1));
-    if (failed(applyPartialConversion(getOperation(), target,
-                                      std::move(patterns))))
-      signalPassFailure();
+    walkAndApplyPatterns(getOperation(), std::move(patterns));
   }
 };
 } // namespace
 
-void mlir::populateSubgroupReduceLoweringPatterns(
-    LLVMTypeConverter &converter, RewritePatternSet &patterns, unsigned subgroupSize, PatternBenefit benefit) {
+void mlir::populateSubgroupReduceLoweringPatterns(RewritePatternSet &patterns,
+                                                  unsigned subgroupSize,
+                                                  PatternBenefit benefit) {
   patterns.add<ScalarSubgroupReduceToShuffles>(
       patterns.getContext(), subgroupSize, /*matchClustered=*/true, benefit);
 }