Remove plier::GlobalIdOp and related code (#174)

This op and it spirv lowering was upstreamed.

Author: Hardcode84

mlir/include/mlir-extensions/dialect/plier_util/PlierUtilOps.td

@@ -17,7 +17,6 @@
 
 include "mlir/IR/OpBase.td"
 include "mlir/Dialect/GPU/GPUBase.td"
-include "mlir/IR/EnumAttr.td"
 include "mlir/Interfaces/ControlFlowInterfaces.td"
 include "mlir/Interfaces/InferTypeOpInterface.td"
 include "mlir/Interfaces/LoopLikeInterface.td"
@@ -275,26 +274,4 @@ def ReleaseContextOp : PlierUtil_Op<"release_context"> {
   let arguments = (ins PlierUtil_OpaqueType:$context);
 }
 
-def GPUEx_Dimension : I32EnumAttr<"GpuDimension",
-    "a dimension, either 'x', 'y', or 'z'",
-    [
-      I32EnumAttrCase<"x", 0>,
-      I32EnumAttrCase<"y", 1>,
-      I32EnumAttrCase<"z", 2>
-    ]>{
-  let genSpecializedAttr = 0;
-  let cppNamespace = "::plier";
-}
-def GPUEx_DimensionAttr : EnumAttr<PlierUtil_Dialect, GPUEx_Dimension, "dim">;
-
-class GPU_IndexOp<string mnemonic, list<Trait> traits = []> :
-    PlierUtil_Op<mnemonic, !listconcat(traits, [NoSideEffect])>,
-    Arguments<(ins GPUEx_DimensionAttr:$dimension)>, Results<(outs Index)> {
-  let assemblyFormat = "$dimension attr-dict";
-}
-
-def GPU_GlobalIdOp : GPU_IndexOp<"gpu_global_id"> {
-  let hasCanonicalizer = 1;
-}
-
 #endif // PLIER_UTIL_OPS

mlir/lib/dialect/plier_util/dialect.cpp

@@ -226,13 +226,57 @@ struct SpirvInputCSE : public mlir::OpRewritePattern<mlir::spirv::LoadOp> {
     return mlir::success();
   }
 };
+
+// TODO: move to separate pass
+struct GenGlobalId : public mlir::OpRewritePattern<mlir::arith::AddIOp> {
+  using OpRewritePattern::OpRewritePattern;
+
+  mlir::LogicalResult
+  matchAndRewrite(mlir::arith::AddIOp op,
+                  mlir::PatternRewriter &rewriter) const override {
+
+    auto getArg = [](auto op, bool rev) -> mlir::Value {
+      return rev ? op.getLhs() : op.getRhs();
+    };
+
+    mlir::gpu::Dimension dim;
+    mlir::arith::MulIOp other;
+    for (auto rev : {false, true}) {
+      auto arg1 = getArg(op, rev);
+      auto arg2 = getArg(op, !rev);
+      if (auto tid = arg1.getDefiningOp<mlir::gpu::ThreadIdOp>()) {
+        dim = tid.dimension();
+        other = arg2.getDefiningOp<mlir::arith::MulIOp>();
+        break;
+      }
+    }
+
+    if (!other)
+      return mlir::failure();
+
+    for (auto rev : {false, true}) {
+      auto arg1 = getArg(other, rev).getDefiningOp<mlir::gpu::BlockIdOp>();
+      auto arg2 = getArg(other, !rev).getDefiningOp<mlir::gpu::BlockDimOp>();
+      if (arg1 && arg2) {
+        if (arg1.dimension() != dim || arg2.dimension() != dim)
+          return mlir::failure();
+
+        rewriter.replaceOpWithNewOp<mlir::gpu::GlobalIdOp>(op, dim);
+        return mlir::success();
+      }
+    }
+
+    return mlir::failure();
+  }
+};
 } // namespace
 
 void PlierUtilDialect::getCanonicalizationPatterns(
     mlir::RewritePatternSet &results) const {
   results.add<DimExpandShape<mlir::tensor::DimOp, mlir::tensor::ExpandShapeOp>,
               DimExpandShape<mlir::memref::DimOp, mlir::memref::ExpandShapeOp>,
-              DimInsertSlice, FillExtractSlice, SpirvInputCSE>(getContext());
+              DimInsertSlice, FillExtractSlice, SpirvInputCSE, GenGlobalId>(
+      getContext());
 }
 
 OpaqueType OpaqueType::get(mlir::MLIRContext *context) {
@@ -1512,56 +1556,6 @@ void TakeContextOp::build(mlir::OpBuilder &b, mlir::OperationState &result,
   build(b, result, allTypes, initFunc, releaseFunc);
 }
 
-namespace {
-struct GenGlobalId : public mlir::OpRewritePattern<mlir::arith::AddIOp> {
-  using OpRewritePattern::OpRewritePattern;
-
-  mlir::LogicalResult
-  matchAndRewrite(mlir::arith::AddIOp op,
-                  mlir::PatternRewriter &rewriter) const override {
-
-    auto getArg = [](auto op, bool rev) -> mlir::Value {
-      return rev ? op.getLhs() : op.getRhs();
-    };
-
-    mlir::gpu::Dimension dim;
-    mlir::arith::MulIOp other;
-    for (auto rev : {false, true}) {
-      auto arg1 = getArg(op, rev);
-      auto arg2 = getArg(op, !rev);
-      if (auto tid = arg1.getDefiningOp<mlir::gpu::ThreadIdOp>()) {
-        dim = tid.dimension();
-        other = arg2.getDefiningOp<mlir::arith::MulIOp>();
-        break;
-      }
-    }
-
-    if (!other)
-      return mlir::failure();
-
-    for (auto rev : {false, true}) {
-      auto arg1 = getArg(other, rev).getDefiningOp<mlir::gpu::BlockIdOp>();
-      auto arg2 = getArg(other, !rev).getDefiningOp<mlir::gpu::BlockDimOp>();
-      if (arg1 && arg2) {
-        if (arg1.dimension() != dim || arg2.dimension() != dim)
-          return mlir::failure();
-
-        rewriter.replaceOpWithNewOp<plier::GlobalIdOp>(
-            op, static_cast<plier::GpuDimension>(dim));
-        return mlir::success();
-      }
-    }
-
-    return mlir::failure();
-  }
-};
-} // namespace
-
-void GlobalIdOp::getCanonicalizationPatterns(::mlir::RewritePatternSet &results,
-                                             ::mlir::MLIRContext *context) {
-  results.insert<GenGlobalId>(context);
-}
-
 } // namespace plier
 
 #include "mlir-extensions/dialect/plier_util/PlierUtilOpsDialect.cpp.inc"

numba_dpcomp/numba_dpcomp/mlir_compiler/lib/pipelines/lower_to_gpu.cpp

@@ -1208,37 +1208,6 @@ class ConvertFunc : public mlir::OpConversionPattern<mlir::FuncOp> {
   }
 };
 
-template <typename SourceOp, mlir::spirv::BuiltIn builtin>
-class LaunchConfigConversion : public mlir::OpConversionPattern<SourceOp> {
-public:
-  LaunchConfigConversion(mlir::TypeConverter &typeConverter,
-                         mlir::MLIRContext *context)
-      : mlir::OpConversionPattern<SourceOp>(typeConverter, context,
-                                            /*benefit*/ 100) {}
-  //  using mlir::OpConversionPattern<SourceOp>::OpConversionPattern;
-
-  mlir::LogicalResult
-  matchAndRewrite(SourceOp op, typename SourceOp::Adaptor adaptor,
-                  mlir::ConversionPatternRewriter &rewriter) const override;
-};
-
-template <typename SourceOp, mlir::spirv::BuiltIn builtin>
-mlir::LogicalResult LaunchConfigConversion<SourceOp, builtin>::matchAndRewrite(
-    SourceOp op, typename SourceOp::Adaptor adaptor,
-    mlir::ConversionPatternRewriter &rewriter) const {
-  auto *typeConverter =
-      this->template getTypeConverter<mlir::SPIRVTypeConverter>();
-  auto indexType = typeConverter->getIndexType();
-
-  // SPIR-V invocation builtin variables are a vector of type <3xi32>
-  auto spirvBuiltin =
-      mlir::spirv::getBuiltinVariableValue(op, builtin, indexType, rewriter);
-  rewriter.replaceOpWithNewOp<mlir::spirv::CompositeExtractOp>(
-      op, indexType, spirvBuiltin,
-      rewriter.getI32ArrayAttr({static_cast<int32_t>(op.dimension())}));
-  return mlir::success();
-}
-
 struct GPUToSpirvPass
     : public mlir::PassWrapper<GPUToSpirvPass,
                                mlir::OperationPass<mlir::ModuleOp>> {
@@ -1293,10 +1262,6 @@ struct GPUToSpirvPass
                 ConvertStoreOp, ConvertAtomicOps, ConvertFunc>(typeConverter,
                                                                context);
 
-    patterns.insert<LaunchConfigConversion<
-        plier::GlobalIdOp, mlir::spirv::BuiltIn::GlobalInvocationId>>(
-        typeConverter, context);
-
     if (failed(
             applyFullConversion(kernelModules, *target, std::move(patterns))))
       return signalPassFailure();