Add a generic "convert-to-llvm" pass delegating to an interface

[mehdi_amini]

Summary

The multiple -convert-XXX-to-llvm passes are really nice testing tools for
individual dialects, but the expectation is that a proper conversion should
assemble the conversion patterns using `populateXXXToLLVMConversionPatterns()
APIs. However most customers just chain the conversion passes by convenience.

This pass makes it composable more transparently to assemble the required
patterns for conversion to LLVM dialect by using an interface.
The Pass will scan the input and collect all the dialect present, and for
those who implement the ConvertToLLVMPatternInterface it will use it to
populate the conversion pattern, and possible the conversion target.

Since these conversions can involve intermediate dialects, or target other
dialects than LLVM (for example AVX or NVVM), this pass can't statically
declare the required getDependentDialects() before the pass pipeline
begins. This is worked around by using an extension in the dialectRegistry
that will be invoked for every new loaded dialects in the context. This
allows to lookup the interface ahead of time and use it to query the
dependent dialects.

This patch demonstrate it with the NVVM dialect, but we should then
apply the same pattern for all the dialect that have conversions to LLVM.

[Mogball]

mlir/lib/Conversion/ConvertToLLVM/ConvertToLLVMPass.cpp

34

nice

mlir/lib/Conversion/ConvertToLLVM/ToLLVMInterface.cpp

20

Ouch. Would it be better to just iterate the loaded dialects in the context?

mlir/lib/Conversion/NVVMToLLVM/NVVMToLLVM.cpp

212

leftover debugging?

220

\n

[mehdi_amini]

Initialize patterns from the dialects loaded in the context

mlir/lib/Conversion/ConvertToLLVM/ToLLVMInterface.cpp

20

I was wondering about this: if we have many functions with only a subset of the dialects, how costly is it to load all the possible patterns that we don't need?

I implemented this alternative now, PTAL :)

[ftynse]

Nice!

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

2

Nit: spurious empty line.

mlir/include/mlir/Conversion/NVVMToLLVM/NVVMToLLVM.h

11

Nit: doesn't look necessary given the forward declaration below.

mlir/include/mlir/Conversion/Passes.td

22

mlir/lib/Conversion/ConvertToLLVM/ConvertToLLVMPass.cpp

31

Nit: a short documentation comment is welcome, here and below.

56–57

Could you comment on unique/shared ptr difference here? Can't the frozen pattern set be shared between copies of the pass?

63

Nit: registry.addExtensions<LoadDependentDialectExtension>() reads better.

74

Should this be also exposed to the interface? Patterns may be producing NVVM dialect and friends, which should be declared legal.

[mehdi_amini]

Address Alex's comments

[mehdi_amini]

mlir/lib/Conversion/ConvertToLLVM/ConvertToLLVMPass.cpp

56–57

I should be able to use shared everywhere now, the mutable state was making me nervous, I added a dependent revision and I'm using of shared_ptr<const T> now

63

I don't think it works? It requires a callback as argument and build an extension class around it? The template arguments are for the dialects to filter on.

74

It is actually :)
populateConvertToLLVMConversionPatterns gets the target for this purpose.

[jpienaar]

One thing I'm wondering: the populate methods allow one to combine it in rather arbitrary ways & hops. But the interface for me signals more a fixed/preferred way. For the lower level dialects there isn't much choice and I don't think we have any multipaths here already, so not a practical concern today, more question about positioning.

mlir/include/mlir/Conversion/ConvertToLLVM/ToLLVMInterface.h

21

Let's keep these sorted

mlir/lib/Conversion/ConvertToLLVM/ConvertToLLVMPass.cpp

36

This feels like something more general ... Especially as this is something dynamic pipelines commonly run into.

[mehdi_amini]

mlir/lib/Conversion/ConvertToLLVM/ConvertToLLVMPass.cpp

36

Here the anchor point is that it delegate some control to the dialects, but I'm not sure how to connect to dynamic pipeline?

The kind of problem I've seen with dynamic pipelines is when people build a complete pipeline where they want for example to delegate the dynamic pipeline to an Op or Attribute interface for example. The dynamic pipeline isn't known until the pass execution starts, since it is queried from the IR.
I'm not sure what's the solution for this though, I think it needs to split the pass pipeline instead and use other means.

[awarzynski]

This will be super helpful, thank you!

Could you add some tests?

mlir/lib/Conversion/ConvertToLLVM/ConvertToLLVMPass.cpp

46

Why not continue?

[mehdi_amini]

Address comments and enable the test

mlir/lib/Conversion/ConvertToLLVM/ConvertToLLVMPass.cpp

46

Ah, I just refactored this out of a lambda, where return was right, we need continue here indeed.

[ftynse]

mlir/lib/Conversion/ConvertToLLVM/ConvertToLLVMPass.cpp

63

I mean this one: https://github.com/llvm/llvm-project/blob/f580901d5d30e37755212f1c09e5b587587fbfeb/mlir/include/mlir/IR/DialectRegistry.h#L199-L203. It only does make_unique inside, so it looks like it should work.

[mehdi_amini]

Address ftynse's comment

[springerm]

mlir/lib/Conversion/NVVMToLLVM/NVVMToLLVM.cpp

204

Can we change this into a type converter? ConversionTarget is an argument of applyPartialConversion etc. and I have not seen it as an argument to populate... methods. ConversionPattern has a type converter field and many populateLLVM... methods require a type converter.

[springerm]

mlir/lib/Conversion/NVVMToLLVM/NVVMToLLVM.cpp

204

Just saw the comment about NVVM above. How about having both?

[nicolasvasilache]

Generally looks good but this seems like a very limited use case for now: the NVVMToLLVM actually does not do conversions, it just applies a single rewrite pattern to lower the ops with PTX builder to inline asm ops.

How do you envision this being extended?
I can see how to compose this with some of @springerm 's work on applying conversions but we still need customizable TypeConverter.
Also some conversions use/require flags (e.g. fp32 reassoc for fmas), I am unclear what the story is there?

Maybe the overarching question is whether you have this connected in a bigger context and have good confidence that you're able to solve the tricky conversion pass compositions we've been discussing?
I'd love to be able to throw some of my local hacks away :)

[mehdi_amini]

In D157183#4568897, @nicolasvasilache wrote:

How do you envision this being extended?

As needed :)
Right now I hope to move the individual passes to be "test-passes", and avoid the issue of unrealized_cast as much as possible.

Also some conversions use/require flags (e.g. fp32 reassoc for fmas), I am unclear what the story is there?

I think we should operate with other means than options: this does not scale well and a pass isn't the best granularity for this in general: data layout and other option may better be in the IR (where you could attach them as attribute to functions or module? Allowing different options for different functions?)

Maybe the overarching question is whether you have this connected in a bigger context and have good confidence that you're able to solve the tricky conversion pass compositions we've been discussing?
I'd love to be able to throw some of my local hacks away :)

mlir/lib/Conversion/ConvertToLLVM/ConvertToLLVMPass.cpp

63

Somehow I missed this! Maybe it is because it has a plural.. :)
Thanks!

[krzysz00]

So, I have some minor thoughts here:

1. It sounds like this'll help with target-specific overrides of type conversion, am I right here?
2. Similarly, it'd be nice if this general pass would take options like triple and chip so you could look up the LLVM data layout etc., since that's logic that'd be a pain to set up everywhere
3. One interesting idea would be to generalize D158287 . I'd like, for example, the notion of "target-specific patterns". That way, for example, when we have these vector-to-llvm patterns, the avx patterns would live in a separate library and be pulled in only if I'm targetting x86. This would allow me to not even build the target-specific dialects (for instance, why does a library targetting AMDGPU need ArmSVE?) and to have that be a general pattern. (As another example, I'm having to pull in NVVM because of some memref patterns.)

[mehdi_amini]

In D157183#4606830, @krzysz00 wrote:

1. It sounds like this'll help with target-specific overrides of type conversion, am I right here?

I don't understand the question?

2. Similarly, it'd be nice if this general pass would take options like triple and chip so you could look up the LLVM data layout etc., since that's logic that'd be a pain to set up everywhere

Should this be handled by adding a data layout and other target information as a module attribute?

3. One interesting idea would be to generalize D158287 . I'd like, for example, the notion of "target-specific patterns". That way, for example, when we have these vector-to-llvm patterns, the avx patterns would live in a separate library and be pulled in only if I'm targetting x86. This would allow me to not even build the target-specific dialects (for instance, why does a library targetting AMDGPU need ArmSVE?) and to have that be a general pattern. (As another example, I'm having to pull in NVVM because of some memref patterns.)

Right now with this pass this would work as long as the pipeline introduces a dependency on a dialect (that implements the interface).

I don't think it is satisfactory for D158287 though: because there can be multiple options to lower to LLVM when multiple vector ISA dialects are available. I'm not sure how to model it, but seems like it comes back to some sort of attributes (even function granularity maybe).

[krzysz00]

1. I'm asking in relation to https://discourse.llvm.org/t/rfc-almost-all-uses-of-typeconverter-should-be-const/72689/9 and the existence of the need to create target-specific overrides of the LLVM type conversion (for instance, AMDGPU lowers bf16 to i16 to match the expectations of the backend). I'm not sure the new pass has much to do with this

Similarly, it'd be nice if this general pass would take options like triple and chip so you could look up the LLVM data layout etc., since that's logic that'd be a pain to set up everywhere

Should this be handled by adding a data layout and other target information as a module attribute?

Yeah, that could probably be attributes. That might actually make for a reasonably clean design overall and make sure thing stay consistent between the lowering to LLVM and any compilation/JIT.

Right now with this pass this would work as long as the pipeline introduces a dependency on a dialect (that implements the interface).

I don't think it is satisfactory for D158287 though: because there can be multiple options to lower to LLVM when multiple vector ISA dialects are available. I'm not sure how to model it, but seems like it comes back to some sort of attributes (even function granularity maybe).

The high-level context on all this is that a (potentially somewhat fossilized) "one big static library" build of our downstream project was tripping up the Windows 4 GB debug symbol limit. That's all worked around, but it sent me looking for ways to slim down our build of MLIR to the dialects we reasonably need (for example, "Wait, why is ArmSVE in here - we're not targeting ARM!").

Currently, the way you'd enable these optional target-specific extensions is either calling a pass with an argument set, or by manually building up a pattern set.

Since this pass is trying to replace (?) building up pattern sets by hand, some way to signal "I'm targeting X, please load all the X-specific patterns", even if you don't have any operations from the X dialect in your code (unless we want some sort of no-op that'll trigger the existing interface search) would be neat.

[mehdi_amini]

The high-level context on all this is that a (potentially somewhat fossilized) "one big static library" build of our downstream project was tripping up the Windows 4 GB debug symbol limit. That's all worked around, but it sent me looking for ways to slim down our build of MLIR to the dialects we reasonably need (for example, "Wait, why is ArmSVE in here - we're not targeting ARM!").
Currently, the way you'd enable these optional target-specific extensions is either calling a pass with an argument set, or by manually building up a pattern set.
Since this pass is trying to replace (?) building up pattern sets by hand, some way to signal "I'm targeting X, please load all the X-specific patterns", even if you don't have any operations from the X dialect in your code (unless we want some sort of no-op that'll trigger the existing interface search) would be neat.

Right, that's why I was thinking about an attribute, on the model of what we did for the GPU pipeline: the attribute could have an externally registered interface that load these patterns. Since it's all opt-in registered that breaks the build dependencies.
It's not obvious to me how to plumb everything together just now though.

mlir/include/mlir/Conversion/ConvertToLLVM/ToLLVMInterface.h

+//===- ToLLVMInterface.h - Conversion to LLVM iface ---*- 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_TOLLVMINTERFACE_H
+#define MLIR_CONVERSION_CONVERTTOLLVM_TOLLVMINTERFACE_H
+
+#include "mlir/IR/DialectInterface.h"
+#include "mlir/IR/MLIRContext.h"
+#include "mlir/Support/LogicalResult.h"
+
+namespace mlir {
+class ConversionTarget;
+class LLVMTypeConverter;
+class MLIRContext;
+class Operation;
+class RewritePatternSet;
+
+/// Base class for dialect interfaces providing translation to LLVM IR.
+/// Dialects that can be translated should provide an implementation of this
+/// interface for the supported operations. The interface may be implemented in
+/// a separate library to avoid the "main" dialect library depending on LLVM IR.
+/// The interface can be attached using the delayed registration mechanism
+/// available in DialectRegistry.
+class ConvertToLLVMPatternInterface
+    : public DialectInterface::Base<ConvertToLLVMPatternInterface> {
+public:
+  ConvertToLLVMPatternInterface(Dialect *dialect) : Base(dialect) {}
+
+  /// Hook for derived dialect interface to load the dialects they
+  /// target. The LLVMDialect is implicitly already loaded, but this
+  /// method allows to load other intermediate dialects used in the
+  /// conversion, or target dialects like NVVM for example.
+  virtual void loadDependentDialects(MLIRContext *context) const {}
+
+  /// Hook for derived dialect interface to provide conversion patterns
+  /// and mark dialect legal for the conversion target.
+  virtual void populateConvertToLLVMConversionPatterns(
+      ConversionTarget &target, RewritePatternSet &patterns) const = 0;
+};
+
+/// Recursively walk the IR and collect all dialects implementing the interface,
+/// and populate the conversion patterns.
+void populateConversionTargetFromOperation(Operation *op,
+                                           ConversionTarget &target,
+                                           RewritePatternSet &patterns);
+
+} // namespace mlir
+
+#endif // MLIR_CONVERSION_CONVERTTOLLVM_TOLLVMINTERFACE_H

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>
+
+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/NVVMToLLVM/NVVMToLLVM.h

 //===- NVVMTOLLVMPass.h - Convert NVVM to LLVM dialect --------*- 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_NVVMTOLLVM_NVVMTOLLVMPASS_H_
 #define MLIR_CONVERSION_NVVMTOLLVM_NVVMTOLLVMPASS_H_
 
 #include <memory>
 
 namespace mlir {
-
+class DialectRegistry;
 class LLVMTypeConverter;
 class RewritePatternSet;
 class Pass;
 
 #define GEN_PASS_DECL_CONVERTNVVMTOLLVMPASS
 #include "mlir/Conversion/Passes.h.inc"
 
 void populateNVVMToLLVMConversionPatterns(RewritePatternSet &patterns);
 
+void registerConvertNVVMToLLVMInterface(DialectRegistry &registry);
+
 } // namespace mlir
 
 #endif // MLIR_CONVERSION_NVVMTOLLVM_NVVMTOLLVMPASS_H_

mlir/include/mlir/Conversion/Passes.h

 #include "mlir/Conversion/BufferizationToMemRef/BufferizationToMemRef.h"
 #include "mlir/Conversion/ComplexToLLVM/ComplexToLLVM.h"
 #include "mlir/Conversion/ComplexToLibm/ComplexToLibm.h"
 #include "mlir/Conversion/ComplexToSPIRV/ComplexToSPIRVPass.h"
 #include "mlir/Conversion/ComplexToStandard/ComplexToStandard.h"
 #include "mlir/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.h"
 #include "mlir/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRV.h"
 #include "mlir/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRVPass.h"
+#include "mlir/Conversion/ConvertToLLVM/ToLLVMPass.h"
 #include "mlir/Conversion/FuncToLLVM/ConvertFuncToLLVMPass.h"
 #include "mlir/Conversion/FuncToSPIRV/FuncToSPIRVPass.h"
 #include "mlir/Conversion/GPUCommon/GPUCommonPass.h"
 #include "mlir/Conversion/GPUToNVVM/GPUToNVVMPass.h"
 #include "mlir/Conversion/GPUToROCDL/GPUToROCDLPass.h"
 #include "mlir/Conversion/GPUToSPIRV/GPUToSPIRVPass.h"
 #include "mlir/Conversion/GPUToVulkan/ConvertGPUToVulkanPass.h"
 #include "mlir/Conversion/IndexToLLVM/IndexToLLVM.h"

mlir/include/mlir/Conversion/Passes.td

 //===-- Passes.td - Conversion pass definition file --------*- tablegen -*-===//
 //
 // 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_PASSES
 #define MLIR_CONVERSION_PASSES
 
 include "mlir/Pass/PassBase.td"
 
+
+//===----------------------------------------------------------------------===//
+// ToLLVM
+//===----------------------------------------------------------------------===//
+
+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.
+  }];
+}
+
 //===----------------------------------------------------------------------===//
 // AffineToStandard
 //===----------------------------------------------------------------------===//
 
 def ConvertAffineToStandard : Pass<"lower-affine"> {
   let summary = "Lower Affine operations to a combination of Standard and SCF "
                 "operations";
   let description = [{

mlir/include/mlir/IR/DialectRegistry.h

 /// required dialects are loaded. The application function can be used to attach
 /// additional functionality to attributes, dialects, operations, types, etc.,
 /// and may also load additional necessary dialects.
 class DialectExtensionBase {
 public:
   virtual ~DialectExtensionBase();
 
   /// Return the dialects that our required by this extension to be loaded
-  /// before applying.
+  /// before applying. If empty then the extension is invoked for every loaded
+  /// dialect indepently.
   ArrayRef<StringRef> getRequiredDialects() const { return dialectNames; }
 
   /// Apply this extension to the given context and the required dialects.
   virtual void apply(MLIRContext *context,
                      MutableArrayRef<Dialect *> dialects) const = 0;
 
   /// Return a copy of this extension.
   virtual std::unique_ptr<DialectExtensionBase> clone() const = 0;
 
 protected:
-  /// Initialize the extension with a set of required dialects. Note that there
-  /// should always be at least one affected dialect.
+  /// Initialize the extension with a set of required dialects.
+  /// If the list is empty, the extension is invoked for every loaded dialect
+  /// independently.
   DialectExtensionBase(ArrayRef<StringRef> dialectNames)
-      : dialectNames(dialectNames.begin(), dialectNames.end()) {
-    assert(!dialectNames.empty() && "expected at least one affected dialect");
-  }
+      : dialectNames(dialectNames.begin(), dialectNames.end()) {}
 
 private:
   /// The names of the dialects affected by this extension.
   SmallVector<StringRef> dialectNames;
 };
 
 /// This class represents a dialect extension anchored on the given set of
 /// dialects. When all of the specified dialects have been loaded, the

mlir/include/mlir/InitAllExtensions.h

 //===- InitAllExtensions.h - MLIR Extension Registration --------*- 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file defines a helper to trigger the registration of all dialect
 // extensions to the system.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_INITALLEXTENSIONS_H_
 #define MLIR_INITALLEXTENSIONS_H_
 
+#include "mlir/Conversion/NVVMToLLVM/NVVMToLLVM.h"
 #include "mlir/Dialect/Func/Extensions/AllExtensions.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) {
   func::registerAllExtensions(registry);
+  registerConvertNVVMToLLVMInterface(registry);
 }
 
 } // namespace mlir
 
 #endif // MLIR_INITALLEXTENSIONS_H_

mlir/lib/Conversion/CMakeLists.txt

 add_subdirectory(AffineToStandard)
 add_subdirectory(AMDGPUToROCDL)
 add_subdirectory(ArithCommon)
 add_subdirectory(ArithToLLVM)
 add_subdirectory(ArithToSPIRV)
 add_subdirectory(ArmNeon2dToIntr)
 add_subdirectory(ArmSMEToSCF)
 add_subdirectory(AsyncToLLVM)
 add_subdirectory(BufferizationToMemRef)
 add_subdirectory(ComplexToLibm)
 add_subdirectory(ComplexToLLVM)
 add_subdirectory(ComplexToSPIRV)
 add_subdirectory(ComplexToStandard)
 add_subdirectory(ControlFlowToLLVM)
 add_subdirectory(ControlFlowToSPIRV)
+add_subdirectory(ConvertToLLVM)
 add_subdirectory(FuncToLLVM)
 add_subdirectory(FuncToSPIRV)
 add_subdirectory(GPUCommon)
 add_subdirectory(GPUToNVVM)
 add_subdirectory(GPUToROCDL)
 add_subdirectory(GPUToSPIRV)
 add_subdirectory(GPUToVulkan)
 add_subdirectory(IndexToLLVM)

mlir/lib/Conversion/ConvertToLLVM/CMakeLists.txt

+set(LLVM_OPTIONAL_SOURCES
+  ConvertToLLVMPass.cpp
+  ToLLVMInterface.cpp
+)
+
+add_mlir_conversion_library(MLIRConvertToLLVMInterface
+  ToLLVMInterface.cpp
+
+  DEPENDS
+
+  LINK_LIBS PUBLIC
+  MLIRIR
+  MLIRSupport
+)
+
+add_mlir_conversion_library(MLIRConvertToLLVMPass
+  ConvertToLLVMPass.cpp
+
+  DEPENDS
+  MLIRConversionPassIncGen
+
+  LINK_LIBS PUBLIC
+  MLIRConvertToLLVMInterface
+  MLIRPass
+  MLIRIR
+  MLIRSupport
+  )

mlir/lib/Conversion/ConvertToLLVM/ConvertToLLVMPass.cpp

+//===- ConvertToLLVMPass.cpp - MLIR LLVM Conversion -----------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Conversion/ConvertToLLVM/ToLLVMInterface.h"
+#include "mlir/Conversion/ConvertToLLVM/ToLLVMPass.h"
+#include "mlir/Conversion/LLVMCommon/ConversionTarget.h"
+#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/IR/PatternMatch.h"
+#include "mlir/Pass/Pass.h"
+#include "mlir/Rewrite/FrozenRewritePatternSet.h"
+#include "mlir/Transforms/DialectConversion.h"
+#include <memory>
+
+#define DEBUG_TYPE "convert-to-llvm"
+
+namespace mlir {
+#define GEN_PASS_DEF_CONVERTTOLLVMPASS
+#include "mlir/Conversion/Passes.h.inc"
+} // namespace mlir
+
+using namespace mlir;
+
+namespace {
+
+/// This DialectExtension can be attached to the context, which will invoke the
+/// `apply()` method for every loaded dialect. If a dialect implements the
+/// `ConvertToLLVMPatternInterface` interface, we load dependent dialects
+/// through the interface. This extension is loaded in the context before
+/// starting a pass pipeline that involves dialect conversion to LLVM.
+class LoadDependentDialectExtension : public DialectExtensionBase {
+public:
+  LoadDependentDialectExtension() : DialectExtensionBase(/*dialectNames=*/{}) {}
+
+  void apply(MLIRContext *context,
+             MutableArrayRef<Dialect *> dialects) const final {
+    LLVM_DEBUG(llvm::dbgs() << "Convert to LLVM extension load\n");
+    for (Dialect *dialect : dialects) {
+      auto iface = dyn_cast<ConvertToLLVMPatternInterface>(dialect);
+      if (!iface)
+        continue;
+      LLVM_DEBUG(llvm::dbgs() << "Convert to LLVM found dialect interface for "
+                              << dialect->getNamespace() << "\n");
+      iface->loadDependentDialects(context);
+    }
+  }
+
+  /// Return a copy of this extension.
+  virtual std::unique_ptr<DialectExtensionBase> clone() const final {
+    return std::make_unique<LoadDependentDialectExtension>(*this);
+  }
+};
+
+/// This is a generic pass to convert to LLVM, it uses the
+/// `ConvertToLLVMPatternInterface` dialect interface to delegate to dialects
+/// the injection of conversion patterns.
+class ConvertToLLVMPass
+    : public impl::ConvertToLLVMPassBase<ConvertToLLVMPass> {
+  std::shared_ptr<const FrozenRewritePatternSet> patterns;
+  std::shared_ptr<const ConversionTarget> target;
+
+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) {}
+
+  LogicalResult initialize(MLIRContext *context) final {
+    RewritePatternSet tempPatterns(context);
+    auto target = std::make_shared<ConversionTarget>(*context);
+    target->addLegalDialect<LLVM::LLVMDialect>();
+    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, tempPatterns);
+    }
+    patterns =
+        std::make_unique<FrozenRewritePatternSet>(std::move(tempPatterns));
+    this->target = target;
+    return success();
+  }
+
+  void runOnOperation() final {
+    if (failed(applyPartialConversion(getOperation(), *target, *patterns)))
+      signalPassFailure();
+  }
+};
+
+} // namespace

mlir/lib/Conversion/ConvertToLLVM/ToLLVMInterface.cpp

+//===- ToLLVMInterface.cpp - MLIR LLVM Conversion -------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Conversion/ConvertToLLVM/ToLLVMInterface.h"
+#include "mlir/IR/Dialect.h"
+#include "mlir/IR/Operation.h"
+#include "llvm/ADT/DenseSet.h"
+
+using namespace mlir;
+
+void mlir::populateConversionTargetFromOperation(Operation *root,
+                                                 ConversionTarget &target,
+                                                 RewritePatternSet &patterns) {
+  DenseSet<Dialect *> dialects;
+  root->walk([&](Operation *op) {
+    Dialect *dialect = op->getDialect();
+    if (!dialects.insert(dialect).second)
+      return;
+    // First time we encounter this dialect: if it implements the interface,
+    // let's populate patterns !
+    auto iface = dyn_cast<ConvertToLLVMPatternInterface>(dialect);
+    if (!iface)
+      return;
+    iface->populateConvertToLLVMConversionPatterns(target, patterns);
+  });
+}

mlir/lib/Conversion/NVVMToLLVM/NVVMToLLVM.cpp

 //===- NVVMToLLVM.cpp - NVVM to LLVM dialect conversion -----------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements a translation NVVM ops which is not supported in LLVM
 // core.
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Conversion/NVVMToLLVM/NVVMToLLVM.h"
 
+#include "mlir/Conversion/ConvertToLLVM/ToLLVMInterface.h"
 #include "mlir/Conversion/LLVMCommon/ConversionTarget.h"
 #include "mlir/Conversion/LLVMCommon/Pattern.h"
 #include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/Dialect/GPU/IR/GPUDialect.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/Dialect/LLVMIR/LLVMTypes.h"
 #include "mlir/Dialect/LLVMIR/NVVMDialect.h"
 #include "mlir/Dialect/MemRef/IR/MemRef.h"
void runOnOperation() override {
     RewritePatternSet pattern(&getContext());
     mlir::populateNVVMToLLVMConversionPatterns(pattern);
     if (failed(
             applyPartialConversion(getOperation(), target, std::move(pattern))))
       signalPassFailure();
   }
 };
 
+/// Implement the interface to convert NNVM to LLVM.
+struct NVVMToLLVMDialectInterface : public ConvertToLLVMPatternInterface {
+  using ConvertToLLVMPatternInterface::ConvertToLLVMPatternInterface;
+  void loadDependentDialects(MLIRContext *context) const final {
+    context->loadDialect<NVVMDialect>();
+  }
+
+  /// Hook for derived dialect interface to provide conversion patterns
+  /// and mark dialect legal for the conversion target.
+  void populateConvertToLLVMConversionPatterns(
+      ConversionTarget &target, RewritePatternSet &patterns) const final {
+    populateNVVMToLLVMConversionPatterns(patterns);
+  }
+};
+
 } // namespace
 
 void mlir::populateNVVMToLLVMConversionPatterns(RewritePatternSet &patterns) {
   patterns.add<PtxLowering>(patterns.getContext());
 }
+
+void mlir::registerConvertNVVMToLLVMInterface(DialectRegistry &registry) {
+  registry.addExtension(+[](MLIRContext *ctx, NVVMDialect *dialect) {
+    dialect->addInterfaces<NVVMToLLVMDialectInterface>();
+  });
+}

mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp

 //
 // The NVVM dialect only contains GPU specific additions on top of the general
 // LLVM dialect.
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/LLVMIR/NVVMDialect.h"
 
+#include "mlir/Conversion/ConvertToLLVM/ToLLVMInterface.h"
 #include "mlir/Dialect/Utils/StaticValueUtils.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/BuiltinAttributes.h"
 #include "mlir/IR/BuiltinTypes.h"
 #include "mlir/IR/DialectImplementation.h"
 #include "mlir/IR/MLIRContext.h"
 #include "mlir/IR/Operation.h"
 #include "mlir/IR/OperationSupport.h"
#include "mlir/Dialect/LLVMIR/NVVMOps.cpp.inc"
   addAttributes<
 #define GET_ATTRDEF_LIST
 #include "mlir/Dialect/LLVMIR/NVVMOpsAttributes.cpp.inc"
       >();
 
   // Support unknown operations because not all NVVM operations are
   // registered.
   allowUnknownOperations();
+  declarePromisedInterface<ConvertToLLVMPatternInterface>();
 }
 
 LogicalResult NVVMDialect::verifyOperationAttribute(Operation *op,
                                                     NamedAttribute attr) {
   StringAttr attrName = attr.getName();
   // Kernel function attribute should be attached to functions.
   if (attrName == NVVMDialect::getKernelFuncAttrName()) {
     if (!isa<LLVM::LLVMFuncOp>(op)) {

mlir/lib/IR/Dialect.cpp

 
 void DialectRegistry::applyExtensions(Dialect *dialect) const {
   MLIRContext *ctx = dialect->getContext();
   StringRef dialectName = dialect->getNamespace();
 
   // Functor used to try to apply the given extension.
   auto applyExtension = [&](const DialectExtensionBase &extension) {
     ArrayRef<StringRef> dialectNames = extension.getRequiredDialects();
+    // An empty set is equivalent to always invoke.
+    if (dialectNames.empty()) {
+      extension.apply(ctx, dialect);
+      return;
+    }
 
     // Handle the simple case of a single dialect name. In this case, the
     // required dialect should be the current dialect.
     if (dialectNames.size() == 1) {
       if (dialectNames.front() == dialectName)
         extension.apply(ctx, dialect);
       return;
     }
void DialectRegistry::applyExtensions(Dialect *dialect) const {
   for (const auto &extension : extensions)
     applyExtension(*extension);
 }
 
 void DialectRegistry::applyExtensions(MLIRContext *ctx) const {
   // Functor used to try to apply the given extension.
   auto applyExtension = [&](const DialectExtensionBase &extension) {
     ArrayRef<StringRef> dialectNames = extension.getRequiredDialects();
+    if (dialectNames.empty()) {
+      auto loadedDialects = ctx->getLoadedDialects();
+      extension.apply(ctx, loadedDialects);
+      return;
+    }
 
     // Check to see if all of the dialects for this extension are loaded.
     SmallVector<Dialect *> requiredDialects;
     requiredDialects.reserve(dialectNames.size());
     for (StringRef dialectName : dialectNames) {
       Dialect *loadedDialect = ctx->getLoadedDialect(dialectName);
       if (!loadedDialect)
         return;

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

 // RUN: mlir-opt --convert-nvvm-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
   llvm.return
 }