[mlir] Support lowering of dialect attributes attached to top-level modules

[skatrak]

Summary

This patch provides an alternative approach to D144883, with the goal of
supporting the processing of dialect attributes attached to top-level module-
type operations during MLIR-to-LLVMIR lowering.

Based on reviewer comments, this approach modifies the
mlir::translateModuleToLLVMIR() function to call
ModuleTranslation::convertOperation() on the top-level operation, after its
body has been lowered. This, in turn, will get the
LLVMTranslationDialectInterface object associated to that operation's dialect
before trying to use it for lowering prior to processing dialect attributes
attached to the operation.

Since there are no LLVMTranslationDialectInterfaces for the builtin and GPU
dialects, which define their own module-type operations, this patch also adds
and registers them. Their purpose is to just succeed when processing module
operations, to allow the lowering process to continue and to prevent the
introduction of failures related to not finding such interfaces.

The mentioned interface is added to the builtin dialect, differently to every
other dialect, on creation. Alternatively, it would require substantial changes
to be registered before any lowering to LLVM IR can be done.

[kiranchandramohan]

mlir/lib/IR/BuiltinDialect.cpp

115 ↗

(On Diff #504622)

I think this is best done in the Flang flow, so as to not make the builtin dialect llvm translation aware. I believe this is what @ftynse also recommended in https://reviews.llvm.org/D144883. We have done this previously for the OpenMP_PointerLikeType interface in https://github.com/llvm/llvm-project/blob/main/flang/include/flang/Tools/PointerModels.h.

We can also use the external model mechanism https://mlir.llvm.org/docs/Interfaces/#external-models-for-attribute-operation-and-type-interfaces so we don't need to make the builtin dialect aware of the LLVM translation interface.

mlir/lib/Target/LLVMIR/Dialect/GPU/GPUToLLVMIRTranslation.cpp

2

Nit: OpenMP -> GPU.

10–11

Nit: OpenMP -> GPU.

[skatrak]

mlir/lib/IR/BuiltinDialect.cpp

115 ↗

(On Diff #504622)

Thank you Kiran for the feedback and for pointing me to an example, as I wasn't sure exactly what that approach would look like in this case.

It makes sense that we wouldn't want the builtin dialect itself to know about an interface for translating it to LLVM IR, but it looks like we can't quite keep it localized to Flang though. The reason is that, after this change, any code intending to lower to LLVM IR must register a LLVMTranslationDialectInterface for at least the builtin dialect and the LLVM dialect. Otherwise, the call to translator.convertOperation(*module, llvmBuilder) in ModuleTranslation.cpp would fail for the builtin.module operation, aborting the translation process. Basically every MLIR lit test outside of Flang that invokes the translation into LLVM IR would fail because of this.

So, I'm currently thinking about incorporating feedback so far into this alternative solution, but I'd like to check whether this is an acceptable way to go:

1. Create a ModuleLikeOp interface, or maybe call it SkippableLLVMIRTranslationOp (listening for any suggestions on a name here), without any methods defined to be used as a tag.
2. Attach that interface to mlir::ModuleOp in mlir::registerLLVMDialectTranslation(), and attach the LLVMTranslationDialectInterface to the builtin dialect.
3. Modify mlir::registerGPUDialectTranslation() to do the same as steps 2 for the mlir::gpu::GPUModuleOp and mlir::gpu::GPUDialect, removing the GPUDialectLLVMIRTranslationInterface I created in this patch.
4. Modify LLVMTranslationDialectInterface::convertOperation() to succeed by default if the op has the ModuleLikeOp interface attached.

This way, by registering the interfaces in mlir::registerLLVMDialectTranslation(), which is called by every spot that needs lowering to LLVM IR, we avoid errors due to not finding these lowering interfaces for the builtin.module without making the dialect itself aware of it. Also, using the ModuleLikeOp interface or similar, we can skip defining an LLVM IR translation interface for dialects such as the GPU one that define their own module operation but are not (at least currently) intended to be lowered to LLVM IR directly. These are two separate ideas, so it could also be that I could just do one of them as well.

The alternative I see to step 2, to avoid mixing dialects in a spot where it might be better not to, would be to create an mlir::registerBuiltinDialectTranslation() that would need to be called in every place in which mlir::registerLLVMDialectTranslation() is called, which I see as very invasive. Maybe the compromise here would be to consolidate both calls into something like mlir::registerDefaultDialectsToLLVMTranslation(), and replace all calls to mlir::registerLLVMDialectTranslation() for it.

[ftynse]

mlir/lib/IR/BuiltinDialect.cpp

115 ↗

(On Diff #504622)

This looks overly complicated IMO. And we still don't want to be touching code in lib/IR for this.

Why can't we have the interface for the Builtin dialect be implemented the same way as any other dialect? Yes, the client will have to call the registration function, but it shouldn't be an issue given how many things they already regsiter.

[skatrak]

Implement translation from builtin dialect to LLVM IR like for all other dialects

[skatrak]

Fix error in comment

[skatrak]

mlir/lib/IR/BuiltinDialect.cpp

115 ↗

(On Diff #504622)

I just made the changes to have the interface for the Builtin dialect be implemented like for the others, as well as registering it by all potential clients to avoid errors during translation. The only reason to avoid doing it this way was that the changes following this approach seem more substantial and error-prone.

[ftynse]

Nice, thanks!

mlir/include/mlir/Target/LLVMIR/Dialect/Builtin/BuiltinToLLVMIRTranslation.h

1–2

Nit: this should be one line.

mlir/include/mlir/Target/LLVMIR/Dialect/GPU/GPUToLLVMIRTranslation.h

1–2

Ditto.

mlir/lib/Target/LLVMIR/Dialect/Builtin/BuiltinToLLVMIRTranslation.cpp

30

Nit: return success(isa<ModuleOp>(op)) is more concise.

[skatrak]

Address last reviewer's comments

[skatrak]

Thank you for all the suggestions!

flang/lib/Frontend/CMakeLists.txt

add_flang_library(flangFrontend
   FIRDialectSupport
   FIRSupport
   FIRBuilder
   FIRCodeGen
   FIRTransforms
   HLFIRDialect
   HLFIRTransforms
   MLIRTransforms
+  MLIRBuiltinToLLVMIRTranslation
   MLIRLLVMToLLVMIRTranslation
   MLIRSCFToControlFlow
   MLIRTargetLLVMIRImport
   ${dialect_libs}
 
   LINK_COMPONENTS
   Passes
   Analysis

flang/lib/Optimizer/CodeGen/CMakeLists.txt

add_flang_library(FIRCodeGen
   FIRDialectSupport
   FIRSupport
   MLIRComplexToLLVM
   MLIRComplexToStandard
   MLIRMathToFuncs
   MLIRMathToLLVM
   MLIRMathToLibm
   MLIROpenMPToLLVM
+  MLIRBuiltinToLLVMIRTranslation
   MLIRLLVMToLLVMIRTranslation
   MLIRTargetLLVMIRExport
 
   LINK_COMPONENTS
   AsmParser
   AsmPrinter
   Remarks
   TargetParser
 )

flang/lib/Optimizer/Dialect/CMakeLists.txt

add_flang_library(FIRDialect
   CanonicalizationPatternsIncGen
   MLIRIR
   FIROpsIncGen
   intrinsics_gen
 
   LINK_LIBS
   FIRDialectSupport
   MLIRArithDialect
+  MLIRBuiltinToLLVMIRTranslation
   MLIROpenMPToLLVM
   MLIRLLVMToLLVMIRTranslation
   MLIRTargetLLVMIRExport
 
   LINK_COMPONENTS
   AsmParser
   AsmPrinter
   Remarks
 )

flang/lib/Optimizer/Support/CMakeLists.txt

 get_property(dialect_libs GLOBAL PROPERTY MLIR_DIALECT_LIBS)
 
 add_flang_library(FIRSupport
   InitFIR.cpp
   InternalNames.cpp
 
   DEPENDS
   FIROpsIncGen
   HLFIROpsIncGen
   MLIRIR
   ${dialect_libs}
 
   LINK_LIBS
   ${dialect_libs}
+  MLIRBuiltinToLLVMIRTranslation
   MLIROpenMPToLLVMIRTranslation
   MLIRLLVMToLLVMIRTranslation
   MLIRTargetLLVMIRExport
 
   LINK_COMPONENTS
   TargetParser
 )

flang/lib/Optimizer/Support/InitFIR.cpp

 //===-- Optimizer/Support/InitFIR.cpp -------------------------------------===//
 //
 // 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 "flang/Optimizer/Support/InitFIR.h"
+#include "mlir/Target/LLVMIR/Dialect/Builtin/BuiltinToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.h"
 
 void fir::support::registerLLVMTranslation(mlir::MLIRContext &context) {
   mlir::DialectRegistry registry;
   // Register OpenMP dialect interface here as well.
-  mlir::registerOpenMPDialectTranslation(registry);
+  registerOpenMPDialectTranslation(registry);
   // Register LLVM-IR dialect interface.
   registerLLVMDialectTranslation(registry);
+  // Register builtin dialect interface.
+  registerBuiltinDialectTranslation(registry);
   context.appendDialectRegistry(registry);
 }

flang/tools/tco/CMakeLists.txt

target_link_libraries(tco PRIVATE
   FIRSupport
   FIRTransforms
   FIRBuilder
   HLFIRDialect
   HLFIRTransforms
   ${dialect_libs}
   MLIRIR
   MLIRLLVMDialect
+  MLIRBuiltinToLLVMIRTranslation
   MLIRLLVMToLLVMIRTranslation
   MLIRTargetLLVMIRExport
   MLIRPass
   MLIRFuncToLLVM
   MLIRTransforms
   MLIRAffineToStandard
   MLIRAnalysis
   MLIRSCFToControlFlow
   MLIRParser
   MLIRSupport
   MLIRVectorToLLVM
 )

mlir/examples/toy/Ch6/CMakeLists.txt

 include_directories(${CMAKE_CURRENT_BINARY_DIR}/include/)
 get_property(dialect_libs GLOBAL PROPERTY MLIR_DIALECT_LIBS)
 get_property(conversion_libs GLOBAL PROPERTY MLIR_CONVERSION_LIBS)
 target_link_libraries(toyc-ch6
   PRIVATE
     ${dialect_libs}
     ${conversion_libs}
     MLIRAnalysis
+    MLIRBuiltinToLLVMIRTranslation
     MLIRCallInterfaces
     MLIRCastInterfaces
     MLIRExecutionEngine
     MLIRIR
     MLIRLLVMCommonConversion
     MLIRLLVMDialect
     MLIRLLVMToLLVMIRTranslation
     MLIRMemRefDialect
     MLIRParser
     MLIRPass
     MLIRSideEffectInterfaces
     MLIRSupport
     MLIRTargetLLVMIRExport
     MLIRTransforms
     )

mlir/examples/toy/Ch6/toyc.cpp

 #include "mlir/IR/AsmState.h"
 #include "mlir/IR/BuiltinOps.h"
 #include "mlir/IR/MLIRContext.h"
 #include "mlir/IR/Verifier.h"
 #include "mlir/InitAllDialects.h"
 #include "mlir/Parser/Parser.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Pass/PassManager.h"
+#include "mlir/Target/LLVMIR/Dialect/Builtin/BuiltinToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Export.h"
 #include "mlir/Transforms/Passes.h"
 
 #include "llvm/ADT/StringRef.h"
 #include "llvm/IR/Module.h"
 #include "llvm/Support/CommandLine.h"
 #include "llvm/Support/ErrorOr.h"
if (!moduleAST)
     return 1;
 
   dump(*moduleAST);
   return 0;
 }
 
 int dumpLLVMIR(mlir::ModuleOp module) {
   // Register the translation to LLVM IR with the MLIR context.
+  mlir::registerBuiltinDialectTranslation(*module->getContext());
   mlir::registerLLVMDialectTranslation(*module->getContext());
 
   // Convert the module to LLVM IR in a new LLVM IR context.
   llvm::LLVMContext llvmContext;
   auto llvmModule = mlir::translateModuleToLLVMIR(module, llvmContext);
   if (!llvmModule) {
     llvm::errs() << "Failed to emit LLVM IR\n";
     return -1;

 
 int runJit(mlir::ModuleOp module) {
   // Initialize LLVM targets.
   llvm::InitializeNativeTarget();
   llvm::InitializeNativeTargetAsmPrinter();
 
   // Register the translation from MLIR to LLVM IR, which must happen before we
   // can JIT-compile.
+  mlir::registerBuiltinDialectTranslation(*module->getContext());
   mlir::registerLLVMDialectTranslation(*module->getContext());
 
   // An optimization pipeline to use within the execution engine.
   auto optPipeline = mlir::makeOptimizingTransformer(
       /*optLevel=*/enableOpt ? 3 : 0, /*sizeLevel=*/0,
       /*targetMachine=*/nullptr);
 
   // Create an MLIR execution engine. The execution engine eagerly JIT-compiles

mlir/examples/toy/Ch7/CMakeLists.txt

 include_directories(${CMAKE_CURRENT_BINARY_DIR}/include/)
 get_property(dialect_libs GLOBAL PROPERTY MLIR_DIALECT_LIBS)
 get_property(conversion_libs GLOBAL PROPERTY MLIR_CONVERSION_LIBS)
 target_link_libraries(toyc-ch7
   PRIVATE
     ${dialect_libs}
     ${conversion_libs}
     MLIRAnalysis
+    MLIRBuiltinToLLVMIRTranslation
     MLIRCallInterfaces
     MLIRCastInterfaces
     MLIRExecutionEngine
     MLIRIR
     MLIRLLVMCommonConversion
     MLIRLLVMToLLVMIRTranslation
     MLIRMemRefDialect
     MLIRParser
     MLIRPass
     MLIRSideEffectInterfaces
     MLIRTargetLLVMIRExport
     MLIRTransforms
     )

mlir/examples/toy/Ch7/toyc.cpp

 #include "mlir/IR/AsmState.h"
 #include "mlir/IR/BuiltinOps.h"
 #include "mlir/IR/MLIRContext.h"
 #include "mlir/IR/Verifier.h"
 #include "mlir/InitAllDialects.h"
 #include "mlir/Parser/Parser.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Pass/PassManager.h"
+#include "mlir/Target/LLVMIR/Dialect/Builtin/BuiltinToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Export.h"
 #include "mlir/Transforms/Passes.h"
 
 #include "llvm/ADT/StringRef.h"
 #include "llvm/IR/Module.h"
 #include "llvm/Support/CommandLine.h"
 #include "llvm/Support/ErrorOr.h"
if (!moduleAST)
     return 1;
 
   dump(*moduleAST);
   return 0;
 }
 
 int dumpLLVMIR(mlir::ModuleOp module) {
   // Register the translation to LLVM IR with the MLIR context.
+  mlir::registerBuiltinDialectTranslation(*module->getContext());
   mlir::registerLLVMDialectTranslation(*module->getContext());
 
   // Convert the module to LLVM IR in a new LLVM IR context.
   llvm::LLVMContext llvmContext;
   auto llvmModule = mlir::translateModuleToLLVMIR(module, llvmContext);
   if (!llvmModule) {
     llvm::errs() << "Failed to emit LLVM IR\n";
     return -1;

 
 int runJit(mlir::ModuleOp module) {
   // Initialize LLVM targets.
   llvm::InitializeNativeTarget();
   llvm::InitializeNativeTargetAsmPrinter();
 
   // Register the translation from MLIR to LLVM IR, which must happen before we
   // can JIT-compile.
+  mlir::registerBuiltinDialectTranslation(*module->getContext());
   mlir::registerLLVMDialectTranslation(*module->getContext());
 
   // An optimization pipeline to use within the execution engine.
   auto optPipeline = mlir::makeOptimizingTransformer(
       /*optLevel=*/enableOpt ? 3 : 0, /*sizeLevel=*/0,
       /*targetMachine=*/nullptr);
 
   // Create an MLIR execution engine. The execution engine eagerly JIT-compiles

mlir/include/mlir/Target/LLVMIR/Dialect/All.h

 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_TARGET_LLVMIR_DIALECT_ALL_H
 #define MLIR_TARGET_LLVMIR_DIALECT_ALL_H
 
 #include "mlir/Target/LLVMIR/Dialect/AMX/AMXToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Dialect/ArmNeon/ArmNeonToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Dialect/ArmSVE/ArmSVEToLLVMIRTranslation.h"
+#include "mlir/Target/LLVMIR/Dialect/Builtin/BuiltinToLLVMIRTranslation.h"
+#include "mlir/Target/LLVMIR/Dialect/GPU/GPUToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMIRToLLVMTranslation.h"
 #include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Dialect/OpenACC/OpenACCToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Dialect/X86Vector/X86VectorToLLVMIRTranslation.h"
 
 namespace mlir {
 class DialectRegistry;
 
 /// Registers all dialects that can be translated to LLVM IR and the
 /// corresponding translation interfaces.
 static inline void registerAllToLLVMIRTranslations(DialectRegistry &registry) {
   registerArmNeonDialectTranslation(registry);
   registerAMXDialectTranslation(registry);
   registerArmSVEDialectTranslation(registry);
+  registerBuiltinDialectTranslation(registry);
+  registerGPUDialectTranslation(registry);
   registerLLVMDialectTranslation(registry);
   registerNVVMDialectTranslation(registry);
   registerOpenACCDialectTranslation(registry);
   registerOpenMPDialectTranslation(registry);
   registerROCDLDialectTranslation(registry);
   registerX86VectorDialectTranslation(registry);
 }
 

mlir/include/mlir/Target/LLVMIR/Dialect/Builtin/BuiltinToLLVMIRTranslation.h

+//===- BuiltinToLLVMIRTranslation.h - Builtin Dialect to LLVM IR --*- 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 provides registration calls for builtin dialect to LLVM IR translation.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_TARGET_LLVMIR_DIALECT_BUILTIN_BUILTINTOLLVMIRTRANSLATION_H
+#define MLIR_TARGET_LLVMIR_DIALECT_BUILTIN_BUILTINTOLLVMIRTRANSLATION_H
+
+namespace mlir {
+
+class DialectRegistry;
+class MLIRContext;
+
+/// Register the translation from the builtin dialect to the LLVM IR in the
+/// given registry.
+void registerBuiltinDialectTranslation(DialectRegistry &registry);
+
+/// Register the translation from the builtin dialect in the registry associated
+/// with the given context.
+void registerBuiltinDialectTranslation(MLIRContext &context);
+
+} // namespace mlir
+
+#endif // MLIR_TARGET_LLVMIR_DIALECT_BUILTIN_BUILTINTOLLVMIRTRANSLATION_H

mlir/include/mlir/Target/LLVMIR/Dialect/GPU/GPUToLLVMIRTranslation.h

+//===- GPUToLLVMIRTranslation.h - GPU Dialect to LLVM IR --*- 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 provides registration calls for GPU dialect to LLVM IR translation.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_TARGET_LLVMIR_DIALECT_GPU_GPUTOLLVMIRTRANSLATION_H
+#define MLIR_TARGET_LLVMIR_DIALECT_GPU_GPUTOLLVMIRTRANSLATION_H
+
+namespace mlir {
+
+class DialectRegistry;
+class MLIRContext;
+
+/// Register the GPU dialect and the translation from it to the LLVM IR in
+/// the given registry;
+void registerGPUDialectTranslation(DialectRegistry &registry);
+
+/// Register the GPU dialect and the translation from it in the registry
+/// associated with the given context.
+void registerGPUDialectTranslation(MLIRContext &context);
+
+} // namespace mlir
+
+#endif // MLIR_TARGET_LLVMIR_DIALECT_GPU_GPUTOLLVMIRTRANSLATION_H

mlir/lib/CAPI/ExecutionEngine/CMakeLists.txt

 set(LLVM_LINK_COMPONENTS
   nativecodegen
   native
 )
   
 # Main API shared library.
 add_mlir_upstream_c_api_library(MLIRCAPIExecutionEngine
   ExecutionEngine.cpp
 
   LINK_LIBS PUBLIC
+  MLIRBuiltinToLLVMIRTranslation
   MLIRExecutionEngine
   MLIRLLVMToLLVMIRTranslation
 )

mlir/lib/CAPI/ExecutionEngine/ExecutionEngine.cpp

 //===- ExecutionEngine.cpp - C API for MLIR JIT ---------------------------===//
 //
 // 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-c/ExecutionEngine.h"
 #include "mlir/CAPI/ExecutionEngine.h"
 #include "mlir/CAPI/IR.h"
 #include "mlir/CAPI/Support.h"
 #include "mlir/ExecutionEngine/OptUtils.h"
+#include "mlir/Target/LLVMIR/Dialect/Builtin/BuiltinToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
 #include "llvm/ExecutionEngine/Orc/Mangling.h"
 #include "llvm/Support/TargetSelect.h"
 
 using namespace mlir;
 
 extern "C" MlirExecutionEngine
 mlirExecutionEngineCreate(MlirModule op, int optLevel, int numPaths,
                           const MlirStringRef *sharedLibPaths,
                           bool enableObjectDump) {
   static bool initOnce = [] {
     llvm::InitializeNativeTarget();
     llvm::InitializeNativeTargetAsmParser(); // needed for inline_asm
     llvm::InitializeNativeTargetAsmPrinter();
     return true;
   }();
   (void)initOnce;
 
-  mlir::registerLLVMDialectTranslation(*unwrap(op)->getContext());
+  auto &ctx = *unwrap(op)->getContext();
+  mlir::registerBuiltinDialectTranslation(ctx);
+  mlir::registerLLVMDialectTranslation(ctx);
 
   auto tmBuilderOrError = llvm::orc::JITTargetMachineBuilder::detectHost();
   if (!tmBuilderOrError) {
     llvm::errs() << "Failed to create a JITTargetMachineBuilder for the host\n";
     return MlirExecutionEngine{nullptr};
   }
   auto tmOrError = tmBuilderOrError->createTargetMachine();
   if (!tmOrError) {

mlir/lib/CAPI/RegisterEverything/CMakeLists.txt

 # Dialect registration.
 get_property(dialect_libs GLOBAL PROPERTY MLIR_DIALECT_LIBS)
 get_property(translation_libs GLOBAL PROPERTY MLIR_TRANSLATION_LIBS)
 get_property(conversion_libs GLOBAL PROPERTY MLIR_CONVERSION_LIBS)
 add_mlir_upstream_c_api_library(MLIRCAPIRegisterEverything
   RegisterEverything.cpp
 
   LINK_LIBS PUBLIC
   ${dialect_libs}
   ${translation_libs}
   ${conversion_libs}
 
+  MLIRBuiltinToLLVMIRTranslation
   MLIRCAPIIR
   MLIRLLVMToLLVMIRTranslation
   MLIRCAPITransforms
 )

mlir/lib/CAPI/RegisterEverything/RegisterEverything.cpp

 //===- RegisterEverything.cpp - Register all MLIR entities ----------------===//
 //
 // 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-c/RegisterEverything.h"
 
 #include "mlir/CAPI/IR.h"
 #include "mlir/InitAllDialects.h"
 #include "mlir/InitAllPasses.h"
+#include "mlir/Target/LLVMIR/Dialect/Builtin/BuiltinToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
 
 void mlirRegisterAllDialects(MlirDialectRegistry registry) {
   mlir::registerAllDialects(*unwrap(registry));
 }
 
 void mlirRegisterAllLLVMTranslations(MlirContext context) {
-  mlir::registerLLVMDialectTranslation(*unwrap(context));
+  auto &ctx = *unwrap(context);
+  mlir::registerBuiltinDialectTranslation(ctx);
+  mlir::registerLLVMDialectTranslation(ctx);
 }
 
 void mlirRegisterAllPasses() { mlir::registerAllPasses(); }

mlir/lib/Dialect/GPU/CMakeLists.txt

add_mlir_dialect_library(MLIRGPUTransforms
   DEPENDS
   MLIRGPUPassIncGen
   MLIRParallelLoopMapperEnumsGen
 
   LINK_LIBS PUBLIC
   MLIRAffineUtils
   MLIRArithDialect
   MLIRAsyncDialect
+  MLIRBuiltinToLLVMIRTranslation
   MLIRDataLayoutInterfaces
   MLIRExecutionEngineUtils
   MLIRGPUOps
   MLIRIR
   MLIRLLVMDialect
+  MLIRGPUToLLVMIRTranslation
   MLIRLLVMToLLVMIRTranslation
   MLIRMemRefDialect
   MLIRPass
   MLIRSCFDialect
   MLIRSideEffectInterfaces
   MLIRSupport
   MLIRTransformUtils
   )

mlir/lib/Dialect/GPU/Transforms/SerializeToBlob.cpp

 //===- SerializeToBlob.cpp - MLIR GPU lowering pass -----------------------===//
 //
 // 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 base class for a pass to serialize a gpu module
 // into a binary blob that can be executed on a GPU. The binary blob is added
 // as a string attribute to the gpu module.
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/GPU/Transforms/Passes.h"
 #include "mlir/Pass/Pass.h"
+#include "mlir/Target/LLVMIR/Dialect/GPU/GPUToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Export.h"
 #include "llvm/IR/LegacyPassManager.h"
 #include "llvm/MC/TargetRegistry.h"
 #include "llvm/Support/TargetSelect.h"
 #include "llvm/Target/TargetMachine.h"
 
 #include <string>
gpu::SerializeToBlobPass::optimizeLlvm(llvm::Module &llvmModule,
                                        llvm::TargetMachine &targetMachine) {
   // TODO: If serializeToCubin ends up defining optimizations, factor them
   // into here from SerializeToHsaco
   return success();
 }
 
 void gpu::SerializeToBlobPass::getDependentDialects(
     DialectRegistry &registry) const {
+  registerGPUDialectTranslation(registry);
   registerLLVMDialectTranslation(registry);
   OperationPass<gpu::GPUModuleOp>::getDependentDialects(registry);
 }
 
 std::unique_ptr<llvm::TargetMachine>
 gpu::SerializeToBlobPass::createTargetMachine() {
   Location loc = getOperation().getLoc();
   std::string error;

mlir/lib/ExecutionEngine/CMakeLists.txt

add_mlir_library(MLIRExecutionEngine
   Vectorize
   TransformUtils
   nativecodegen
   IPO
   Passes
   ${LLVM_JIT_LISTENER_LIB}
 
   LINK_LIBS PUBLIC
+  MLIRBuiltinToLLVMIRTranslation
   MLIRExecutionEngineUtils
   MLIRLLVMDialect
   MLIRLLVMToLLVMIRTranslation
   MLIRTargetLLVMIRExport
   )
 
 get_property(dialect_libs GLOBAL PROPERTY MLIR_DIALECT_LIBS)
 add_mlir_library(MLIRJitRunner

mlir/lib/Target/LLVMIR/CMakeLists.txt

 
 add_mlir_translation_library(MLIRToLLVMIRTranslationRegistration
   ConvertToLLVMIR.cpp
 
   LINK_LIBS PUBLIC
   MLIRArmNeonToLLVMIRTranslation
   MLIRArmSVEToLLVMIRTranslation
   MLIRAMXToLLVMIRTranslation
+  MLIRBuiltinToLLVMIRTranslation
+  MLIRGPUToLLVMIRTranslation
   MLIRX86VectorToLLVMIRTranslation
   MLIRLLVMToLLVMIRTranslation
   MLIRNVVMToLLVMIRTranslation
   MLIROpenACCToLLVMIRTranslation
   MLIROpenMPToLLVMIRTranslation
   MLIRROCDLToLLVMIRTranslation
   )
 

mlir/lib/Target/LLVMIR/Dialect/Builtin/BuiltinToLLVMIRTranslation.cpp

+//===- BuiltinToLLVMIRTranslation.cpp - Translate builtin to LLVM IR ------===//
+//
+// 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 between the MLIR builtin dialect and LLVM
+// IR.
+//
+//===----------------------------------------------------------------------===//
+#include "mlir/Target/LLVMIR/Dialect/Builtin/BuiltinToLLVMIRTranslation.h"
+#include "mlir/IR/BuiltinDialect.h"
+#include "mlir/IR/BuiltinOps.h"
+#include "mlir/Target/LLVMIR/LLVMTranslationInterface.h"
+
+using namespace mlir;
+
+namespace {
+
+class BuiltinDialectLLVMIRTranslationInterface
+    : public LLVMTranslationDialectInterface {
+public:
+  using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+
+  LogicalResult
+  convertOperation(Operation *op, llvm::IRBuilderBase &builder,
+                   LLVM::ModuleTranslation &moduleTranslation) const override {
+    return isa<ModuleOp>(op) ? success() : failure();
+  }
+};
+
+} // namespace
+
+void mlir::registerBuiltinDialectTranslation(DialectRegistry &registry) {
+  registry.addExtension(+[](MLIRContext *ctx, BuiltinDialect *dialect) {
+    dialect->addInterfaces<BuiltinDialectLLVMIRTranslationInterface>();
+  });
+}
+
+void mlir::registerBuiltinDialectTranslation(MLIRContext &context) {
+  DialectRegistry registry;
+  registerBuiltinDialectTranslation(registry);
+  context.appendDialectRegistry(registry);
+}

mlir/lib/Target/LLVMIR/Dialect/Builtin/CMakeLists.txt

+add_mlir_translation_library(MLIRBuiltinToLLVMIRTranslation
+  BuiltinToLLVMIRTranslation.cpp
+
+  LINK_COMPONENTS
+  Core
+
+  LINK_LIBS PUBLIC
+  MLIRIR
+  MLIRLLVMDialect
+  MLIRSupport
+  MLIRTargetLLVMIRExport
+  )

mlir/lib/Target/LLVMIR/Dialect/CMakeLists.txt

 add_subdirectory(ArmNeon)
 add_subdirectory(ArmSVE)
 add_subdirectory(AMX)
+add_subdirectory(Builtin)
+add_subdirectory(GPU)
 add_subdirectory(LLVMIR)
 add_subdirectory(NVVM)
 add_subdirectory(OpenACC)
 add_subdirectory(OpenMP)
 add_subdirectory(ROCDL)
 add_subdirectory(X86Vector)

mlir/lib/Target/LLVMIR/Dialect/GPU/CMakeLists.txt

+add_mlir_translation_library(MLIRGPUToLLVMIRTranslation
+  GPUToLLVMIRTranslation.cpp
+
+  LINK_COMPONENTS
+  Core
+
+  LINK_LIBS PUBLIC
+  MLIRIR
+  MLIRGPUOps
+  MLIRLLVMDialect
+  MLIRSupport
+  MLIRTargetLLVMIRExport
+  )

mlir/lib/Target/LLVMIR/Dialect/GPU/GPUToLLVMIRTranslation.cpp

+//===- GPUToLLVMIRTranslation.cpp - Translate GPU dialect to LLVM IR ------===//
+//
+// 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 between the MLIR GPU dialect and LLVM IR.
+//
+//===----------------------------------------------------------------------===//
+#include "mlir/Target/LLVMIR/Dialect/GPU/GPUToLLVMIRTranslation.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Target/LLVMIR/LLVMTranslationInterface.h"
+
+using namespace mlir;
+
+namespace {
+
+class GPUDialectLLVMIRTranslationInterface
+    : public LLVMTranslationDialectInterface {
+public:
+  using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+
+  LogicalResult
+  convertOperation(Operation *op, llvm::IRBuilderBase &builder,
+                   LLVM::ModuleTranslation &moduleTranslation) const override {
+    return isa<gpu::GPUModuleOp>(op) ? success() : failure();
+  }
+};
+
+} // namespace
+
+void mlir::registerGPUDialectTranslation(DialectRegistry &registry) {
+  registry.insert<gpu::GPUDialect>();
+  registry.addExtension(+[](MLIRContext *ctx, gpu::GPUDialect *dialect) {
+    dialect->addInterfaces<GPUDialectLLVMIRTranslationInterface>();
+  });
+}
+
+void mlir::registerGPUDialectTranslation(MLIRContext &context) {
+  DialectRegistry registry;
+  registerGPUDialectTranslation(registry);
+  context.appendDialectRegistry(registry);
+}

mlir/lib/Target/LLVMIR/ModuleTranslation.cpp

for (Operation &o : getModuleBody(module).getOperations()) {
     if (!isa<LLVM::LLVMFuncOp, LLVM::GlobalOp, LLVM::GlobalCtorsOp,
              LLVM::GlobalDtorsOp, LLVM::MetadataOp>(&o) &&
         !o.hasTrait<OpTrait::IsTerminator>() &&
         failed(translator.convertOperation(o, llvmBuilder))) {
       return nullptr;
     }
   }
 
+  // Convert module itself.
+  if (failed(translator.convertOperation(*module, llvmBuilder)))
+    return nullptr;
+
   if (llvm::verifyModule(*translator.llvmModule, &llvm::errs()))
     return nullptr;
 
   return std::move(translator.llvmModule);
 }

mlir/tools/mlir-cpu-runner/CMakeLists.txt

 set(LLVM_LINK_COMPONENTS
   Core
   Support
   nativecodegen
   native
   )
 
 add_mlir_tool(mlir-cpu-runner
   mlir-cpu-runner.cpp
   )
 llvm_update_compile_flags(mlir-cpu-runner)
 target_link_libraries(mlir-cpu-runner PRIVATE
   MLIRAnalysis
+  MLIRBuiltinToLLVMIRTranslation
   MLIRExecutionEngine
   MLIRIR
   MLIRJitRunner
   MLIRLLVMDialect
   MLIRLLVMToLLVMIRTranslation
   MLIRToLLVMIRTranslationRegistration
   MLIRParser
   MLIRTargetLLVMIRExport
   MLIRSupport
   )

mlir/tools/mlir-spirv-cpu-runner/CMakeLists.txt

if (MLIR_ENABLE_SPIRV_CPU_RUNNER)
   llvm_update_compile_flags(mlir-spirv-cpu-runner)
 
   get_property(conversion_libs GLOBAL PROPERTY MLIR_CONVERSION_LIBS)
 
   target_link_libraries(mlir-spirv-cpu-runner PRIVATE
     ${conversion_libs}
     MLIRAnalysis
     MLIRArithDialect
+    MLIRBuiltinToLLVMIRTranslation
     MLIRExecutionEngine
     MLIRFuncDialect
     MLIRGPUOps
     MLIRIR
     MLIRJitRunner
     MLIRLLVMDialect
     MLIRLLVMToLLVMIRTranslation
     MLIRMemRefDialect
     MLIRParser
     MLIRSPIRVDialect
     MLIRTargetLLVMIRExport
     MLIRTransforms
     MLIRTranslateLib
     MLIRSupport
   )
 endif()

mlir/tools/mlir-vulkan-runner/CMakeLists.txt

target_link_libraries(vulkan-runtime-wrappers
     ${Vulkan_LIBRARY}
   )
 
   get_property(conversion_libs GLOBAL PROPERTY MLIR_CONVERSION_LIBS)
   set(LIBS
     ${conversion_libs}
     MLIRAnalysis
     MLIRArithDialect
+    MLIRBuiltinToLLVMIRTranslation
     MLIRExecutionEngine
     MLIRFuncDialect
     MLIRGPUOps
     MLIRIR
     MLIRJitRunner
     MLIRLLVMDialect
     MLIRLLVMCommonConversion
     MLIRLLVMToLLVMIRTranslation

mlir/unittests/ExecutionEngine/Invoke.cpp

 #include "mlir/ExecutionEngine/CRunnerUtils.h"
 #include "mlir/ExecutionEngine/ExecutionEngine.h"
 #include "mlir/ExecutionEngine/MemRefUtils.h"
 #include "mlir/ExecutionEngine/RunnerUtils.h"
 #include "mlir/IR/MLIRContext.h"
 #include "mlir/InitAllDialects.h"
 #include "mlir/Parser/Parser.h"
 #include "mlir/Pass/PassManager.h"
+#include "mlir/Target/LLVMIR/Dialect/Builtin/BuiltinToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Export.h"
 #include "llvm/Support/TargetSelect.h"
 #include "llvm/Support/raw_ostream.h"
 
 #include "gmock/gmock.h"
 
 // SPARC currently lacks JIT support.
TEST(MLIRExecutionEngine, SKIP_WITHOUT_JIT(AddInteger)) {
   std::string moduleStr = R"mlir(
   func.func @foo(%arg0 : i32) -> i32 attributes { llvm.emit_c_interface } {
     %res = arith.addi %arg0, %arg0 : i32
     return %res : i32
   }
   )mlir";
   DialectRegistry registry;
   registerAllDialects(registry);
+  registerBuiltinDialectTranslation(registry);
   registerLLVMDialectTranslation(registry);
   MLIRContext context(registry);
   OwningOpRef<ModuleOp> module =
       parseSourceString<ModuleOp>(moduleStr, &context);
   ASSERT_TRUE(!!module);
   ASSERT_TRUE(succeeded(lowerToLLVMDialect(*module)));
   auto jitOrError = ExecutionEngine::create(*module);
   ASSERT_TRUE(!!jitOrError);
TEST(MLIRExecutionEngine, SKIP_WITHOUT_JIT(SubtractFloat)) {
   std::string moduleStr = R"mlir(
   func.func @foo(%arg0 : f32, %arg1 : f32) -> f32 attributes { llvm.emit_c_interface } {
     %res = arith.subf %arg0, %arg1 : f32
     return %res : f32
   }
   )mlir";
   DialectRegistry registry;
   registerAllDialects(registry);
+  registerBuiltinDialectTranslation(registry);
   registerLLVMDialectTranslation(registry);
   MLIRContext context(registry);
   OwningOpRef<ModuleOp> module =
       parseSourceString<ModuleOp>(moduleStr, &context);
   ASSERT_TRUE(!!module);
   ASSERT_TRUE(succeeded(lowerToLLVMDialect(*module)));
   auto jitOrError = ExecutionEngine::create(*module);
   ASSERT_TRUE(!!jitOrError);
TEST(NativeMemRefJit, SKIP_WITHOUT_JIT(ZeroRankMemref)) {
   func.func @zero_ranked(%arg0 : memref<f32>) attributes { llvm.emit_c_interface } {
     %cst42 = arith.constant 42.0 : f32
     memref.store %cst42, %arg0[] : memref<f32>
     return
   }
   )mlir";
   DialectRegistry registry;
   registerAllDialects(registry);
+  registerBuiltinDialectTranslation(registry);
   registerLLVMDialectTranslation(registry);
   MLIRContext context(registry);
   auto module = parseSourceString<ModuleOp>(moduleStr, &context);
   ASSERT_TRUE(!!module);
   ASSERT_TRUE(succeeded(lowerToLLVMDialect(*module)));
   auto jitOrError = ExecutionEngine::create(*module);
   ASSERT_TRUE(!!jitOrError);
   auto jit = std::move(jitOrError.get());
func.func @one_ranked(%arg0 : memref<?xf32>) attributes { llvm.emit_c_interface } {
     %cst42 = arith.constant 42.0 : f32
     %cst5 = arith.constant 5 : index
     memref.store %cst42, %arg0[%cst5] : memref<?xf32>
     return
   }
   )mlir";
   DialectRegistry registry;
   registerAllDialects(registry);
+  registerBuiltinDialectTranslation(registry);
   registerLLVMDialectTranslation(registry);
   MLIRContext context(registry);
   auto module = parseSourceString<ModuleOp>(moduleStr, &context);
   ASSERT_TRUE(!!module);
   ASSERT_TRUE(succeeded(lowerToLLVMDialect(*module)));
   auto jitOrError = ExecutionEngine::create(*module);
   ASSERT_TRUE(!!jitOrError);
   auto jit = std::move(jitOrError.get());
func.func @rank2_memref(%arg0 : memref<?x?xf32>, %arg1 : memref<?x?xf32>) attributes { llvm.emit_c_interface } {
     %cst42 = arith.constant 42.0 : f32
     memref.store %cst42, %arg0[%y, %x] : memref<?x?xf32>
     memref.store %cst42, %arg1[%x, %y] : memref<?x?xf32>
     return
   }
   )mlir";
   DialectRegistry registry;
   registerAllDialects(registry);
+  registerBuiltinDialectTranslation(registry);
   registerLLVMDialectTranslation(registry);
   MLIRContext context(registry);
   OwningOpRef<ModuleOp> module =
       parseSourceString<ModuleOp>(moduleStr, &context);
   ASSERT_TRUE(!!module);
   ASSERT_TRUE(succeeded(lowerToLLVMDialect(*module)));
   auto jitOrError = ExecutionEngine::create(*module);
   ASSERT_TRUE(!!jitOrError);
TEST(NativeMemRefJit, MAYBE_JITCallback) {
   func.func @caller_for_callback(%arg0: memref<?x?xf32>, %coefficient: i32) attributes { llvm.emit_c_interface } {
     %unranked = memref.cast %arg0: memref<?x?xf32> to memref<*xf32>
     call @callback(%arg0, %coefficient) : (memref<?x?xf32>, i32) -> ()
     return
   }
   )mlir";
   DialectRegistry registry;
   registerAllDialects(registry);
+  registerBuiltinDialectTranslation(registry);
   registerLLVMDialectTranslation(registry);
   MLIRContext context(registry);
   auto module = parseSourceString<ModuleOp>(moduleStr, &context);
   ASSERT_TRUE(!!module);
   ASSERT_TRUE(succeeded(lowerToLLVMDialect(*module)));
   auto jitOrError = ExecutionEngine::create(*module);
   ASSERT_TRUE(!!jitOrError);
   auto jit = std::move(jitOrError.get());