[mlir][Conversion] Implement ConvertToLLVMPatternInterface for FuncDialect

[springerm]

Summary

Also a new pass option ConvertToLLVMPass to populate only patterns from the specified dialects. This is needed because the existing test cases expect that only ops from certain dialects are lowered. (E.g., "arith-to-llvm" expects that only "arith" ops are lowered but not "func" ops.)

[springerm]

error checking

[mehdi_amini]

mlir/include/mlir/Conversion/Passes.td

30

What about calling this "filterDialects"? That's more descriptive of what it does right?

mlir/lib/Conversion/ConvertToLLVM/ConvertToLLVMPass.cpp

90

The error message does not check for the interface, or if the dialect is registered, just that it is loaded.

101–102

121

What's this comment about here?

mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp

797

This one is implicitly loaded already, the interface documents it as a guarantee

[springerm]

address comments

mlir/lib/Conversion/ConvertToLLVM/ConvertToLLVMPass.cpp

121

deleted, this is from an earlier version of this code

mlir/include/mlir/Conversion/ConvertToLLVM/ToLLVMPass.h

 //===- ToLLVMPass.h - Conversion to LLVM pass ---*- C++ -*-===================//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_CONVERSION_CONVERTTOLLVM_TOLLVM_PASS_H
 #define MLIR_CONVERSION_CONVERTTOLLVM_TOLLVM_PASS_H
 
 #include <memory>
 
+#include "mlir/Pass/Pass.h"
+
+#define GEN_PASS_DECL_CONVERTTOLLVMPASS
+#include "mlir/Conversion/Passes.h.inc"
+
 namespace mlir {
-class Pass;
 
 /// Create a pass that performs dialect conversion to LLVM  for all dialects
 /// implementing `ConvertToLLVMPatternInterface`.
 std::unique_ptr<Pass> createConvertToLLVMPass();
 
 } // namespace mlir
 
 #endif // MLIR_CONVERSION_CONVERTTOLLVM_TOLLVM_PASS_H

mlir/include/mlir/Conversion/FuncToLLVM/ConvertFuncToLLVM.h

 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_CONVERSION_FUNCTOLLVM_CONVERTFUNCTOLLVM_H
 #define MLIR_CONVERSION_FUNCTOLLVM_CONVERTFUNCTOLLVM_H
 
 namespace mlir {
 
+class DialectRegistry;
 class LLVMTypeConverter;
 class RewritePatternSet;
 
 /// Collect the default pattern to convert a FuncOp to the LLVM dialect. If
 /// `emitCWrappers` is set, the pattern will also produce functions
 /// that pass memref descriptors by pointer-to-structure in addition to the
 /// default unpacked form.
 void populateFuncToLLVMFuncOpConversionPattern(LLVMTypeConverter &converter,
                                                RewritePatternSet &patterns);
 
 /// Collect the patterns to convert from the Func dialect to LLVM. The
 /// conversion patterns capture the LLVMTypeConverter and the LowerToLLVMOptions
 /// by reference meaning the references have to remain alive during the entire
 /// pattern lifetime.
 void populateFuncToLLVMConversionPatterns(LLVMTypeConverter &converter,
                                           RewritePatternSet &patterns);
 
+void registerConvertFuncToLLVMInterface(DialectRegistry &registry);
+
 } // namespace mlir
 
 #endif // MLIR_CONVERSION_FUNCTOLLVM_CONVERTFUNCTOLLVM_H

mlir/include/mlir/Conversion/Passes.td

 
 def ConvertToLLVMPass : Pass<"convert-to-llvm"> {
   let summary = "Convert to LLVM via dialect interfaces found in the input IR";
   let description = [{
     This is a generic pass to convert to LLVM, it uses the
     `ConvertToLLVMPatternInterface` dialect interface to delegate to dialects
     the injection of conversion patterns.
   }];
+
+  let constructor = "mlir::createConvertToLLVMPass()";
+  let options = [
+    ListOption<"filterDialects", "filter-dialects", "std::string",
+               "Test conversion patterns of only the specified dialects">,
+  ];
 }
 
 //===----------------------------------------------------------------------===//
 // AffineToStandard
 //===----------------------------------------------------------------------===//
 
 def ConvertAffineToStandard : Pass<"lower-affine"> {
   let summary = "Lower Affine operations to a combination of Standard and SCF "

mlir/include/mlir/InitAllExtensions.h

 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_INITALLEXTENSIONS_H_
 #define MLIR_INITALLEXTENSIONS_H_
 
 #include "mlir/Conversion/ArithToLLVM/ArithToLLVM.h"
 #include "mlir/Conversion/ComplexToLLVM/ComplexToLLVM.h"
 #include "mlir/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.h"
+#include "mlir/Conversion/FuncToLLVM/ConvertFuncToLLVM.h"
 #include "mlir/Conversion/IndexToLLVM/IndexToLLVM.h"
 #include "mlir/Conversion/MathToLLVM/MathToLLVM.h"
 #include "mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h"
 #include "mlir/Conversion/NVVMToLLVM/NVVMToLLVM.h"
 #include "mlir/Conversion/UBToLLVM/UBToLLVM.h"
 #include "mlir/Dialect/Func/Extensions/AllExtensions.h"
 #include "mlir/Target/LLVM/NVVM/Target.h"
 
 #include <cstdlib>
 
 namespace mlir {
 
 /// This function may be called to register all MLIR dialect extensions with the
 /// provided registry.
 /// If you're building a compiler, you generally shouldn't use this: you would
 /// individually register the specific extensions that are useful for the
 /// pipelines and transformations you are using.
 inline void registerAllExtensions(DialectRegistry &registry) {
   arith::registerConvertArithToLLVMInterface(registry);
   registerConvertComplexToLLVMInterface(registry);
   cf::registerConvertControlFlowToLLVMInterface(registry);
   func::registerAllExtensions(registry);
+  registerConvertFuncToLLVMInterface(registry);
   index::registerConvertIndexToLLVMInterface(registry);
   registerConvertMathToLLVMInterface(registry);
   registerConvertMemRefToLLVMInterface(registry);
   registerConvertNVVMToLLVMInterface(registry);
   ub::registerConvertUBToLLVMInterface(registry);
   registerNVVMTarget(registry);
 }
 
 } // namespace mlir
 
 #endif // MLIR_INITALLEXTENSIONS_H_

mlir/lib/Conversion/ConvertToLLVM/ConvertToLLVMPass.cpp

 
 public:
   using impl::ConvertToLLVMPassBase<ConvertToLLVMPass>::ConvertToLLVMPassBase;
   void getDependentDialects(DialectRegistry &registry) const final {
     registry.insert<LLVM::LLVMDialect>();
     registry.addExtensions<LoadDependentDialectExtension>();
   }
 
-  ConvertToLLVMPass(const ConvertToLLVMPass &other)
-      : ConvertToLLVMPassBase(other), patterns(other.patterns),
-        target(other.target), typeConverter(other.typeConverter) {}
-
   LogicalResult initialize(MLIRContext *context) final {
     RewritePatternSet tempPatterns(context);
     auto target = std::make_shared<ConversionTarget>(*context);
     target->addLegalDialect<LLVM::LLVMDialect>();
     auto typeConverter = std::make_shared<LLVMTypeConverter>(context);
+
+    if (!filterDialects.empty()) {
+      // Test mode: Populate only patterns from the specified dialects. Produce
+      // an error if the dialect is not loaded or does not implement the
+      // interface.
+      for (std::string &dialectName : filterDialects) {
+        Dialect *dialect = context->getLoadedDialect(dialectName);
+        if (!dialect)
+          return emitError(UnknownLoc::get(context))
+                 << "dialect not loaded: " << dialectName << "\n";
+        auto iface = dyn_cast<ConvertToLLVMPatternInterface>(dialect);
+        if (!iface)
+          return emitError(UnknownLoc::get(context))
+                 << "dialect does not implement ConvertToLLVMPatternInterface: "
+                 << dialectName << "\n";
+        iface->populateConvertToLLVMConversionPatterns(*target, *typeConverter,
+                                                       tempPatterns);
+      }
+    } else {
+      // Normal mode: Populate all patterns from all dialects that implement the
+      // interface.
       for (Dialect *dialect : context->getLoadedDialects()) {
         // First time we encounter this dialect: if it implements the interface,
         // let's populate patterns !
         auto iface = dyn_cast<ConvertToLLVMPatternInterface>(dialect);
         if (!iface)
           continue;
         iface->populateConvertToLLVMConversionPatterns(*target, *typeConverter,
                                                        tempPatterns);
       }
-    patterns =
+    }
+
+    this->patterns =
         std::make_unique<FrozenRewritePatternSet>(std::move(tempPatterns));
     this->target = target;
     this->typeConverter = typeConverter;
     return success();
   }
 
   void runOnOperation() final {
     if (failed(applyPartialConversion(getOperation(), *target, *patterns)))
       signalPassFailure();
   }
 };
 
 } // namespace
+
+std::unique_ptr<Pass> mlir::createConvertToLLVMPass() {
+  return std::make_unique<ConvertToLLVMPass>();
+}

mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp

 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Conversion/FuncToLLVM/ConvertFuncToLLVMPass.h"
 
 #include "mlir/Analysis/DataLayoutAnalysis.h"
 #include "mlir/Conversion/ArithToLLVM/ArithToLLVM.h"
 #include "mlir/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.h"
+#include "mlir/Conversion/ConvertToLLVM/ToLLVMInterface.h"
 #include "mlir/Conversion/FuncToLLVM/ConvertFuncToLLVM.h"
 #include "mlir/Conversion/LLVMCommon/ConversionTarget.h"
 #include "mlir/Conversion/LLVMCommon/Pattern.h"
 #include "mlir/Conversion/LLVMCommon/VectorPattern.h"
 #include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/Dialect/LLVMIR/FunctionCallUtils.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/Dialect/LLVMIR/LLVMTypes.h"
void runOnOperation() override {
     if (failed(applyPartialConversion(m, target, std::move(patterns))))
       signalPassFailure();
 
     m->setAttr(LLVM::LLVMDialect::getDataLayoutAttrName(),
                StringAttr::get(m.getContext(), this->dataLayout));
   }
 };
 } // namespace
+
+//===----------------------------------------------------------------------===//
+// ConvertToLLVMPatternInterface implementation
+//===----------------------------------------------------------------------===//
+
+namespace {
+/// Implement the interface to convert Func to LLVM.
+struct FuncToLLVMDialectInterface : public ConvertToLLVMPatternInterface {
+  using ConvertToLLVMPatternInterface::ConvertToLLVMPatternInterface;
+  /// Hook for derived dialect interface to provide conversion patterns
+  /// and mark dialect legal for the conversion target.
+  void populateConvertToLLVMConversionPatterns(
+      ConversionTarget &target, LLVMTypeConverter &typeConverter,
+      RewritePatternSet &patterns) const final {
+    populateFuncToLLVMConversionPatterns(typeConverter, patterns);
+  }
+};
+} // namespace
+
+void mlir::registerConvertFuncToLLVMInterface(DialectRegistry &registry) {
+  registry.addExtension(+[](MLIRContext *ctx, func::FuncDialect *dialect) {
+    dialect->addInterfaces<FuncToLLVMDialectInterface>();
+  });
+}

mlir/test/Conversion/ArithToLLVM/arith-to-llvm.mlir

 // RUN: mlir-opt -pass-pipeline="builtin.module(func.func(convert-arith-to-llvm))" %s -split-input-file | FileCheck %s
 
 // Same below, but using the `ConvertToLLVMPatternInterface` entry point
 // and the generic `convert-to-llvm` pass.
-// RUN: mlir-opt --convert-to-llvm --split-input-file %s | FileCheck %s
+// RUN: mlir-opt --convert-to-llvm="filter-dialects=arith" --split-input-file %s | FileCheck %s
 
 // CHECK-LABEL: @vector_ops
 func.func @vector_ops(%arg0: vector<4xf32>, %arg1: vector<4xi1>, %arg2: vector<4xi64>, %arg3: vector<4xi64>) -> vector<4xf32> {
 // CHECK-NEXT:  %0 = llvm.mlir.constant(dense<4.200000e+01> : vector<4xf32>) : vector<4xf32>
   %0 = arith.constant dense<42.> : vector<4xf32>
 // CHECK-NEXT:  %1 = llvm.fadd %arg0, %0 : vector<4xf32>
   %1 = arith.addf %arg0, %0 : vector<4xf32>
 // CHECK-NEXT:  %2 = llvm.sdiv %arg2, %arg2 : vector<4xi64>

mlir/test/Conversion/ComplexToLLVM/convert-to-llvm.mlir

 // RUN: mlir-opt %s -convert-complex-to-llvm | FileCheck %s
 
 // Same below, but using the `ConvertToLLVMPatternInterface` entry point
 // and the generic `convert-to-llvm` pass.
-// RUN: mlir-opt --convert-to-llvm --split-input-file %s | FileCheck %s
+// RUN: mlir-opt --convert-to-llvm="filter-dialects=complex" --split-input-file %s | FileCheck %s
 
 // CHECK-LABEL: func @complex_create
 // CHECK-SAME:    (%[[REAL0:.*]]: f32, %[[IMAG0:.*]]: f32)
 // CHECK-NEXT:    %[[CPLX0:.*]] = llvm.mlir.undef : !llvm.struct<(f32, f32)>
 // CHECK-NEXT:    %[[CPLX1:.*]] = llvm.insertvalue %[[REAL0]], %[[CPLX0]][0] : !llvm.struct<(f32, f32)>
 // CHECK-NEXT:    %[[CPLX2:.*]] = llvm.insertvalue %[[IMAG0]], %[[CPLX1]][1] : !llvm.struct<(f32, f32)>
 func.func @complex_create(%real: f32, %imag: f32) -> complex<f32> {
   %cplx2 = complex.create %real, %imag : complex<f32>

mlir/test/Conversion/ControlFlowToLLVM/assert.mlir

 // RUN: mlir-opt %s -convert-cf-to-llvm='use-opaque-pointers=1' | FileCheck %s
 
 // Same below, but using the `ConvertToLLVMPatternInterface` entry point
 // and the generic `convert-to-llvm` pass.
-// RUN: mlir-opt --convert-to-llvm --split-input-file %s | FileCheck %s
+// RUN: mlir-opt --convert-to-llvm="filter-dialects=cf" --split-input-file %s | FileCheck %s
 
 func.func @main() {
   %a = arith.constant 0 : i1
   cf.assert %a, "assertion foo"
   return
 }
 
 // CHECK: llvm.func @puts(!llvm.ptr)
 
 // CHECK-LABEL: @main
 // CHECK: llvm.cond_br %{{.*}}, ^{{.*}}, ^[[FALSE_BRANCH:[[:alnum:]]+]]
 
 // CHECK: ^[[FALSE_BRANCH]]:
 // CHECK: %[[ADDRESS_OF:.*]] = llvm.mlir.addressof @{{.*}} : !llvm.ptr{{$}}
 // CHECK: %[[GEP:.*]] = llvm.getelementptr %[[ADDRESS_OF]][0] : (!llvm.ptr) -> !llvm.ptr, !llvm.array<{{[0-9]+}} x i8>
 // CHECK: llvm.call @puts(%[[GEP]]) : (!llvm.ptr) -> ()

mlir/test/Conversion/FuncToLLVM/func-to-llvm.mlir

-// RUN: mlir-opt -pass-pipeline="builtin.module(func.func(convert-math-to-llvm,convert-arith-to-llvm),convert-func-to-llvm{use-opaque-pointers=1},reconcile-unrealized-casts)" %s -split-input-file | FileCheck %s
+// RUN: mlir-opt -pass-pipeline="builtin.module(func.func(convert-math-to-llvm,convert-arith-to-llvm),convert-func-to-llvm{use-opaque-pointers=1},reconcile-unrealized-casts)" %s | FileCheck %s
 
-// RUN: mlir-opt -pass-pipeline="builtin.module(func.func(convert-math-to-llvm,convert-arith-to-llvm{index-bitwidth=32}),convert-func-to-llvm{index-bitwidth=32 use-opaque-pointers=1},reconcile-unrealized-casts)" %s -split-input-file | FileCheck --check-prefix=CHECK32 %s
+// RUN: mlir-opt -pass-pipeline="builtin.module(func.func(convert-math-to-llvm,convert-arith-to-llvm{index-bitwidth=32}),convert-func-to-llvm{index-bitwidth=32 use-opaque-pointers=1},reconcile-unrealized-casts)" %s | FileCheck --check-prefix=CHECK32 %s
 
 // RUN: mlir-opt -test-transform-dialect-interpreter %s | FileCheck --check-prefix=CHECK32 %s
 
+// Same below, but using the `ConvertToLLVMPatternInterface` entry point
+// and the generic `convert-to-llvm` pass.
+// RUN: mlir-opt --convert-to-llvm="filter-dialects=arith,cf,func,math" %s | FileCheck %s
+
 // CHECK-LABEL: func @empty() {
 // CHECK-NEXT:  llvm.return
 // CHECK-NEXT: }
 func.func @empty() {
 ^bb0:
   return
 }
 

mlir/test/Conversion/IndexToLLVM/index-to-llvm.mlir

 // RUN: mlir-opt %s -convert-index-to-llvm | FileCheck %s
 // RUN: mlir-opt %s -convert-index-to-llvm=index-bitwidth=32 | FileCheck %s --check-prefix=INDEX32
 // RUN: mlir-opt %s -convert-index-to-llvm=index-bitwidth=64 | FileCheck %s --check-prefix=INDEX64
 
 // Same below, but using the `ConvertToLLVMPatternInterface` entry point
 // and the generic `convert-to-llvm` pass.
-// RUN: mlir-opt --convert-to-llvm --split-input-file %s | FileCheck %s
+// RUN: mlir-opt --convert-to-llvm="filter-dialects=index" --split-input-file %s | FileCheck %s
 
 // CHECK-LABEL: @trivial_ops
 func.func @trivial_ops(%a: index, %b: index) {
   // CHECK: llvm.add
   %0 = index.add %a, %b
   // CHECK: llvm.sub
   %1 = index.sub %a, %b
   // CHECK: llvm.mul

mlir/test/Conversion/MathToLLVM/math-to-llvm.mlir

 // RUN: mlir-opt %s -split-input-file -pass-pipeline="builtin.module(func.func(convert-math-to-llvm))" | FileCheck %s
 
 // Same below, but using the `ConvertToLLVMPatternInterface` entry point
 // and the generic `convert-to-llvm` pass.
-// RUN: mlir-opt --convert-to-llvm --split-input-file %s | FileCheck %s
+// RUN: mlir-opt --convert-to-llvm="filter-dialects=math" --split-input-file %s | FileCheck %s
 
 // CHECK-LABEL: @ops
 func.func @ops(%arg0: f32, %arg1: f32, %arg2: i32, %arg3: i32, %arg4: f64) {
   // CHECK: = llvm.intr.exp(%{{.*}}) : (f32) -> f32
   %0 = math.exp %arg0 : f32
   // CHECK: = llvm.intr.exp2(%{{.*}}) : (f32) -> f32
   %1 = math.exp2 %arg0 : f32
   // CHECK: = llvm.intr.sqrt(%{{.*}}) : (f32) -> f32

mlir/test/Conversion/MemRefToLLVM/memref-to-llvm.mlir

 // RUN: mlir-opt -finalize-memref-to-llvm='use-opaque-pointers=1' %s -split-input-file | FileCheck %s
 // RUN: mlir-opt -finalize-memref-to-llvm='index-bitwidth=32 use-opaque-pointers=1' %s -split-input-file | FileCheck --check-prefix=CHECK32 %s
 
 // Same below, but using the `ConvertToLLVMPatternInterface` entry point
 // and the generic `convert-to-llvm` pass. This produces slightly different IR
 // because the conversion target is set up differently. Only one test case is
 // checked.
-// RUN: mlir-opt --convert-to-llvm --split-input-file %s | FileCheck --check-prefix=CHECK-INTERFACE %s
+// RUN: mlir-opt --convert-to-llvm="filter-dialects=memref" --split-input-file %s | FileCheck --check-prefix=CHECK-INTERFACE %s
 
 // CHECK-LABEL: func @view(
 // CHECK: %[[ARG0F:.*]]: index, %[[ARG1F:.*]]: index, %[[ARG2F:.*]]: index
 func.func @view(%arg0 : index, %arg1 : index, %arg2 : index) {
   // CHECK: %[[ARG2:.*]] = builtin.unrealized_conversion_cast %[[ARG2F:.*]]
   // CHECK: %[[ARG0:.*]] = builtin.unrealized_conversion_cast %[[ARG0F:.*]]
   // CHECK: %[[ARG1:.*]] = builtin.unrealized_conversion_cast %[[ARG1F:.*]]
   // CHECK: llvm.mlir.constant(2048 : index) : i64

mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir

 // RUN: mlir-opt --convert-nvvm-to-llvm --convert-arith-to-llvm --split-input-file %s | FileCheck %s
+
 // Same below, but using the `ConvertToLLVMPatternInterface` entry point
 // and the generic `convert-to-llvm` pass.
 // RUN: mlir-opt --convert-to-llvm --split-input-file %s | FileCheck %s
 
 // CHECK-LABEL : @init_mbarrier_arrive_expect_tx
 llvm.func @init_mbarrier_arrive_expect_tx(%barrier : !llvm.ptr<3>, %txcount : i32) {
   //CHECK : llvm.inline_asm has_side_effects asm_dialect = att "mbarrier.arrive.expect_tx.shared.b64 _, [$0], $1;", "r,r" 
   nvvm.mbarrier.arrive.expect_tx.shared %barrier, %txcount : !llvm.ptr<3>, i32

mlir/test/Conversion/UBToLLVM/ub-to-llvm.mlir

 // RUN: mlir-opt -pass-pipeline="builtin.module(func.func(convert-ub-to-llvm))" %s -split-input-file | FileCheck %s
 
 // Same below, but using the `ConvertToLLVMPatternInterface` entry point
 // and the generic `convert-to-llvm` pass.
-// RUN: mlir-opt --convert-to-llvm --split-input-file %s | FileCheck %s
+// RUN: mlir-opt --convert-to-llvm="filter-dialects=ub" --split-input-file %s | FileCheck %s
 
 // CHECK-LABEL: @check_poison
 func.func @check_poison() {
 // CHECK: {{.*}} = llvm.mlir.poison : i64
   %0 = ub.poison : index
 // CHECK: {{.*}} = llvm.mlir.poison : i16
   %1 = ub.poison : i16
 // CHECK: {{.*}} = llvm.mlir.poison : f64
   %2 = ub.poison : f64
   return
 }

utils/bazel/llvm-project-overlay/mlir/BUILD.bazel

 cc_library(
     name = "AllExtensions",
     hdrs = ["include/mlir/InitAllExtensions.h"],
     deps = [
         ":ArithToLLVM",
         ":ComplexToLLVM",
         ":ControlFlowToLLVM",
         ":FuncExtensions",
+        ":FuncToLLVM",
         ":IndexToLLVM",
         ":MathToLLVM",
         ":MemRefToLLVM",
         ":NVVMTarget",
         ":NVVMToLLVM",
         ":UBToLLVM",
     ],
 )
hdrs = [
         "include/mlir/Conversion/FuncToLLVM/ConvertFuncToLLVMPass.h",
     ],
     includes = ["include"],
     deps = [
         ":Analysis",
         ":ArithToLLVM",
         ":ControlFlowToLLVM",
         ":ConversionPassIncGen",
+        ":ConvertToLLVM",
         ":DataLayoutInterfaces",
         ":DialectUtils",
         ":FuncDialect",
         ":FuncTransforms",
         ":IR",
         ":LLVMCommonConversion",
         ":LLVMDialect",
         ":MathDialect",